Computing motion of pixels in an image

Abstract

Apparatuses, systems, and techniques for computing motion of one or more pixels in a region in an image. In at least one embodiment, motion is computed based on motion of one or more pixels in different regions of the image that overlap the region, with one or more algorithms represented in CUDA code, e.g., for efficient execution on a GPU.

Description

[0001] CLAIM

[0002] This application incorporates by reference the entirety of [[U.S. Application Number Not Yet Assigned]] titled “SUPERPIXEL GENERATION AND USE” filed on April 5, 2021. TECHNICAL FIELD

[0003] At least one embodiment relates to processing resources for performing one or more compute unified device architecture (CUDA) programs. For example, at least one embodiment relates to a processor or computing system for performing one or more CUDA programs that compute motion of one or more pixels in one or more images according to various new techniques described herein. BACKGROUND

[0004] In various cases, it can be difficult to compute motion of pixels in an image. Often, neural networks are used to compute motion of pixels from an image, but the computational cost of using neural networks can be high. Using CUDA code to configure a GPU to perform computational operations in parallel can improve the amount of memory, time, or computational resources used to compute motion of pixels in an image. BRIEF DESCRIPTION OF DRAWINGS

[0005] FIG. 1 An example framework to compute motion of one or more pixels in a plurality of regions of an image is shown in accordance with at least one embodiment;

[0006] FIG. 2 A block diagram of a system to group one or more pixels in an image together based on shared features is shown in accordance with at least one embodiment;

[0007] FIG. 3 An example to group one or more pixels in an image based on shared features is shown in accordance with at least one embodiment;

[0008] FIG. 4 An example to interpolate optical flow of one or more pixels in a region of an image based on one or more pixels in an overlapping region is shown in accordance with at least one embodiment;

[0009] FIG. 5 An example to interpolate optical flow of one or more pixels in a first region and a second overlapping region of an image is shown in accordance with at least one embodiment;

[0010] FIG. 6 An example to interpolate optical flow of one or more pixels by scanning an image and one or more rows and one or more columns of a cost map is shown in accordance with at least one embodiment;

[0011] FIG. 7 An example of updating one or more pixels in an image using a cost map is shown in accordance with at least one embodiment;

[0012] FIG. 8 An example of sparse interpolation to dense interpolation of optical flow is shown in accordance with at least one embodiment;

[0013] FIG. 9 An example of a process for a framework to perform sparse interpolation to dense interpolation of optical flow is shown in accordance with at least one embodiment;

[0014] FIG. 10 An example data center is shown in accordance with at least one embodiment;

[0015] FIG. 11 A processing system is shown in accordance with at least one embodiment;

[0016] FIG. 12 A computer system is shown in accordance with at least one embodiment;

[0017] FIG. 13 A system is shown in accordance with at least one embodiment;

[0018] FIG. 14 An example integrated circuit is shown in accordance with at least one embodiment;

[0019] FIG. 15 A computing system is shown in accordance with at least one embodiment;

[0020] FIG. 16 An APU is shown in accordance with at least one embodiment;

[0021] FIG. 17 A CPU is shown in accordance with at least one embodiment;

[0022] FIG. 18 An example accelerator integration slice is shown in accordance with at least one embodiment;

[0023] FIG. 19A And 19B An example graphics processor is shown in accordance with at least one embodiment;

[0024] FIG. 20A A graphics core is shown in accordance with at least one embodiment;

[0025] FIG. 20B A GPGPU is shown in accordance with at least one embodiment;

[0026] FIG. 21A A parallel processor is shown in accordance with at least one embodiment;

[0027] FIG. 21B A processing cluster is shown in accordance with at least one embodiment;

[0028] FIG. 21C A graphics multiprocessor is shown in accordance with at least one embodiment;

[0029] FIG. 22 A graphics processor is shown in accordance with at least one embodiment;

[0030] FIG. 23 A processor is shown in accordance with at least one embodiment;

[0031] FIG. 24 A processor is shown in accordance with at least one embodiment;

[0032] FIG. 25 A graphics processor core is shown in accordance with at least one embodiment;

[0033] FIG. 26 A PPU is shown in accordance with at least one embodiment;

[0034] FIG. 27 A GPC is shown in accordance with at least one embodiment;

[0035] FIG. 28 A streaming multiprocessor is shown in accordance with at least one embodiment;

[0036] FIG. 29 A software stack of a programming platform is shown in accordance with at least one embodiment;

[0037] FIG. 30 A CUDA implementation of the software stack of FIG. 29 is shown in accordance with at least one embodiment;

[0038] FIG. 31 A ROCm implementation of the software stack of FIG. 29 is shown in accordance with at least one embodiment;

[0039] FIG. 32 An OpenCL implementation of the software stack of FIG. 29 is shown in accordance with at least one embodiment;

[0040] FIG. 33 Software supported by a programming platform is shown in accordance with at least one embodiment;

[0041] FIG. 34 Compiled code executed on a programming platform of FIGS. 29-32 is shown in accordance with at least one embodiment;

[0042] FIG. 35 Compiled code executed onFIGS. 29-32 More detailed compiled code executed on the programming platform;

[0043] FIG. 36 This illustrates the transformation of source code before compilation, according to at least one embodiment;

[0044] FIG. 37A A system configured to compile and execute CUDA source code using different types of processing units, according to at least one embodiment, is shown;

[0045] FIG. 37B The diagram illustrates a configuration, according to at least one embodiment, to compile and execute using a CPU and a CUDA-enabled GPU. FIG. 37A The system of CUDA source code;

[0046] FIG. 37C The diagram illustrates a configuration, according to at least one embodiment, to compile and execute using a CPU and a GPU with CUDA disabled. FIG. 37A The system of CUDA source code;

[0047] FIG. 38 The diagram illustrates a method according to at least one embodiment. FIG. 37C An example kernel converted by the CUDA to HIP conversion tool;

[0048] FIG. 39 A more detailed description is provided according to at least one embodiment. FIG. 37C GPUs without CUDA enabled;

[0049] FIG. 40 This illustrates how threads of an exemplary CUDA grid, according to at least one embodiment, are mapped to... FIG. 39 Different computing units; and

[0050] FIG. 41 This paper illustrates how to migrate existing CUDA code to data-parallel C++ code according to at least one embodiment. Detailed Implementation

[0051] In the following description, numerous specific details are set forth to provide a more thorough understanding of at least one embodiment. However, it will be apparent to those skilled in the art that the concepts of the invention can be practiced without one or more of these specific details.

[0052] In at least one embodiment, one or more circuits, which can be part of one or more processors in a computer system, compute motion of one or more pixels of an image. In at least one embodiment, the one or more circuits compute motion of one or more pixels in a first region of an image using motion information of one or more pixels in a second region that overlaps the first region of the image. In at least one embodiment, the one or more circuits compute motion of one or more pixels of an image by obtaining an image and a flow graph such that one or more pixels are grouped together, where the group of pixels share one or more characteristics corresponding to at least one of intensity and gradient, and iteratively convert one or more invalid pixels in the group of pixels to valid pixels using one or more valid pixels from the group of pixels based on the flow graph.

[0053] In at least one embodiment, optical flow (e.g., an optical flow graph) is used in computer vision tasks (e.g., image segmentation and object detection) to demonstrate motion of objects between consecutive frames of a video caused by relative motion between objects and cameras. While it is somewhat feasible to perform optical flow and improve resolution of tracked objects in video frames, doing so often requires a system to perform object detection and semantic segmentation, which is difficult to implement in parallel computation and requires multiple iterations. Moreover, in computing optical flow vectors (i.e., vectors representing movement of pixels from one frame to the next), often there are many pixels for which vectors cannot be computed in a straightforward manner. This is due to changes in lighting conditions between frames, objects being covered by other objects, and other factors.

[0054] In at least one embodiment, a framework is created to compute missing optical flow vectors by filling in the missing optical flow vectors using previously computed optical flow vectors. In at least one embodiment, a first image and a sparse optical flow map corresponding to the first image and a second image (e.g., two subsequent frames in a video) are obtained. In at least one embodiment, the framework can then select a candidate pixel from one or more pixels of an object in the first image and cause a group of neighboring pixels in a first region of the first image to be grouped together (e.g., form a superpixel) representing the object. In at least one embodiment, the group of pixels is determined based on how each of these pixels is similar in gradient and intensity to the candidate pixel. In at least one embodiment, the framework determines whether a pixel from the group of pixels is a valid flow pixel (e.g., valid pixel) or an invalid flow pixel (e.g., invalid pixel) based on the sparse optical flow map by applying a cost function. In at least one embodiment, a valid flow pixel moves from a first frame (e.g., first image) to a subsequent frame (e.g., second image) in the same way as when the pixel moves backward from the subsequent frame back to the first frame, while an invalid flow pixel does not. In at least one embodiment, the framework knows which pixels in the group of pixels are not part of the object (e.g., edge pixels) in order to exclude the edge pixels from the computation. In at least one embodiment, the framework then uses at least one valid flow pixel to fix neighboring pixels with invalid flow by interpolation. In at least one embodiment, in multiple iterations over different regions in the image, a new candidate pixel is selected and its neighboring optical flow pixels are interpolated based on at least one valid flow pixel (determined based on shared features such as gradient and / or intensity). In at least one embodiment, neighboring pixels are determined based on depth of the image. In at least one embodiment, neighboring pixels are determined based on at least one of: depth of the image, gradient, and intensity. In at least one embodiment, after multiple iterations, the sparse optical flow map is modified to a dense optical flow map. In at least one embodiment, a median filter is applied to the dense optical flow map for refinement.

[0055] In at least one embodiment, the techniques described herein achieve various technical advantages, including but not limited to improvements in computational performance (e.g., speed, processing power, and memory consumption), and are also configured to be compatible with parallel computing. In at least one embodiment, the techniques described herein use a framework to address problems related to sparse-to-dense interpolation in real-time optics. Dense optical flow is used in applications such as frame rate upconversion (accelerating clips to 2x, 4x, 8x speeds), video coding, robot navigation, and / or video classification for data center applications. In some frameworks, sparse-to-dense interpolation is performed by oversegmenting the entire scene / image, using a Random Sample Consensus (RANSAC) flow process to estimate the model for each segment, and further using model propagation to refine the flow graph. In at least one embodiment, the techniques described herein segment around invalid locations, perform iterative nearest neighbor and scanline filling, and further use a median filter to refine the flow graph.

[0056] In at least one embodiment, instead of using image segmentation, model generation and propagation for each segment, and variational refinement to achieve sparse-to-dense interpolation, the technique described herein uses a framework to fix all invalid flow locations and replace them with valid flows in a manner compatible with parallel computing. In at least one embodiment, the framework uses CUDA code to perform sparse-to-dense interpolation (e.g., calculating the motion of pixels in the image based on other pixels in the image) to configure the GPU to perform computational operations related to the computational motion of pixels in the image in parallel.

[0057] FIG. 1 An example framework 100 is illustrated for calculating the motion of one or more pixels in a plurality of regions of an image according to at least one embodiment. In at least one embodiment, the example framework 100 includes one or more computing devices that receive one or more input data (e.g., image 1, image 2, sparse flow graph, gradient of image 1) via a network to perform sparse-to-dense interpolation. In at least one embodiment, the one or more computing devices include one or more processors (e.g., GPUs supporting Computing Unified Device Architecture (CUDA)) configured to execute CUDA code to calculate the motion of one or more pixels in an image (e.g., image 1). In at least one embodiment, the framework 100 is a single computing device configured to execute CUDA code to calculate the motion of one or more pixels in the image. In at least one embodiment, the components of the framework 100 can communicate with each other using any of a variety of different network types, including but not limited to local area networks (LANs) and / or wide area networks (WANs), via wired and / or wireless communication protocols.

[0058] In at least one embodiment, one or more computing devices receive input data. In at least one embodiment, input data includes one or more images (e.g., image 1 and image 2) and a sparse flow map. In at least one embodiment, a sparse optical flow map is computed between image 1 and image 2. In at least one embodiment, a sparse optical flow map includes a sparse set of features of pixels that track the velocity vectors (e.g., motion) of the pixels. In at least one embodiment, an optical flow map can be generated using two images and computed using the Lucas Kanade method, which provides an estimate of the movement of features of interest in successive images of a scene. In at least one embodiment, the sparse optical flow map is the same shape as the image in image 1 and image 2. In at least one embodiment, in a sparse flow map, there are some locations (coordinates) that have valid flow pixels and some that have invalid flow pixels. In at least one embodiment, invalid flow pixels have an invalid value (MAX POSSIBLE_FLOW + 1).

[0059] In at least one embodiment, one or more computing devices execute one or more programs (e.g., CUDA or other such programs) at least in part to compute motion of one or more pixels in a first region of an image based at least in part on motion of one or more pixels in a second region of the image that overlaps the first region of the image. In at least one embodiment, one or more computing devices execute CUDA code to determine a set of pixels (sometimes referred to herein as superpixels) in a first region of an image. In at least one embodiment, one or more computing devices receive input data and generate at least one superpixel 102. In at least one embodiment, a framework for generating superpixels, also referred to herein as Edgeaware superpixels, reads a sparse flow map and image 1. In at least one embodiment, one or more computing devices perform one or more operations to generate superpixels for invalid flow locations (by grouping pixels that share common features such as intensity and gradients). In at least one embodiment, a sparse flow map is used to identify invalid flow locations and image 1 is used for superpixel generation.

[0060] In at least one embodiment, a framework for interpolating flow vectors for each pixel in an image, also referred to herein as performing Edgeaware-Infill 104, reads the superpixels for each invalid location and checks for validity of the superpixels (a superpixel is valid if it contains a subset of pixels with valid flow locations). In at least one embodiment, the best flow between valid superpixels is chosen by inflling_cost = (0.2 * pixel_distance + 0.4 * pixel_gradient_diff + 0.4 * pixel_intensity_diff). In at least one embodiment, the flow vector that minimizes a cost function such as “infilling_cost” is chosen. In at least one embodiment, performing Edgeaware-Infill 104 includes iteratively interpolating flow vectors 106 in the image. In at least one embodiment, the number of iterations is fixed at 4; however, more or fewer iterations can also be used. In at least one embodiment, each iteration fixes some points / superpixels, which ultimately helps the next iteration to fix other invalid points / superpixels. In at least one embodiment, Edgeaware-Infill can not fix all locations, but it helps to reduce the density of invalid pixels. In at least one embodiment, one or more computing devices perform detection 108 of whether one or more pixels in the image are edge pixels (e.g., edge pixels are pixels that do not represent an object). In at least one embodiment, Edgeaware-Infill 104 is performed by excluding any detected pixels that are edge pixels. In at least one embodiment, one or more computing devices perform flow scanline fill 110. In at least one embodiment, a valid vector is dragged in four directions (left, right, up, down) until another valid point is detected. In at least one embodiment, once another valid point is detected, neighboring pixels with invalid locations can be fixed based on the detected valid points during the traversal. In at least one embodiment, flow vectors are chosen among multiple flows based on a cost function “infilling_cost” as described above. In at least one embodiment, the output of this step is a dense flow map. In at least one embodiment, one or more computing devices run a 5x5 median filter 112 on the dense flow map for refinement. In at least one embodiment, other sizes and / or filters besides a 5x5 median filter can also be used.

[0061] FIG. 2A block diagram illustrating a system 200 that groups one or more pixels in an image together based on shared features, in accordance with at least one embodiment, is shown. In at least one embodiment, a pixel grouping (e.g., superpixel generation) technique is performed by system 200 that traverses neighboring pixels starting from a center location until an edge point or a specified boundary (e.g., a 2S x 2S window boundary, where S is a window radius) is reached. In at least one embodiment, by stopping at an edge point, computation can be reduced. In at least one embodiment, system 200 is executed by a computing system or by a vision system having a GPU (graphics processing unit), a CPU (central processing unit). In at least one embodiment, system 200 can include hardware (e.g., a processing device, a circuit, a special purpose logic, programmable logic, microcode, a hardware of a device, an integrated circuit, etc.), software (e.g., a device running or executing instructions on a processing device), or a combination thereof.

[0062] In at least one embodiment, system 200 includes one or more modules: down-samplers 201-203, forward flow generator 205, backward flow generator 207, flow comparator 209, invalid flow coordinate extractor / compressor 211, edge detector 213, superpixel generator 215, and up-samplers 217. In at least one embodiment, some modules 201-217 are integrated modules. In at least one embodiment, forward flow generator 205 and backward flow generator 207 can be integrated modules. In at least one embodiment, modules include one or more electronic circuits that are located on a circuit board that provide functionality within a computer. In at least one embodiment, system 200 includes, without limitation, one or more processors communicatively coupled to any suitable number or kind of components, peripherals, modules, or devices.

[0063] In at least one embodiment, down-samplers 201-203 can receive images of a sequence of images (I1’ and I2’) and down-sample the images by a scale factor to generate images (I1 and I2). In at least one embodiment, images I1’ and I2’ are digital images and / or color images, such as RGB (red-green-blue) color images. In at least one embodiment, images I1’ and I2’ can be still images, video images, or real-time captured images from a vision system. In at least one embodiment, down-sampling an image scales the image to a manageable size by reducing pixel size of the image. In at least one embodiment, if an image is scaled less than one, then traversing neighboring pixels to generate superpixels as discussed further below can use pixels that are far from a center of a superpixel. In at least one embodiment, if an image has a size of 800x800 pixels, then down-sampling by a factor of 1 / 2 provides an image with 400x400 pixels. In at least one embodiment, an image can be down-sampled by periodically dropping some rows / columns from the image. In at least one embodiment, other techniques for down-sampling can include bicubic down-sampling, average down-sampling, etc.

[0064] In at least one embodiment, forward flow generator 205 can receive images (I1 and I2, or I1’ and I2’ if down-sampling is not performed) and obtain forward optical flow data (F12) using images I1 and I2 (or I1’ and I2’). In at least one embodiment, optical flow data (also referred to as flow data) includes optical flow vectors (also referred to as flow vectors) that indicate direction and magnitude of movement of image pixel data from a first image to a second image. In at least one embodiment, optical flow data indicates direction and magnitude of movement of image pixel data in a two-dimensional (2D) image plane. In at least one embodiment, pixel movement is tracked by assuming that a pixel maintains the same intensity and / or color in two image frames and moves by a small amount (e.g., a few pixels). In at least one embodiment, a particular color / intensity of a particular pixel or object edge in a first frame can move to a different pixel in a second frame. In at least one embodiment, the color / intensity is associated with an object in a scene that moves between the first and second image frames. In at least one embodiment, a camera position used to capture the scene can change (pan) between the first and second frames, such that the object appears to move between the image frames. In at least one embodiment, forward optical flow data represents optical flow vectors from I1 to I2 (or I1’ to I2’), where I1 (and I1’) is an image at time t and I2 (and I2’) is an image at time t+1.

[0065] In at least one embodiment, backward generator 207 receives images I1 and I2 (or I1’ and I2’) and obtains backward optical flow data (F21) using images I1 and I2 (or I1’ and I2’). In at least one embodiment, backward optical flow data represents optical flow vectors from I2 to I1 (or I2’ to I1’), where I1 (and I1’) is an image at time t and I2 (and I2’) is an image at time t+1.2-1 )) In at least one embodiment, the backward optical flow data represents the optical flow vector from I2 to I1 (or I2' to I1'), where I1 and I1' are images at time t, and I2 and I2' are images at time t+1.

[0066] In at least one embodiment, the forward and backward optical flow data are obtained using forward-backward constraint equations, assuming that the intensity of a pixel is constant, e.g., the intensity of a pixel is not affected by the motion of the pixel. In at least one embodiment, the continuity equation for a pixel at (x, y) assumes constant intensity as follows:

[0067] I(x, y, t) - I(x + dx, y + dy, t + dt) = 0 (1)

[0068] where I(x, y, t) is the intensity of the pixel at time t, and I(x + dx, y + dy, t + dt) is the intensity of the same pixel assumed to be at (x + dx, y + dy) at time t + dt.

[0069] From equation (1), the forward and backward constraint motion equations can be expressed as:

[0070] I t (x, y) = I t+dt (x + dx2, y + dy2), and (2)

[0071] I t+dt (x, y) = I t (x - dx1, y - dy1). (3)

[0072] The Taylor series expansion of equations (2) and (3) provides:

[0073]

[0074]

[0075] Equations (4) and (5) can be rewritten as:

[0076]

[0077]

[0078] or,

[0079]

[0080]

[0081] where and

[0082] Here, and are the derivatives of the image at (x, y, t); while and are the derivatives of the image at (x, y, t+1), (u2, v2) is the forward motion vector, and (ul, vl) is the backward motion vector.

[0083] Equations (8) and (9) can be simply rewritten as:

[0084]

[0085]

[0086] In at least one embodiment, equations (10) and (11) can be solved by numerical methods to estimate the forward and backward motion (or flow) vectors, where the forward motion vector represents the amount a pixel has moved from the first image (time t) to the second image (time t+dt), and the backward motion vector represents the amount a pixel has moved from the second image (time t+dt) to the first image (time t).

[0087] In at least one embodiment, the flow comparator 209 receives the forward and backward optical flow data and performs a comparison between the forward and backward optical flow data for each optical flow vector in the forward flow data to determine consistent and inconsistent optical flow data. In at least one embodiment, consistent optical flow data has flow vectors that match (or approximately match) the field magnitude and field direction between the forward optical flow data F 1-2 and the backward optical flow data F 2-1 In at least one embodiment, inconsistent flow data has at least a mismatch in at least one of the field magnitude or field direction between the forward optical flow data F 1-2 and the backward optical flow data F 2-1

[0088] In at least one embodiment, due to various reasons (e.g., false or edge detection of objects, such as occlusions, blurry images, noisy images, or thin objects), the flow vectors in the backward optical data F 2-1 compared to the forward optical data F 1-2 ​may not be consistent. For example, for occlusion, if there is no object in time t frame II, and the object enters frame I2 at time t+1, the flow vector corresponding to the object from frame II to I2 can not be consistent with the flow vector corresponding to the object from frame I2 to II because II lacks pixel information of the object. Thus, the forward and backward optical flow data are two independent sets of optical flow information and can not be consistent. Optical flow data that is inconsistent between the forward optical flow data and the backward optical flow data is unreliable.

[0089] In at least one embodiment, the invalid optical flow data can be replaced or “filled-in” with valid optical flow data from surrounding pixels having similar features. In at least one embodiment, the surrounding pixels having similar features are pixels in the same superpixel.

[0090] In at least one embodiment, the invalid flow coordinate extractor / compressor 211 extracts or compresses pixel coordinates corresponding to optical flow locations having unreliable / inconsistent optical flow data (e.g., invalid optical flow coordinates). In at least one embodiment, the invalid flow coordinate extractor / compressor 211 generates an invalid optical flow coordinate map including pixel coordinates associated with invalid optical flow data. In at least one embodiment, the invalid optical flow coordinate map can include (x, y) coordinates of pixel locations corresponding to invalid optical flow data. In at least one embodiment, the invalid flow coordinate extractor / compressor 211 performs parallel flow compression to remove valid flow coordinates and keep invalid flow coordinates.

[0091] In at least one embodiment, the edge detector 213 receives the image (II or I1’) and performs edge detection on the image II or I1’ to generate a gradient map or other edge detection map. In at least one embodiment, the gradient map is generated using a Laplacian operator of a Gaussian edge detection algorithm.

[0092] In at least one embodiment, the superpixel generator 215 receives the gradient map or other edge detection data, the invalid flow coordinates (e.g., the map of invalid optical flow coordinates), and the images II and I2 (or images I1’ and I2’) to generate one or more superpixels. In at least one embodiment, a superpixel corresponds to a group of cells sharing one or more common features (e.g., pixel intensity). In at least one embodiment, the superpixel generator 215 outputs one or more images (O1) including superpixels. In at least one embodiment, O1 is a two-dimensional (2D) array structure whose coordinates correspond to the pixel size of the input image II (or I1’). In at least one embodiment, each cell in the 2D array contains an integer value identifying the superpixel to which the cell belongs.

[0093] In at least one embodiment, superpixel generator 215 uses the invalid flow coordinates as the center / candidate coordinates for superpixel generation. In at least one embodiment, generating a superpixel includes traversing neighboring pixels until a radius (S) around the center location of the superpixel. In at least one embodiment, a neighboring pixel is a coordinate with a valid and / or invalid flow vector.

[0094] In at least one embodiment, superpixel generator 215 traverses neighboring pixels to determine whether the neighboring pixels belong to the same superpixel as the center pixel, e.g., the neighboring pixels are similar to the center pixel in certain characteristics. In at least one embodiment, a neighboring pixel is similar to the center pixel if a predetermined criterion is satisfied. In at least one embodiment, the predetermined criterion is satisfied if an absolute difference between intensities of the neighboring pixel and the center pixel is below a first predetermined threshold and an absolute difference between gradients of the neighboring pixel and the center pixel is below a second predetermined threshold.

[0095] |I Center -I Neighbor |<I Threshold &|G Center -G Neighbor |<GT hreshold →S Member

[0096] where I center is the intensity of the center pixel, I neighbor is the intensity of the neighboring pixel under consideration, G center is the gradient of the center pixel, G neighbor is the gradient of the neighboring pixel under consideration, and S member denotes that the neighboring pixel is a member of the superpixel.

[0097] In at least one embodiment, the image is pre-processed and converted to CIELAB color space. In at least one embodiment, a neighboring pixel is similar to the center pixel if a predetermined criterion is satisfied, where the predetermined criterion can be:

[0098] G Neighbor (x, y) < G Threshold , G(x, y)

[0099] = ||I(x+1, y) - I(x-1, y) || 2 + ||I(x, y+1) - I(x, y-1) || 2 → S Member

[0100] where I(x, y) is the lab vector corresponding to the pixel at location (x, y) and ||.|| is the L2 norm.

[0101] In at least one embodiment, if the above criteria are met, superpixel generator 215 can assign one or more neighboring pixels to the same superpixel as a center pixel. In at least one embodiment, pixel coordinates belonging to the same superpixel are assigned the same integer value in an array data structure storing superpixel information. In at least one embodiment, superpixel generator 215 iterates through neighboring pixels using a separate processing thread to obtain one or more superpixel centers / candidate coordinates. In at least one embodiment, a superpixel candidate can be selected as an invalid stream coordinate. In at least one embodiment, superpixel generator 215 processes multiple invalid stream coordinates in parallel.

[0102] In at least one embodiment, rather than using a queue to store locations of neighboring pixels for iteration, superpixel generator 215 uses bits of one or more buffers to store neighboring coordinates. In at least one embodiment, one or more 64-bit buffers are used to store neighboring coordinates. In at least one embodiment, 8-bit buffers, 16-bit buffers, 32-bit buffers, 128-bit buffers, or buffers of other sizes are used. Bits of at least one buffer (e.g., 64-bit buffer) can reference an index in a constant array data structure of neighboring coordinates that map to one or more neighboring pixels. In at least one embodiment, bits of at least one buffer reference index in a data structure of neighboring coordinates map to S = 3 neighboring pixels of a pixel, where S is a radius of a superpixel iteration pattern.

[0103] In at least one embodiment, upscaler 217 upsample an image (01) containing superpixel information by a scaling factor (SF) corresponding to downscalers 301-303, where Ol has superpixel information. In at least one embodiment, if images II’ and I2’ of size 800 x 800 pixels are downsampled with SF = 1 / 2, an output image Ol’ can be correspondingly upsampled from size 400 x 400 to 800 x 800 pixels to generate image Ol. In this case, superpixel iteration at output Ol’ can cover twice the radius (S) compared to Ol to determine superpixel clusters with similar characteristics.

[0104] FIG. 3 An example 300 of grouping one or more pixels in an image 302 based on shared characteristics (e.g., intensity and / or gradients) is shown, in accordance with at least one embodiment. In at least one embodiment, a candidate pixel in image 302 has a corresponding set of neighboring pixels, some of which are part of a superpixel and some of which are not. In at least one embodiment, candidate pixel and neighboring pixels (as FIG. 3A superpixel (described above as "superpixel valid stream" and "superpixel invalid stream") is a portion of a group of pixels. In at least one embodiment, a superpixel includes pixels that are 5x5 in size; however, other sizes can be used, such as 7x7. The above description of superpixels is with respect to a 2D image. In at least one embodiment, a 3D image is used, and a superpixel is a portion of a group of voxels. FIG. 1 and 2 The generation and grouping of pixels together to form superpixels is described above. In at least one embodiment, pixels that are outside the boundary are not considered to be part of a group of pixels (e.g., a superpixel). In at least one embodiment, pixels in the image 302 are evaluated to determine if they are pixels that are part of a superpixel with respect to a candidate pixel. In at least one embodiment, if a pixel is not a superpixel, then pixels connected to those pixels do not need to be evaluated. In at least one embodiment, one or more 64-bit buffers are used to traverse neighboring coordinates from some candidate pixel coordinates to generate superpixels, where bits of the one or more 64-bit buffers corresponding to the neighboring coordinates are used for the traversal. In at least one embodiment, there is no queue used for the neighbor coordinate traversal. In at least one embodiment, valid pixels are evaluated against superpixel criteria, but pixels connected to valid pixels are not evaluated against superpixel criteria. In at least one embodiment, if a pixel is not a superpixel, then neighboring pixels of such a pixel are not evaluated against superpixel criteria.

[0105] FIG. 4An example 400 of interpolating optical flow for one or more pixels in a region in image 402 based on one or more pixels in an overlap region is shown, according to at least one embodiment. In at least one embodiment, some or all of the processes described herein, such as 400 (or any other process described herein, or variations and / or combinations thereof), are performed under the control of one or more computer systems configured with computer-executable instructions and are implemented as code (e.g., computer-executable instructions, one or more computer programs or one or more applications, CUDA code) executing collectively on one or more processors, by hardware or combinations thereof. In at least one embodiment, a system performing at least a portion of the processes in example 400 includes executable code to perform sparse-to-dense interpolation of optical flow using superpixels and infilling techniques in one iteration. In at least one embodiment, energy minimization for all valid superpixels is minimized by infilling cost = (a * warp cost + b * intensity diff + c * pixel distance), where a, b, c = range(0, 1) and a + b + c = 1. In at least one embodiment, warp cost is a difference in current pixel intensity (of a first image, “image I1”) and converging pixel intensity (of a second image, “image I2”). In at least one embodiment, interpolation in one iteration is performed with some fixed invalid flow pixels for some points / superpixels, which ultimately helps the next iteration to fix other invalid flow points / superpixels. In at least one embodiment, traditional scanline infilling (e.g., scanline algorithm) is performed in multiple iterations, which ultimately uses valid flow pixels to fix all pixels with invalid points / superpixels.

[0106] FIG. 5 An example 500 of interpolating optical flow for one or more pixels in a first region and a second overlap region in image 502 is shown, according to at least one embodiment. FIG. 5 An example process described in FIG. 4 continues in multiple iterations. In at least one embodiment, in multiple iterations (e.g., second and third iterations), all invalid flow pixels of a portion of a superpixel are ultimately fixed based on at least one or more valid flow pixels of the superpixel, while edge pixels (not part of the superpixel) are not processed, computed, or fixed during the interpolation process.

[0107] FIG. 6Example 600 of optical flow for interpolating one or more pixels by scanning rows and columns of an image and a cost map is shown. In at least one embodiment, iterative edge-aware filling may not fix all positions, but it helps reduce the density of invalid pixels. In at least one embodiment, a fast scanline method (e.g., scanline rendering) is used to fill most of the invalid flow pixels. In at least one embodiment, a fast scanline method operates on invalid boundary pixels. In at least one embodiment, Example 600 shows an image where the fast scanline method performs optical flow interpolation on 9 pixels. In at least one embodiment, a cost map is used where for all invalid pixels, each cost = INF. In at least one embodiment, the cost can be modified after each scanline operation and candidate pixel is filled with flow, which minimizes the cost.

[0108] FIG. 7 Example 700 of updating one or more pixels in an image using a cost map is shown. In at least one embodiment, and as described above with respect to FIG. 1 As FIG. 1 shown, Dijkstra's algorithm (or other similar shortest path algorithms) is used to minimize the FIG. 6 cost described in

[0109] FIG. 8 Example 800 of sparse-to-dense interpolation of optical flow is shown. As shown in FIG. 8 a coordinate map with valid and invalid optical flow is shown, and a coordinate map with optical flow filled using generated superpixels is shown. In at least one embodiment, Image 802 represents a coordinate map with one or more pixels having invalid and valid flow, where white patches (unfilled pixels) are coordinates with invalid flow and filled pixels are coordinates with valid flow. In at least one embodiment, Image 804 is a filled version of Image 802. In at least one embodiment, Image 804 does not show invalid optical flow coordinates because pixels with invalid optical flow have been filled with valid optical flow data using filling techniques (as described in reference to FIG. 1 )

[0110] FIG. 9An example of a sparse-to-dense interpolation process 900 for performing optical flow according to at least one embodiment is shown. In at least one embodiment, some or all of process 900 (or any other process described herein, or variations and / or combinations thereof) is executed under the control of one or more computer systems configured with computer-executable instructions and is implemented as code (e.g., computer-executable instructions, one or more computer programs or one or more application programs, CUDA code) executed jointly on one or more processors via hardware, software, or a combination thereof. In at least one embodiment, the code is stored in the form of a computer program on a computer-readable storage medium comprising a plurality of computer-readable instructions executable by one or more processors. In at least one embodiment, the computer-readable storage medium is a non-transitory computer-readable medium. In at least one embodiment, at least some of the computer-readable instructions available for executing process 900 are not stored using only transient signals (e.g., propagating transient electrical or electromagnetic transmissions). In at least one embodiment, the non-transitory computer-readable medium does not necessarily include non-transitory data storage circuitry (e.g., buffers, caches, and queues) within a transceiver of transient signals. In at least one embodiment, process 900 is executed at least partially on a computer system, such as those described elsewhere in this disclosure. In at least one embodiment, process 900 is executed by one or more circuits to calculate the motion of one or more pixels in the first region of the image based at least partially on the motion of one or more pixels in a second region of the image that overlaps with the first region.

[0111] In at least one embodiment, the system performing at least a portion of process 900 includes executable code for obtaining input data (e.g., a first image, a second image, and a first flow graph) at 902. In at least one embodiment, the input data is video data captured by one or more cameras. In at least one embodiment, the first image is a frame from the video data, and the second image is a subsequent frame from the video data. In at least one embodiment, the first flow graph is a sparse optical flow graph, which includes vectors indicating the movement of pixels from one frame to the next.

[0112] In at least one embodiment, the system performing at least a portion of process 900 includes executable code to group a set of pixels of an image object 904 based on similar features (e.g., similar pixel intensity and / or gradient). In at least one embodiment, the system combines one or more pixels of an image together based on similar features through one or more processes, such as combining... FIG. 2 The superpixels generated as described in the text.

[0113] In at least one embodiment, a system performing at least a portion of process 900 includes executable code to compute 906 motion of a group of pixels of a group based on motion of a second group of overlapping pixels to generate a second flow map. In at least one embodiment, a framework for sparse-to-dense interpolation computes missing optical flow vectors by first selecting overlapping boxes in an entire image, where each box contains similar pixels (e.g., with similar intensities and gradients). In at least one embodiment, a framework for sparse-to-dense interpolation then selects boxes for which some optical flow vectors have already been computed. In at least one embodiment, the already computed optical flow vectors are used to compute remaining optical flow vectors for a box. In at least one embodiment, because the boxes overlap, computing optical flow vectors for one box results in computing some optical flow vectors for overlapping boxes. In at least one embodiment, any remaining optical flow vectors for overlapping boxes are then computed. In at least one embodiment, this process is iteratively repeated until optical flow vectors for all boxes are computed and a second flow map is generated. In at least one embodiment, a system performs one or more pixel interpolations by one or more processes such as those described in connection with FIG. 1 framework 100.

[0114] In at least one embodiment, a system performing at least a portion of process 900 includes executable code to perform 908 scanline filling. In at least one embodiment, scanline filling includes traversing one or more rows or one or more columns of the second flow map to update the second flow map by converting one or more invalid pixels to valid pixels using one or more valid pixels in the one or more rows or one or more columns. In at least one embodiment, an updated second flow map is a dense optical flow map. In at least one embodiment, dense optical flow provides flow for all pixels throughout an entire image (as opposed to tracking only a group of feature pixels of interest used in sparse optical flow). In at least one embodiment, dense optical flow computes an optical flow vector for each pixel of each frame. In at least one embodiment, a system performs scanline filling by one or more processes such as those described in connection with FIG. 1those described in the framework 100 of FIG. 1. In at least one embodiment, the iterative edge-aware filling can not fix all positions, but it helps to reduce the density of invalid pixels, so the remaining amount of invalid regions is filled using the scanline filling method. In at least one embodiment, the scanline method works on invalid boundary pixels by using a cost map, where for all invalid pixels, each cost = INF. In at least one embodiment, the cost is modified after each scanline operation and the candidate pixels are filled with a flow that minimizes the cost. In at least one embodiment, Dijkstra’s algorithm (e.g., an algorithm that finds the shortest path between two locations) is used to minimize the cost, where the cost depends at least on travel distance, convergence cost, and intensity difference. In at least one embodiment, if (newcost < c(a,b)), then update candidate flow and cost.

[0115] In at least one embodiment, a system performing at least a portion of process 900 includes executable code to apply 910 a 5x5 median filter to the second flow map for refinement. In at least one embodiment, a system performs applying a filter to a flow map through one or more processes such as those described in the framework 100 of FIG. 1. FIG. 1

[0116] In at least one embodiment, one or more processes of process 900 are performed in any order, such as sequential, parallel, and / or variations thereof, and process 900 can include other processes not depicted in FIG. 9. Processes of process 900 can be omitted. FIG. 9 FIG. 9 ​​the various processes described herein. In at least one embodiment, another process that can be performed includes filling in missing portions or parts of an image based on neighboring pixels of the image. In at least one embodiment, similar steps described herein for performing sparse to dense interpolation of optical flow can also be used to perform operations with respect to filling in missing portions of an image (e.g., image inpainting). In at least one embodiment, instead of a sparse flow map, image inpainting uses a distance map to identify neighboring pixels around a candidate pixel. In at least one embodiment, a distance map is used to identify pixels closest to a candidate pixel to form a superpixel. In at least one embodiment, any missing portions in a superpixel can then be fixed based on neighboring pixels in the candidate pixel and / or the superpixel. In at least one embodiment, some or all of the processes for performing image inpainting are performed under control of one or more computer systems configured with computer-executable instructions and are implemented as code (e.g., computer-executable instructions, one or more computer programs, or one or more applications, CUDA code), collectively referred to as“computer-readable media,” that is, a hardware description language, that is executed by one or more processors on hardware.

[0117] data center

[0118] FIG. 10 An example data center 1000 is shown in accordance with at least one embodiment. In at least one embodiment, data center 1000 includes, without limitation, a data center infrastructure layer 1010, a framework layer 1020, a software layer 1030, and an application layer 1040.

[0119] In at least one embodiment, as FIG. 10As shown, the data center infrastructure layer 1010 can include a resource orchestrator 1012, grouped computing resources 1014, and node computing resources (“node C.R.s”) 1016(1)-1016(N), where “N” represents any whole, positive integer. In at least one embodiment, node C.R.s 1016(1)-1016(N) can include, but are not limited to, any number of central processing units (“CPUs” or “processors”), memory devices (e.g., dynamic read-only memory), storage devices (e.g., solid state or disk drives), network input / output (“NW I / O”) devices, network switches, power modules, cooling modules, virtual machines (“VMs”), and / or the like. In at least one embodiment, one or more node C.R.s of node C.R.s 1016(1)-1016(N) can be a server having one or more of the above-described computing resources.

[0120] In at least one embodiment, grouped computing resources 1014 can include individual groups of node C.R.s housed within one or more racks (not shown), or housed within a number of racks (also not shown) within various geographic locations. Individual groups of node C.R.s within grouped computing resources 1014 can include groups of computing, network, memory, or storage resources that can be configured or allocated to support one or more workloads. In at least one embodiment, several node C.R.s including CPUs or processors can be grouped within one or more racks to provide computing resources to support one or more workloads. In at least one embodiment, one or more racks can also include any number of power modules, cooling modules, and network switches, in any combination.

[0121] In at least one embodiment, resource orchestrator 1012 can configure or otherwise control one or more node C.R.s 1016(1)-1016(N) and / or grouped computing resources 1014. In at least one embodiment, resource orchestrator 1012 can include a software design infrastructure (“SDI”) management entity for data center 1000. In at least one embodiment, resource orchestrator 1012 can comprise hardware, software, or some combination thereof.

[0122] In at least one embodiment, as FIG. 10As shown, the framework layer 1020 includes, without limitation, a job scheduler 1032, a configuration manager 1034, a resource manager 1036, and a distributed file system 1038. In at least one embodiment, the framework layer 1020 can include a framework that supports the software 1052 of the software layer 1030 and / or one or more applications 1042 of the application layer 1040. In at least one embodiment, the software 1052 or applications 1042 can include, respectively, web-based service software or applications, such as services or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. In at least one embodiment, the framework layer 1020 can be, without limitation, a type of free and open-source software web application framework such as Apache Spark™ (hereinafter “Spark”) that can utilize the distributed file system 1038 for large-scale data processing (e.g., “big data”). In at least one embodiment, the job scheduler 1032 can include a Spark driver to facilitate scheduling workloads supported by various layers of the data center 1000. In at least one embodiment, the configuration manager 1034 can be capable of configuring different layers, such as the software layer 1030 and the framework layer 1020 including Spark and the distributed file system 1038 for supporting large-scale data processing. In at least one embodiment, the resource manager 1036 can be capable of managing clustered or grouped computing resources mapped to or allocated for supporting the distributed file system 1038 and the job scheduler 1032. In at least one embodiment, the clustered or grouped computing resources can include grouped computing resources 1014 on the data center infrastructure layer 1010. In at least one embodiment, the resource manager 1036 can coordinate with the resource orchestrator 1012 to manage these mapped or allocated computing resources.

[0123] In at least one embodiment, the software 1052 included in the software layer 1030 can include software used by at least a portion of the node C.R.s 1016(1)-1016(N), the grouped computing resources 1014, and / or the distributed file system 1038 of the framework layer 1020. One or more types of software can include, without limitation, Internet web page search software, email virus scanning software, database software, and streaming video content software.

[0124] In at least one embodiment, one or more application programs 1042 included in application layer 1040 can include one or more types of application programs used by at least portions of node C.R. 1016(1)-1016(N), grouped computing resources 1014, and / or distributed file system 1038 of framework layer 1020. One or more types of application programs can include, but are not limited to, CUDA application programs.

[0125] In at least one embodiment, any of configuration manager 1034, resource manager 1036, and resource orchestrator 1012 can implement any number and type of self-modifying actions based on any quantity and type of data acquired in any technically feasible fashion. In at least one embodiment, self-modifying actions can relieve data center operators of data center 1000 from making possibly poor configuration decisions and can avoid underutilization and / or poor performing portions of a data center.

[0126] Computer-based system

[0127] The following figures set forth, without limitation, exemplary computer-based systems that can be used to implement at least one embodiment.

[0128] FIG. 11 A processing system 1100, in accordance with at least one embodiment, is shown. In at least one embodiment, system 1100 includes one or more processor(s) 1102 and one or more graphics processing unit(s) 1108, and can be a single processor desktop system, a multiprocessor workstation system, or a server system having many processors 1102 or processor cores 1107. In at least one embodiment, processing system 1100 is a processing platform incorporated within a system- on-a-chip (SoC) integrated circuit for use in mobile, handheld, or embedded devices.

[0129] In at least one embodiment, processing system 1100 can include, or be incorporated within a server-based gaming platform, a game console, a media console, a mobile gaming console, a handheld game console, or an online game console. In at least one embodiment, processing system 1100 is a mobile phone, a smart phone, a tablet device, or a mobile internet device. In at least one embodiment, processing system 1100 can also include, couple with, or be integrated within a wearable device, such as a smart watch wearable device, a smart eyewear device, an augmented reality device, or a virtual reality device. In at least one embodiment, processing system 1100 is a television or set-top box device having one or more processors 1102 and a graphical interface generated by one or more graphics processors 1108.

[0130] In at least one embodiment, one or more processors 1102 each include one or more processor cores 1107 to process instructions which, when executed, perform operations such as operations for systems and user software. In at least one embodiment, each of the one or more processor cores 1107 is configured to process a specific instruction set 1109. In at least one embodiment, instruction set 1109 can facilitate complex instruction set computing (CISC), reduced instruction set computing (RISC), or computing via a very long instruction word (VLIW). In at least one embodiment, multiple processor cores 1107 can each process a different instruction set 1109, which can include instructions to facilitate emulation of other instruction sets. In at least one embodiment, processor core 1107 can also include other processing devices, such as a digital signal processor (DSP).

[0131] In at least one embodiment, processor 1102 includes cache memory 1104. In at least one embodiment, processor 1102 can have single-level or multi-level internal caches. In at least one embodiment, cache memory is shared among multiple components of processor 1102. In at least one embodiment, processor 1102 also uses an external cache (e.g., a level three (L3) cache or last level cache (LLC)) (not shown), which can be shared among processor cores 1107 using known cache coherency techniques. In at least one embodiment, additional included in processor 1102 is a register file 1106, which processor 1102 can include different types of registers such as integer registers, floating point registers, status registers, and instruction pointer registers. In at least one embodiment, register file 1106 can include a general register file or another register.

[0132] In at least one embodiment, one or more processors 1102 are coupled with one or more interface buses 1110 for passing control information and receiving data signals between processor 1102 and other components of system 1100. In at least one embodiment, interface bus 1110 can be versions of the Peripheral Component Interconnect (PCI) bus, including PCI, PCI Extended (PCI-X), and PCI Express (PCIe). In at least one embodiment, interface bus 1110 is not a PCI bus implementation, and can include one or more bus implementations including, but not limited to, Industry Standard Architecture (ISA), Micro Channel Architecture (MCA), Digital Electronic Bus (DEB), MicroBlaze, Advanced eXtensible Interface (AXI), Advanced Microcontroller Bus Architecture (AMBA), and / or any other bus implementation that provides for communication to one or more processors 1102.

[0133] In at least one embodiment, storage devices 1120 can be Dynamic Random Access Memory (DRAM) devices, Static Random Access Memory (SRAM) devices, flash memory devices, and / or other memory devices. In at least one embodiment, storage devices 1120 can be used as the main memory of processor 1102, to store data and instructions for use by one or more processors 1102. In at least one embodiment, memory controller 1116 also interfaces with an optional external graphics processor 1112, which can communicate with one or more graphics processors 1108 in processor 1102 to perform graphics and media operations.

[0134] In at least one embodiment, platform controller hub 1130 enables peripherals to connect to storage device 1120 and processor 1102 via a high-speed I / O bus. In at least one embodiment, I / O peripherals include, but are not limited to, audio controller 1146, network controller 1134, firmware interface 1128, wireless transceiver 1126, touch sensors 1125, data storage device 1124 (e.g., hard disk drive, flash memory, etc.). In at least one embodiment, data storage device 1124 can connect to the storage interface via a storage interface bus, e.g., a SATA, or a peripheral bus, e.g., a peripheral component interconnect bus (such as PCI, PCI Express). In at least one embodiment, touch sensors 1125 can include touch screen sensors, pressure sensors, or fingerprint sensors. In at least one embodiment, wireless transceiver 1126 can be a Wi-Fi transceiver, a Bluetooth transceiver, or a mobile network transceiver such as a 3G, 4G, or Long Term Evolution (LTE) transceiver. In at least one embodiment, firmware interface 1128 enables communication with system firmware, and can be, for example, a unified extensible firmware interface (UEFI) according to at least one embodiment. In at least one embodiment, network controller 1134 can enable network connectivity to one or more wired networks. In at least one embodiment, a high-performance network controller (not shown) is coupled

[0135] In at least one embodiment, memory controller 1116 and instances of platform controller hub 1130 can be integrated into a discrete external graphics processor, such as external graphics processor 1112. In at least one embodiment, platform controller hub 1130 and / or memory controller 1116 can be external to one or more processor(s) 1102. For example, in at least one embodiment, processing system 1100 can include an external memory controller 1116 and platform controller hub 1130, which can be configured as a memory controller hub and a peripheral controller hub in a system chipset that is in communication with processor(s) 1102.

[0136] FIG. 12A computer system 1200 according to at least one embodiment is shown. In at least one embodiment, computer system 1200 can be a system-on-a-chip, SOC, or some combination of a system and a processor. In at least one embodiment, computer system 1200 is formed from a processor 1202 that can include execution units to execute an instruction, in at least one embodiment, computer system 1200 can include, without limitation, a component, such as a processor 1202, to employ execution units including logic to perform algorithms for process data. In at least one embodiment, computer system 1200 can include processors such as a Pentium®, Core™, Xeon™, XScale™, and / or StrongARM™, Core TM or Nervana TM microprocessors, although other systems (including PCs, workstations, set-top boxes, etc. with other microprocessors) can also be used. In at least one embodiment, computer system 1200 can execute a version of the WINDOWS operating system available from Microsoft Corporation of Redmond, Wash., although other operating systems, embedded software, and / or graphical user interfaces can also be used.

[0137] In at least one embodiment, computer system 1200 can be used in other devices such as handheld devices and embedded applications. Some examples of handheld devices include cellular phones, Internet Protocol devices, digital cameras, personal digital assistants (“PDAs”), and handheld PCs. In at least one embodiment, embedded applications can include a microcontroller, a digital signal processor (“DSP”), a System-on-a-chip (“SOC”), a network computer (“NetPC”), a set-top box, a network hub, a wide area network (“WAN”) switch, or another system that can perform one or more instructions in accordance with at least one embodiment.

[0138] In at least one embodiment, computer system 1200 can include, without limitation, a processor 1202 that can include, without limitation, one or more execution units 1208 that can be configured to execute a compute unified device architecture (“CUDA”) (registered trademark) algorithm. In at least one embodiment, computer system 1200 can include, without limitation, a processor 1202 that can include, without limitation, one or more execution units 1208 that can be configured to execute an OpenCL algorithm. a CUDA program. In at least one embodiment, a CUDA program is at least a portion of a software application written in the CUDA programming language. In at least one embodiment, computer system 1200 is a single processor desktop or server system. In at least one embodiment, computer system 1200 can be a multiprocessor system. In at least one embodiment, processor 1202 can include, without limitation, a CISC microprocessor, a RISC microprocessor, a VLIW microprocessor, a processor implementing a combination of instruction sets, or any other processor device, such as a digital signal processor. In at least one embodiment, processor 1202 can be coupled to a processor bus 1210 that can transmit data signals between processor 1202 and other components in computer system 1200.

[0139] In at least one embodiment, processor 1202 can include, without limitation, a level 1 (“L1”) internal cache memory (“cache”) 1204. In at least one embodiment, processor 1202 can have a single internal cache or multiple levels of internal caches. In at least one embodiment, cache memory can reside in the processor 1202’s external. In at least one embodiment, processor 1202 can include a combination of internal and external caches. In at least one embodiment, register file 1206 can store different types of data within various registers including, without limitation, integer registers, floating point registers, status registers, and instruction pointer registers.

[0140] In at least one embodiment, execution unit 1208, including, without limitation, logic to perform integer and floating point operations, also resides in processor 1202. Processor 1202 can also include microcode (“ucode”) read-only memory (“ROM”) that stores microcode for certain macro instructions. In at least one embodiment, execution unit 1208 can include logic to handle a packed instruction set 1209. In at least one embodiment, by including the packed instruction set 1209 in the instruction set of a general-purpose processor 1202, along with associated circuitry to execute the instructions, the general-purpose processor 1202 can be used to perform the operations of many multimedia applications faster and more effectively than a general-purpose processor without such support.

[0141] In at least one embodiment, execution unit 1208 can also be used in microcontrollers, embedded processors, graphics devices, DSPs, and other types of logic circuits. In at least one embodiment, computer system 1200 can include, but not limited to, memory 1220. In at least one embodiment, memory 1220 can be implemented as a DRAM device, SRAM device, flash memory device, or other memory device. Memory 1220 can store data and / or instructions (e.g., software) that can be executed by processor 1202.

[0142] In at least one embodiment, a system logic chip can be coupled to processor bus 1210 and memory 1220. In at least one embodiment, system logic chip can include, without limitation, a memory controller hub (“MCH”) 1216, and processor 1202 can communicate with MCH 1216 via processor bus 1210. In at least one embodiment, MCH 1216 can provide a high bandwidth memory path 1218 to memory 1220 for instruction and data storage and for storage of graphics commands, data, and textures. In at least one embodiment, MCH 1216 can direct data signals between processor 1202, memory 1220, and other components in computer system 1200, and can

[0143] In at least one embodiment, computer system 1200 can use system I / O 1222 as a proprietary hub interface bus to couple MCH 1216 to I / O controller hub (“ICH”) 1230. In at least one embodiment, ICH 1230 can provide direct connections to some I / O devices and a high-speed I / O bus to connect to other I / O devices. In at least one embodiment, the high-speed I / O bus can include, without limitation, a PCI Express bus or a revved version thereof. Examples can include, without limitation, audio controller 1229, firmware hub (“Flash BIOS”) 1228, wireless transceiver 1226, data storage 1224, legacy I / O controller 1223 containing user input 1225 and keyboard interface, serial expansion port 1227 (e.g., USB), and network controller 1234. Data storage 1224 can include a hard disk drive, a floppy disk drive, a CD-ROM device, a flash memory device, or other mass storage device.

[0144] In at least one embodiment, FIG. 12 A system including interconnected hardware devices or “chips” is shown. In at least one embodiment, FIG. 12 An exemplary SoC can be shown. In at least one embodiment, FIG. 12 Devices shown in FIG. 13 can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, one or more components of system 1300 are interconnected using a compute express link (CXL) interconnect.

[0145] FIG. 13 A system 1300 according to at least one embodiment is shown. In at least one embodiment, system 1300 is an electronic device that utilizes a processor 1310. In at least one embodiment, system 1300 can be, for example and without limitation, a laptop personal computer, a tower server, a rack server, a blade server, an edge device communicatively coupled with one or more local or cloud service providers, a laptop computer, a desktop computer, a tablet, a mobile device, a phone, an embedded computer, or any other suitable electronic device.

[0146] In at least one embodiment, system 1300 can include, without limitation, a processor 1310 communicatively coupled to any suitable number or kind of components, peripherals, modules, or devices. In at least one embodiment, processor 1310 is coupled using a bus or interface, such as I 2a C bus, a system management bus (“SMBus”), a low pin count (LPC) bus, a serial peripheral interface (“SPI”), a high-definition audio (“HDA”) bus, a serial advanced technology attachment (“SATA”) bus, a USB (versions 1, 2, 3), or a universal asynchronous receiver / transmitter (“UART”) bus. In at least one embodiment, FIG. 13 A system is shown that includes interconnected hardware devices or “chips.” In at least one embodiment, FIG. 13 An exemplary SoC can be shown. In at least one embodiment, FIG. 13 Devices shown in FIG. 13B can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, FIG. 13 One or more components of FIG. 13B are interconnected using compute express link (CXL) interconnects.

[0147] In at least one embodiment, FIG. 13 may include a display 1324, a touchscreen 1325, a touchpad 1330, a near field communication unit (“NFC”) 1345, a sensor hub 1340, a thermal sensor 1346, an express chipset (“EC”) 1335, a trusted platform module (“TPM”) 1338, a BIOS / firmware / flash memory (“BIOS, FW Flash”) 1322, a DSP 1360, a solid state disk (“SSD”) or hard disk drive (“HDD”) 1320, a wireless local area network unit (“WLAN”) 1350, a Bluetooth unit 1352, a wireless wide area network unit (“WWAN”) 1356, a global positioning system (GPS) 1355, a camera (“USB 3.0 camera”) 1354 (e.g., a USB 3.0 camera), or a low power double data rate (“LPDDR”) memory unit (“LPDDR3”) 1315 implemented in, for example, LPDDR3 standard. These components can each be implemented in any suitable manner.

[0148] In at least one embodiment, other components can be communicatively coupled to processor 1310 by components discussed above. In at least one embodiment, an accelerometer 1341, an ambient light sensor (“ALS”) 1342, a compass 1343, and a gyroscope 1344 can be communicatively coupled to a sensor hub 1340. In at least one embodiment, a thermal sensor 1339, a fan 1337, a keyboard 1336, and a touchpad 1330 can be communicatively coupled to a EC 1335. In at least one embodiment, a speaker 1363, a headphone 1364, and a microphone (“mic”) 1365 can be communicatively coupled to an audio unit (“audio codec and class D amplifier”) 1362, which in turn can be communicatively coupled to a DSP 1360. In at least one embodiment, audio unit 1362 can include, for example and without limitation, an audio coder / decoder (“codec”) and a class D amplifier. In at least one embodiment, a SIM card (“SIM”) 1357 can be communicatively coupled to a WWAN unit 1356. In at least one embodiment, components such as WLAN unit 1350 and Bluetooth unit 1352, as well as WWAN unit 1356, can be implemented in a Next Generation Form Factor (NGFF).

[0149] FIG. 14 An exemplary integrated circuit 1400, in accordance with at least one embodiment, is shown. In at least one embodiment, exemplary integrated circuit 1400 is a SoC, which can be fabricated using one or more IP cores. In at least one embodiment, integrated circuit 1400 includes one or more application processors 1405 (e.g., CPUs), at least one graphics processor 1410, and can additionally include an image processor 1415 and / or a video processor 1420, any of which can be a modular IP core. In at least one embodiment, integrated circuit 1400 includes peripheral or bus logic including a USB controller 1425, a UART controller 1430, an SPI / SDIO controller 1435, and an I 2 S / I 2 C controller 1440. In at least one embodiment, integrated circuit 1400 can include a display device 1445 coupled to one or more of a high-definition multimedia interface (HDMI) controller 1450 and a mobile industry processor interface (MIPI) display interface 1455. In at least one embodiment, storage can be provided by a flash memory subsystem 1460, including flash memory and a flash memory controller. In at least one embodiment, a memory interface can be provided via a memory controller 1465 for access to SDRAM or SRAM memory devices. In at least one embodiment, some integrated circuits also include an embedded security engine 1470.

[0150] FIG. 15A computing system 1500 is shown in accordance with at least one embodiment. In at least one embodiment, computing system 1500 includes a processing subsystem 1501 having one or more processor(s) 1502 and a system memory 1504 communicating via an interconnection path 1505 that can include a memory hub 1505. In at least one embodiment, memory hub 1505 can be a separate component coupled with one or more processors 1502 via individual communication links 1507A-1507N, or can be integrated within one or more processors 1502. In at least one embodiment, memory hub 1505 couples with system memory 1504, which can include one or more memory modules 1506A-1506N. In at least one embodiment, one or more memory modules 1506A-1506N have a common memory component design, thus leading to a common form factor used to implement system memory 1504. In at least one embodiment, one or more memory modules 1506A-1506N are low power, high-density memory modules.

[0151] In at least one embodiment, processing subsystem 1501 includes one or more parallel processor(s) 1512 coupled with memory hub 1505 via a bus or other communication link 1513. In at least one embodiment, communication link 1513 can be one of many such links, which can be established as a peripheral component interconnect (PCI) bus using a PCI local bus protocol or a PCI Express bus using a PCI-Express local bus protocol, among other protocols. In at least one embodiment, one or more parallel processor(s) 1512 form a computationally intensive, parallel- processing processing system that can include a number of processor cores, each of which can be configured to execute a large number of threads.

[0152] In at least one embodiment, system storage 1514 can connect to I / O hub 1507 to provide storage mechanisms for computing system 1500. In at least one embodiment, I / O switches 1516 can be used to provide interface mechanisms to enable connections between I / O hub 1507 and other components such as network adapter 1518 and / or wireless network adapter 1519 that can be integrated into a platform, as well as various other devices that can be added via one or more add-in devices 1520. In at least one embodiment, network adapter 1518 can be an Ethernet adapter or another wired

[0153] In at least one embodiment, computing system 1500 can include other components not explicitly shown, including USB or other port connections, optical storage drives, video capture devices, and so on, that can also be connected to I / O hub 1507. In at least one embodiment, communication paths interconnecting various components in FIG. 15 Communication paths interconnecting various components in at least one embodiment can use any suitable protocols including, for example, PCI (Peripheral Component Interconnect) based protocols (e.g., PCI Express), or other bus or point-to-point communication interfaces and / or protocols (e.g., NVLink high-speed interconnect, or interconnect protocols).

[0154] In at least one embodiment, parallel processor(s) 1512 include circuitry optimized for graphics and video processing, including, for example, video output circuitry. In at least one embodiment, parallel processor(s) 1512 include circuitry optimized for general purpose processing. In at least one embodiment, components of computing system 1500 can be integrated with one or more other system elements on a single integrated circuit. For example, in at least one embodiment, parallel processor(s) 1512, memory hub 1505, processor(s) 1502, and I / O hub 1507 can be integrated into a system on a chip (SoC) integrated circuit. In at least one embodiment, components of computing system 1500 can be integrated into a single package to form a system in a package (SIP) configuration. In at least one embodiment, at least a portion of components of computing system 1500 can be integrated into a multichip module (MCM), which can be interconnected with other multichip modules into a modular computing system. In at least one embodiment, I / O subsystem 1511 and display device 1510B are omitted from computing system 1500.

[0155] Processing system

[0156] The following figures illustrate, but are not limited to, exemplary processing systems that can be used to implement at least one embodiment.

[0157] FIG. 16 An accelerated processing unit (“APU”) 1600 according to at least one embodiment is illustrated. In at least one embodiment, the APU 1600 was developed by AMD Inc. of Santa Clara, California. In at least one embodiment, the APU 1600 can be configured to execute applications, such as CUDA programs. In at least one embodiment, the APU 1600 includes, but is not limited to, a core complex 1610, a graphics complex 1640, an architecture 1660, an I / O interface 1670, a memory controller 1680, a display controller 1692, and a multimedia engine 1694. In at least one embodiment, the APU 1600 can be, but is not limited to, any combination of any number of core complexes 1610, any number of graphics complexes 1640, any number of display controllers 1692, and any number of multimedia engines 1694. For illustrative purposes, multiple instances of similar objects are indicated herein by reference numerals, wherein the reference numerals identify the object, and the numbers in parentheses identify the desired instances.

[0158] In at least one embodiment, the core complex 1610 is a CPU, the graphics complex 1640 is a GPU, and the APU 1600 is a processing unit that is not limited to 1610 and 1640 integrated onto a single chip. In at least one embodiment, some tasks may be assigned to the core complex 1610, while other tasks may be assigned to the graphics complex 1640. In at least one embodiment, the core complex 1610 is configured to execute main control software associated with the APU 1600, such as an operating system. In at least one embodiment, the core complex 1610 is the main processor of the APU 1600, which controls and coordinates the operation of other processors. In at least one embodiment, the core complex 1610 issues commands to control the operation of the graphics complex 1640. In at least one embodiment, the core complex 1610 may be configured to execute host executable code derived from CUDA source code, and the graphics complex 1640 may be configured to execute device executable code derived from CUDA source code.

[0159] In at least one embodiment, the core complex 1610 includes, but is not limited to, cores 1620(1)-1620(4) and L3 cache 1630. In at least one embodiment, the core complex 1610 may include, but is not limited to, any combination of any number of cores 1620 and any number and type of cache. In at least one embodiment, the cores 1620 are configured to execute instructions of a specific instruction set architecture (“ISA”). In at least one embodiment, each core 1620 is a CPU core.

[0160] In at least one embodiment, each core 1620 includes, without limitation, a fetch / decode unit 1622, an integer execution engine 1624, a floating point execution engine 1626, and an L2 cache 1628. In at least one embodiment, fetch / decode unit 1622 fetches instructions, decodes such instructions, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1624 and floating point execution engine 1626. In at least one embodiment, fetch / decode unit 1622 can dispatch one micro-instruction to integer execution engine 1624 and another micro-instruction to floating point execution engine 1626 simultaneously. In at least one embodiment, integer execution engine 1624 executes, without limitation, integer and memory operations. In at least one embodiment, floating point engine 1626 executes, without limitation, floating point and vector operations. In at least one embodiment, fetch-decode unit 1622 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1624 and floating point execution engine 1626.

[0161] In at least one embodiment, each core 1620(i) has access to an L2 cache 1628(i) included in core 1620(i), where i is an integer representing a particular instance of core 1620. In at least one embodiment, each core 1620 included in core complex 1610(j) is connected to other cores 1620 included in core complex 1610(j) via an L3 cache 1630(j) included in core complex 1610(j), where j is an integer representing a particular instance of core complex 1610. In at least one embodiment, cores 1620 included in core complex 1610(j) have access to all L3 caches 1630(j) included in core complex 1610(j), where j is an integer representing a particular instance of core complex 1610. In at least one embodiment, L3 cache 1630 can include, without limitation, any number of slices.

[0162] In at least one embodiment, graphics complex 1640 can be configured to perform compute operations in a highly parallel manner. In at least one embodiment, graphics complex 1640 is configured to perform graphics pipeline operations such as draw commands, pixel operations, geometric calculations, and other operations associated with rendering images to a display. In at least one embodiment, graphics complex 1640 is configured to perform operations that are not graphics related. In at least one embodiment, graphics complex 1640 is configured to perform graphics related operations and operations that are not graphics related.

[0163] In at least one embodiment, graphics processing engine 1630 includes, without limitation, any number of SIMD units 1632 and shared memory 1634. In at least one embodiment, SIMD units 1632 share shared memory 1634. In at least one embodiment, shared memory 1634 is partitioned. In at least one embodiment, graphics processing engine 1630 includes, without limitation, any number of SIMD units 1632 and any number (including zero) and type of cache. In at least one embodiment, graphics processing engine 1630 includes, without limitation, any number of specialized graphics hardware.

[0164] In at least one embodiment, each compute unit 1650 includes, without limitation, any number of SIMD units 1652 and shared memory 1654. In at least one embodiment, each SIMD unit 1652 implements a SIMD architecture and is configured to perform operations in parallel. In at least one embodiment, each compute unit 1650 can execute any number of thread blocks, but each thread block executes on a single compute unit 1650. In at least one embodiment, a thread block includes, without limitation, any number of execution threads. In at least one embodiment, a workgroup is a thread block. In at least one embodiment, each SIMD unit 1652 executes a different thread warp. In at least one embodiment, a thread warp is a group of threads (e.g., 16 threads), where each thread in a thread warp belongs to a single thread block and is configured to process a different set of data based on a single instruction set. In at least one embodiment, one or more threads in a thread warp can be disabled using predication. In at least one embodiment, a lane is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a thread warp. In at least one embodiment, different wavefronts in a thread block can be synchronized together and communicate via shared memory 1654.

[0165] In at least one embodiment, fabric 1660 is a system interconnect that facilitates data and control transmissions across core complex 1610, graphics complex 1640, I / O interface 1670, memory controllers 1680, display controller 1692, and multimedia engine 1694. In at least one embodiment, APU 1600 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1660 that facilitate data and control transmissions across any number and type of directly or indirectly linked components that can be internal or external to APU 1600. In at least one embodiment, I / O interface 1670 represents any number and type of I / O interface (e.g., PCI, PCI-Extended (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to I / O interface 1670. In at least one embodiment, peripheral devices coupled to I / O interface 1670 can include, without limitation, a keyboard, a mouse, a printer, a scanner, a joystick or other type of game controller, a media recording device, an external storage device, a network interface card, etc.

[0166] In at least one embodiment, display controller 1692 displays images on one or more display devices, such as liquid crystal display (“LCD”) devices. In at least one embodiment, multimedia engine 1694 includes, without limitation, any number and type of multimedia-related circuitry, such as a video decoder, a video encoder, an image signal processor, etc. In at least one embodiment, memory controllers 1680 facilitate data transfers between APU 1600 and unified system memory 1690. In at least one embodiment, core complex 1610 and graphics complex 1640 share unified system memory 1690.

[0167] In at least one embodiment, APU 1600 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1680 and memory devices (e.g., shared memory 1654) that can be dedicated to one component or shared among multiple components. In at least one embodiment, APU 1600 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1628, L3 cache 1630, and L2 cache 1642), each of which can be private to a component or shared among any number of components (e.g., core 1620, core complex 1610, SIMD unit 1652, compute unit 1650, and graphics complex 1640).

[0168] FIG. 17A CPU 1700 is shown, in accordance with at least one embodiment. In at least one embodiment, CPU 1700 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, CPU 1700 can be configured to execute an application program. In at least one embodiment, CPU 1700 is configured to execute a main control software, such as an operating system. In at least one embodiment, CPU 1700 issues commands that control the operation of an external GPU (not shown). In at least one embodiment, CPU 1700 can be configured to execute host executable code derived from CUDA source code, and an external GPU can be configured to execute device executable code derived from such CUDA source code. In at least one embodiment, CPU 1700 includes, without limitation, any number of core complexes 1710, fabric 1760, I / O interfaces 1770, and memory controllers 1780.

[0169] In at least one embodiment, core complex 1710 includes, without limitation, cores 1720(1)-1720(4) and L3 cache 1730. In at least one embodiment, core complex 1710 can include, without limitation, any number of cores 1720 and any combination and type of caches. In at least one embodiment, cores 1720 are configured to execute instructions of a particular ISA. In at least one embodiment, each core 1720 is a CPU core.

[0170] In at least one embodiment, each core 1720 includes, without limitation, a fetch / decode unit 1722, an integer execution engine 1724, a floating point execution engine 1726, and an L2 cache 1728. In at least one embodiment, fetch / decode unit 1722 fetches instructions, decodes them, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1724 and floating point execution engine 1726. In at least one embodiment, fetch / decode unit 1722 can dispatch one micro-instruction to integer execution engine 1724 and another micro-instruction to floating point execution engine 1726 simultaneously. In at least one embodiment, integer execution engine 1724 executes, without limitation, integer and memory operations. In at least one embodiment, floating point engine 1726 executes, without limitation, floating point and vector operations. In at least one embodiment, fetch-decode unit 1722 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1724 and floating point execution engine 1726.

[0171] In at least one embodiment, each core 1720(i) has access to an L2 cache 1728(i) included in the core 1720(i), where i is an integer representing a particular instance of a core 1720. In at least one embodiment, each core 1720 included in a core complex 1710(j) is connected to other cores 1720 in the core complex 1710(j) via an L3 cache 1730(j) included in the core complex 1710(j), where j is an integer representing a particular instance of a core complex 1710. In at least one embodiment, a core 1720 included in a core complex 1710(j) has access to all L3 caches 1730(j) included in the core complex 1710(j), where j is an integer representing a particular instance of a core complex 1710. In at least one embodiment, an L3 cache 1730 can include, without limitation, any number of slices.

[0172] In at least one embodiment, fabric 1760 is a system interconnect that facilitates data and control transfers across core complexes 1710(1)-1710(N) (where N is an integer greater than zero), I / O interface 1770, and memory controllers 1780. In at least one embodiment, CPU 1700 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1760 that facilitate data and control transfers across any number and type of directly or indirectly linked components that can be internal or external to CPU 1700. In at least one embodiment, I / O interface 1770 represents any number and type of I / O interface (e.g., PCI, PCI-X, PCIe, GBE, USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to I / O interface 1770. In at least one embodiment, peripheral devices coupled to I / O interface 1770 can include, without limitation, a display, a keyboard, a mouse, a printer, a scanner, a joystick or other types of game controller, a media recording device, an external storage device, a network interface card, etc.

[0173] In at least one embodiment, memory controller 1780 facilitates communication between CPU 1700 and system memory 1790. In at least one embodiment, core complex 1710 and graphics complex 1740 share system memory 1790. In at least one embodiment, CPU 1700 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1780 and memory devices that can be dedicated to one component or shared among multiple components. In at least one embodiment, CPU 1700 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1728 and L3 cache 1730), each of which can be private to a component or shared among any number of components (e.g., core 1720 and core complex 1710).

[0174] FIG. 18 An exemplary accelerator integration slice 1890 is shown in accordance with at least one embodiment. As used herein, a “slice” includes a specified portion of processing resources of an accelerator integration circuit. In at least one embodiment, an accelerator integration circuit provides cache management, memory access, environment management, and interrupt management services on behalf of multiple graphics processing engines that are part of graphics acceleration modules. Graphics processing engines can each comprise a separate GPU. Alternatively, graphics processing engines can include different types of graphics processing engines within a GPU, such as graphics execution units, media processing engines (e.g., video encoders / decoders), samplers, and blit engines. In at least one embodiment, a graphics acceleration module can be a GPU with a plurality of graphics processing engines. In at least one embodiment, a graphics processing engine can be a separate GPU integrated on a common package, line card, or chip as the CPU.

[0175] Application effective address space 1882 within system memory 1814 stores process elements 1883. In one embodiment, process elements 1883 are stored in response to GPU invocations 1881 from applications 1880 executing on processor 1807. Process elements 1883 contain processing state for corresponding applications 1880. Work descriptors (WDs) 1884 contained in process elements 1883 can be individual jobs requested by applications or can contain pointers to queues of jobs. In at least one embodiment, WDs 1884 are pointers to job request queues in application effective address space 1882.

[0176] Graphics acceleration module 1846 and / or individual graphics processing engines can be shared by all or a subset of processes in a system. In at least one embodiment, there can be included infrastructure for setting up processing state and sending WDs 1884 to graphics acceleration module 1846 to start a job in a virtualized environment.

[0177] In at least one embodiment, a dedicated process programming model is implemented. In this model, a single process owns a graphics acceleration module 1846 or individual graphics processing engines. As graphics acceleration module 1846 is owned by a single process, a hypervisor initializes the accelerator integration circuit for the owning partition and an operating system initializes the accelerator integration circuit for the owning partition when assigning graphics acceleration module 1846.

[0178] In operation, a WD fetch unit 1891 in accelerator integration slice 1890 fetches a next WD 1884 including an indication of work to be completed by one or more graphics processing engines of graphics acceleration module 1846. Data from WD 1884 can be stored in registers 1845 used by memory management unit (MMU) 1839, interrupt management circuit 1847, and / or environment management circuit 1848, as shown. For example, one embodiment of MMU 1839 includes segment / page walk circuitry to access segment / page tables 1886 within an OS virtual address space 1885. Interrupt management circuit 1847 can handle interrupt events (INTs) 1892 received from graphics acceleration module 1846. Effective addresses 1893 produced by graphics processing engines, when executing graphics operations, are translated to real addresses by MMU 1839.

[0179] In one embodiment, a same set of registers 1845 is replicated for each graphics processing engine and / or graphics acceleration module 1846 and can be initialized by a system hypervisor or operating system. Each of these replicated registers can be included in accelerator integration slice 1890. Exemplary registers that can be initialized by a hypervisor are shown in Table 1.

[0180] Table 1 - Hypervisor Initialized Registers

[0181]

[0182]

[0183] Exemplary registers that can be initialized by an operating system are shown in Table 2.

[0184] Table 2 - Operating System Initialized Registers

[0185] 1 Process and thread identification 2 Effective address (EA) environment save / restoration pointer 3 Virtual address (VA) accelerator utilization record pointer 4 Virtual address (VA) storage segment table pointer 5 Authority mask 6 Work descriptor

[0186] In one embodiment, each WD 1884 is specific to a particular graphics processing module 1846 and / or a particular graphics processing engine. It contains all information the graphics processing engine needs to work or work needed, or it can be a pointer to a memory location where an application has set up a command queue of work to be done.

[0187] FIGS. 19A-19B An exemplary graphics processor according to at least one embodiment is shown. In at least one embodiment, any of the exemplary graphics processors can be fabricated as one or more IP cores. In addition to the graphics processors illustrated, in at least one embodiment, other logic and circuits can be included, including additional graphics processors cores or special-purpose processing cores, peripheral interface controllers or general-purpose processor cores. In at least one embodiment, graphics processor is used in a SoC implementation.

[0188] FIG. 19A An exemplary graphics processor 1910 of a SoC integrated circuit, according to at least one embodiment, is shown, which can be fabricated using one or more IP cores. FIG. 19B An additional exemplary graphics processor 1940 of a SoC integrated circuit, according to at least one embodiment, is shown, which can be fabricated using one or more IP cores. In at least one embodiment, FIG. 19A Graphics processor 1910 of FIG. 19 is a low-power graphics processor core. In at least one embodiment, graphics processor 1910 of FIG. 19 is a low-power graphics processor core. FIG. 19B Graphics processor 1940 of FIG. 19 is a higher performance graphics processor core. In at least one embodiment, each graphics processor 1910, 1940 can be a FIG. 14 Variants of graphics processor 1410 of FIG. 14.

[0189] In at least one embodiment, graphics processor 1910 includes a vertex processor 1905 and one or more fragment processor(s) 1915A-1915N (e.g., 1915A, 1915B, 1915C, 1915D, through 1915N-1, and 1915N). In at least one embodiment, graphics processor 1910 can execute different shader programs via separate logic for vertex processing and / or for fragment / pixel processing. In at least one embodiment, vertex processor 1905 is optimized to execute operations on vertex data that frequently process complex, high-precision calculations for geometry transformations, motion blur, and dynamic shading. In at least one embodiment, one or more fragment processor(s) 1915A-1915N can be optimized to execute more SIMT-based calculations for rendering an image to a display device based on vertex data. In at least one embodiment, image / texture data are stored in memory. In at least one embodiment, vertex processor 1905 and / or one or more fragment processor(s) 1915A-1915N can access and process the image / texture data in memory that is stored in a high-speed cache memory area.

[0190] In at least one embodiment, graphics processor 1910 additionally includes one or more MMU(s) 1920A-1920B, cache memory 1925A-1925B, and circuit interconnect 1930A-1930B. In at least one embodiment, one or more MMU(s) 1920A-1920B provide for virtual to physical address mapping for graphics processor 1910, including for vertex processor 1905 and / or for fragment processor(s) 1915A-1915N, which can reference vertex or image / texture data stored in memory, in addition to vertex or image / texture data stored in one or more cache memories 1925A-1925B. In at least one embodiment, one or more MMU(s) 1920A-1920B can be synchronized with other MMUs within FIG. 14 one or more application processor(s) 1405, image processors 1415, and / or video processors 1420 such that each processor 1405-1420 can participate in a shared or unified virtual memory system. In at least one embodiment, one or more circuit interconnects 1930A-1930B enable graphics processor 1910 to interface with other IP cores within a SoC, via an internal bus, or via a direct connection.

[0191] In at least one embodiment, graphics processor 1940 includes FIG. 19Aone or more MMUs 1920A-1920B, caches 1925A-1925B, and circuit interconnect 1930A-1930B of graphics processor 1910. In at least one embodiment, graphics processor 1940 includes one or more shader cores 1955A-1955N (e.g., 1955A, 1955B, 1955C, 1955D, 1955E, 1955F, through 1955N-1, and 1955N), which provide for a unified shader core architecture in which a single core or type or core can execute all types of programmable shader code, including shader program code to implement vertex shaders, fragment shaders, and / or compute shaders. In at least one embodiment, a number of shader cores can vary from one implementation to another. In at least one embodiment, graphics processor 1940 includes an inter-core task manager 1945, which acts as a thread dispatcher to distribute execution threads to one or more shader cores 1955A-1955N and a tiling unit 1958 to accelerate tiling operations for tile-based rendering in which rendering operations for a scene are subdivided in image space, e.g., to

[0192] FIG. 20A A graphics core 2000 according to at least one embodiment is shown. In at least one embodiment, graphics core 2000 can be included within graphics processor 1410. FIG. 14 A graphics core 2000 according to at least one embodiment is shown. In at least one embodiment, graphics core 2000 can be included within graphics processor 1410. FIG. 19B In at least one embodiment, graphics core 2000 includes shared instruction cache 2002, texture unit 2018, and cache / shared memory 2020, which are common to execution resources within graphics core 2000. In at least one embodiment, graphics core 2000 can include multiple slices 2001A-2001N or partitions of each core, and graphics processor can include multiple instances of graphics core 2000. Slices 2001A-2001N can include support logic including a local instruction cache 2004A-2004N, a thread scheduler 2006A-2006N, a thread dispatcher 2008A-2008N, and a set of registers 2010A-2010N. In at least one embodiment, slices 2001A-2001N can include a set of additional functional units (AFUs) 2012A-2012N, floating point units (FPUs) 2014A-2014N, integer arithmetic logic units (ALUs) 2016A-2016N, address computation units (ACUs) 2013A-2013N, double precision floating point units (DPFPUs) 2015A-2015N, and matrix processing units (MPUs) 2017A-2017N.

[0193] In one embodiment, FPUs 2014A-2014N can perform single-precision (32-bit) and half-precision (16-bit) floating point operations, while DPFPUs 2015A-2015N can perform double-precision (64-bit) floating point operations. In at least one embodiment, ALUs 2016A-2016N can perform variable precision integer operations at 8-bit, 16-bit, and 32-bit precision, and can be configured for mixed precision operations. In at least one embodiment, MPUs 2017A-2017N can also be configured for mixed precision matrix operations, including half-precision floating point operations and 8-bit integer operations. In at least one embodiment, MPUs 2017A-2017N can perform various matrix operations to accelerate CUDA programs, including enabling support for accelerated General Matrix to Matrix multiplication (GEMM). In at least one embodiment, AFUs 2012A-2012N can perform additional logical operations not supported by floating point or integer units, including trigonometric operations (e.g., Sine, Cosine, etc.).

[0194] FIG. 20B A general purpose graphics processing unit (GPGPU) 2030 in at least one embodiment is shown. In at least one embodiment, GPGPU 2030 is highly parallel and suitable for deployment on a multi-chip module. In at least one embodiment, GPGPU 2030 can be configured to enable highly parallel compute operations to be performed by a GPU array. In at least one embodiment, GPGPU 2030 can be directly linked to other instances of GPGPU 2030 to create a multi-GPU cluster to improve execution time for CUDA programs. In at least one embodiment, GPGPU 2030 includes a host interface 2032 to enable connection to a host processor. In at least one embodiment, host interface 2032 is a PCIe interface. In at least one embodiment, host interface 2032 can be a vendor-specific communications interface or communications fabric. In at least one embodiment, GPGPU 2030 receives commands from a host processor and uses a global scheduler 2034 to dispatch execution threads associated with those commands to a group of compute clusters 2036A-2036H. In at least one embodiment, compute clusters 2036A-2036H share a cache memory 2038. In at least one embodiment, cache memory 2038 can be used as an upper level cache for cache memory within compute clusters 2036A-2036H.

[0195] In at least one embodiment, GPGPU 2030 includes memory 2044A-2044B coupled to compute clusters 2036A-2036H via a set of memory controllers 2042A-2042B. In at least one embodiment, memory 2044A-2044B can include various types of memory devices including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory.

[0196] In at least one embodiment, compute clusters 2036A-2036H each include a group of graphics cores, such as graphics core 2000, which can include multiple types of integer and floating point logic units that can perform computational operations at various precisions, including suitable for computations related to CUDA programs. For example, in at least one embodiment, at least a subset of floating point units in each compute cluster 2036A-2036H can be configured to perform 16- or 32-bit floating point operations, while a different subset of floating point units can be configured to perform 64-bit floating point operations. FIG. 20A

[0197] In at least one embodiment, multiple instances of GPGPU 2030 can be configured to operate as compute clusters. Compute clusters 2036A-2036H can implement any technically feasible communication technology for synchronization and data exchange. In at least one embodiment, multiple instances of GPGPU 2030 communicate over host interface 2032. In at least one embodiment, GPGPU 2030 includes I / O hub 2039 that couples GPGPU 2030 with GPU link 2040, enabling a direct connection to other instances of GPGPU 2030. In at least one embodiment, GPU link 2040 is coupled to a specialized GPU-to-GPU bridge that enables communication and synchronization between multiple instances of GPGPU 2030. In at least one embodiment, GPU link 2040 is coupled with a high-speed interconnect to transmit and receive data to other GPGPUs or parallel processors. In at least one embodiment, multiple instances of GPGPU 2030 are located in separate data processing systems and communicate over a network device accessible via host interface 2032. In at least one embodiment, GPU link 2040 can be configured to connect to a host processor, in addition to or in place of host interface 2032. In at least one embodiment, GPGPU 2030 can be configured to execute CUDA programs.

[0198] FIG. 21A ​A parallel processor 2100, in accordance with at least one embodiment, is shown. In at least one embodiment, various components of parallel processor 2100 can be implemented using one or more integrated circuit devices, such as programmable processors, application specific integrated circuits (ASICs), or FPGAs.

[0199] In at least one embodiment, parallel processor 2100 includes a parallel processing unit 2102. In at least one embodiment, parallel processing unit 2102 includes an I / O unit 2104 that enables communication with other devices, including other instances of parallel processing unit 2102. In at least one embodiment, I / O unit 2104 can be directly connected to the other devices. In at least one embodiment, I / O unit 2104 connects with other devices using a hub or switch interface, such as memory hub 2105. In at least one embodiment, connections between memory hub 2105 and I / O unit 2104 form a communication link. In at least one embodiment, I / O unit 2104 connects with a host interface 2106 and a memory crossbar switch 2116, where host interface 2106 receives commands directed to the processing operations and memory crossbar switch 2116 receives commands directed to memory operations.

[0200] In at least one embodiment, when host interface 2106 receives a command buffer via I / O unit 2104, host interface 2106 can direct work operations to execute those commands to front end 2108. In at least one embodiment, front end 2108 couples with a scheduler 2110, which is configured to assign commands or other work items to processing arrays 2112. In at least one embodiment, scheduler 2110 ensures that processing arrays 2112 are properly configured and in an active state before assigning tasks to processing arrays 2112 of processing arrays 2112. In at least one embodiment, scheduler 2110 is implemented by firmware logic executing on a microcontroller. In at least one embodiment, microcontroller- implemented scheduler 2110 is configurable to perform complex scheduling and work distribution operations at a coarse and fine grain level to achieve a low-latency access to threads executing on processing arrays 2112. In at least one embodiment, host software can prove a workload for scheduling on processing arrays 2112 through one of a number of graphics processing doorbells. In at least one embodiment, workload can then be automatically distributed on processing arrays 2112 by scheduler 2110 logic within microcontroller that includes scheduler 2110.

[0201] In at least one embodiment, processing array 2112 can include up to “N” processing clusters (e.g., cluster 2114A, 2114B through 2114N). In at least one embodiment, each cluster 2114A-2114N of processing array 2112 can execute a large number of concurrent threads. In at least one embodiment, scheduler 2110 can allocate work to clusters 2114A-2114N of processing array 2112 using various scheduling and / or work distribution algorithms, which can be determined at least in part by workload arriving at processing array 2112, capacity of processing array 2112, and / or capacities of individual clusters 2114A-2114N within processing array 2112. In at least one embodiment, scheduling can be handled dynamically by scheduler 2110, or can be aided in part by compiler logic during compilation of program logic configured for execution by processing array 2112. In at least one embodiment, different clusters 2114A-2114N of processing array 2112 can be allocated for processing different types of programs or for performing different types of computations.

[0202] In at least one embodiment, processing array 2112 can be configured to perform various types of parallel processing operations. In at least one embodiment, processing array 2112 is configured to perform general purpose parallel compute operations. For example, in at least one embodiment, processing array 2112 can include logic to perform processing tasks including filtering of video and / or audio data, performing modeling operations including physical operations, and performing data transformations.

[0203] In at least one embodiment, processing array 2112 is configured to perform parallel graphics processing operations. In at least one embodiment, processing array 2112 can include additional logic to support performance of such graphics processing operations, including but not limited to texture sampling logic to perform texture operations, tessellation logic, and other vertex processing logic. In at least one embodiment, processing array 2112 can be configured to execute graphics processing related shader programs, such as, but not limited to, vertex shaders, tessellation shaders, geometry shaders, and pixel

[0204] In at least one embodiment, when parallel processing unit 2102 is used to perform graphics processing, scheduler 2110 can be configured to divide incoming workloads into tasks of approximately equal size to better enable distribution of graphics processing operations across multiple clusters 2114A-2114N of processing array 2112. In at least one embodiment, portions of processing array 2112 can be configured to perform different types of processing. For example, in at least one embodiment, a first portion can be configured to perform vertex shading and topology generation, a second portion can be configured to perform tessellation and geometry shading, and a third portion can be configured to perform pixel shading or other screen space operations to produce a rendered image for display on a display device. In at least one embodiment, intermediate data produced by one or more of clusters 2114A-2114N can be stored in buffers to allow transmission of intermediate data between clusters 2114A-2114N for further processing.

[0205] In at least one embodiment, processing array 2112 can receive processing tasks to be executed from scheduler 2110, which receives commands defining the processing tasks from front end 2108. In at least one embodiment, a processing task can include an index into data to be processed, such as can include surface (patch) data, raw data, vertex data, and / or pixel data, as well as state parameters and commands defining how the data is to be processed (e.g., what program is to be executed). In at least one embodiment, scheduler 2110 can be configured to fetch the index corresponding to a task, or can receive the index from front end 2108. In at least one embodiment, front end 2108 can be configured to ensure that processing array 2112 is configured in an effective state before launching a workload specified by an incoming command buffer (e.g., a batch-buffer, a push buffer, etc.).

[0206] In at least one embodiment, each of one or more instances of parallel processing unit 2102 can be coupled to a parallel processor memory 2122. In at least one embodiment, parallel processor memory 2122 can be accessed by the processing array 2112, as well as the I / O unit 2104, via a memory crossbar 2116. In at least one embodiment, memory crossbar 2116 can be used to access memory banks 2124A-2124N of parallel processor memory 2122 by the processing array 2112, I / O unit 2104, and possibly other processing units or processing arrays in at least one embodiment. In at least one embodiment, memory crossbar 2116 can be configured to couple to a high-speed data bus such as a high-bandwidth memory (HBM) bus.

[0207] In at least one embodiment, memory units 2124A-2124N can include various types of memory devices including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM). In at least one embodiment, memory units 2124A-2124N can also include 3D stacked memory including but not limited to high bandwidth memory (HBM). In at least one embodiment, rendering targets such as frame buffers or texture maps can be stored across memory units 2124A-2124N allowing partition units 2120A-2120N to write portions of each rendering target in parallel to effectively use available bandwidth of parallel processor memory 2122. In at least one embodiment, local instances of parallel processor memory 2122 can be excluded from a unified memory design that utilizes system memory in combination with local cache memory.

[0208] In at least one embodiment, any of clusters 2114A-2114N of processing array 2112 can process data that is to be written into any of memory units 2124A-2124N within parallel processor memory 2122. In at least one embodiment, memory crossbar 2116 can be configured to transmit outputs of each cluster 2114A-2114N to any partition unit 2120A-2120N or another cluster 2114A-2114N, which can perform other processing operations on the outputs. In at least one embodiment, each cluster 2114A-2114N can communicate with memory interface 2118 through memory crossbar 2116 to read from or write to various external memory devices. In at least one embodiment, memory crossbar 2116 has a connection to memory interface 2118 to communicate with I / O unit 2104, as well as a local instance of parallel processor memory 2122, to enable processing units within different processing clusters 2114A-2114N to communicate with system memory or other memory that is not local to the parallel processing units 2102. In at least one embodiment, memory crossbar 2116 can use virtual channels to separate traffic streams between clusters 2114A-2114N and partition units 2120A-2120N.

[0209] In at least one embodiment, multiple instances of parallel processing unit 2102 can be provided on a single add-in card, or multiple add-in cards can be interconnected. In at least one embodiment, different instances of parallel processing unit 2102 can be configured to operate together as a single parallel processing unit 2102, even if the different instances have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences. In at least one embodiment, some instances of parallel processing unit 2102 can include higher precision floating point units relative to other instances. In at least one embodiment, a system incorporating one or more instances of parallel processing unit 2102 or parallel processor 2100 can be implemented in a variety of configurations and form factors, including but not limited to desktop, laptop, or handheld personal computers, servers, workstations, game consoles, and / or embedded systems.

[0210] FIG. 21BA processing cluster 2194 is shown in accordance with at least one embodiment. In at least one embodiment, processing cluster 2194 is included in a parallel processing unit. In at least one embodiment, processing cluster 2194 is an instance of one of processing clusters 2114A-2114N of FIG. 21. In at least one embodiment, processing cluster 2194 can be configured to execute many threads in parallel, where the term“thread” refers to an instance of a particular program executing on a particular set of input data. In at least one embodiment, Single Instruction Multiple Data (SIMD) instruction issue techniques are used to support parallel execution of a large number of threads with no or negligible specification impact. In at least one embodiment, Single Instruction Multiple Thread (SIMT) techniques are used to support parallel execution of a large number of generally synchronous threads using a common instruction unit configured to issue instructions to a set of processing engines within each processing cluster 2194.

[0211] In at least one embodiment, operation of processing cluster 2194 can be controlled via a pipeline manager 2132 that is assigned to SIMT parallel processor. In at least one embodiment, pipeline manager 2132 receives instructions from scheduler 2110 of FIG. 21, and manages execution of those instructions via graphics multiprocessor 2134 and / or texture unit 2136. In at least one embodiment, graphics multiprocessor 2134 is an exemplary instance of a SIMT parallel processor. However, in at least one embodiment, various types of SIMT parallel processors of differing architectures can be included within processing cluster 2194. In at least one embodiment, one or more instances of graphics multiprocessor 2134 can be included within processing cluster 2194. In at least one embodiment, graphics multiprocessor 2134 can process data, and a data crossbar 2140 can be used to distribute processed data to one of a number of possible destinations, including other shader units. In at least one embodiment, pipeline manager 2132 can facilitate distribution by specifying destinations for processed data as a function of its origin.

[0212] In at least one embodiment, each graphics multiprocessor 2134 within processing cluster 2194 can include an identical set of functional execution logic (e.g., arithmetic logic units, load store units (LSUs), etc.). In at least one embodiment, functional execution logic is configured in a pipelined manner in which new instructions can be issued before previous instructions are complete. In at least one embodiment, functional execution logic supports a variety of operations including integer and floating point arithmetic, comparison operations, Boolean operations, shift operations, and the like. In at least one embodiment, same functional-unit hardware is utilized to perform different operations in response to different instruction sets being issued to the functional units.

[0213] In at least one embodiment, instructions sent to the processing cluster 2194 constitute threads. In at least one embodiment, a group of threads executed across a set of parallel processing engines is a thread group. In at least one embodiment, the thread group executes programs on different input data. In at least one embodiment, each thread within the thread group may be assigned to a different processing engine within the graphics multiprocessor 2134. In at least one embodiment, the thread group may include fewer threads than the number of processing engines within the graphics multiprocessor 2134. In at least one embodiment, when the number of threads included in the thread group is less than the number of processing engines, one or more processing engines may be idle during a loop that is processing the thread group. In at least one embodiment, the thread group may also include more threads than the number of processing engines within the graphics multiprocessor 2134. In at least one embodiment, when the thread group includes more threads than the number of processing engines within the graphics multiprocessor 2134, processing can be performed in consecutive clock cycles. In at least one embodiment, multiple thread groups can be executed simultaneously on the graphics multiprocessor 2134.

[0214] In at least one embodiment, the graphics multiprocessor 2134 includes an internal cache memory for performing load and store operations. In at least one embodiment, the graphics multiprocessor 2134 may forgo the internal cache and use a cache memory within the processing cluster 2194 (e.g., L1 cache 2148). In at least one embodiment, each graphics multiprocessor 2134 may also access partition units (e.g., FIG. 21A The L2 cache is located within partition units 2120A-2120N, which are shared among all processing clusters 2194 and can be used to transfer data between threads. In at least one embodiment, the graphics multiprocessor 2134 can also access off-chip global memory, which may include one or more of local parallel processor memory and / or system memory. In at least one embodiment, any memory outside of the parallel processing unit 2102 can be used as global memory. In at least one embodiment, the processing cluster 2194 includes multiple instances of the graphics multiprocessor 2134, which can share common instructions and data that can be stored in the L1 cache 2148.

[0215] In at least one embodiment, each processing cluster 2194 may include an MMU 2145 configured to map virtual addresses to physical addresses. In at least one embodiment, one or more instances of MMU 2145 may reside within the memory interface 2118 of FIG. 21. In at least one embodiment, MMU 2145 includes a set of page table entries (PTEs) for mapping virtual addresses to physical addresses of tiles (more information about tiles) and optionally to cache line indices. In at least one embodiment, MMU 2145 may include an address translation back buffer (TLB) or a cache that may reside within the graphics multiprocessor 2134, L1 cache 2148, or processing cluster 2194. In at least one embodiment, physical addresses are processed to allocate surface data access locality for efficient request interleaving between partition units. In at least one embodiment, cache line indices may be used to determine whether a request for a cache line is a hit or a miss.

[0216] In at least one embodiment, the processing cluster 2194 may be configured such that each graphics multiprocessor 2134 is coupled to a texture unit 2136 to perform texture mapping operations, such as determining texture sample locations, reading texture data, and filtering texture data. In at least one embodiment, texture data is read as needed from an internal texture L1 cache (not shown) or from an L1 cache within the graphics multiprocessor 2134, and texture data is also retrieved from an L2 cache, local parallel processor memory, or system memory. In at least one embodiment, each graphics multiprocessor 2134 outputs a processed task to a data crossbar switch 2140 to provide the processed task to another processing cluster 2194 for further processing or to store the processed task in an L2 cache, local parallel processor memory, or system memory via a memory crossbar switch 2116. In at least one embodiment, the pre-raster operation unit (preROP) 2142 is configured to receive data from the graphics multiprocessor 2134 and direct the data to a ROP unit that may be located together with partitioning units described herein (e.g., partitioning units 2120A-2120N of FIG. 21). In at least one embodiment, the PreROP 2142 unit may perform optimizations for color blending, organize pixel color data, and perform address translation.

[0217] FIG. 21C A graphics multiprocessor 2196 according to at least one embodiment is illustrated. In at least one embodiment, the graphics multiprocessor 2196 is FIG. 21Bgraphics multiprocessor 2134. In at least one embodiment, graphics multiprocessor 2196 is coupled with pipeline manager 2132 of processing cluster 2194. In at least one embodiment, graphics multiprocessor 2196 has a graphics processing unit (GPU) execution pipeline that includes, without limitation, an instruction cache 2152, an instruction unit 2154, an address mapping unit 2156, a register file 2158, one or more GPGPU cores 2162, and one or more LSU’s 2166. GPGPU cores 2162 and LSUs 2166 are coupled with cache memory 2172 and shared memory 2170 via a memory and cache interconnect 2168.

[0218] In at least one embodiment, instruction cache 2152 receives a stream of instructions 2130 to execute from pipeline manager 2132. In at least one embodiment, instructions are cached in instruction cache 2152 and dispatched for execution by instruction unit 2154. In one embodiment, instruction unit 2154 can dispatch instructions to the threads of a thread group, with each thread of the thread group allocated a different execution unit within GPGPU cores 2162. In at least one embodiment, instructions can access any of the local, shared, or global address spaces. In at least one embodiment, address mapping unit 2156 can be used to convert addresses from the uniform memory address space into addresses that can be accessed by LSU 2166.

[0219] In at least one embodiment, register file 2158 provides a set of registers for functional units of graphics multiprocessor 2196. In at least one embodiment, register file 2158 provides temporary storage for operands of the data

[0220] In at least one embodiment, GPGPU cores 2162 can each include FPUs and / or ALUs for executing instructions for graphics processing. GPGPU cores 2162 can be similar to each other in architecture or can differ from each other in architecture. In at least one embodiment, a first portion of GPGPU cores 2162 includes single precision FPUs and integer ALUs, while a second portion of GPGPU cores includes double precision FPUs. In at least one embodiment, FPUs can implement IEEE 754-2008 standard for floating point arithmetic or enable variable precision floating point arithmetic. In at least one embodiment, graphics multi-processor 2196 can further include one or more fixed function or special-purpose processing units to perform specific functions in addition to the general purpose FLUs in GPGPU cores 2162. In at least one embodiment, one or more of GPGPU cores 2162 can also include fixed or special-purpose logic.

[0221] In at least one embodiment, GPGPU cores 2162 include SIMD logic capable of

[0222] In at least one embodiment, memory and cache interconnect 2168 is an interconnect network that connects each functional unit of graphics multiprocessor 2196 to register file 2158 and shared memory 2170. In at least one embodiment, memory and cache interconnect 2168 is a crossbar interconnect that allows LSUs 2166 to implement load and store operations between shared memory 2170 and register file 2158. In at least one embodiment, register file 2158 can operate at same frequency as GPGPU cores 2162, making the latency to transfer data between GPGPU cores 2162 and register file 2158 very low. In at least one embodiment, shared memory 2170 can be used to enable

[0223] In at least one embodiment, parallel processor or GPGPU as described herein is communicatively coupled to host / processor cores to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general purpose GPU (GPGPU) functions. In at least one embodiment, GPU can be communicatively coupled to host processor / cores over a bus or other interconnect (e.g., a high-speed

[0224] FIG. 22A graphics processor 2200 according to at least one embodiment is shown. In at least one embodiment, graphics processor 2200 includes a ring interconnect 2202, a pipeline front-end 2204, a media engine 2237, and graphics cores 2280A-2280N. In at least one embodiment, ring interconnect 2202 couples graphics processor 2200 to other processing units including other graphics processors or one or more general-purpose processor cores. In at least one embodiment, graphics processor 2200 is one of many processors integrated within a multi-core processing system.

[0225] In at least one embodiment, graphics processor 2200 receives batches of commands via ring interconnect 2202. In at least one embodiment, incoming commands are interpreted by a command streamer 2203 in pipeline front-end 2204. In at least one embodiment, graphics processor 2200 includes scalable execution logic to perform 3D geometry processing and media processing via the graphics cores 2280A-2280N. In at least one embodiment, for 3D geometry processing commands, command streamer 2203 supplies commands to geometry pipeline 2236. In at least one embodiment, for at least some media processing commands, command streamer 2203 supplies commands to a video front end 2234, which couples with a media engine 2237. In at least one embodiment, media engine 2237 includes a video quality engine (VQE) 2230 for video and image post-processing, and a multi-format encode / decode (MFX) 2233 engine to provide hardware-accelerated

[0226] In at least one embodiment, graphics processor 2200 includes a scalable thread execution resource with modular graphics cores 2280A-2280N (sometimes referred to as core slices) that each include multiple sub-cores 2250A-2250N, 2260A-2260N (sometimes referred to as core sub-slices). In at least one embodiment, graphics processor 2200 can have any number of graphics cores 2280A-2280N. In at least one embodiment, graphics processor 2200 includes graphics core 2280A having at least a first sub-core 2250A and a second sub-core 2260A. In at least one embodiment, graphics processor 2200 is a low power processor with a single sub-core (e.g., 2250A). In at least one embodiment, graphics processor 2200 includes multiple graphics cores 2280A-2280N each including a set of first sub-cores 2250A-2250N and a set of second sub-cores 2260A-2260N. In at least one embodiment, each of the first sub-cores 2250A-2250N includes at least a first set of execution units (EUs) 2252A-2252N and a media / texture sampler 2254A-2254N. In at least one embodiment, each of the second sub-cores 2260A-2260N includes at least a second set of execution units (EUs) 2262A-2262N and a sampler 2264A-2264N. In at least one embodiment, each of the sub-cores 2250A-2250N, 2260A-2260N share a set of shared resources 2270A-2270N. In at least one embodiment, shared resources include shared cache memory and pixel operation logic.

[0227] FIG. 23 A processor 2300 according to at least one embodiment is shown. In at least one embodiment, processor 2300 can include, without limitation, a logic circuit that executes instructions. In at least one embodiment, processor 2300 can execute instructions including x86 instructions, ARM instructions, specialized instructions for ASICs, and the like. In at least one embodiment, processor 2310 can include registers to store packed data, such as 64-bit wide MMX™ registers in microprocessors enabled with MMX technology by Intel Corporation of Santa Clara, California. In at least one embodiment, MMX registers available in integer and floating point form can operate with packed data elements that accompany SIMD and Streaming SIMD Extensions (“SSE”) instructions. In at least one embodiment, 128-bit wide XMM registers related to SSE2, SSE3, SSE4, AVX, or higher (“SSEx” in general) technology can hold such packed data operands. In at least one embodiment, processor 2310 can execute instructions to accelerate CUDA programs.

[0228] In at least one embodiment, processor 2300 includes an in-order front-end (“front-end”) 2301 to fetch instructions to be executed and prepare instructions to be used later in processor pipeline. In at least one embodiment, front-end 2301 can include several units. In at least one embodiment, instruction prefetcher 2326 fetches instructions from memory and provides instructions to instruction decoder 2328 which in turn decodes or interprets instructions. For example, in at least one embodiment, instruction decoder 2328 decodes a received instruction into one or more operations called “micro-instructions” or “micro-operations” (also called “micro ops” or “uops”) for execution. In at least one embodiment, instruction decoder 2328 parses the instruction into an operation code that it uses to identify an appropriate one of a set of executable micro ops and data and control fields that can be used by micro ops to execute the instruction. In at least one embodiment, a trace cache 2330 can assemble decoded micro ops into program ordered sequences or traces in a micro op queue 2334 for execution. In at least one embodiment, when trace cache 2330 encounters a complex instruction, microcode ROM 2332 provides micro ops needed to implement the instruction.

[0229] In at least one embodiment, some instructions can be converted into a single micro op, while others can need several micro ops to complete the full operation. In at least one embodiment, if more than four micro ops are needed to complete an instruction, then instruction decoder 2328 can access microcode ROM 2332 to execute the instruction. In at least one embodiment, instructions can be decoded into a small number of micro ops for processing at instruction decoder 2328. In at least one embodiment, if multiple micro ops are needed to complete an operation, then the instruction can be stored in microcode ROM 2332. In at least one embodiment, trace cache 2330 references a entry point programmable logic array (“PLA”) to determine the correct micro op pointer for reading microcode sequences from microcode ROM 2332 to complete one or more instructions, in accordance with at least one embodiment. In at least one embodiment, after microcode ROM 2332 completes the micro op ordering for an instruction, front-end 2301 of the machine can resume fetching micro ops from trace cache 2330.

[0230] In at least one embodiment, out-of-order execution engine (“out-of-order engine”) 2303 can prepare instructions for execution. In at least one embodiment, out-of-order execution logic has multiple buffers to smooth and reorder instruction flow to optimize performance as instructions are pipelined down and dispatched for execution. Out-of-order execution engine 2303 includes, without limitation, an allocator / register renamer 2340, a memory micro instruction queue 2342, an integer / floating point micro instruction queue 2344, a memory scheduler 2346, a fast scheduler 2302, a slow / general floating point scheduler (“slow / general FP scheduler”) 2304, and a simple floating point scheduler (“simple FP scheduler”) 2306. In at least one embodiment, fast scheduler 2302, slow / general floating point scheduler 2304, and simple floating point scheduler 2306 are also collectively referred to as “micro instruction schedulers 2302, 2304, 2306.” Allocator / register renamer 2340 allocates machine buffers and resources needed for each micro instruction to execute in order. In at least one embodiment, allocator / register renamer 2340 renames logical registers to entries in a register file. In at least one embodiment, allocator / register renamer 2340 also allocates entries for each micro instruction in one of two micro instruction queues, memory micro instruction queue 2342 for memory operations and integer / floating point micro instruction queue 2344 for non-memory operations, in front of memory scheduler 2346 and micro instruction schedulers 2302, 2304, 2306. In at least one embodiment, micro instruction schedulers 2302, 2304, 2306 determine when micro instructions are ready to execute based on readiness of their dependent input register operand sources and availability of execution resource micro instructions needed to complete. In at least one embodiment, fast scheduler 2302 of at least one embodiment can schedule on every half of a main clock cycle, while slow / general floating point scheduler 2304 and simple floating point scheduler 2306 can schedule once per main processor clock cycle. In at least one embodiment, micro instruction schedulers 2302, 2304, 2306 arbitrate for a dispatch port to dispatch micro instructions for execution.

[0231] In at least one embodiment, execution block 2311 includes, without limitation, integer register file / bypass network 2308, floating point register file / bypass network (“FP register file / bypass network”) 2310, address generation units (“AGUs”) 2312 and 2314, fast ALUs 2316 and 2318, slow ALU 2320, floating point ALU (“FP ALU”) 2322, and floating point move unit (“FP move”) 2324. In at least one embodiment, integer register file / bypass network 2308 and floating point register file / bypass network 2310 are also referred to herein as “register files 2308, 2310.” In at least one embodiment, AGUs 2312 and 2314, fast ALUs 2316 and 2318, slow ALU 2320, floating point ALU 2322, and floating point move unit 2324 are also referred to herein as “execution units 2312, 2314, 2316, 2318, 2320, 2322, and 2324.” In at least one embodiment, execution block can include, without limitation, any number (including zero) and type of register files, bypass networks, address generation units, and execution units (in any combination).

[0232] In at least one embodiment, register files 2308, 2310 can be arranged between microinstruction schedulers 2302, 2304, 2306 and execution units 2312, 2314, 2316, 2318, 2320, 2322, and 2324. In at least one embodiment, integer register file / bypass network 2308 performs integer operations. In at least one embodiment, floating point register file / bypass network 2310 performs floating point operations. In at least one embodiment, each of register files 2308, 2310 can include, without limitation, a bypass network that can bypass or forward a just-completed result that has not yet been written into a register file to a new dependee. In at least one embodiment, register files 2308, 2310 can communicate data with each other. In at least one embodiment, integer register file / bypass network 2308 can include, without limitation, two separate register files, one for lower 32 bits of data and a second for upper 32 bits of data. In at least one embodiment, floating point register file / bypass network 2310 can include, without limitation, 128 bit wide entries, as floating point instructions typically have operands that are 64 to 128 bits wide.

[0233] In at least one embodiment, execution units 2312, 2314, 2316, 2318, 2320, 2322, 2324 can execute instructions. In at least one embodiment, register files 2308, 2310 store integer and floating point data operand values upon which microinstructions require execution. In at least one embodiment, processor 2300 can include, without limitation, any number of execution units 2312, 2314, 2316, 2318, 2320, 2322, 2324 and combinations thereof. In at least one embodiment, floating point ALU 2322 and floating point move unit 2324 can execute floating point, MMX, SIMD, AVX and SSE, or other operations, including specialized machine learning instructions. In at least one embodiment, floating point ALU 2322 can include, without limitation, a 64 bit by 64 bit floating point divider to execute divide, square root, and remainder micro-ops. In at least one embodiment, instructions for dealing with floating point values can be handled with floating point hardware. In at least one embodiment, ALU operations can be passed to fast ALUs 2316, 2318. In at least one embodiment, fast ALUS 2316, 2318 can execute fast operations with an effective latency of half a clock cycle. In at least one embodiment, most complex integer operations enter slow ALU 2320 as slow ALU 2320 can include, without limitation, integer execution hardware for long latency type operations such as multiplies, shifts, flag logic, and branch processing. In at least one embodiment, memory load / store operations can be executed by AGUS 2312, 2314. In at least one embodiment, fast ALU 2316, fast ALU 2318, and slow ALU 2320 can execute integer operations on 64 bit data operands. In at least one embodiment, fast ALU 2316, fast ALU 2318, and slow ALU 2320 can be implemented to support a variety of data bit sizes including 16, 32, 128, 256, etc. In at least one embodiment, floating point ALU 2322 and floating point move unit 2324 can be implemented to support a range of operands having bits of various widths. In at least one embodiment, floating point ALU 2322 and floating point move unit 2324 can operate on 128 bits wide packed data operands in conjunction with SIMD and multimedia instructions.

[0234] In at least one embodiment, micro-instruction schedulers 2302, 2304, 2306 schedule dependent operations before parent load completion execution. In at least one embodiment, because micro-instructions can be speculatively scheduled and executed in processor 2300, processor 2300 can also include logic to handle memory misses. In at least one embodiment, if a data load in a data cache misses, there can be dependent operations running in the pipeline that cause the scheduler to temporarily have incorrect data. In at least one embodiment, a replay mechanism tracks and re-executes instructions that use incorrect data. In at least one embodiment, dependent operations can need to be replayed and independent operations can be allowed to complete. In at least one embodiment, a scheduler and replay mechanism of at least one embodiment of a processor can also be designed to capture instruction sequences for text string compare operations.

[0235] In at least one embodiment, the term “register” can refer to an on-board processor storage location that can be used as part of an instruction that identifies an operand. In at least one embodiment, a register can be one that can be used from outside of a processor (from a programmer’s perspective). In at least one embodiment, a register can not be limited to a particular type of circuit. Rather, in at least one embodiment, a register can store data, provide data, and perform functions described herein. In at least one embodiment, registers described herein can be implemented by circuitry within a processor using a variety of different techniques, such as dedicated physical registers, physical registers dynamically allocated using register renaming, a combination of dedicated and dynamically allocated physical registers, etc. In at least one embodiment, an integer register stores 32-bit integer data. A register file of at least one embodiment also contains eight multimedia SIMD registers for packed data.

[0236] FIG. 24 A processor 2400 according to at least one embodiment is shown. In at least one embodiment, processor 2400 includes, without limitation, one or more processor cores (cores) 2402A-2402N, an integrated memory controller 2414, and an integrated graphics processor 2408. In at least one embodiment, processor 2400 can include additional cores up to and including the additional processor cores 2402N represented by the dashed box. In at least one embodiment, each processor core 2402A-2402N includes one or more internal cache units 2404A-2404N. In at least one embodiment, each processor core can also access one or more shared cache units 2406.

[0237] In at least one embodiment, internal cache units 2404A-2404N and shared cache unit 2406 represent a cache memory hierarchy within processor 2400. In at least one embodiment, cache memory units 2404A-2404N can include one or more levels of cache, such as L2, L3, 4th level (L4), or other levels of cache, within each processor core and in a shared on-chip cache, where the highest level of cache is classified as an LLC before off-chip memory. In at least one embodiment, cache coherence logic maintains coherence between various cache units 2406 and 2404A-2404N.

[0238] In at least one embodiment, processor 2400 can also include a set of one or more bus controller units 2416 and a system agent core 2410. In at least one embodiment, one or more bus controller units 2416 manage a set of peripheral buses, such as one or more PCI or PCI Express buses. In at least one embodiment, system agent core 2410 provides management functionality for various processor components. In at least one embodiment, system agent core 2410 includes one or more integrated memory controllers 2414 to manage access to various external memory devices (not shown), such as one or more dynamic random access memory DRAM or static RAM (SRAM) devices.

[0239] In at least one embodiment, one or more processor cores 2402A-2402N include support for simultaneous multi-threading. In at least one embodiment, system agent core 2410 includes components for coordination and operation of processor cores 2402A-2402N during multi-threaded processing. In at least one embodiment, system agent core 2410 can additionally include a power control unit (PCU) including logic and components to regulate one or more power states of processor cores 2402A-2402N and graphics processor 2408.

[0240] In at least one embodiment, processor 2400 additionally includes graphics processor 2408 to perform graphics processing operations. In at least one embodiment, graphics processor 2408 couples with shared cache unit 2406 and system agent core 2410, including one or more integrated memory controllers 2414. In at least one embodiment, system agent core 2410 also includes a display controller 2411 to drive graphics processor output to one or more coupled displays. In at least one embodiment, display controller 2411 can also be a separate module coupled with graphics processor 2408 via at least one interconnect, or can be integrated within graphics processor 2408.

[0241] In at least one embodiment, ring-based interconnect unit 2412 is used to couple internal components of processor 2400. In at least one embodiment, alternative interconnect units, such as point-to-point interconnects, switched interconnects, or other technologies, may be used. In at least one embodiment, graphics processor 2408 is coupled to ring interconnect 2412 via I / O link 2413.

[0242] In at least one embodiment, I / O link 2413 represents at least one of a variety of I / O interconnects, including packaged I / O interconnects that facilitate communication between various processor components and high-performance embedded memory module 2418 (e.g., eDRAM module). In at least one embodiment, each of processor cores 2402A-2402N and graphics processor 2408 uses embedded memory module 2418 as a shared LLC.

[0243] In at least one embodiment, processor cores 2402A-2402N are homogeneous cores executing a common instruction set architecture. In at least one embodiment, processor cores 2402A-2402N are heterogeneous in terms of the instruction set architecture (ISA), with one or more processor cores 2402A-2402N executing a common instruction set, while one or more other processor cores 2402A-2402N execute a common instruction set or a subset of a different instruction set. In at least one embodiment, processor cores 2402A-2402N are heterogeneous in terms of microarchitecture, with one or more cores having relatively high power consumption coupled to one or more power cores having lower power consumption. In at least one embodiment, processor 2400 can be implemented on one or more chips or implemented as a SoC integrated circuit.

[0244] FIG. 25 A graphics processing unit (GPU) core 2500 according to at least one embodiment described is illustrated. In at least one embodiment, the GPU core 2500 is included within a GPU core array. In at least one embodiment, the GPU core 2500 (sometimes referred to as a core slice) may be one or more GPU cores within a modular GPU. In at least one embodiment, the GPU core 2500 is an example of a GPU core slice, and the GPU described herein may include multiple GPU core slices based on target power and performance envelopes. In at least one embodiment, each GPU core 2500 may include a fixed-function block 2530, also referred to as a sub-slice, coupled to a plurality of sub-cores 2501A-2501F, which includes modular blocks of general-purpose and fixed-function logic.

[0245] In at least one embodiment, fixed function block 2530 includes geometry / fixed function pipeline 2536, which can be shared by all of the sub-cores in graphics processor 2500, for example, in a lower performance and / or lower power graphics processor implementation. In at least one embodiment, geometry / fixed function pipeline 2536 includes a 3D fixed function pipeline, a video front-end unit, a thread generator and thread dispatcher, and a unified return buffer manager which manages a unified return buffer.

[0246] In at least one embodiment, fixed function block 2530 also includes a graphics SoC interface 2537, a graphics microcontroller 2538, and a media pipeline 2539. Graphics SoC interface 2537 provides an interface between graphics core 2500 and other processor cores within a SoC integrated circuit. In at least one embodiment, graphics microcontroller 2538 is a programmable sub-processor that is configurable to manage various functions of graphics processor 2500, including thread dispatch, scheduling, and pre-emption.

[0247] In at least one embodiment, SoC interface 2537 enables graphics core 2500 to communicate with general application processor cores within a SoC, including memory hierarchy elements such as a shared LLC cache, system RAM, and / or embedded on-chip or package DRAM. In at least one embodiment, SoC interface 2537 can also enable communication with fixed function devices within SoC, such as camera imaging pipelines, and enable use of and / or implementation of global memory atoms that can be shared between graphics core 2500 and a CPU within SoC. In at least one embodiment, SoC interface 2537 can also implement power management controls for graphics core 2500 and enable an interface between a clock domain of graphics core 2500 with other clock domains within SoC. In at least one embodiment, SoC interface 2537 enables receiving command buffers from a command streamer and global thread dispatcher that are configured to provide commands and instructions to each of one or more graphics cores within graphics processor 2500. In at least one embodiment, commands and instructions can be dispatched from a command streamer and global thread dispatcher to media pipeline 2539 when media operations are to be performed, or to the geometry and fixed function pipeline (e.g., geometry and fixed function pipeline 2536, geometry and fixed function pipeline 2514) when graphics processing operations are to be performed.

[0248] In at least one embodiment, graphics microcontroller 2538 can be configured to perform various scheduling and management tasks for graphics core 2500. In at least one embodiment, graphics microcontroller 2538 can perform graphics and / or compute workload scheduling on various graphics processing engines within execution unit (EU) arrays 2502A-2502F, 2504A-2504F of sub-cores 2501A-2501F. In at least one embodiment, host software executing on a CPU core of an SoC including graphics core 2500 can submit workloads for one of graphics processor bellman, which invokes scheduling operations on appropriate graphics engines. In at least one embodiment, scheduling operations include determining which workload to run next, submitting a workload to a command streamer, pre-empting existing workloads running on an engine, monitoring progress of a workload, and notifying host software when a workload completes. In at least one embodiment, graphics microcontroller 2538 can also facilitate low power or idle state conditioning of graphics core 2500, providing the ability to save and restore registers across low power state transitions of graphics core 2500 independent of operating system and / or graphics driver software on the system.

[0249] In at least one embodiment, graphics core 2500 can have more or less than the illustrated sub-cores 2501A-2501F, up to N modular sub-cores. For each set of N sub-cores, graphics core 2500 can also include shared function logic 2510, shared and / or cache memory 2512, geometry / fixed function pipeline 2514, and additional fixed function logic 2516 to accelerate various graphics and compute processing operations, in at least one embodiment. Shared function logic 2510 can include logic units (e.g., samplers, math, and / or inter-thread communication logic) that can be shared by each N sub-core within graphics core 2500, in at least one embodiment. Shared and / or cache memory 2512 can be an LLC for N sub-cores 2501A-2501F within graphics core 2500, and can also serve as shared memory accessible by multiple sub-cores, in at least one embodiment. Geometry / fixed function pipeline 2514 can be included in place of geometry / fixed function pipeline 2536 within fixed function block 2530, and can include the same or similar logic units, in at least one embodiment.

[0250] In at least one embodiment, graphics core 2500 includes additional fixed function logic 2516 which can include various fixed function acceleration logic to be utilized by graphics core 2500. In at least one embodiment, additional fixed function logic 2516 includes additional geometry pipeline for use in position only shading. In position only shading, there are at least two geometry pipelines, a full geometry pipeline and a cull pipeline within geometry / fixed function pipeline 2516, 2536, which is an additional geometry pipeline that can be included in additional fixed function logic 2516. In at least one embodiment, the cull pipeline is a trimmed down version of the full geometry pipeline. In at least one embodiment, the full pipeline and the cull pipeline can execute different instances of an application, each with separate state. In at least one embodiment, position only shading can hide long cull runs of triangles that are discarded, such that shading can be completed earlier in some cases. For example, in at least one embodiment, cull pipeline logic in additional fixed function logic 2516 can execute position shaders in parallel with a main application, and often generate critical results faster than the full pipeline because the cull pipeline takes and shades the position attributes of vertices without needing to perform rasterization and render pixels to a frame buffer. In at least one embodiment, the cull pipeline can use generated critical results to compute visibility information for all triangles, regardless of whether those triangles are culled. In at least one embodiment, the full pipeline, which can be referred to in this case as a replay pipeline, can consume the visibility information to skip culled triangles to only shade visible triangles that are finally passed to a rasterization stage.

[0251] In at least one embodiment, additional fixed function logic 2516 can also include general-purpose processing acceleration logic, such as fixed function matrix multiplication logic for implementing a reduced CUAD program.

[0252] In at least one embodiment, each graphics sub-core 2501A-2501F includes a set of execution resources within that can be used to perform graphics, media, and compute operations in response to requests by graphics pipeline, media pipeline, or shader program. In at least one embodiment, graphics sub-cores 2501A-2501F include multiple arrays of EU’s 2502A-2502F, 2504A-2504F, thread dispatch and inter-thread communication (TD / IC) logic 2503A-2503F, 3D (e.g., texture) samplers 2505A-2505F, media samplers 2506A-2506F, shader processors 2507A-2507F, and shared local memory (SLM) 2508A-2508F. Arrays of EU’s 2502A-2502F, 2504A-2504F each include multiple execution units (EUs), which are GPGPUs capable of performing floating-point and integer / fixed-point logic operations in service of a graphics, media, or compute operation, including graphics, media, or compute shader programs. In at least one embodiment, TD / IC logic 2503A-2503F performs local thread dispatch and thread control operations for execution units within a sub-core and facilitate communication between threads executing on execution units of a sub-core. In at least one embodiment, 3D samplers 2505A-2505F can read texture or other 3D graphics related data into memory. In at least one embodiment, 3D samplers can read texture data differently based on a configured sample state and texture format associated with a given texture. In at least one embodiment, media samplers 2506A-2506F can perform similar read operations based on a type and format associated with media data. In at least one embodiment, each graphics sub-core 2501A-2501F can alternatively include unified 3D and media samplers. In at least one embodiment, threads executing on execution units within each sub-core 2501A-2501F can make use of shared local memory 2508A-2508F within each sub-core for storage of thread private data and / or shared memory for thread group data sharing.

[0253] FIG. 26A parallel processing unit (“PPU”) 2600, in accordance with at least one embodiment, is shown. In at least one embodiment, PPU 2600 is configured with machine-readable code that, if executed by PPU 2600, causes PPU 2600 to perform some or all of the processes and techniques described herein. In at least one embodiment, PPU 2600 is a multi-threaded processor implemented on one or more integrated circuit devices and utilizes multi-threading as a latency-hiding technique designed to process computer-readable instructions (also referred to as machine-readable instructions or simply instructions) that are executed in parallel on multiple threads. In at least one embodiment, a thread refers to an execution thread and is an instance of a set of instructions configured to be executed by PPU 2600. In at least one embodiment, PPU 2600 is a graphics processing unit (“GPU”) configured to implement a graphics rendering pipeline for processing three-dimensional (“3D”) graphics data in order to generate two-dimensional (“2D”) image data for display on a display device, such as an LCD device. In at least one embodiment, PPU 2600 is used to perform computations, such as linear algebra operations and machine learning operations. FIG. 26 The example parallel processor is shown for illustrative purposes only and should be construed as a non-limiting example of a processor architecture implemented in at least one embodiment.

[0254] In at least one embodiment, one or more PPUs 2600 are configured to accelerate high-performance computing (“HPC”), data center, and machine learning applications. In at least one embodiment, one or more PPUs 2600 are configured to accelerate CUDA programs. In at least one embodiment, the PPU 2600 includes, but is not limited to, I / O unit 2606, front-end unit 2610, scheduler unit 2612, job allocation unit 2614, hub 2616, crossbar (“Xbar”) 2620, one or more general-purpose processing clusters (“GPC”) 2618, and one or more partitioning units (“memory partitioning units”) 2622. In at least one embodiment, the PPU 2600 is connected to a host processor or other PPU 2600 via one or more high-speed GPU interconnects (“GPU interconnects”) 2608. In at least one embodiment, the PPU 2600 is connected to a host processor or other peripheral device via a system bus or interconnect 2602. In one embodiment, the PPU 2600 is connected to local memory including one or more memory devices (“memory”) 2604. In at least one embodiment, the memory device 2604 includes, but is not limited to, one or more dynamic random access memory (“DRAM”) devices. In at least one embodiment, the one or more DRAM devices are configured and / or configurable as a high bandwidth memory (“HBM”) subsystem, and multiple DRAM dies are stacked within each device.

[0255] In at least one embodiment, the high-speed GPU interconnect 2608 may refer to a wire-based multi-channel communication link used by the system for scaling, and includes one or more PPUs 2600s (“CPUs”) coupled with one or more CPUs, supporting cache coherency between the PPUs 2600s and the CPUs, as well as CPU master control. In at least one embodiment, the high-speed GPU interconnect 2608 transmits data and / or commands to other units of the PPU 2600, such as one or more copy engines, video encoders, video decoders, power management units, and / or other components, via a hub 2616. FIG. 26 Other components that may not be explicitly shown.

[0256] In at least one embodiment, the I / O unit 2606 is configured to access the host processor via the system bus 2602. FIG. 26The I / O unit 2606 sends and receives communications (e.g., commands, data) to and from the system bus 2602 (e.g., over a wired or wireless connection). In at least one embodiment, the I / O unit 2606 communicates directly with the host processor(s) via the system bus 2602 or through one or more intermediate devices such as a memory bridge. In at least one embodiment, the I / O unit 2606 can communicate with one or more other processors (e.g., one or more PPUs 2600) via the system bus 2602. In at least one embodiment, the I / O unit 2606 implements a PCIe interface for communications over a PCIe bus. In at least one embodiment, the I / O unit 2606 implements interfaces for communicating with external devices.

[0257] In at least one embodiment, the I / O unit 2606 decodes packets received via the system bus 2602. In at least one embodiment, at least some packets represent commands configured to cause the PPU 2600 to perform various operations. In at least one embodiment, the I / O unit 2606 transmits decoded commands to various other units of the PPU 2600 as designated by the commands. In at least one embodiment, commands are transmitted to the front-end unit 2610 and / or to the hub 2616 or other units of the PPU 2600 such as one or more copy engines, video encoders, video decoders, power management units, etc. FIG. 26 In at least one embodiment, the I / O unit 2606 is configured to route communications between various logical units of the PPU 2600 (not explicitly shown). In at least one embodiment, the I / O unit 2606 is configured to provide a translation look-aside buffer (TLB) and a branch prediction buffer.

[0258] In at least one embodiment, a program executed by the host processor encodes a command stream in a buffer that provides a workload to the PPU 2600 for processing. In at least one embodiment, the workload includes instructions and data to be processed by those instructions. In at least one embodiment, the buffer is a region in memory that is accessible (e.g., read / write) by both the host processor and the PPU 2600 - the host interface unit can be configured to access memory requests transmitted via the I / O unit 2606 over the system bus 2602 to connect to a system memory buffer in the system memory via the system bus 2602. In at least one embodiment, the host processor writes the command stream to the buffer and then sends a pointer to the beginning of the command stream to the PPU 2600, causing the front-end unit 2610 to receive the one or more command stream pointers and manage one or more command streams, reading commands from the command stream and forwarding the commands to various units of the PPU 2600.

[0259] In at least one embodiment, front-end unit 2610 is coupled to a scheduler unit 2612 which configures various GPCs 2618 to process tasks defined by one or more command streams. In at least one embodiment, scheduler unit 2612 is configured to track state information related to various tasks managed by scheduler unit 2612, where state information can indicate which task is assigned to which GPC 2618, whether task is active or inactive, priority of task associated with it, and so forth. In at least one embodiment, scheduler unit 2612 manages multiple tasks that are executed on one or more GPCs 2618.

[0260] In at least one embodiment, scheduler unit 2612 is coupled to a work distribution unit 2614, which is configured to dispatch tasks for execution on GPCs 2618. In at least one embodiment, work distribution unit 2614 tracks a number of scheduled tasks received from scheduler unit 2612 and work distribution unit 2614 manages a pending task pool and an active task pool for each GPC 2618. In at least one embodiment, pending task pool includes a number of slots (e.g., 32 slots) that hold tasks assigned to be processed by a particular GPC 2618; active task pool can include a number of slots (e.g., 4 slots) for tasks that are actively being processed by GPC 2618, such that as one of GPCs 2618 completes processing of a task, the task is evicted from active task pool of GPC 2618 and one of other tasks from pending task pool is selected and scheduled for execution on GPC 2618. In at least one embodiment, if an active task is idle, for example, while waiting for a data dependency to resolve, the active task is evicted from GPC 2618 and returned to pending task pool while another task from pending task pool is selected and scheduled for execution on GPC 2618.

[0261] In at least one embodiment, work distribution unit 2614 communicates with one or more GPCs 2618 via XBar 2620. In at least one embodiment, XBar 2620 is an interconnect network that couples many units of PPU 2600 to other units of PPU 2600 and can be configured to couple work distribution unit 2614 to a particular GPC 2618. In at least one embodiment, other units of one or more PPUs 2600 can also be connected to XBar 2620 via hub 2616.

[0262] In at least one embodiment, tasks are managed by a scheduler unit 2612 and dispatched to one of GPCs 2618 by a work distribution unit 2614. GPCs 2618 are configured to process tasks and generate results. In at least one embodiment, results can be consumed by other tasks within GPC 2618, routed to different GPCs 2618 over the XBar 2620, or stored in memory 2604. In at least one embodiment, results can be written to memory 2604 by a partition unit 2622, which implements a memory interface for reading and writing data to memory 2604. In at least one embodiment, results can be transmitted over a high-speed GPU interconnect 2608 to another PPU 2600 or CPU. In at least one embodiment, PPU 2600 includes, without limitation, U partition units 2622 equal to the number of separate and distinct memory devices 2604 coupled to PPU 2600.

[0263] In at least one embodiment, a host processor executes a driver core that implements an application programming interface (API) that enables one or more applications executing on host processor to schedule operations to be performed on PPU 2600. In one embodiment, multiple compute applications are executed simultaneously by PPU 2600 and PPU 2600 provides isolation, quality of service (“QoS”), and independent address spaces for multiple compute applications. In at least one embodiment, an application generates instructions (e.g., in the form of API calls) that cause a driver core to generate one or more tasks for execution by PPU 2600 and driver core outputs tasks to one or more streams that are processed by PPU 2600. In at least one embodiment, each task includes one or more related thread groups, which can be referred to as warps. In at least one embodiment, a warp includes a plurality of related threads (e.g., 32 threads) that can be executed in parallel. In at least one embodiment, a cooperative thread can refer to a plurality of threads including instructions for performing a task and exchanging data via shared memory.

[0264] FIG. 27 A GPC 2700 according to at least one embodiment is shown. In at least one embodiment, GPC 2700 is a FIG. 26GPC 2618. In at least one embodiment, each GPC 2700 includes, without limitation, a plurality of hardware units for processing tasks and each GPC 2700 includes, without limitation, a pipeline manager 2702, a pre-raster operations unit (“PROP”) 2704, a raster engine 2708, a work distribution crossbar (“WDX”) 2716, a memory management unit (“MMU”) 2718, one or more data processing clusters (“DPCs”) 2706, and any suitable combination of such components.

[0265] In at least one embodiment, operation of GPC 2700 is controlled by pipeline manager 2702. In at least one embodiment, pipeline manager 2702 manages configuration of one or more DPCs 2706 to process tasks assigned to GPC 2700. In at least one embodiment, pipeline manager 2702 configures at least one of one or more DPCs 2706 to implement at least a portion of a graphics rendering pipeline. In at least one embodiment, DPC 2706 is configured to execute vertex shader programs on a programmable streaming multi-processor (“SM”) 2714. In at least one embodiment, pipeline manager 2702 is configured to route packets received from a work distribution unit to appropriate logical units within GPC 2700, and in at least one embodiment, some packets can be routed to fixed function hardware units in PROP 2704 and / or raster engine 2708 while other packets can be routed to DPCs 2706 for processing by a primitive engine 2712 or SM 2714. In at least one embodiment, pipeline manager 2702 configures at least one of DPCs 2706 to implement a neural network model and / or compute pipeline. In at least one embodiment, pipeline manager 2702 configures at least one of DPCs 2706 to execute at least a portion of a CUDA program.

[0266] In at least one embodiment, PROP unit 2704 is configured to route data generated by raster engine 2708 and DPCs 2706 to a raster operations (“ROP”) unit in a partition unit for processing. In at least one embodiment, PROP unit 2704 includes, without limitation, a cache memory for storing data used by PROP unit 2704 to perform one or more pre-rasterization operations on primitive. FIG. 25The memory partition unit 2522, among other things, is described in greater detail. In at least one embodiment, the PROP unit 2704 is configured to perform optimizations for color blending, organize pixel data, perform address translations, and the like. In at least one embodiment, the raster engine 2708 includes, without limitation, a number of fixed function hardware units configured to perform various raster operations, and in at least one embodiment, the raster engine 2708 includes, without limitation, a setup engine, a coarse raster engine, a cull engine, a clip engine, a fine raster engine, a tile aggregation engine, and any suitable combinations thereof. In at least one embodiment, the setup engine receives transformed vertices and generates plane equations associated with geometric primitives defined by the vertices; the plane equations are passed to the coarse raster engine to generate coverage information (e.g., x, y coverage masks for tiles) for the primitive; output from the coarse raster engine is passed to the cull engine, where fragments associated with primitives that fail a z-test are culled, and to the clip engine, where fragments that are outside a viewing frustum are clipped. In at least one embodiment, the culled and clipped fragments are passed to the fine raster engine to generate attributes of pixel fragments based on the plane equations generated by the setup engine. In at least one embodiment, output from the raster engine 2708 includes fragments that are to be processed by any suitable entity, such as by a fragment shader implemented within the DPC 2706.

[0267] In at least one embodiment, each DPC 2706 included in GPC 2700 includes, without limitation, an M-Pipe Controller (“MPC”) 2710; a primitive engine 2712; one or more SMs 2714; and any suitable combination thereof. In at least one embodiment, MPC 2710 controls operation of DPC 2706, routing incoming packets from pipeline manager 2702 to appropriate units in DPC 2706. In at least one embodiment, packets associated with vertices are routed to primitive engine 2712, which is configured to fetch vertex attributes associated with the vertices from memory; in contrast, data packets associated with a shader program can be transmitted to SM 2714.

[0268] In at least one embodiment, the SM 2714 includes, but is not limited to, a programmable streaming processor configured to process tasks represented by multiple threads. In at least one embodiment, the SM 2714 is multithreaded and configured to execute multiple threads (e.g., 32 threads) from a specific thread group concurrently, and implements a Single Instruction, Multiple Data (“SIMD”) architecture, wherein each thread in a group of threads (e.g., a thread bundle) is configured to process a different dataset based on the same instruction set. In at least one embodiment, all threads in the thread group execute the same instructions. In at least one embodiment, the SM 2714 implements a Single Instruction, Multiple Thread (“SIMT”) architecture, wherein each thread in a group of threads is configured to process a different dataset based on the same instruction set, but wherein individual threads in the thread group are allowed to diverge during execution. In at least one embodiment, a program counter, call stack, and execution state are maintained for each thread bundle, thereby achieving concurrency between the thread bundle and serial execution within the thread bundle when threads in the thread bundle diverge. In another embodiment, a program counter, call stack, and execution state are maintained for each individual thread, thereby ensuring equal concurrency among all threads within and between thread bundles. In at least one embodiment, an execution state is maintained for each individual thread, and threads executing the same instructions can be converged and executed in parallel to improve efficiency. The following is in conjunction with... FIG. 28 At least one embodiment of SM 2714 is described in more detail.

[0269] In at least one embodiment, the MMU 2718 is integrated with the GPC 2700 and memory partitioning unit (e.g., FIG. 25 The MMU 2718 provides an interface between partition units 2522 and provides virtual address to physical address translation, memory protection, and memory request arbitration. In at least one embodiment, the MMU 2718 provides one or more translation back buffers (“TLBs”) for performing virtual address to physical address translation in memory.

[0270] FIG. 28 A streaming multiprocessor (“SM”) 2800 according to at least one embodiment is illustrated. In at least one embodiment, the SM 2800 is FIG. 27SM 2714. In at least one embodiment, SM 2800 includes, without limitation, an instruction cache 2802; one or more scheduler units 2804; a register file 2808; one or more processing cores (“cores”) 2810; one or more special-function units (“SFUs”) 2812; one or more load / store units (“LSUs”) 2814; an interconnect network 2816; shared memory / level-one (“LI”) cache 2818; and any suitable combination thereof. In at least one embodiment, a work distribution unit dispatches tasks for execution on general processing clusters (“GPCs”) of parallel processing units (“PPUs”) and each task is assigned a specific data processing cluster (“DPC”) within a GPC and, if task is associated with a shader program, to one of SMs 2800. In at least one embodiment, scheduler units 2804 receive tasks from work distribution unit and manage dispatch of instructions to one or more thread blocks for execution on SM 2800. In at least one embodiment, scheduler units 2804 schedule thread blocks to be executed to be dispatched as thread warps of parallel threads, where at least one thread block is assigned to each thread warp. In at least one embodiment, each thread warp executes a thread. In at least one embodiment, scheduler units 2804 manage a plurality of different thread blocks, allocate thread warps to different thread blocks, and then dispatch instructions from different ones of the plurality of cooperation groups to various functional units (e.g., processing cores 2810, SFUs 2812, and LSUs 2814) during each clock cycle.

[0271] In at least one embodiment, a “cooperative group” can refer to a programming model for organizing groups of communication threads that allows developers to express the granularity at which threads are communicating, enabling richer, more efficient parallel decomposition. In at least one embodiment, a cooperative launch API supports synchronization between thread blocks to execute parallel algorithms. In at least one embodiment, an API of a conventional programming model provides a single, simple construct for synchronizing cooperating threads: a barrier across all threads of a thread block (e.g., a syncthreads() function). However, in at least one embodiment, a programmer can define thread groups at less than a thread block granularity and synchronize within defined groups to achieve higher performance, design flexibility, and software reuse in the form of collective group-wide function interfaces. In at least one embodiment, cooperative groups enable programmers to explicitly define thread groups at sub-block and multi-block granularities and perform collective operations, such as synchronizing threads in a cooperative group. In at least one embodiment, sub-block granularity is as small as a single thread. In at least one embodiment, a programming model supports clean composition across software boundaries, so that library and utility functions can safely synchronize in their local environment without having to make assumptions about convergence. In at least one embodiment, cooperative group primitives enable new patterns of cooperative parallelism, including but not limited to producer-consumer parallelism, opportunistic parallelism, and global synchronization across a grid of thread blocks.

[0272] In at least one embodiment, a dispatch unit 2806 is configured to send instructions to one or more of the functional units, and a scheduler unit 2804 includes, without limitation, two dispatch units 2806 that enable two different instructions from the same thread to be dispatched in each clock cycle. In at least one embodiment, each scheduler unit 2804 includes a single dispatch unit 2806 or an additional dispatch unit 2806.

[0273] In at least one embodiment, each SM 2800 includes, without limitation, a register file 2808 that provides a set of registers for functional units of SM 2800. In at least one embodiment, register file 2808 is partitioned between functional units, with each functional unit being allocated a dedicated portion of register file 2808. In at least one embodiment, register file 2808 is partitioned between different warps executed by SM 2800, with register file 2808 providing temporary storage for operands of the operations calculated and performed by functional units. In at least one embodiment, each SM 2800 includes, without limitation, a plurality L of processing cores 2810. In at least one embodiment, SM 2800 includes, without limitation, a large number (e.g., 128 or more) of different processing cores 2810. In at least one embodiment, each processing core 2810 includes, without limitation, a full-pipe, single precision, double precision, and / or mixed precision processing unit including, without limitation, a floating point arithmetic logic unit and an integer arithmetic logic unit. In at least one embodiment, floating point arithmetic logic units implement IEEE 754-2008 standard for floating point arithmetic. In at least one embodiment, processing cores 2810 include, without limitation, 64 single-precision (32-bit) floating point cores, 64 integer cores, 32 double-precision (64-bit) floating point cores, and 8 tensor cores.

[0274] In at least one embodiment, tensor cores are configured to perform matrix operations. In at least one embodiment, one or more tensor cores are included in processing cores 2810. In at least one embodiment, tensor cores are configured to perform deep learning matrix arithmetic, such as convolution operations for neural network training and inferencing. In at least one embodiment, each tensor core operates on 4x4 matrices and performs matrix multiplication and accumulation operations D = A x B + C, where A, B, C, and D are 4x4 matrices.

[0275] In at least one embodiment, matrix multiplication inputs A and B are 16-bit floating point matrices, and accumulation matrices C and D are 16-bit floating point or 32-bit floating point matrices. In at least one embodiment, a tensor core performs 32-bit floating point accumulation operations on 16-bit floating point input data. In at least one embodiment, 16-bit floating point multiplication uses 64 operations and results in a full precision product, which is then accumulated with other intermediate products using 32-bit floating point addition for 4x4x4 matrix multiplication. In at least one embodiment, tensor cores are used to perform larger two-dimensional or higher dimensional matrix operations composed of these smaller elements. In at least one embodiment, an API such as a CUDA-C++ API exposes specialized matrix load, matrix multiply and accumulate, and matrix store operations to efficiently use tensor cores from a CUDA-C++ program. In at least one embodiment, at a CUDA level, a warp level interface assumes a 16x16 size matrix across all 32 warp threads.

[0276] In at least one embodiment, each SM 2800 includes, without limitation, M SFUs 2812 to perform special functions (e.g., certain math functions, atomics, bit endcomplement, etc.). In at least one embodiment, SFUs 2812 include, without limitation, tree traversal units configured to traverse a hierarchical tree data structure. In at least one embodiment, SFUs 2812 include, without limitation, texture units configured to perform texture mapping operations. In at least one embodiment, texture units are configured to load a texture map (e.g., a 2D array of texture pixels) from memory and sample the texture map to produce sampled texture values for use by a shader program executed by SM 2800. In at least one embodiment, texture maps are stored in shared memory / L1 cache 2818. In at least one embodiment, texture units use mip-maps (e.g., different levels of detail for a texture map) to perform texture operations such as filtering operations. In at least one embodiment, each SM 2800 includes, without limitation, two texture units.

[0277] In at least one embodiment, each SM 2800 includes, without limitation, N LSUs 2814 that implement load and store operations between shared memory / L1 cache 2818 and register file 2808. In at least one embodiment, each SM 2800 includes, without limitation, interconnect network 2816 that connects each of the functional units to register file 2808 and LSUs 2814 to register file 2808 and shared memory / L1 cache 2818. In at least one embodiment, interconnect network 2816 is a cross-bar switch that can be configured to connect any of the functional units to any of the registers in register file 2808, and to connect LSUs 2814 to registers in register file 2808 and to memory locations in shared memory / L1 cache 2818.

[0278] In at least one embodiment, shared memory / L1 cache 2818 is an array of on-chip memory that, in at least one embodiment, allows data storage and communication between SMs 2800 and the graphics processing engine and between threads within SMs 2800. In at least one embodiment, shared memory / L1 cache 2818 includes, without limitation, 128 KB of storage and is located on a path from SMs 2800 to partition units. In at least one embodiment, shared memory / L1 cache 2818 is used for caching reads and writes, in at least one embodiment. In at least one embodiment, one or more of shared memory / L1 cache 2818, L2 cache, and memory are backing stores.

[0279] In at least one embodiment, combining data cache and shared memory functionality into a single memory block provides improved performance for both types of memory access. In at least one embodiment, capacity is used by programs that do not use shared memory or used as a cache, e.g., if shared memory is configured to use half of capacity, then textures and load / store operations can use remaining capacity. According to at least one embodiment, integration within shared memory / L1 cache 2818 enables shared memory / L1 cache 2818 to be used as a high-throughput pipeline for streaming data while providing high bandwidth and low latency access to frequently reused data. In at least one embodiment, when configured for general purpose parallel computation, a simpler configuration can be used compared to graphics processing. In at least one embodiment, fixed function GPU is bypassed, creating a more straightforward programming model. In at least one embodiment, in a general purpose parallel computation configuration, work distribution unit allocates and distributes blocks of threads directly to DPCs. In at least one embodiment, threads in a block execute the same program, use unique thread IDs in a computation to ensure that each thread generates unique results, use SM 2800 to execute program and perform computation, use shared memory / L1 cache 2818 to communicate between threads, and use LSU 2814 to read and write global memory through shared memory / L1 cache 2818 and memory partition unit. In at least one embodiment, when configured for general purpose parallel computation, SM 2800 writes commands to scheduler unit 2704 that can be used to launch new work on DPCs.

[0280] In at least one embodiment, PPU is included in a desktop computer, laptop computer, tablet computer, server computer, supercomputer, smart- phone (e.g., a wireless, hand-held device), PDA, digital camera, vehicle, head mounted display, hand-held electronic device, etc. or is coupled to such devices. In at least one embodiment, PPU is implemented on a single semiconductor

[0281] In at least one embodiment, PPU can be included on a graphics card that includes one or more memory devices. In at least one embodiment, graphics card can be configured to interface with a PCIe slot on a motherboard of a desktop computer. In at least one embodiment, PPU can be an integrated GPU (“iGPU”) included in a chipset of a motherboard.

[0282] Software constructs for general purpose computation

[0283] The following figures set forth, without limitation, example software constructs for implementing at least one embodiment.

[0284] FIG. 29 A software stack of a programming platform is shown, in accordance with at least one embodiment. In at least one embodiment, a programming platform is a platform for utilizing hardware on a computing system to accelerate computational tasks. In at least one embodiment, a software developer can access a programming platform through libraries, compiler directives, and / or extensions to a programming language. In at least one embodiment, a programming platform can be, but is not limited to, CUDA, Radeon Open Compute Platform (“ROCm”), OpenCL (OpenCL TM ), SYCL, or Intel One API.

[0285] In at least one embodiment, software stack 2900 of a programming platform provides an execution environment for application 2901. In at least one embodiment, application 2901 can include any computer software capable of launching on software stack 2900. In at least one embodiment, application 2901 can include, but is not limited to, artificial intelligence (“AI”) / machine learning (“ML”) applications, high performance computing (“HPC”) applications, virtual desktop infrastructure (“VDI”), or data center workloads.

[0286] In at least one embodiment, application 2901 and software stack 2900 run on hardware 2907. In at least one embodiment, hardware 2907 can include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices that support a programming platform. In at least one embodiment, for example with CUDA, software stack 2900 can be vendor specific and only compatible with devices from a particular vendor. In at least one embodiment, for example in OpenCL, software stack 2900 can be used with devices from different vendors. In at least one embodiment, hardware 2907 includes a host connected to one or more devices that can be accessed via an application programming interface (API) call to perform computational tasks. In at least one embodiment, in contrast to a host within hardware 2907, which can include, but is not limited to, a CPU (but can also include a computing device) and its memory, a device within hardware 2907 can include, but is not limited to, a GPU, FPGA, AI engine, or other computing device (but can also include a CPU) and its memory.

[0287] In at least one embodiment, software stack 2900 of a programming platform includes, without limitation, a plurality of libraries 2903, a runtime 2905, and a device kernel driver 2906. In at least one embodiment, each of libraries 2903 can include data and programming code that can be used by computer programs and leveraged during software development. In at least one embodiment, libraries 2903 can include, without limitation, pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. In at least one embodiment, libraries 2903 include functions that are optimized for execution on one or more types of devices. In at least one embodiment, libraries 2903 can include, without limitation, functions for performing mathematical, deep learning, and / or other types of operations on a device. In at least one embodiment, libraries 2903 are associated with corresponding APIs 2902, which can include one or more APIs that expose functions implemented in libraries 2903.

[0288] In at least one embodiment, application 2901 is written as source code that is compiled into executable code, as discussed in more detail below with respect to FIG. 29B. FIGS. 34-36 In at least one embodiment, executable code of application 2901 can run, at least partially, on an execution environment provided by software stack 2900. In at least one embodiment, during execution of application 2901, code can be derived that needs to run on a device (as opposed to a host). In such a case, in at least one embodiment, runtime 2905 can be invoked to load and launch the necessary code on a device. In at least one embodiment, runtime 2905 can include any technically feasible runtime system capable of supporting execution of application 2901.

[0289] In at least one embodiment, runtime 2905 is implemented as one or more runtime libraries associated with corresponding APIs (which are shown as APIs 2904). In at least one embodiment, one or more such runtime libraries can include, without limitation, functions for memory management, execution control, device management, error handling, and / or synchronization, among others. In at least one embodiment, memory management functions can include, without limitation, functions for allocating, deallocating, and copying device memory, as well as transferring data between host memory and device memory. In at least one embodiment, execution control functions can include, without limitation, functions for launching functions on a device (sometimes referred to as “kernels” when functions are global functions that are callable from a host), and functions for setting attribute values in buffers maintained by a runtime library for a given function to be executed on a device.

[0290] In at least one embodiment, the runtime library and the corresponding API 2904 can be implemented in any technically feasible manner. In at least one embodiment, one (or any number of) APIs may expose a low-level set of functions for fine-grained control of the device, while another (or any number of) APIs may expose such a higher-level set of functions. In at least one embodiment, a high-level runtime API can be built on top of the low-level APIs. In at least one embodiment, one or more runtime APIs may be language-specific APIs layered on top of language-independent runtime APIs.

[0291] In at least one embodiment, device kernel driver 2906 is configured to facilitate communication with the underlying device. In at least one embodiment, device kernel driver 2906 can provide APIs such as API 2904 and / or low-level functions upon which other software depends. In at least one embodiment, device kernel driver 2906 can be configured to compile intermediate representation (“IR”) code into binary code at runtime. In at least one embodiment, for CUDA, device kernel driver 2906 can compile non-hardware-specific parallel thread execution (“PTX”) IR code into binary code for a specific target device (cached compiled binary code), sometimes also referred to as “final” code. In at least one embodiment, doing so allows the final code to run on the target device, which may not exist when the source code was initially compiled into PTX code. Alternatively, in at least one embodiment, the device source code can be compiled into binary code offline, without requiring device kernel driver 2906 to compile the IR code at runtime.

[0292] FIG. 30 The illustration shows an embodiment according to at least one of the embodiments. FIG. 29 The software stack 2900 is a CUDA implementation. In at least one embodiment, the CUDA software stack 3000 on which an application 3001 can be launched includes a CUDA library 3003, a CUDA runtime 3005, a CUDA driver 3007, and a device kernel driver 3008. In at least one embodiment, the CUDA software stack 3000 executes on hardware 3009, which may include a CUDA-enabled GPU developed by NVIDIA Corporation of Santa Clara, California.

[0293] In at least one embodiment, application 3001, CUDA runtime 3005, and device kernel driver 3008 can respectively perform functions similar to those of application 2801, runtime 2805, and device kernel driver 2906, in combination with the above. FIG. 29CUDA Driver 3007, in at least one embodiment, includes a library (libcuda.so) that implements a CUDA Driver API 3006. In at least one embodiment, similar to a CUDA Runtime API 3004 implemented by a CUDA runtime library (cudart), the CUDA Driver API 3006 can expose, without limitation, functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability, etc. In at least one embodiment, the CUDA Driver API 3006 differs from the CUDA Runtime API 3004 in that the CUDA Runtime API 3004 simplifies device code management by providing implicit initialization, context (similar to a process) management, and module (similar to a dynamically loaded library) management. In contrast to the high-level CUDA Runtime API 3004, in at least one embodiment, the CUDA Driver API 3006 is a low-level API that provides more fine-grained control over the device, particularly with respect to context and module loading. In at least one embodiment, the CUDA Driver API 3006 can expose functions for context management that are not exposed by the CUDA Runtime API 3004. In at least one embodiment, the CUDA Driver API 3006 is also language agnostic and supports, for example, OpenCL in addition to supporting the CUDA Runtime API 3004. Further, in at least one embodiment, development libraries including the CUDA Runtime 3005 can be considered separate from driver components, including the user-mode CUDA Driver 3007 and the kernel-mode device driver 3008 (sometimes also referred to as a “display” driver).

[0294] In at least one embodiment, CUDA libraries 3003 can include, without limitation, mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries that can be utilized by parallel computing applications, such as application 3001. In at least one embodiment, CUDA libraries 3003 can include mathematical libraries such as a cuBLAS library that is an implementation of basic linear algebra subprograms (“BLAS”) for performing linear algebra operations, a cuFFT library for computing fast Fourier transforms (“FFTs”), a cuRAND library for generating random numbers, etc. In at least one embodiment, CUDA libraries 3003 can include deep learning libraries such as a cuDNN library for primitives of deep neural networks and a TensorRT platform for high-performance deep learning inference, etc.

[0295] FIG. 31 FIG. 30 shows a diagram of a system including a training data preparation component, in accordance with at least one embodiment FIG. 29ROCm implementation of the software stack 2900. In at least one embodiment, the ROCm software stack 3100 on which the application 3101 can launch includes a language runtime 3103, a system runtime 3105, a thunk 3107, and a ROCm kernel driver 3108. In at least one embodiment, the ROCm software stack 3100 executes on hardware 3109, which can include a GPU that supports ROCm, developed by AMD Corporation of Santa Clara, California.

[0296] In at least one embodiment, the application 3101 can perform similar functionality as the application 2901 discussed above in conjunction with FIG. 29 In at least one embodiment, the language runtime 3103 and the system runtime 3105 can perform similar functionality as the runtime 2905 discussed above in conjunction with FIG. 29 In at least one embodiment, the language runtime 3103 and the system runtime 3105 differ in that the system runtime 3105 is a language-agnostic runtime that implements the ROCr system runtime API 3104 and utilizes a Heterogeneous System Architecture (“HSA”) runtime API. In at least one embodiment, the HSA runtime API is a thin user-mode API that exposes interfaces for accessing and interacting with AMD GPUs, including functions for memory management, execution control dispatching of kernels through the architecture, error handling, system and agent information, and runtime initialization and shutdown, among others. In at least one embodiment, the language runtime 3103 is an implementation of a language-specific runtime API 3102 layered on top of the ROCr system runtime API 3104 as compared to the system runtime 3105. In at least one embodiment, a language runtime API can include, without limitation, a Heterogeneous Compute Interface for Portability (“HIP”) language runtime API, a Heterogeneous Compute Compiler (“HCC”) language runtime API, or an OpenCL API, among others. In particular, the HIP language is an extension of the C++ programming language with functionally similar versions of CUDA mechanisms, and in at least one embodiment, the HIP language runtime API includes functions similar to the CUDA runtime API 3004 discussed above in conjunction with FIG. 30

[0297] ​In at least one embodiment, thunk (ROCt) 3107 is an interface 3106 that can be used to interface with underlying ROCm drivers 3108. In at least one embodiment, ROCm drivers 3108 are ROCk drivers, which are a combination of AMDGPU drivers and HSA kernel drivers (amdkfd). In at least one embodiment, AMDGPU drivers are device kernel drivers for GPUs developed by AMD that perform similar functions to those discussed above in connection with FIG. 29 HSA kernel drivers 2906. In at least one embodiment, HSA kernel drivers are drivers that allow different types of processors to more efficiently share system resources via hardware features.

[0298] In at least one embodiment, various libraries (not shown) can be included in ROCm software stack 3100 above language runtime 3103 and provide similar functionality to CUDA libraries 3003 discussed above in connection with FIG. 30 In at least one embodiment, various libraries can include, but are not limited to, math, deep learning, and / or other libraries such as a hipBLAS library that implements similar functions to CUDA cuBLAS, a rocFFT library similar to CUDA cuFFT for computing FFTs, etc.

[0299] FIG. 32 FIG. 32 illustrates an OpenCL implementation of software stack 2900 of FIG. 29 In at least one embodiment, OpenCL software stack 3200 on which application 3201 can be launched includes an OpenCL framework 3210, an OpenCL runtime 3206, and drivers 3207. In at least one embodiment, OpenCL software stack 3200 executes on hardware 3009 that is not vendor-specific. In at least one embodiment, because devices developed by different vendors support OpenCL, specific OpenCL drivers can be needed to interoperate with hardware from such vendors.

[0300] In at least one embodiment, application 3201, OpenCL runtime 3206, device kernel drivers 3207, and hardware 3208 can perform similar functions to those discussed above in connection with FIG. 29 application 2901, runtime 2905, device kernel drivers 2906, and hardware 2907. In at least one embodiment, application 3201 also includes OpenCL kernels 3202 that have code to be executed on a device.

[0301] In at least one embodiment, OpenCL defines a “platform” that allows a host to control devices connected to that host. In at least one embodiment, OpenCL framework provides a platform layer API and a runtime API, shown as platform API 3203 and runtime API 3205. In at least one embodiment, runtime API 3205 uses a context to manage execution of kernels on a device. In at least one embodiment, each identified device can be associated with a respective context, which runtime API 3205 can use to manage a command queue, program and kernel objects, shared memory objects, etc. for that device. In at least one embodiment, platform API 3203 exposes functions that allow a device context to be used to select and initialize a device, submit work to a device via a command queue, enable data transfers to and from a device, etc. Additionally, in at least one embodiment, OpenCL framework provides various built-in functions (not shown), including mathematical functions, relational functions, and image processing functions, etc.

[0302] In at least one embodiment, compiler 3204 is also included in OpenCL framework 3210. In at least one embodiment, source code can be compiled offline before an application is executed or online during execution of an application. In contrast to CUDA and ROCm, OpenCL applications in at least one embodiment can be compiled online by compiler 3204, which is included to represent any number of compilers that can be used to compile source and / or IR code (e.g., Standard Portable Intermediate Representation (“SPIR-V”) code) into binary code. Alternatively, in at least one embodiment, OpenCL applications can be compiled offline before such applications are executed.

[0303] FIG. 33 Software supported by a programming platform in accordance with at least one embodiment is shown. In at least one embodiment, programming platform 3304 is configured to support various programming models 3303, middleware and / or libraries 3302, and frameworks 3301 that an application 3300 can rely on. In at least one embodiment, application 3300 can be an AI / ML application implemented using, for example, a deep learning framework (e.g., MXNet, PyTorch, or TensorFlow) that can rely on libraries such as cuDNN, NVIDIA Collective Communications Library (“NCCL”), and / or NVIDIA Developer Data Loading Library (“DALI”) CUDA libraries to provide accelerated computation on underlying hardware.

[0304] In at least one embodiment, programming platform 3304 can be a combination of the above described programming platforms 3102 and 3204, respectively. In at least one embodiment, programming platform 3304 can be a combination of the above described programming platforms 3102 and 3204, respectively.FIG. 30 , FIG. 31 and FIG. 32 One of the CUDA, ROCm, or OpenCL platforms described above. In at least one embodiment, programming platform 3304 supports multiple programming models 3303, which are abstractions of the underlying computing system that allow for expression of algorithms and data structures. In at least one embodiment, programming models 3303 can expose features of the underlying hardware in order to improve performance. In at least one embodiment, programming models 3303 can include, but are not limited to, CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++ AMP”), Open Multi-Processing (“OpenMP”), Open Accelerators (“OpenACC”), and / or Vulcan Compute.

[0305] In at least one embodiment, libraries and / or middleware 3302 provide implementations of abstractions of programming models 3304. In at least one embodiment, such libraries include data and programming code that can be used by computer programs and utilized during software development. In at least one embodiment, such middleware includes software that provides services to applications in addition to those that can be obtained from programming platform 3304. In at least one embodiment, libraries and / or middleware 3302 can include, but are not limited to, cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Additionally, in at least one embodiment, libraries and / or middleware 3302 can include NCCL and ROCm Communication Collectives Library (“RCCL”) libraries, which provide communication routines for GPUs, MIOpen libraries for deep learning acceleration, and / or Eigen libraries for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0306] In at least one embodiment, application frameworks 3301 rely on libraries and / or middleware 3302. In at least one embodiment, each application framework 3301 is a software framework used to implement a standard structure for application software. Returning to the AI / ML example discussed above, in at least one embodiment, an AI / ML application can be implemented using a framework such as a Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet deep learning framework.

[0307] FIG. 34 shows compiled code to run on FIGS. 29-32on one of the programming platforms. In at least one embodiment, compiler 3401 receives source code 3400, which includes both host code as well as device code. In at least one embodiment, compiler 3401 is configured to convert source code 3400 into host executable code 3402 for execution on a host and device executable code 3403 for execution on a device. In at least one embodiment, source code 3400 can be compiled offline prior to execution of an application, or online during execution of an application.

[0308] In at least one embodiment, source code 3400 can include code in any programming language supported by compiler 3401, such as C++, C, Fortran, etc. In at least one embodiment, source code 3400 can include a single-source file with a mix of host code and device code, with locations of device code indicated therein. In at least one embodiment, the single-source file can be a.cu file including CUDA code or a.hip.cpp file including HIP code. Alternatively, in at least one embodiment, source code 3400 can include multiple source code files, rather than a single-source file, with host code and device code separated.

[0309] In at least one embodiment, compiler 3401 is configured to compile source code 3400 into host executable code 3402 for execution on a host and device executable code 3403 for execution on a device. In at least one embodiment, compiler 3401 performs operations including parsing source code 3400 into an abstract syntax tree (AST), performing optimizations, and generating executable code. In at least one embodiment where source code 3400 includes a single-source file, compiler 3401 can separate device code from host code in such single-source file, compile device code and host code into device executable code 3403 and host executable code 3402, respectively, and link device executable code 3403 and host executable code 3402 together in a single file, as discussed in more detail below with respect to FIG. 4. FIG. 35 discussed in more detail.

[0310] In at least one embodiment, host executable code 3402 and device executable code 3403 can be in any suitable format, such as binary code and / or IR code. In the case of CUDA, in at least one embodiment, host executable code 3402 can include native object code, while device executable code 3403 can include PTX intermediate representation code. In the case of ROCm, in at least one embodiment, both host executable code 3402 and device executable code 3403 can include object binary code.

[0311] FIGS. 29-32 is a more detailed illustration of compiling code to execute on one of the programming platforms of FIG. 36 In at least one embodiment, compiler 3501 is configured to receive source code 3500, compile source code 3500, and output executable 3510. In at least one embodiment, source code 3500 is a single source file, such as a.cu file, a.hip.cpp file, or other format of file, that includes both host code and device code. In at least one embodiment, compiler 3501 can be, without limitation, an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code in a.cu file, or an HCC compiler for compiling HIP code in a.hip.cpp file.

[0312] In at least one embodiment, compiler 3501 includes a compiler front end 3502, a host compiler 3505, a device compiler 3506, and a linker 3509. In at least one embodiment, compiler front end 3502 is configured to separate device code 3504 from host code 3503 in source code 3500. In at least one embodiment, device code 3504 is compiled by device compiler 3506 into device executable code 3508, which can include binary code or IR code, as described. In at least one embodiment, host code 3503 is separately compiled by host compiler 3505 into host executable code 3507. In at least one embodiment, for NVCC, host compiler 3505 can be, without limitation, a general C / C++ compiler that outputs native object code, while device compiler 3506 can be, without limitation, a low-level virtual machine (“LLVM”) based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code. In at least one embodiment, for HCC, both host compiler 3505 and device compiler 3506 can be, without limitation, LLVM based compilers that output target binary code.

[0313] In at least one embodiment, after source code 3500 is compiled into host executable code 3507 and device executable code 3508, linker 3509 links the host and device executable code 3507 and 3508 together in executable 3510. In at least one embodiment, native object code for the host and PTX or binary code for the device can be linked together in an executable and linkable format (“ELF”) file, which is a container format for storing object code.

[0314] FIG. 34Conversion of source code prior to compilation is shown in accordance with at least one embodiment. In at least one embodiment, source code 3600 is passed through a conversion tool 3601 that converts source code 3600 to converted source code 3602. In at least one embodiment, a compiler 3603 is used to compile converted source code 3602 to host executable code 3604 and device executable code 3605, a process similar to that of compiler 3401 compiling source code 3400 to host executable code 3402 and device executable code 3403, as described above in conjunction with FIG. 3. In at least one embodiment, compiler 3603 can be a different compiler than compiler 3401, or can be the same compiler as compiler 3401. FIGS. 37A-38

[0315] In at least one embodiment, the conversion performed by conversion tool 3601 is used to port source code 3600 to execute in a different environment than originally intended. In at least one embodiment, conversion tool 3601 can include, without limitation, a HIP translator that is used to “hipify” CUDA code for a CUDA platform to HIP code that can be compiled and executed on a ROCm platform. In at least one embodiment, conversion of source code 3600 can include parsing source code 3600 and converting calls to APIs provided by one programming model (e.g., CUDA) to corresponding calls to APIs provided by another programming model (e.g., HIP), as discussed in more detail below in conjunction with FIG. 36. FIG. 37A In at least one embodiment, automatic conversion performed by conversion tool 3601 can sometimes be incomplete, requiring additional human effort to fully port source code 3600.

[0316] Configuring a GPU for general-purpose computing

[0317] The following figures set forth, without limitation, exemplary architectures for compiling and executing compute source code, in accordance with at least one embodiment.

[0318] FIG. 37B ​A system 3700 configured to compile and execute CUDA source code 3710 using different types of processing units is shown, in accordance with at least one embodiment. In at least one embodiment, system 3700 includes, without limitation, CUDA source code 3710, a CUDA compiler 3750, host executable code 3770(1), host executable code 3770(2), CUDA device executable code 3784, a CPU 3790, a CUDA-enabled GPU 3794, a GPU 3792, a CUDA to HIP translation tool 3720, a HIP source code 3730, a HIP compiler driver 3740, an HCC 3760, and HCC device executable code 3782.

[0319] In at least one embodiment, CUDA source code 3710 is a collection of human-readable code in the CUDA programming language. In at least one embodiment, CUDA code is human-readable code in the CUDA programming language. In at least one embodiment, the CUDA programming language is an extension of the C++ programming language that includes, without limitation, mechanisms to define device code and to distinguish between device code and host code. In at least one embodiment, device code is source code that, after compilation, can be executed in parallel on a device. In at least one embodiment, a device can be a processor optimized for parallel instruction processing, such as CUDA-enabled GPU 3790, GPU 37192, or another GPGPU, etc. In at least one embodiment, host code is source code that, after compilation, can be executed on a host. In at least one embodiment, a host is a processor optimized for sequential instruction processing, such as CPU 3790.

[0320] In at least one embodiment, CUDA source code 3710 includes, without limitation, any number (including zero) of global functions 3712, any number (including zero) of device functions 3714, any number (including zero) of host functions 3716, and any number (including zero) of host / device functions 3718. In at least one embodiment, global functions 3712, device functions 3714, host functions 3716, and host / device functions 3718 can be intermixed in CUDA source code 3710. In at least one embodiment, each global function 3712 is executable on a device and is callable from a host. Thus, in at least one embodiment, one or more of global functions 3712 can serve as an entry point for a device. In at least one embodiment, each global function 3712 is a kernel. In at least one embodiment, and in a technique known as dynamic parallelism, one or more global functions 3712 define a kernel that is executable on a device and is callable from such a device. In at least one embodiment, a kernel is executed N times in parallel by N different threads on a device during execution (where N is any positive integer).

[0321] In at least one embodiment, each device function 3714 executes on a device and is callable only from such a device. In at least one embodiment, each host function 3716 executes on a host and is callable only from such a host. In at least one embodiment, each host / device function 3716 defines both a host version of a function executable on a host and callable only from such a host, as well as a device version of a function executable on a device and callable only from such a device.

[0322] In at least one embodiment, CUDA source code 3710 can also include, without limitation, any number of calls to any number of functions defined by CUDA runtime API 3702. In at least one embodiment, CUDA runtime API 3702 can include, without limitation, any number of functions that execute on a host for allocating and deallocating device memory, transferring data between host memory and device memory, managing systems with multiple devices, etc. In at least one embodiment, CUDA source code 3710 can also include any number of calls to any number of functions specified in any number of other CUDA APIs. In at least one embodiment, a CUDA API can be any API designed to be used by CUDA code. In at least one embodiment, CUDA APIs include, without limitation, CUDA runtime API 3702, CUDA driver APIs, APIs for any number of CUDA libraries, etc. In at least one embodiment and relative to CUDA runtime API 3702, CUDA driver APIs are lower-level APIs, but can provide more fine-grained control over a device. In at least one embodiment, examples of CUDA libraries include, without limitation, cuBLAS, cuFFT, cuRAND, cuDNN, etc.

[0323] In at least one embodiment, CUDA compiler 3750 compiles input CUDA code (e.g., CUDA source code 3710) to generate host executable code 3770(1) and CUDA device executable code 3784. In at least one embodiment, CUDA compiler 3750 is NVCC. In at least one embodiment, host executable code 3770(1) is a compiled version of host code included in input source code that is executable on CPU 3790. In at least one embodiment, CPU 3790 can be any processor optimized for sequential instruction processing.

[0324] In at least one embodiment, CUDA device executable code 3784 is a compiled version of device code included in input source code executable on a CUDA-enabled GPU 3794. In at least one embodiment, CUDA device executable code 3784 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3784 includes, without limitation, IR code, such as PTX code, which is further compiled into binary code for a particular target device (e.g., CUDA-enabled GPU 3794) at runtime by a device driver. In at least one embodiment, CUDA-enabled GPU 3794 can be any processor optimized for parallel instruction processing and supporting CUDA. In at least one embodiment, CUDA-enabled GPU 3794 is developed by NVIDIA Corporation of Santa Clara, CA.

[0325] In at least one embodiment, CUDA to HIP translation tool 3720 is configured to translate CUDA source code 3710 into functionally similar HIP source code 3730. In at least one embodiment, HIP source code 3730 is a collection of human-readable code in the HIP programming language. In at least one embodiment, HIP code is human-readable code in the HIP programming language. In at least one embodiment, the HIP programming language is an extension of the C++ programming language that includes, without limitation, functionally similar versions of CUDA mechanisms for defining device code and distinguishing device code from host code. In at least one embodiment, the HIP programming language can include a subset of functionality of the CUDA programming language. In at least one embodiment, for example, the HIP programming language includes, without limitation, mechanisms to define global functions 3712, but such a HIP programming language can lack support for dynamic parallelism, and thus global functions 3712 defined in HIP code are only callable from a host.

[0326] In at least one embodiment, HIP source code 3730 includes, without limitation, any number (including zero) of global functions 3712, any number (including zero) of device functions 3714, any number (including zero) of host functions 3716, and any number (including zero) of host / device functions 3718. In at least one embodiment, HIP source code 3730 can also include any number of calls to any number of functions specified in a HIP runtime API 3732. In one embodiment, HIP runtime API 3732 includes, without limitation, functionally similar versions of a subset of functions included in CUDA runtime API 3702. In at least one embodiment, HIP source code 3730 can also include any number of calls to any number of functions specified in any number of other HIP APIs. In at least one embodiment, a HIP API can be any API designed for use with HIP code and / or ROCm. In at least one embodiment, HIP APIs include, without limitation, HIP runtime API 3732, a HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, and the like.

[0327] In at least one embodiment, CUDA to HIP translation tool 3720 translates each kernel call in CUDA code from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in CUDA code to any number of other functionally similar HIP calls. In at least one embodiment, a CUDA call is a call to a function specified in a CUDA API, and a HIP call is a call to a function specified in a HIP API. In at least one embodiment, CUDA to HIP translation tool 3720 translates any number of calls to functions specified in CUDA runtime API 3702 to any number of calls to functions specified in HIP runtime API 3732.

[0328] In at least one embodiment, CUDA to HIP translation tool 3720 is a tool known as hipify-perl, which performs a text-based translation process. In at least one embodiment, CUDA to HIP translation tool 3720 is a tool known as hipify-clang, which performs a more complex and robust translation process relative to hipify-perl, involving parsing of CUDA code using clang (a compiler frontend), followed by translation of resulting symbols. In at least one embodiment, in addition to modifications performed by CUDA to HIP translation tool 3720, proper translation of CUDA code to HIP code can also require modifications (e.g., manual edits).

[0329] In at least one embodiment, HIP compiler driver 3740 is a front end that determines target device 3746 and then configures a compiler compatible with target device 3746 to compile HIP source code 3730. In at least one embodiment, target device 3746 is a processor optimized for parallel instruction processing. In at least one embodiment, HIP compiler driver 3740 can determine target device 3746 in any technically feasible manner.

[0330] In at least one embodiment, if target device 3746 is compatible with CUDA (e.g., a CUDA-enabled GPU 3794), then HIP compiler driver 3740 generates HIP / NVCC compilation commands 3742. In at least one embodiment and in conjunction with FIG. 37C In more detail, HIP / NVCC compilation commands 3742 configure CUDA compiler 3750 to compile HIP source code 3730 using, without limitation, a HIP-to-CUDA translation header and a CUDA runtime library. In at least one embodiment and in response to HIP / NVCC compilation commands 3742, CUDA compiler 3750 generates host executable code 3770(1) and CUDA device executable code 3784.

[0331] In at least one embodiment, if target device 3746 is not compatible with CUDA, then HIP compiler driver 3740 generates HIP / HCC compilation commands 3744. In at least one embodiment and in conjunction with FIG. 37A In more detail, HIP / HCC compilation commands 3744 configure HCC 3760 to compile HIP source code 3730 using a HCC header and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3744, HCC 3760 generates host executable code 3770(2) and HCC device executable code 3782. In at least one embodiment, HCC device executable code 3782 is a compiled version of device code contained in HIP source code 3730 that is executable on GPU 3792. In at least one embodiment, GPU 3792 can be any processor optimized for parallel instruction processing that is not compatible with CUDA and is compatible with HCC. In at least one embodiment, GPU 3792 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, GPU 3792 is a GPU 3792 that is not CUDA-enabled.

[0332] For illustrative purposes only, in FIG. 37BCUDA source code 3710 to execute on CPU 3790 and a different device in at least one embodiment. In at least one embodiment, a direct CUDA flow compiles CUDA source code 3710 to execute on CPU 3790 and a CUDA-enabled GPU 3794 without converting CUDA source code 3710 to HIP source code 3730. In at least one embodiment, an indirect CUDA flow converts CUDA source code 3710 to HIP source code 3730 and then compiles HIP source code 3730 to execute on CPU 3790 and a CUDA-enabled GPU 3794. In at least one embodiment, a CUDA / HCC flow converts CUDA source code 3710 to HIP source code 3730 and then compiles HIP source code 3730 to execute on CPU 3790 and a GPU 3792.

[0333] A direct CUDA flow that can be implemented in at least one embodiment can be depicted by the dashed line and series of bubble annotations Al-A3. In at least one embodiment, and as shown by bubble annotation Al, a CUDA compiler 3750 receives CUDA source code 3710 and a CUDA compile command 3748 that configures CUDA compiler 3750 to compile CUDA source code 3710. In at least one embodiment, CUDA source code 3710 used in a direct CUDA flow is written in a CUDA programming language that is based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In at least one embodiment, and in response to CUDA compile command 3748, CUDA compiler 3750 generates host executable code 3770(1) and CUDA device executable code 3784 (denoted with bubble annotation A2). In at least one embodiment and as shown with bubble annotation A3, host executable code 3770(1) and CUDA device executable code 3784 can be executed on CPU 3790 and CUDA-enabled GPU 3794, respectively. In at least one embodiment, CUDA device executable code 3784 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3784 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0334] The indirect CUDA flow that can be implemented in at least one embodiment can be described by the dashed line and series of bubble annotations Bl- B6. In at least one embodiment and as shown by bubble annotation Bl, the CUDA to HIP translation tool 3720 receives CUDA source code 3710. In at least one embodiment and as shown by bubble annotation B2, the CUDA to HIP translation tool 3720 translates the CUDA source code 3710 to HIP source code 3730. In at least one embodiment and as shown by bubble annotation B3, the HIP compiler driver 3740 receives the HIP source code 3730 and determines whether the target device 3746 has CUDA enabled.

[0335] In at least one embodiment and as shown by bubble annotation B4, the HIP compiler driver 3740 generates HIP / NVCC compilation commands 3742 and sends both the HIP / NVCC compilation commands 3742 and the HIP source code 3730 to the CUDA compiler 3750. In at least one embodiment and as described in greater detail below in connection with FIG. 37, the HIP / NVCC compilation commands 3742 configure the CUDA compiler 3750 to compile the HIP source code 3730 using, without limitation, a HIP to CUDA translation header and a CUDA runtime library. FIG. 37C In at least one embodiment and in response to the HIP / NVCC compilation commands 3742, the CUDA compiler 3750 generates host executable code 3770(1) and CUDA device executable code 3784 (represented by bubble annotation B5). In at least one embodiment and as shown by bubble annotation B6, the host executable code 3770(1) and the CUDA device executable code 3784 can be executed on a CPU 3790 and a CUDA-enabled GPU 3794, respectively. In at least one embodiment, the CUDA device executable code 3784 includes, without limitation, binary code. In at least one embodiment, the CUDA device executable code 3784 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0336] The CUDA / HCC flow that can be implemented in at least one embodiment can be described by the solid line and series of bubble annotations Cl- C6. In at least one embodiment and as shown by bubble annotation Cl, the CUDA to HIP translation tool 3720 receives CUDA source code 3710. In at least one embodiment and as shown by bubble annotation C2, the CUDA to HIP translation tool 3720 translates the CUDA source code 3710 to HIP source code 3730. In at least one embodiment and as shown by bubble annotation C3, the HIP compiler driver 3740 receives the HIP source code 3730 and determines that the target device 3746 does not have CUDA enabled.

[0337] In at least one embodiment, HIP compiler driver 3740 generates HIP / HCC compilation commands 3744 and sends both HIP / HCC compilation commands 3764 and HIP source code 3730 to HCC 3760 (represented with bubble annotation C4). In at least one embodiment and as described in greater detail below, HIP / HCC compilation commands 3764 configure HCC 3760 to compile HIP source code 3730 using, without limitation, HCC headers and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3744, HCC 3760 generates host executable code 3770(2) and HCC device executable code 3782 (represented with bubble annotation C5). In at least one embodiment and as represented by bubble annotation C6, host executable code 3770(2) and HCC device executable code 3782 can be executed on CPU 3790 and GPU 3792, respectively. FIG. 37B In at least one embodiment and as described in greater detail below, HIP / HCC compilation commands 3764 configure HCC 3760 to compile HIP source code 3730 using, without limitation, HCC headers and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3744, HCC 3760 generates host executable code 3770(2) and HCC device executable code 3782 (represented with bubble annotation C5). In at least one embodiment and as represented by bubble annotation C6, host executable code 3770(2) and HCC device executable code 3782 can be executed on CPU 3790 and GPU 3792, respectively.

[0338] In at least one embodiment, after CUDA source code 3710 is converted to HIP source code 3730, HIP compiler driver 3740 can then be used to generate executable code for a CUDA-enabled GPU 3794 or GPU 3792 without CUDA being re-implemented as a HIP conversion tool 3720. In at least one embodiment, CUDA-to-HIP conversion tool 3720 converts CUDA source code 3710 to HIP source code 3730, which is then stored in memory. In at least one embodiment, HIP compiler driver 3740 then configures HCC 3760 to generate host executable code 3770(2) and HCC device executable code 3782 based on HIP source code 3730. In at least one embodiment, HIP compiler driver 3740 then configures CUDA compiler 3750 to generate host executable code 3770(1) and CUDA device executable code 3784 based on stored HIP source code 3730.

[0339] FIG. 37A A system 3704 configured to compile and execute CUDA source code 3710 using CPU 3790 and a CUDA-enabled GPU 3794 is shown, in accordance with at least one embodiment. In at least one embodiment, system 3704 includes, without limitation, CUDA source code 3710, CUDA-to-HIP conversion tool 3720, HIP source code 3730, HIP compiler driver 3740, CUDA compiler 3750, host executable code 3770(1), CUDA device executable code 3784, CPU 3790, and CUDA-enabled GPU 3794. FIG. 37A A system 3704 configured to compile and execute CUDA source code 3710 using CPU 3790 and a CUDA-enabled GPU 3794 is shown, in accordance with at least one embodiment. In at least one embodiment, system 3704 includes, without limitation, CUDA source code 3710, CUDA-to-HIP conversion tool 3720, HIP source code 3730, HIP compiler driver 3740, CUDA compiler 3750, host executable code 3770(1), CUDA device executable code 3784, CPU 3790, and CUDA-enabled GPU 3794.

[0340] In at least one embodiment, and as previously described herein in connection with FIG. 37C In at least one embodiment, and as previously described herein in connection with

[0341] In at least one embodiment, CUDA to HIP translation tool 3720 translates CUDA source code 3710 into HIP source code 3730. In at least one embodiment, CUDA to HIP translation tool 3720 translates each kernel call in CUDA source code 3710 from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in CUDA source code 3710 to any number of other functionally similar HIP calls.

[0342] In at least one embodiment, HIP compiler driver 3740 determines that target device 3746 is CUDA-enabled, and generates HIP / NVCC compilation command 3742. In at least one embodiment, HIP compiler driver 3740 then configures CUDA compiler 3750 via HIP / NVCC compilation command 3742 to compile HIP source code 3730. In at least one embodiment, as part of configuring CUDA compiler 3750, HIP compiler driver 3740 provides access to a HIP to CUDA translation header 3752. In at least one embodiment, HIP to CUDA translation header 3752 translates any number of mechanisms (e.g., functions) specified in any number of HIP APIs to any number of mechanisms specified in any number of CUDA APIs. In at least one embodiment, CUDA compiler 3750 uses HIP to CUDA translation header 3752 in conjunction with a CUDA runtime library 3754 corresponding to CUDA runtime API 3702 to generate host executable code 3770(1) and CUDA device executable code 3784. In at least one embodiment, host executable code 3770(1) and CUDA device executable code 3784 can then be executed on CPU 3790 and CUDA-enabled GPU 3794, respectively. In at least one embodiment, CUDA device executable code 3784 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3784 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0343] FIG. 37A A system 3706 configured to compile and execute CUDA source code 3710 using a CPU 3790 and a non-CUDA-enabled GPU 3792 is shown, in accordance with at least one embodiment. FIG. 37A In at least one embodiment, system 3706 includes, without limitation, CUDA source code 3710, a CUDA to HIP translation tool 3720, HIP source code 3730, a HIP compiler driver 3740, an HCC 3760, host executable code 3770(2), HCC device executable code 3782, CPU 3790, and GPU 3792.

[0344] In at least one embodiment, and as previously described herein in conjunction with FIG. 38 CUDA source code 3710 includes, without limitation, any number (including zero) of global functions 3712, any number (including zero) of device functions 3714, any number (including zero) of host functions 3716, and any number (including zero) of host / device functions 3718, in at least one embodiment. In at least one embodiment, CUDA source code 3710 also includes, without limitation, any number of calls to any number of functions specified in any number of CUDA APIs.

[0345] In at least one embodiment, CUDA to HIP translation tool 3720 translates CUDA source code 3710 into HIP source code 3730. In at least one embodiment, CUDA to HIP translation tool 3720 translates each kernel call in CUDA source code 3710 from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in source code 3710 to any number of other functionally similar HIP calls.

[0346] In at least one embodiment, HIP compiler driver 3740 then determines that target device 3746 is not CUDA-enabled and generates HIP / HCC compilation commands 3744. In at least one embodiment, HIP compiler driver 3740 then configures HCC 3760 to perform HIP / HCC compilation commands 3744, thereby compiling HIP source code 3730. In at least one embodiment, HIP / HCC compilation commands 3744 configure HCC 3760 to use, without limitation, HIP / HCC runtime library 3758 and HCC headers 3756 to generate host executable code 3770(2) and HCC device executable code 3782. In at least one embodiment, HIP / HCC runtime library 3758 corresponds to HIP runtime API 3732. In at least one embodiment, HCC headers 3756 include, without limitation, any number and type of interoperability mechanisms for HIP and HCC. In at least one embodiment, host executable code 3770(2) and HCC device executable code 3782 can be executed on CPU 3790 and GPU 3792, respectively.

[0347] FIG. 37C FIG. 19 illustrates an exemplary kernel converted by CUDA to HIP conversion tool 3720 of FIG. 18, in accordance with at least one embodiment. FIG. 38 In at least one embodiment, CUDA source code 3710 divides the overall problem that a given kernel is designed to solve into relatively coarse sub-problems that can be solved independently using thread blocks. In at least one embodiment, each thread block includes, without limitation, any number of threads. In at least one embodiment, each sub-problem is divided into relatively fine pieces that can be solved in parallel by threads in a thread block in collaboration. In at least one embodiment, threads within a thread block can collaborate by sharing data through shared memory and by synchronizing execution to coordinate memory access.

[0348] In at least one embodiment, CUDA source code 3710 organizes thread blocks associated with a given kernel into a one-, two-, or three-dimensional grid of thread blocks. In at least one embodiment, each thread block includes, without limitation, any number of threads, and the grid includes, without limitation, any number of thread blocks.

[0349] In at least one embodiment, a kernel is a function in device code defined with a “__global__” declaration specifier. In at least one embodiment, a CUDA kernel launch syntax 3810 is used to specify a size of a grid of threads to execute a kernel for a given kernel call and an associated stream. In at least one embodiment, a CUDA kernel launch syntax 3810 is specified as “KernelName<<<GridSize, BlockSize, SharedMemorySize, Stream>>>(KernelArguments);”. In at least one embodiment, an execution configuration syntax is a “<<<...>>>” construct that is inserted between a kernel name (“KernelName”) and a parenthetical list of kernel arguments (“KernelArguments”). In at least one embodiment, a CUDA kernel launch syntax 3810 includes, without limitation, a CUDA launch function syntax instead of an execution configuration syntax.

[0350] In at least one embodiment, “GridSize” is of type dim3 and specifies a size and dimension of a grid. In at least one embodiment, type dim3 is a CUDA-defined structure that includes, without limitation, unsigned integers x, y, and z. In at least one embodiment, if z is not specified, it defaults to 1. In at least one embodiment, if y is not specified, it defaults to 1. In at least one embodiment, a number of thread blocks in a grid is equal to a product of GridSize.x, GridSize.y, and GridSize.z. In at least one embodiment, “BlockSize” is of type dim3 and specifies a size and dimension of each thread block. In at least one embodiment, a number of threads per thread block is equal to a product of BlockSize.x, BlockSize.y, and BlockSize.z. In at least one embodiment, each thread executing a kernel has a unique thread ID that is accessible within a kernel through a built-in variable such as “threadIdx”.

[0351] In at least one embodiment, with respect to CUDA kernel launch syntax 3810, “SharedMemorySize” is an optional parameter that specifies the number of bytes of shared memory to dynamically allocate for each thread block in addition to statically allocated memory for a given kernel call. In at least one embodiment and with respect to CUDA kernel launch syntax 3810, SharedMemorySize defaults to zero. In at least one embodiment and with respect to CUDA kernel launch syntax 3810, “stream” is an optional parameter that specifies an associated stream and defaults to zero to specify the default stream. In at least one embodiment, a stream is a sequence of commands (which can be issued by different host threads) that are executed in-order. In at least one embodiment, different streams can execute commands out-of-order or concurrently with respect to each other.

[0352] In at least one embodiment, CUDA source code 3710 includes, without limitation, a kernel definition and a main function for exemplary kernel “MatAdd”. In at least one embodiment, the main function is host code that executes on a host and includes, without limitation, a kernel call to cause kernel MatAdd to execute on a device. In at least one embodiment, as shown, kernel MatAdd adds two matrices A and B of size NxN, where N is a positive integer, and stores the result in matrix C. In at least one embodiment, the main function defines a threadsPerBlock variable to be 16x16 and a numBlocks variable to be N / 16 x N / 16. In at least one embodiment, the main function then specifies a kernel call “MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);”. In at least one embodiment, and in accordance with CUDA kernel launch syntax 3810, kernel MatAdd is executed using a thread block grid of size N / 16 x N / 16, with each thread block of size 16 x 16. In at least one embodiment, each thread block includes 256 threads, a grid with enough blocks is created so that each matrix element has one thread, and each thread in that grid executes kernel MatAdd to perform one pairwise addition.

[0353] In at least one embodiment, while converting CUDA source code 3710 to HIP source code 3730, CUDA to HIP translation tool 3720 converts each kernel call in CUDA source code 3710 from CUDA kernel launch syntax 3810 to HIP kernel launch syntax 3820, and converts any number of other CUDA calls in source code 3710 to any number of other functionally similar HIP calls. In at least one embodiment, HIP kernel launch syntax 3820 is specified as “hipLaunchKernelGGL(KernelName, GridSize, BlockSize, SharedMemorySize, Stream, KernelArguments);”. In at least one embodiment, each of KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments has the same meaning in HIP kernel launch syntax 3820 as in CUDA kernel launch syntax 3810 (described previously herein). In at least one embodiment, parameters SharedMemorySize and Stream are required in HIP kernel launch syntax 3820, but are optional in CUDA kernel launch syntax 3810.

[0354] In at least one embodiment, in addition to the kernel call to cause kernel MatAdd to execute on a device, FIG. 38 The portion of HIP source code 3730 depicted in FIG. 39 is the same as the portion of CUDA source code 3710 depicted in In at least one embodiment, kernel MatAdd is defined in HIP source code 3730 with the same “__global__” declaration specifier as kernel MatAdd is defined in CUDA source code 3710. In at least one embodiment, the kernel call in HIP source code 3730 is “hipLaunchKernelGGL(MatAdd, numBlocks, threadsPerBlock, 0, 0, A, B, C);”, while the corresponding kernel call in CUDA source code 3710 is “MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);”.

[0355] FIG. 37C is shown in more detail in accordance with at least one embodiment FIG. 40CUDA-enabled GPU 3792. In at least one embodiment, GPU 3792 is developed by AMD Corporation of Santa Clara, CA. In at least one embodiment, GPU 3792 can be configured to perform compute operations in a highly parallel manner. In at least one embodiment, GPU 3792 is configured to perform graphics pipeline operations such as draw commands, pixel operations, geometric calculations, and other operations associated with rendering images to a display. In at least one embodiment, GPU 3792 is configured to perform operations that are not graphics-related. In at least one embodiment, GPU 3792 is configured to perform both graphics-related operations and operations that are not graphics-related. In at least one embodiment, GPU 3792 can be configured to execute device code included in HIP source code 3730.

[0356] In at least one embodiment, GPU 3792 includes, without limitation, any number of programmable processing units 3920, a command processor 3910, an L2 cache 3922, a memory controller 3970, a DMA engine 3980(1), a system memory controller 3982, a DMA engine 3980(2), and a GPU controller 3984. In at least one embodiment, each programmable processing unit 3920 includes, without limitation, a workload manager 3930 and any number of compute units 3940. In at least one embodiment, command processor 3910 reads commands from one or more command queues (not shown) and distributes commands to workload managers 3930. In at least one embodiment, for each programmable processing unit 3920, a relevant workload manager 3930 distributes work to compute units 3940 included in programmable processing unit 3920. In at least one embodiment, each compute unit 3940 can execute any number of thread blocks, but each thread block executes on a single compute unit 3940. In at least one embodiment, a workgroup is a thread block.

[0357] In at least one embodiment, each compute unit 3940 includes, without limitation, any number of SIMD units 3950 and a shared memory 3960. In at least one embodiment, each SIMD unit 3950 implements a SIMD architecture and is configured to execute operations in parallel. In at least one embodiment, each SIMD unit 3950 includes, without limitation, a vector ALU 3952 and a vector register file 3954. In at least one embodiment, each SIMD unit 3950 executes a different thread bundle. In at least one embodiment, a thread bundle is a group of threads (e.g., 16 threads), where each thread in a thread bundle belongs to a single thread block and is configured to process a different set of data based on a single instruction set. In at least one embodiment, one or more threads in a thread bundle can be disabled using predication. In at least one embodiment, a lane is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a thread bundle. In at least one embodiment, different wavefronts in a thread block can be synchronized together and communicate via shared memory 3960.

[0358] In at least one embodiment, programmable processing units 3920 are referred to as “shader processing units.” In at least one embodiment, each programmable processing unit 3920 includes, without limitation, any number of geometry processors, any number of rasterizers, any number of render back-ends, workload manager 3930, and any number of compute units 3940, in addition to compute units 3940.

[0359] In at least one embodiment, compute units 3940 share L2 cache 3922. In at least one embodiment, L2 cache 3922 is partitioned. In at least one embodiment, all compute units 3940 in GPU 3792 have access to GPU memory 3990. In at least one embodiment, memory controllers 3970 and system memory controllers 3982 facilitate data transfers between GPU 3792 and a host, and DMA engines 3980(1) enable asynchronous memory transfers between GPU 3792 and such host. In at least one embodiment, memory controllers 3970 and GPU controllers 3984 facilitate data transfers between GPU 3792 and other GPUs 3792, and DMA engines 3980(2) enable asynchronous memory transfers between GPU 3792 and other GPUs 3792.

[0360] In at least one embodiment, GPU 3792 includes, without limitation, any number and type of system interconnects that facilitate data and control transmission between any number and type of directly or indirectly linked components within or external to GPU 3792. In at least one embodiment, GPU 3792 includes, without limitation, any number and type of I / O interfaces (e.g., PCIe) that couple to any number and type of peripheral devices. In at least one embodiment, GPU 3792 can include, without limitation, any number (including zero) of display engines and any number (including zero) of multimedia engines. In at least one embodiment, GPU 3792 implements a memory subsystem that includes, without limitation, any number and type of memory controllers (e.g., memory controller 3970 and system memory controller 3982) and memory devices that are dedicated to one component or shared between multiple components (e.g., shared memory 3960). In at least one embodiment, GPU 3792 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 3922), each of which can be private or shared between any number of components (e.g., SIMD units 3950, compute units 3940, and programmable processing units 3920).

[0361] FIG. 39 threads of an exemplary CUDA grid 4020 are mapped to different compute units 3940 of a GPU 3792 in accordance with at least one embodiment are shown. FIG. 40 In at least one embodiment, and for purposes of illustration only, grid 4020 has a GridSize of BX by BY by 1 and a BlockSize of TX by TY by 1. Thus, in at least one embodiment, grid 4020 includes, without limitation, (BX*BY) thread blocks 4030, each of which includes, without limitation, (TX*TY) threads 4040. Threads 4040 are depicted as wavy arrows in FIG. 39

[0362] In at least one embodiment, grid 4020 is mapped to programmable processing unit 3920(1), which includes, without limitation, compute units 3940(1)-3940(C). In at least one embodiment and as shown, (BJ*BY) thread blocks 4030 are mapped to compute unit 3940(1), and the remaining thread blocks 4030 are mapped to compute unit 3940(2). In at least one embodiment, each thread block 4030 can include, without limitation, any number of thread warps, and each thread warp is mapped to FIG. 41 different SIMD units 3950 of a GPU 3792. ​

[0363] In at least one embodiment, thread warps in a given thread block 4030 can be synchronized together and communicate through shared memory 3960 included in associated compute units 3940. For example, and in at least one embodiment, thread warps in thread block 4030(BJ, 1) can be synchronized together and communicate through shared memory 3960(1). For example, and in at least one embodiment, thread warps in thread block 4030(BJ+1, 1) can be synchronized together and communicate through shared memory 3960(2).

[0364] FIG. 41 Migration of existing CUDA code to data parallel C++ code is shown, according to at least one embodiment. Data Parallel C++ (DPC++) can refer to an open, standards-based alternative to single-architecture proprietary languages that allows developers to reuse code across hardware targets (CPUs and accelerators, such as GPUs and FPGAs), and also perform custom tuning for specific accelerators. DPC++ uses similar and / or the same C and C++ constructs according to ISO C++ that developers can be familiar with. DPC++ incorporates the SYCL standard by The Khronos Group to support data parallelism and heterogeneous programming. SYCL refers to an abstraction layer across platforms that builds on the underlying concepts, portability, and efficiency of OpenCL, which enables code for heterogeneous processors to be written in “single-source” style using standard C++. SYCL enables single-source development, where C++ template functions can contain both host and device code to build complex algorithms using OpenCL acceleration, and then reuse them throughout a source code of different types of data.

[0365] In at least one embodiment, a DPC++ compiler is used to compile DPC++ source code that can be deployed across a variety of hardware targets. In at least one embodiment, a DPC++ compiler is used to generate DPC++ applications that can be deployed across a variety of hardware targets, and DPC++ compatibility tools can be used to migrate CUDA applications to multi-platform programs in DPC++. In at least one embodiment, a DPC++ foundation toolset includes: a DPC++ compiler to deploy applications across a variety of hardware targets; DPC++ libraries to improve productivity and performance of CPUs, GPUs, and FPGAs; DPC++ compatibility tools to migrate CUDA applications to multi-platform applications; and any suitable combination thereof.

[0366] In at least one embodiment, a DPC++ programming model is used to simplify one or more aspects related to programming CPUs and accelerators by using modern C++ features to express parallelism with a programming language called Data Parallel C++ (DPC++). The DPC++ programming language can be used for code reuse for both hosts (e.g., CPUs) and accelerators (e.g., GPUs or FPGAs) using a single source language, and to clearly convey execution and memory dependencies. Mappings within DPC++ code can be used to convert an application to run on hardware or a set of hardware devices that best accelerate a workload. Hosts can be used to simplify development and debugging of device code even on platforms where no accelerators are available.

[0367] In at least one embodiment, CUDA source code 4100 is provided as input to a DPC++ compatibility tool 4102 to generate human readable DPC++ 4104. In at least one embodiment, human readable DPC++ 4104 includes inline comments generated by DPC++ compatibility tool 4102 that guide a developer how and / or where to modify DPC++ code to complete coding and tuning to desired performance 4106, resulting in DPC++ source code 4108.

[0368] In at least one embodiment, CUDA source code 4100 is or includes a collection of human readable source code in a CUDA programming language. In at least one embodiment, CUDA source code 4100 is human readable source code in a CUDA programming language. In at least one embodiment, a CUDA programming language is an extension of a C++ programming language that includes, without limitation, mechanisms to define device code and to distinguish between device code and host code. In at least one embodiment, device code is source code that, when compiled, can be executed on a device (e.g., a GPU or FPGA), and can include one or more parallelizable workstreams that can be executed on one or more processor cores of a device. In at least one embodiment, a device can be a processor that is optimized for parallel instruction processing, such as a CUDA-enabled GPU, a GPU, or another GPGPU, etc. In at least one embodiment, host code is source code that, when compiled, can be executed on a host. In at least one embodiment, some or all of host code and device code can be executed in parallel across a CPU and a GPU / FPGA. In at least one embodiment, a host is a processor that is optimized for sequential instruction processing, such as a CPU. In conjunction with ​ Described CUDA source code 4100 can be consistent with that discussed elsewhere in this document.

[0369] In at least one embodiment, DPC++ compatibility tool 4102 refers to an executable tool, program, application, or any other suitable type of tool for facilitating migration of CUDA source code 4100 to DPC++ source code 4108. In at least one embodiment, DPC++ compatibility tool 4102 is a command-line based code migration tool that is available as part of a DPC++ toolchain for porting existing CUDA sources to DPC++. In at least one embodiment, DPC++ compatibility tool 4102 converts some or all of source code of a CUDA application from CUDA to DPC++ and generates a result file written at least partially in DPC++, referred to as human-readable DPC++ 4104. In at least one embodiment, human-readable DPC++ 4104 includes annotations generated by DPC++ compatibility tool 4102 to indicate places where user intervention can be necessary. In at least one embodiment, user intervention is necessary when CUDA source code 4100 calls a CUDA API that does not have a similar DPC++ API; other examples of where user intervention is needed are discussed in greater detail later.

[0370] In at least one embodiment, a workflow for migrating CUDA source code 4100 (e.g., an application or portions thereof) includes creating one or more compilation database files; using DPC++ compatibility tool 4102 to migrate CUDA to DPC++; completing the migration and verifying correctness, resulting in DPC++ source code 4108; and compiling DPC++ source code 4108 using a DPC++ compiler to generate a DPC++ application. In at least one embodiment, the compatibility tool provides a utility that intercepts commands used when Makefiles are executed and stores them in a compilation database file. In at least one embodiment, the file is stored in JSON format. In at least one embodiment, intercepting build commands converts Makefile commands to DPC compatibility commands.

[0371] In at least one embodiment, intercept-build is a utility script that intercepts the build process to capture compilation options, macro definitions, and include paths, and writes this data to a compilation database file. In at least one embodiment, the compilation database file is a JSON file. In at least one embodiment, DPC++ compatibility tool 4102 parses the compilation database and applies the options when migrating input sources. In at least one embodiment, use of intercept-build is optional, but is highly recommended for Make or CMake based environments. In at least one embodiment, a migration database includes commands, directories, and files: commands can include necessary compilation flags; directories can include paths to header files; and files can include paths to CUDA files.

[0372] In at least one embodiment, DPC++ compatibility tool 4102 migrates CUDA code (e.g., applications) written in CUDA to DPC++ by generating DPC++ wherever possible. In at least one embodiment, DPC++ compatibility tool 4102 is available as part of a toolchain. In at least one embodiment, DPC++ toolchain includes an intercept-build tool. In at least one embodiment, intercept-build tool creates a compilation database that captures compilation commands to migrate CUDA files. In at least one embodiment, DPC++ compatibility tool 4102 uses the compilation database generated by intercept-build tool to migrate CUDA code to DPC++. In at least one embodiment, non-CUDA C++ code and files are migrated as-is. In at least one embodiment, DPC++ compatibility tool 4102 generates human-readable DPC++ 4104, which can be DPC++ code as generated by DPC++ compatibility tool 4102, cannot be compiled by a DPC++ compiler and requires additional pipeline to validate portions of code that were not migrated correctly, and can involve manual intervention, e.g., by a developer. In at least one embodiment, DPC++ compatibility tool 4102 provides hints or tools embedded in code to help a developer manually migrate additional code that could not be automatically migrated. In at least one embodiment, migration is a one-time activity for a source file, project, or application.

[0373] In at least one embodiment, DPC++ compatibility tool 41002 is able to successfully migrate all portions of CUDA code to DPC++ and can simply exist as an optional step for manual verification and adjustment of performance of generated DPC++ source code. In at least one embodiment, DPC++ compatibility tool 4102 directly generates DPC++ source code 4108 that is compiled by a DPC++ compiler without the need or utilization of human intervention to modify DPC++ code generated by DPC++ compatibility tool 4102. In at least one embodiment, DPC++ compatibility tool generates DPC++ code that is compilable, which a developer can selectively adjust for performance, readability, maintainability, and other various considerations or any combination thereof.

[0374] In at least one embodiment, one or more CUDA source files are migrated to DPC++ source files using, at least in part, DPC++ compatibility tool 4102. In at least one embodiment, CUDA source code includes one or more header files, which can include CUDA header files. In at least one embodiment, CUDA source files include a <cuda.h> header file and a <stdio.h> header file that can be used to print text. In at least one embodiment, a portion of a VectorAddKernel CUDA source file can be written or related to:

[0375]

[0376] In at least one embodiment, and in connection with the CUDA source file presented above, DPC++ compatibility tool 4102 parses CUDA source code and replaces header files with appropriate DPC++ and SYCL header files. In at least one embodiment, DPC++ header files include helper declarations. In CUDA, there is a concept of thread IDs, accordingly, in DPC++ or SYCL, there is a local identifier for each element.

[0377] In at least one embodiment, and in connection with the CUDA source file presented above, there are two vectors A and B that are initialized and a vector addition result is put into vector C as part of VectorAddKernel(). In at least one embodiment, as part of migrating CUDA code to DPC++ code, DPC++ compatibility tool 4102 converts CUDA thread IDs used to index work elements via local IDs to SYCL standard addressing of work elements. In at least one embodiment, DPC++ code generated by DPC++ compatibility tool 4102 can be optimized - for example, by reducing the dimensionality of nd_item, thereby increasing memory and / or processor utilization.

[0378] In at least one embodiment, and in connection with the CUDA source file presented above, memory allocation is migrated. In at least one embodiment, depending on SYCL concepts such as platform, device, context, and queue, cudaMalloc() is migrated to a unified shared memory SYCL call malloc_device() to which a device and context are passed. In at least one embodiment, a SYCL platform can have multiple devices (e.g., a host and a GPU device); a device can have multiple queues to which jobs can be submitted; each device can have a context; and a context can have multiple devices and manage shared memory objects.

[0379] In at least one embodiment and in conjunction with the CUDA source file presented above, the main() function invokes or calls VectorAddKernel() to add two vectors A and B and store the result in vector C. In at least one embodiment, the CUDA code that calls VectorAddKernel() is replaced with DPC++ code to submit the kernel to a command queue for execution. In at least one embodiment, the command group handler cgh passes data submitted to the queue, synchronization, and computation, and parallel_for is called for the multiple global elements and multiple work items in this workgroup that invoke VectorAddKernel().

[0380] In at least one embodiment and in conjunction with the CUDA source file presented above, CUDA calls to copy device memory and then free memory of vectors A, B, and C are migrated to corresponding DPC++ calls. In at least one embodiment, C++ code (e.g., standard ISO C++ code to print a vector of floating point variables) is migrated as-is without modification by DPC++ compatibility tools 4102. In at least one embodiment, DPC++ compatibility tools 4102 modify CUDA APIs for memory setup and / or host calls to execute a kernel on an accelerated device. In at least one embodiment and in conjunction with the CUDA source file presented above, a corresponding human-readable DPC++ 4104 (e.g., compilable) is written as or in relation to:

[0381]

[0382]

[0383] In at least one embodiment, human-readable DPC++ 4104 refers to output generated by DPC++ compatibility tool 4102 and can be optimized in one way or another. In at least one embodiment, human-readable DPC++ 4104 generated by DPC++ compatibility tool 4102 can be manually edited by a developer after migration to make it more maintainable, performant, or for other considerations. In at least one embodiment, DPC++ code (e.g., disclosed DPC++) generated by DPC++ compatibility tool 41002 can be optimized by removing repeated calls to get_current_device() and / or get_default_context() for each malloc_device() call. In at least one embodiment, DPC++ code generated above uses a 3-dimensional nd_range, which can be refactored to use only a single dimension, reducing memory usage. In at least one embodiment, a developer can manually edit DPC++ code generated by DPC++ compatibility tool 4102 to replace use of uniform shared memory with an accessor. In at least one embodiment, DPC++ compatibility tool 4102 has options to change how it migrates CUDA code to DPC++ code. In at least one embodiment, DPC++ compatibility tool 4102 is verbose because it uses generic templates to migrate CUDA code to DPC++ code, which works for a large number of cases.

[0384] In at least one embodiment, a CUDA to DPC++ migration workflow includes the following steps: preparing migration using an intercept-build script; performing migration of CUDA project to DPC++ using DPC++ compatibility tool 4102; reviewing and editing migrated source files to ensure their integrity and correctness; and compiling final DPC++ code to generate a DPC++ application. In at least one embodiment, human review of DPC++ source code can be required in one or more scenarios, including but not limited to: migrated APIs do not return error codes (CUDA code can return error codes, which can then be used by an application, but SYCL uses exceptions to report errors, so error codes are not used to surface errors); DPC++ does not support CUDA compute capability related logic; cannot remove statements. In at least one embodiment, scenarios where DPC++ code requires human intervention can include, but are not limited to: error code logic replaced with (*, 0) code or commented out; equivalent DPC++ API is not available; CUDA compute capability related logic; hardware related API (clock()); missing feature unsupported APIs; execution time measurement logic; handling built-in vector type conflicts; migration of cuBLAS APIs; and more.

[0385] In at least one embodiment, one or more techniques described herein utilize a oneAPI programming model. In at least one embodiment, oneAPI programming model refers to a programming model for interacting with various compute accelerator architectures. In at least one embodiment, oneAPI refers to an application programming interface (API) designed for interacting with various compute accelerator architectures. In at least one embodiment, oneAPI programming model utilizes a Data Parallel C++ (DPC++) programming language. In at least one embodiment, DPC++ programming language refers to a high-level language for data parallel programming productivity. In at least one embodiment, DPC++ programming language is based at least in part on C and / or C++ programming languages. In at least one embodiment, oneAPI programming model is a programming model such as those developed by Intel Corporation of Santa Clara, California.

[0386] In at least one embodiment, oneAPI and / or oneAPI programming model is used to interact with various accelerators, GPUs, processors, and / or variations thereof, architectures. In at least one embodiment, oneAPI includes a set of libraries that implement various functionality. In at least one embodiment, oneAPI includes at least oneAPI DPC++ library, oneAPI math kernel library, oneAPI data analytics library, oneAPI deep neural network library, oneAPI collective communications library, oneAPI threading building blocks library, oneAPI video processing library, and / or variations thereof.

[0387] In at least one embodiment, oneAPI DPC++ library, also referred to as oneDPL, is a library that implements algorithms and functions to accelerate DPC++ kernel programming. In at least one embodiment, oneDPL implements one or more standard template library (STL) functions. In at least one embodiment, oneDPL implements one or more parallel STL functionalities. In at least one embodiment, oneDPL provides a set of library classes and functions such as parallel algorithms, iterators, functor classes, range-based APIs, and / or variations thereof. In at least one embodiment, oneDPL implements one or more classes and / or functions of a C++ standard library. In at least one embodiment, oneDPL implements one or more random number generator functionalities.

[0388] In at least one embodiment, oneAPI Math Kernel Library, also referred to as oneMKL, is a library that implements various optimizations and parallelization routines for various mathematical functions and / or operations. In at least one embodiment, oneMKL implements one or more Basic Linear Algebra Subprograms (BLAS) and / or Linear Algebra Package (LAPACK) dense linear algebra routines. In at least one embodiment, oneMKL implements one or more sparse BLAS linear algebra routines. In at least one embodiment, oneMKL implements one or more random number generators (RNGs). In at least one embodiment, oneMKL implements one or more vector math (VM) routines for mathematical operations on vectors. In at least one embodiment, oneMKL implements one or more Fast Fourier Transform (FFT) functions.

[0389] In at least one embodiment, oneAPI Data Analytics Library, also referred to as oneDAL, is a library that implements various data analytics applications and distributed computing. In at least one embodiment, oneDAL implements various algorithms for data analytics in batch, online, and distributed computing processing modes for pre-processing, transformation, analysis, modeling, validation, and decision making. In at least one embodiment, oneDAL implements various C++ and / or Java APIs and various connectors to one or more data sources. In at least one embodiment, oneDAL implements DPC++ API extensions to traditional C++ interfaces and enables GPUs for various algorithms.

[0390] In at least one embodiment, oneAPI Deep Neural Network Library, also referred to as oneDNN, is a library that implements various deep learning functionality. In at least one embodiment, oneDNN implements various neural network, machine learning, and deep learning functionality, algorithms, and / or variations thereof.

[0391] In at least one embodiment, oneAPI Collective Communications Library, also referred to as oneCCL, is a library that implements various applications for deep learning and machine learning workloads. In at least one embodiment, oneCCL builds on lower-level communication middleware such as Message Passing Interface (MPI) and libfabrics. In at least one embodiment, oneCCL enables a set of deep learning specific optimizations such as priority, persistent operations, out-of-order execution, and / or variations thereof. In at least one embodiment, oneCCL implements various CPU and GPU functionality.

[0392] In at least one embodiment, oneAPI threading building blocks library, also referred to as oneTBB, is a library that implements various parallelization processes for various applications. In at least one embodiment, oneTBB is used for task-based shared parallel programming on a host. In at least one embodiment, oneTBB implements generic parallel algorithms. In at least one embodiment, oneTBB implements concurrent containers. In at least one embodiment, oneTBB implements a scalable memory allocator. In at least one embodiment, oneTBB implements a work-stealing task scheduler. In at least one embodiment, oneTBB implements low-level synchronization primitives. In at least one embodiment, oneTBB is independent of a compiler and is usable with various processors, such as GPUs, PPUs, CPUs, and / or variations thereof.

[0393] In at least one embodiment, oneAPI video processing library, also referred to as oneVPL, is a library for accelerating video processing in one or more applications. In at least one embodiment, oneVPL implements various video decoding, encoding, and processing functions. In at least one embodiment, oneVPL implements various functions for media pipelines on CPUs, GPUs, and other accelerators. In at least one embodiment, oneVPL implements device discovery and selection in media-centric and video analytics workloads. In at least one embodiment, oneVPL implements API primitives for zero-copy buffer sharing.

[0394] In at least one embodiment, oneAPI programming model utilizes a DPC++ programming language. In at least one embodiment, a DPC++ programming language is a programming language that includes, without limitation, functionally similar versions of CUDA mechanisms to define device code and differentiate between device code and host code. In at least one embodiment, a DPC++ programming language can include a functional subset of a CUDA programming language. In at least one embodiment, one or more CUDA programming model operations are performed using a DPC++ programming language using oneAPI programming model.

[0395] Other variations are within the spirit of the present disclosure. Thus, while the disclosed technology is susceptible to various modifications and alternative constructions, certain illustrated embodiments thereof are shown in the drawings and have been described above in detail. It should be understood, however, that there is no intention to limit the disclosure to the specific form or forms disclosed, but on the contrary, the intention is to cover all modifications, alternative constructions, and equivalents falling within the spirit and scope of the disclosure, as defined in the appended claims.

[0396] Unless otherwise indicated or contradicted by context, the use of the terms "a" and "an" and "the" and similar referents in the context of describing the disclosed embodiments (especially in the context of the following claims) are to be construed to cover both the singular and the plural, unless otherwise indicated or contradicted by context. The terms "comprising," "having," "including," and "containing" are to be construed as open-ended terms (meaning "including, but not limited to") unless otherwise noted or contradicted by context. The term "connected" (when used without modification) is to be construed as partly or wholly encompassed, attached to, or joined together, even if there are some intervening items. Unless otherwise indicated herein, a reference to a range of values herein is merely intended to serve as a shorthand method of referring individually to each separate value falling within the range, and each separate value is incorporated in the specification as if it were individually recited herein. Unless otherwise indicated or contradicted by context, the use of the term "set" (e.g., "set of items") or "subset" is to be construed as a non-empty set of one or more members. Also, unless otherwise indicated or contradicted by context, the term "subset" of a corresponding set does not necessarily denote a proper subset of the corresponding set, but rather the subset and the corresponding set can be equal.

[0397] Unless explicitly indicated otherwise or contradicted by context, conjunction language such as phrases in the form "at least one of A, B, and C" or "at least one of A, B, and C" is to be construed in context as generally used to mean that the item, term, etc. can be A or B or C, or any non-empty subset of the set of A and B and C. For example, in the illustrative example of a set having three members, the conjunction phrases "at least one of A, B, and C" and "at least one of A, B, and C" refer to any of the following sets: {A}, {B}, {C}, {A, B}, {A, C}, {B, C}, {A, B, C}. Thus, such conjunction language is not generally intended to imply that certain embodiments require the existence of at least one of A, at least one of B, and at least one of C. Additionally, unless otherwise indicated or contradicted by context, the term "plurality" denotes a plural state (e.g., "a plurality of items" denotes multiple items). The number of items in a plurality of items is at least two, but can be more if explicitly indicated or indicated by context. Furthermore, unless otherwise indicated or clear from context, the phrase "based on" means "based, at least in part, on" rather than "based solely on."

[0398] The operations of a process described herein can be performed in any suitable order unless otherwise indicated herein or otherwise clearly contradicted by context. In at least one embodiment, a process, such as those described herein (or variations and / or combinations thereof) is performed under the control of one or more computer systems configured with executable instructions to perform the operations of the process, and the process is implemented as code (e.g., executable instructions, one or more computer programs or one or more applications) that is collectively executed by a hardware or combination of hardware and / or software. In at least one embodiment, the code is stored on a computer-readable storage medium, such as a computer program product, which is readable by a computer system including one or more processors. In at least one embodiment, a computer-readable storage medium is a non-transitory computer-readable storage medium that excludes transitory signals (e.g., a propagating transient electric or electromagnetic transmission) but includes non-transitory data storage circuitry (e.g., buffers, cache, and queues). In at least one embodiment, the code (e.g., executable instructions or source code) is stored on a set of one or more non-transitory computer-readable storage media (or other memory for storing executable instructions) having stored thereon executable instructions that, as a result of being executed by one or more processors of a computer system (i.e., as a result of being executed), cause the computer system to perform operations described herein. In at least one embodiment, a set of non-transitory computer-readable storage media includes multiple non-transitory computer-readable storage media, and one or more of the individual non-transitory storage media in the multiple non-transitory computer-readable storage media lack all of the code, with the multiple non-transitory computer-readable storage media collectively storing the entire code. In at least one embodiment, executable instructions are executed by different processors, e.g., a non-transitory computer-readable storage medium stores instructions and a main central processing unit (“CPU”) executes some instructions, while a graphics processing unit (“GPU”) executes other instructions. In at least one embodiment, different components of a computer system have separate processors and different processors execute different subsets of the instructions.

[0399] Accordingly, in at least one embodiment, a computer system is configured to implement one or more services that individually or collectively perform operations of processes described herein, and such a computer system is configured with applicable hardware and / or software to enable implementation of the operations. Moreover, a computer system implementing at least one embodiment of the present disclosure is a single device, and in another embodiment is a distributed computer system including multiple devices operating in different manners such that the distributed computer system performs operations described herein and such that a single device does not perform all of the operations.

[0400] The use of any and all examples, or exemplary language (e.g., "such as") provided herein is intended merely to better illuminate embodiments of the disclosure and does not pose a limitation on the scope of the disclosure unless otherwise claimed. No language in the specification should be construed as indicating any non-claimed element as essential to the practice of the disclosure.

[0401] At least one embodiment of the present disclosure can be described with consideration of the following clauses:

[0402] 1. A processor comprising:

[0403] one or more circuits to calculate motion of one or more pixels in a first region of an image based at least in part on motion of one or more pixels in a second region overlapping the first region of the image.

[0404] 2. The processor of clause 1, wherein the one or more circuits are further to:

[0405] obtain the image and a flow map, wherein the flow map is calculated based at least in part on the image and a second image;

[0406] cause a subset of the one or more pixels in the second region to form a set of pixels, wherein the subset of the one or more pixels share one or more characteristics corresponding to at least one of an intensity and a gradient;

[0407] perform an operation based on the flow map to indicate whether the set of pixels includes one or more valid or invalid pixels; and

[0408] iteratively convert one or more invalid pixels in the set of pixels to valid pixels using the one or more valid pixels.

[0409] 3. The processor of clause 1 or 2, wherein the one or more circuits are further to:

[0410] iteratively convert one or more invalid pixels in the first region to one or more valid pixels using one or more valid pixels from the second region; and

[0411] update a flow map to generate a second flow map based at least in part on the conversion.

[0412] 4. The processor of clause 3, wherein the one or more circuits are further to traverse one or more rows or one or more columns of the second flow map and update the second flow map by converting the one or more invalid pixels to valid pixels using the one or more invalid pixels in the one or more rows or the one or more columns.

[0413] 5. The processor of any of clauses 1-4, wherein the one or more circuits are further to generate a dense optical flow map from the sparse optical flow map.

[0414] 6. The processor of clause 5, wherein the one or more circuits are further to apply a median filter to the dense optical flow map.

[0415] 7. A system comprising:

[0416] one or more processors to compute motion of one or more pixels in a first region of an image based at least in part on motion of one or more pixels in a second region overlapping the first region of the image.

[0417] 8. The system of clause 7, wherein the one or more processors are further to:

[0418] identify a subset of pixels from the one or more pixels in the second region, wherein the subset of pixels is identified based at least in part on at least one of an intensity, a gradient, and a depth of the image;

[0419] perform an operation based at least in part on the obtained flow map corresponding to the image and a second image to indicate whether the subset of pixels includes valid pixels;

[0420] interpolate optical flow for the subset of pixels using the valid pixels to generate a subset of valid pixels in the second region; and

[0421] interpolate optical flow for the one or more pixels in the first region using the subset of valid pixels in the second region to generate a subset of valid pixels in the first region.

[0422] 9. The system of clause 7 or 8, wherein the one or more processors are further to update a flow map to generate a second flow map based at least in part on the subset of valid pixels in the first region and the subset of valid pixels in the second region.

[0423] 10. The system of any of clauses 7-9, wherein the one or more processors are further to:

[0424] scan one or more rows of pixels in the second region of the image; and

[0425] interpolate optical flow for pixels in one or more rows of pixels based at least in part on valid pixels in the one or more rows of pixels to generate interpolated optical flow pixels.

[0426] 11. The system of any of clauses 7-10, wherein the one or more processors are further to:

[0427] scan one or more columns of pixels in the second region of the image; and

[0428] interpolate optical flow for pixels in the one or more columns of pixels based at least in part on valid pixels in the one or more columns of pixels to generate interpolated optical flow pixels.

[0429] 12. The system of clause 7, wherein the one or more processors are further to compute motion of the one or more pixels in the first region of the image for frame rate upconversion, video encoding, robot navigation, or image classification.

[0430] 13. A method comprising:

[0431] computing motion of one or more pixels in a first region of an image based at least in part on motion of one or more pixels in a second region of the image that overlaps the first region.

[0432] 14. The method of clause 13, further comprising:

[0433] obtaining a first flow map corresponding to the image and a subsequent image;

[0434] determining validity of one or more pixels in the image using the first flow map;

[0435] grouping a subset of the one or more pixels in the second region;

[0436] identifying valid pixels in the subset of the one or more pixels in the second region using the first flow map; and

[0437] interpolating optical flow for the subset of the one or more pixels in the second region based at least in part on the valid pixels to generate interpolated optical flow for the subset of the one or more pixels in the second region.

[0438] 15. The method of clause 14, wherein computing motion of the one or more pixels in the first region of the image comprises interpolating optical flow for one or more pixels in the first region using interpolated optical flow for the subset of the one or more pixels in the second region.

[0439] 16. The method of any of clauses 13-15, applying a cost function to determine validity of the one or more pixels in the image.

[0440] 17. The method of any of clauses 13-16, further comprising:

[0441] generating an updated flow map based at least in part on the computed motion of the one or more pixels in the one of the first region and the second region of the image.

[0442] 18. The method of clause 17, further comprising:

[0443] applying a filter to refine the updated flow map.

[0444] 19. A machine-readable medium having stored thereon a set of instructions, which if executed by one or more processors, cause the one or more processors to at least:

[0445] compute motion of one or more pixels in a first region of an image based at least in part on motion of one or more pixels in a second region overlapping the first region of the image.

[0446] 20. The machine-readable medium of clause 19, wherein the set of instructions further comprises instructions which, if executed by the one or more processors, cause the one or more processors to:

[0447] obtain a sparse flow map based on the image and a second image;

[0448] generate a subset of the one or more pixels in the second region, wherein at least one pixel in the subset is a valid pixel;

[0449] replace, in one or more iterations, one or more invalid pixels from the subset of the one or more pixels in the second region with the valid pixel; and

[0450] replace, in one or more additional iterations, one or more invalid pixels from the one or more pixels in the first region with the replaced one or more pixels from the second region to generate a set of replacement pixels in the first region.

[0451] 21. The machine-readable medium of clause 19 or 20, wherein the set of instructions further comprises instructions which, if executed by the one or more processors, cause the one or more processors to:

[0452] scan, in one or more iterations, one or more rows and one or more columns in the first region and the second region;

[0453] by replacing one or more pixels in the first region and the second region using one or more valid pixels in the first region or the second region; and

[0454] generating a dense flow map based on the replacement of one or more pixels in the first region and the second region.

[0455] 22. The machine-readable medium of clause 21, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to update the dense flow map using a median filter.

[0456] 23. The machine-readable medium of any of clauses 19-22, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to use a sparse flow map to indicate whether a pixel in one or more pixels in the first region and the second region is valid or invalid.

[0457] 24. The machine-readable medium of any of clauses 19-23, wherein the one or more pixels in the second region do not include computing motion for the one or more pixels labeled as background.

[0458] 25. A system comprising:

[0459] one or more computers having one or more processors to compute motion of one or more pixels in a first region of an image based at least in part on motion of one or more pixels in a second region overlapping the first region of the image.

[0460] 26. The system of clause 25, wherein the one or more processors are further to:

[0461] receive a flow map based on the image and a second image;

[0462] identify invalid pixels in the one or more pixels of the second region;

[0463] identify one or more neighboring pixels from the identified invalid pixels including at least one valid pixel;

[0464] replace the invalid pixels and one or more invalid pixels from the one or more neighboring pixels using at least one valid pixel to generate replacement pixels in the second region; and

[0465] replace one or more invalid pixels in the first region using the replacement pixels in the second region.

[0466] 27. The system of clause 26, wherein the one or more processors are further to:

[0467] replace one or more pixels in a third region of the image using the replacement pixels in the first region or the second region, wherein the first region or the second region overlaps the third region.

[0468] 28. The system of any of clauses 25-27, wherein the one or more processors are further to traverse the pixels in the first region via one or more rows or one or more columns and replace invalid pixels with valid pixels.

[0469] 29. The system of any of clauses 25-28, wherein the image is captured using a video of a semi-autonomous vehicle.

[0470] 30. The system of any of clauses 25-29, wherein the computed motion of one or more pixels in the first region is used as part of performing a frame rate up-conversion on frames in a video.

[0471] All references, including publications, patent applications, and patents, cited herein are hereby incorporated by reference to the same extent as if each reference were individually and specifically indicated to be incorporated by reference and were set forth in its entirety herein.

[0472] Throughout this specification and claims, the terms “coupled” and “connected,” along with derivatives thereof, can be used. It should be understood that these terms are not intended as synonyms for each other. Rather, in particular examples, “connected” or “coupled” is used to indicate that two or more elements are in direct or indirect physical or electrical contact with each other. “Coupled” can also mean that two or more elements are not in direct contact with each other, but yet still co-operate or interact with each other.

[0473] Unless specifically stated otherwise, it can be appreciated that throughout this specification discussions using terms such as “processing,” “computing,” “calculating,” “determining,” or the like can refer to actions and / or processes of a computer or computing system, or similar electronic computing device, that manipulate and / or transform data represented as physical, such as electronic, quantities within the computer or computing system's registers and / or memories into other data similarly represented as physical quantities within the computer or computing system's memories, registers or other such information storage, transmission or display devices.

[0474] In a similar manner, the term“processor” can refer to any device or portion of a device that processes electronic data from registers and / or memory to transform that electronic data into other electronic data that can be stored in registers and / or memory. As a non-limiting example,“processor” can refer to a CPU or GPU. A“computing platform” can include one or more processors. As used herein, a“software” process can include, for example, software and / or hardware entities such as tasks, threads, and intelligent agents that perform work over time. Likewise, each process can refer to multiple processes that execute instructions sequentially or in parallel, continuously or intermittently. The terms“system” and“method” can be used interchangeably herein so long as a system can embody a method or methods and a method can be considered a system.

[0475] In this document, reference can be made to obtaining, acquiring, receiving, or inputting analog or digital data into a subsystem, computer system, or computer-implemented machine. The process of obtaining, acquiring, receiving, or inputting analog and digital data can be accomplished in a variety of ways, such as by receiving data as a parameter of a function call or a call to an application programming interface. In some implementations, the process of obtaining, acquiring, receiving, or inputting analog or digital data can be accomplished by transmitting data via a serial or parallel interface. In another implementation, the process of obtaining, acquiring, receiving, or inputting analog or digital data can be accomplished by transmitting data from a providing entity to an acquiring entity via a computer network. Reference can also be made to providing, outputting, transmitting, sending, or presenting analog or digital data. In various examples, the process of providing, outputting, transmitting, sending, or presenting analog or digital data can be accomplished by transmitting data as an input or output parameter of a function call, a parameter of an application programming interface, or an inter-process communication mechanism.

[0476] Although the above discussion sets forth example implementations of the described technology, other architectures can be used to implement the described functionality and are intended to be within the scope of this disclosure. Moreover, although specific distributions of responsibilities are defined above for purposes of discussion, the various functions and responsibilities can be distributed and divided among various architectures in different ways, depending on circumstances.

[0477] Further, although the subject matter has been described in language specific to structural features and / or methodological acts, it is to be understood that the subject matter defined in the appended claims is not necessarily limited to the specific features or acts described. Rather, the specific features and acts are disclosed as example forms of implementing the claims.

Claims

1. A processor comprising circuitry to: identify at least a first region of an image and a second region of the image, the first region and the second region overlapping in an overlap portion; for a pixel within the overlap portion that lacks a corresponding optical flow vector, use one or more optical flow vectors of the first region to compute motion to convert the pixel to a valid pixel; and for each pixel within the second region that lacks a corresponding optical flow vector, iteratively use one or more valid pixels within the overlap portion, including the converted pixel, to compute motion to convert each pixel.

2. The processor of claim 1, wherein the circuitry is further to: obtain the image and a flow map, wherein the flow map is computed based at least in part on the image and a second image; and determine each pixel within the second region that lacks a corresponding optical flow vector.

3. The processor of claim 1, wherein the circuitry is further to: update a flow map based at least in part on the conversion of each pixel within the second region that lacks a corresponding optical flow vector to generate a second flow map.

4. The processor of claim 3, wherein the circuitry is further to traverse one or more rows or one or more columns of the second flow map and update the second flow map to convert each pixel within the second region that lacks a corresponding optical flow vector to a valid pixel using the one or more valid pixels in the one or more rows or the one or more columns.

5. The processor of claim 1, wherein the circuitry is further to generate a dense optical flow map from a sparse optical flow map.

6. The processor of claim 5, wherein the circuitry is further to apply a median filter to the dense optical flow map.

7. A system comprising: one or more processors to: identify at least a first region of an image and a second region of the image, the first region and the second region overlapping in an overlap portion; for a pixel within the overlap portion that lacks a corresponding optical flow vector, use one or more optical flow vectors of the first region to compute motion to convert the pixel to a valid pixel; and for each pixel within the second region that lacks a corresponding optical flow vector, iteratively use one or more valid pixels within the overlap portion, including the converted pixel, to compute motion to convert each pixel.

8. The system of claim 7, wherein the one or more processors are further to: obtain a flow map corresponding to the image and a second image to indicate each pixel within the second region that lacks a corresponding optical flow vector.

9. The system of claim 7, wherein the one or more processors are further to update a flow map based at least in part on the conversion of each pixel that lacks a corresponding optical flow vector to generate a second flow map.

10. The system of claim 7, wherein the one or more processors are further to: interpolate one or more valid pixels for each pixel that lacks a corresponding optical flow vector.

11. The system of claim 10, wherein the one or more processors are further to: scan one or more columns or one or more rows of pixels in the second region of the image for interpolation.

12. The system of claim 7, wherein the one or more processors are further to compute motion of the one or more pixels in the first region of the image for frame rate up-conversion, video encoding, robot navigation, or image classification.

13. A method comprising: identifying at least a first region of an image and a second region of the image, the first region and the second region overlapping in an overlap portion; for pixels within the overlap portion that lack a corresponding optical flow vector, using one or more optical flow vectors of the first region to compute motion to convert the pixels to valid pixels; and for each pixel within the second region that lacks a corresponding optical flow vector, iteratively using one or more valid pixels within the overlap portion, including converted pixels, to compute motion to convert each pixel.

14. The method of claim 13, further comprising: obtaining a first flow map corresponding to the image and a subsequent image; and using the first flow map to determine whether each pixel within the second region lacks a corresponding optical flow vector.

15. The method of claim 14, wherein iteratively using the one or more valid pixels comprises performing one or more interpolation operations to convert each pixel.

16. The method of claim 13, applying a cost function to determine validity of one or more pixels in the image.

17. The method of claim 13, further comprising: generating an updated flow map based at least in part on computed motion of one or more pixels in the first region and the second region of the image.

18. The method of claim 17, further comprising: applying a filter to refine the updated flow map.

19. A non-transitory machine-readable medium having stored thereon a set of instructions, which if executed by one or more processors, cause the one or more processors to at least: identify at least a first region of an image and a second region of the image, the first region and the second region overlapping in an overlap portion; for pixels within the overlap portion that lack a corresponding optical flow vector, using one or more optical flow vectors of the first region to compute motion to convert the pixels to valid pixels; and for each pixel within the second region that lacks a corresponding optical flow vector, iteratively using one or more valid pixels within the overlap portion, including converted pixels, to compute motion to convert each pixel.

20. The non-transitory machine-readable medium of claim 19, wherein the set of instructions further comprises instructions, which if executed by the one or more processors, cause the one or more processors to: obtain a sparse flow map based on the image and a second image, wherein the sparse flow map is used to identify each pixel within the second region that lacks a corresponding optical flow vector.

21. The non-transitory machine-readable medium of claim 19, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to: in one or more iterations, scan one or more rows and one or more columns in the first region and the second region to identify each pixel within the second region that lacks a corresponding flow vector; and generate a dense flow map based on the conversion of each pixel.

22. The non-transitory machine-readable medium of claim 21, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to update the dense flow map using a median filter.

23. The non-transitory machine-readable medium of claim 19, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to use a sparse flow map to indicate the one or more valid pixels.

24. The non-transitory machine-readable medium of claim 19, wherein the one or more pixels in the second region do not include computing motion for the one or more pixels labeled as background.

25. A system comprising: one or more computers having one or more processors configured to: identify at least a first region of an image and a second region of the image, the first region and the second region overlapping in an overlap portion; for a pixel within the overlap portion that lacks a corresponding optical flow vector, use one or more optical flow vectors of the first region to compute motion to convert the pixel to a valid pixel; and for each pixel within the second region that lacks a corresponding optical flow vector, iteratively use one or more valid pixels within the overlap portion, including the converted pixel, to compute motion to convert each pixel.

26. The system of claim 25, wherein the one or more processors are further configured to update a flow map to generate a second flow map, wherein the flow map is obtained based at least in part on the image and a second image.

27. The system of claim 25, wherein the one or more processors are further configured to: replace one or more pixels in a third region of the image with at least one converted pixel within the second region, wherein the first region or the second region overlaps with the third region.

28. The system of claim 25, wherein the one or more processors are further configured to traverse pixels in the first region via one or more rows or one or more columns to enable conversion of each pixel that lacks a corresponding optical flow vector.

29. The system of claim 25, wherein the image is captured using a video of a semi-autonomous vehicle.

30. The system of claim 25, wherein the computed motion of one or more pixels in the first region is used as part of performing a frame rate up-conversion on frames in the video.

Images