Tone management using tone gain function for high dynamic range imaging applications

Abstract

To minimize negative effects of tone and contrast enhancements.SOLUTION: The invention provides apparatuses, systems and techniques to perform effective tone management for image data. In an embodiment, a set of contrast gain curves are generated corresponding to a set of tonal ranges of an input image. An output image may then be generated by at least applying corresponding contrast gain curves to tonal ranges of the input image.SELECTED DRAWING: Figure 1

Description

[Technical Field]

[0001] This invention relates to tone control using a tone gain function for high dynamic range imaging applications. [Background technology]

[0002] High-Dynamic-Range Imaging (HDRI) involves techniques for generating high-dynamic-range (HDR) image data. Specifically, HDRI provides an increased ratio between the largest possible pixel value (e.g., the largest possible luminosity value) and the smallest possible pixel value (e.g., the smallest possible luminosity value). Pixels with smaller luminosity values ​​render as darker (e.g., blacker) areas of the encoded image, while pixels with larger luminosity values ​​render as brighter (e.g., whiter) areas of the image. Due to the increased ratio between the largest and smallest luminosity values, and when displayed on a device capable of rendering an increased dynamic range (DR), HDR images can offer greater detail and contrast, and therefore appear more realistic and natural to the human viewer. Furthermore, the increased contrast and detail of these HDR images can improve performance on various image processing tasks, such as identifying entities or objects in an image. For example, when rendered properly, the scene image encoded in HDR image data may have higher detail and contrast, potentially enabling more accurate object detection in the scene by image processing software.

[0003] Conventional methods for tone-mapping HDR image data can result in image saturation and / or clipping as a result of inappropriate tone and contrast treatment, which can lead to data loss and / or detection failure. Furthermore, when controlling the brightness of HDR images and / or images encoded by tone-mapped HDR image data, various image processing techniques can reduce much of the visual benefits associated with HDRI, rendering HDR images and / or tone-mapped HDR images as, for example, "washed-out," unrealistic, or otherwise lacking in detail. [Prior art documents] [Patent Documents]

[0004] [Patent Document 1] U.S. Patent Application No. 16 / 526,902, titled "ENHANCED HIGH-DYNAMIC-RANGE IMAGING AND TONE MAPPING," filed on July 30, 2019. [Overview of the project] [Means for solving the problem]

[0005] Various techniques are explained while referring to the drawings. [Brief explanation of the drawing]

[0006] [Figure 1] This figure shows an exemplary method for applying separate tone gain functions to individual tonal regions of an input image, according to at least one embodiment. [Figure 2] This figure shows an exemplary method for applying separate tone gain functions to individual tonal regions of an input image, according to at least one embodiment. [Figure 3]A diagram showing an exemplary histogram generated at least partially based on an input image according to at least one embodiment. [Figure 4] A diagram showing an exemplary method for calculating contrast values using different scales according to at least one embodiment. [Figure 5] A diagram showing an exemplary contrast image according to at least one embodiment. [Figure 6] A diagram showing an exemplary method for determining a contrast gain curve for individual tonal regions of an input image according to at least one embodiment. [Figure 7] A diagram showing an exemplary contrast gain curve for individual tonal regions of an input image according to at least one embodiment. [Figure 8] A diagram showing an exemplary method for contrast mapping according to at least one embodiment. [Figure 9] A diagram showing an exemplary global tone mapping curve according to at least one embodiment. [Figure 10] A diagram showing an exemplary method for generating an output image according to at least one embodiment. [Figure 11] A diagram showing an exemplary data center according to at least one embodiment. [Figure 12] A diagram showing a processing system according to at least one embodiment. [Figure 13] A diagram showing a computer system according to at least one embodiment. [Figure 14] A diagram showing a system according to at least one embodiment. [Figure 15] A diagram showing an exemplary integrated circuit according to at least one embodiment. [Figure 16] A diagram showing a computing system according to at least one embodiment. [Figure 17] A diagram showing an APU according to at least one embodiment. [Figure 18] A diagram showing a CPU according to at least one embodiment. [Figure 19] A diagram showing an exemplary accelerator integration slice according to at least one embodiment. [Figure 20A] A diagram showing an exemplary graphics processor according to at least one embodiment. [Figure 20B] A diagram showing an exemplary graphics processor according to at least one embodiment. [Figure 21A] A diagram showing a graphics core according to at least one embodiment. [Figure 21B] A diagram showing a GPGPU according to at least one embodiment. <( [Figure 22A] A diagram showing a parallel processor according to at least one embodiment. [Figure 22B] A diagram showing a processing cluster according to at least one embodiment. [Figure 22C] A diagram showing a graphics multiprocessor according to at least one embodiment. [Figure 23] A diagram showing a graphics processor according to at least one embodiment. [Figure 24] A diagram showing a processor according to at least one embodiment. [Figure 25] A diagram showing a processor according to at least one embodiment. [Figure 26] A diagram showing a graphics processor core according to at least one embodiment. [Figure 27] A diagram showing a PPU according to at least one embodiment. [Figure 28] A diagram showing a GPC according to at least one embodiment. [Figure 29] A diagram showing a streaming multiprocessor according to at least one embodiment. [Figure 30]This figure shows the software stack of a programming platform, based on at least one embodiment. [Figure 31] This figure shows a CUDA implementation of the software stack in Figure 30, based on at least one embodiment. [Figure 32] This figure shows the ROCm implementation of the software stack in Figure 30, according to at least one embodiment. [Figure 33] This figure shows an OpenCL implementation of the software stack in Figure 30, based on at least one embodiment. [Figure 34] This figure shows software supported by a programming platform, according to at least one embodiment. [Figure 35] This figure shows the compilation of code for execution on the programming platform shown in Figures 30-33, using at least one embodiment. [Figure 36] This figure shows in more detail how to compile code for execution on the programming platform shown in Figures 30-33, using at least one embodiment. [Figure 37] This figure shows, in at least one embodiment, that the source code is translated before it is compiled. [Figure 38A] This figure shows a system configured to compile and execute CUDA source code using different types of processing units, according to at least one embodiment. [Figure 38B] This figure shows a system configured to compile and execute the CUDA source code in Figure 38A using a CPU and a CUDA-enabled GPU, according to at least one embodiment. [Figure 38C] This figure shows a system configured to compile and run the CUDA source code in Figure 38A using a CPU and a non-CUDA-enabled GPU, according to at least one embodiment. [Figure 39]This figure shows an exemplary kernel translated by the CUDA-to-HIP translation tool shown in Figure 38C, according to at least one embodiment. [Figure 40] This figure shows a more detailed view of the CUDA-incompatible GPU in Figure 38C, based on at least one embodiment. [Figure 41] This figure shows how threads in an exemplary CUDA grid are mapped to different compute units in Figure 40, according to at least one embodiment. [Figure 42] This diagram shows how to migrate existing CUDA code to Data Parallel C++ code, using at least one example. [Modes for carrying out the invention]

[0007] Embodiments of this disclosure relate to effective tone management using multiple tone gain functions. A system and method are disclosed for determining a set of regions in HDR image data, wherein a separate tone gain function is applied to a specific region of the set of regions.

[0008] In contrast to conventional systems, such as those described above, applying separate tone gain functions to individual tonal regions of HDR image data minimizes adverse effects on tone and contrast enhancement, such as saturation, clipping, and / or lack of contrast. In various embodiments, the optimal tone gain function is determined for multiple tone keys representing tonal ranges within the HDR image data. In one example, these tonal ranges include shadow ranges, midtone ranges, and highlight ranges. These tonal ranges may represent local areas within the image, and a specific tone gain function may be applied to these local areas.

[0009] In various embodiments, the HDR image processing pipeline includes a statistical analyzer, a contrast calculator, a contrast-gain curve generator, contrast mapping, and global tone mapping. The operation of the HDR image processing pipeline may be carried out in parallel, serially, or in combination thereof, as will be described in more detail below. For example, local contrast processing, which includes the statistical analyzer, contrast calculator, contrast-gain curve generator, and local contrast mapping, may be performed either before or after the global tone mapping.

[0010] The statistical analyzers described in more detail below can process HDR image data to determine various statistics of the HDR image data. In one embodiment, the statistical analyzer processes a histogram representing the HDR image data to generate multiple tonal ranges, a mean for each tonal range, and a variance for each tonal range. In one example, the HDR image data is divided into three tonal ranges: a shadow range, a midtone range, and a highlight range, where the shadow range represents the bottom 10 percent (e.g., the number of pixels represented in the histogram), the highlight range represents the top 10 percent, and the midtone range represents the portion of the HDR image data between the shadow range and the highlight range. In various embodiments, the multiple ranges and / or associated mean are provided as inputs to a contrast-gain curve generator, as described in more detail below.

[0011] In one embodiment, the contrast calculator determines the tonal difference between an input image and a reference tone image scaled by a smoothing factor in Gaussian scale space. Furthermore, multiple contrast images may be generated based at least partially on different scales. In various embodiments, luminance values ​​for pixels in the input image are determined, and these luminance values ​​are then smoothed using a Gaussian kernel to generate smoothed luminance values. Contrast images may then be generated by comparing the luminance values ​​with the smoothed luminance values, at least. In one example, the contrast images include both positive and negative values.

[0012] In various embodiments, the contrast gain curve generator produces a set of contrast gain curves for a tonal range generated by a statistical analyzer. Furthermore, various control parameters can be applied to the contrast gain curves within the set, as well as to a global contrast gain curve, which may be a combination of the contrast gain curves within the set. In addition, in various embodiments, these control factors are adjusted based on scene conditions (e.g., nighttime or low-light image). Returning to the example described above, three gain curves are generated: a shadow contrast gain curve, a highlight contrast gain curve, and a midtone contrast gain curve. In various embodiments, the contrast gain curve is determined at least partially based on the mean value for a tonal range and a sigma representing a Gaussian curve generated by a statistical analyzer according to the tonal range distribution. Generally, the contrast gain curve is a function of the input tone value, producing the corresponding gain for that tone value.

[0013] In various embodiments, contrast mapping applies a contrast gain curve to a contrast image generated by a contrast calculator to produce an adjusted contrast image. In one embodiment, these adjusted contrast images are combined to create a final contrast image. In some embodiments, the final contrast image is merged with the output of global tone mapping. In other embodiments, global tone mapping is applied to the input image before local tone mapping is performed. The output image (e.g., the final contrast image, or the final contrast image merged with the global tone-mapped image) can be used as input to an object detection algorithm or other image processing algorithm. In other embodiments, the output image is simply generated to create a more appealing image for human users.

[0014] Figure 1 illustrates Method 100, in which a separate tone gain function is applied to individual tonal regions of an HDR image and / or HDR image data to generate an output image. In the example shown in Figure 1, the input image 102 is processed using several functions of the process, including, but not limited to, a stat analyzer 106, a contrast gain curve generator 108, a contrast calculator 110, a contrast mapping 112, and a global tone mapping 104, to produce an output image 120. The input image 102 may be data representing something to be displayed, such as a scene captured by an image capture device (e.g., a digital camera), something generated by an automated process (e.g., for display in a video game or animated movie), and / or other data to be displayed. In one embodiment, the input image 102 is formatted as a bitmap file, but the techniques of this disclosure can be adapted for input images in other formats, such as Joint Photographic Experts Group (JPEG), Graphics Interchange Format (GIF), and Portable Network Graphics (PNG). The methods and operations described herein are applicable to two-dimensional and three-dimensional images (e.g., animations, movies, video games, and other image data). Furthermore, as shown in Figure 1, global tone mapping 104 is applied to the input image 102 in parallel with the application of separate tone gain functions to individual tonal regions of the input image 102, and the two resulting images (e.g., the result of global tone mapping 104 and local tone mapping) are combined (represented by a plus sign in Figure 1) to produce an output image 120.As described in this disclosure, combining images involves a variety of operations, such as merging images, applying mathematical transformations to image pixel values ​​(e.g., luminance values ​​and / or RGB values), adding, subtracting, or otherwise modifying pixel values, or other methods of combing two or more images. Alternatively, as will be described in more detail below with respect to Figure 2, global tone mapping 104 is applied to the input image 102 before local tone mapping is performed. As described in this disclosure, local tone mapping may refer to the process of applying separate tone gain functions to individual tonal regions of the input image 102.

[0015] The methods described herein, including Figures 1, 2, 4, 6, 8, and 10, include computing processes that can be implemented using any combination of hardware, firmware, and / or software. For example, various functions can be performed by a processor that executes instructions stored in memory. In various embodiments, the elements of these figures include source code or other executable instructions stored in memory that, when executed by the processor of the computing device, cause the computing device to perform the operations described below. Each of the methods may also be embodied as computer-usable instructions stored in a computer storage medium. The methods may be provided, to name a few examples, as standalone applications, as a service or a hosted service (standalone or in combination with another hosted service), or as a plug-in to another product. However, these methods may be implemented by any one system, or any combination of systems, including, but not limited to, those described herein, as additional or alternative. Furthermore, the operations of the methods described above may be omitted, performed in different orders, and performed in parallel or in a combination of serial and parallel.

[0016] Returning to Figure 1, the statistical analyzer 106 determines the statistical data and key values ​​used in contrast gain curve generation 108 and global tone mapping 104. As will be described in more detail below with respect to Figure 3, in one embodiment, the statistical analyzer 106 generates a histogram based at least partially on the input image 102. The contrast calculator 110 determines a contrast image including a set of local contrast values, where the local contrast values ​​in the set of local contrast values ​​correspond to specific pixel locations in the input image 102. In one example, each pixel location with respect to the input image is assigned a local contrast value by the contrast calculator 110 based at least partially on the image data obtained from the input image. Various operations of the contrast calculator 110 are described in more detail below with respect to Figure 4.

[0017] In one embodiment, the contrast gain curve generator 108 generates a contrast gain transfer function applied to a contrast image obtained from the contrast calculator 110. Furthermore, data obtained from the statistical analyzer 106 (e.g., a histogram of the input image 102) may be used to determine the contrast gain curve generated during the contrast gain curve generator 108. In one embodiment, during contrast mapping 112, individual contrast gain curves are applied to a contrast image (e.g., the output of the contrast calculator 110). As will be described in more detail below with respect to Figure 8, in various embodiments, during contrast mapping 112, individual tone functions may be applied to different tonal regions, at least partially based on representative local pixels.

[0018] In one embodiment, the statistical analyzer 106 includes a process for analyzing statistical data of the input image 102. In one example, the statistical data of the input image 102 is estimated and presented in a histogram. In another example, the statistical data is calculated directly from the input image 102 or other statistical representatives. In various embodiments, the statistical data and / or key statistical data include the midtone mean, the highlight mean, the shadow mean, and the variance. In one example, these values ​​are defined by the following formulas.

number

[0019] In the above formula, R is the set of histogram edge indices. For example, R S This defines the range of shadow tones. Furthermore, the variables m, h, and s correspond to the mean midtone, mean highlight, and mean shadow, respectively. Although the above equation includes three distinct ranges (e.g., midtone, highlight, and shadow), the statistical analyzer 106 can subdivide the input image 102 into any number of ranges and generate corresponding key statistical data. Returning to the exemplary equation above, e[n] and k[n] represent the edge value for index n and the number of pixels on the histogram edge of e[n]. As will be explained in more detail below with respect to Figure 3, this key statistical data generated by the statistical analyzer 106 is illustrated in graph form in various embodiments, where the base curve is a representative Gaussian function with std being 3 times sigma and the mean midtone (m). Furthermore, the key statistical data can be used in the contrast gain curve generation 108 with different sigma and mean values, as will be explained in more detail below with respect to Figure 6.

[0020] Figure 2 illustrates Method 200, in which a separate tone gain function is applied to individual tonal regions of an HDR image and / or HDR image data to generate an output image. In the example shown in Figure 2, the input image 202 is processed using several functions of the process, including, but not limited to, a statistical analyzer 206, a contrast gain curve generator 208, a contrast calculator 210, a contrast mapping 212, and a global tone mapping 204, to produce an output image 220. Furthermore, as shown in Figure 2, the global tone mapping 204 is applied before the separate tone gain function is applied to individual tonal regions of the result (or output) of the global tone mapping 204.

[0021] In various embodiments, after global tone mapping 204 is applied to the input image 202, the resulting image is provided as input to a statistical analyzer 206 and a contrast calculator 210, and contrast mapping 212 is performed. Furthermore, the resulting image (for example, after performing global tone mapping 204 on the input image 202) is combined with the results of applying a separate tone gain function (for example, a contrast gain curve, which will be described in more detail below with respect to Figure 6) to individual tonal regions of the resulting image. In other words, the contrast gain curve is determined and applied to the first result of global tone mapping 204 to produce a second result (for example, an image to which local tone mapping has been applied), and the output image 220 is produced by combining at least the first and second results, as shown in Figure 2 with a plus sign.

[0022] Figure 3 is an exemplary graph 300 showing statistical data and / or key statistical data generated by a statistical analyzer based at least partially on an input image, according to one embodiment. In one embodiment, curve 302 represents the histogram of the input image, curve 306 represents the cumulative histogram of the input image, curve 304 represents the Gaussian function, area 308 below curve 302 represents the bottom 10 percent of the input image (for example, based on the luminance values ​​for the pixels of the input image), and area 310 below curve 302 represents the top 10 percent of the input image (for example, based on the luminance values ​​for the pixels of the input image). Furthermore, between 308 and 310 (for example, the value e[N]) s ] and e[N h The area between ] represents the rest of the input image. As shown in the figure, Graph 300 uses 10 percent values ​​for the upper and lower ranges, but these values ​​may be modified in accordance with this disclosure. For example, different lighting conditions (e.g., night images, low-light images, high-light images, reflective objects in the image, moving objects in the image, etc.), application, sensitivity, or other factors may require these values ​​to be adjusted (e.g., an even distribution of ranges). Furthermore, the input image may be divided into four or more regions.

[0023] As shown in Graph 300, s represents the average shadow tone, m represents the average midtone, and h represents the average highlight. In various embodiments, these values ​​(e.g., x-axis) represent luminance values ​​for pixels in the input image. In various embodiments, these luminance values ​​are determined at least partially based on color values ​​(e.g., RGB) for pixels in the input image. In the example in Figure 3, 0 represents the darker parts of the input image and 1 represents the brighter parts of the input image. As will be explained in more detail below, the s, m, and h values ​​include key statistical data determined by the statistical analyzer at least partially based on the input image.

[0024] Figure 4 shows Method 400, according to one embodiment, in which a contrast calculator is applied to an input image to generate a tonal image. Contrast is defined as the tonal difference between the input and a reference tone image scaled by a smoothing factor in Gaussian-scale space. As shown in Figure 4, an input image 402 is provided as input to the contrast calculator. The input image 402 may include an HDR image, as described above. Furthermore, in various embodiments, a tone image 404 is generated based at least in part on the input image 402. For example, the luminance values ​​for pixels in the input image 402 are determined at least in part on the formula, such as BT.709 color coding for the RGB values ​​of the input image, which is expressed by the following formula. Y(x,y)=0.2126×RGB(x,y,1)+0.7152×RGB(x,y,2)+0.0722×RGB(x,y,3).

[0025] However, various other representations of the luminance of an image may be used in relation to this disclosure. As shown below with respect to Figure 5, the reference tone image was transformed in various embodiments into scale-space representations 406A and 406B based at least partially on different t-values. The formulas for scale-space representations 406A and 406B are described in one embodiment as follows: Y t (x,y)=Y(x,y,t)=Y(x,y)*g(x,y,t) In the above formula, Y tThese are scale-space representations (e.g., 406A and 406B) defined by the convolution of Y(x,y) with a Gaussian kernel g(x,y,t) having a standard deviation (t). In various embodiments, the scale-space representations 406A and 406B have their luminance values ​​smoothed, as shown in Figure 5. In this way, multiple contrast images are generated at different scales to regenerate multiple scale-space representations 406A and 406B. In various embodiments, these scale-space representations 406A and 406B are compared to generate sets of contrast maps 408A and 408B (for example, the luminance values ​​are subtracted and shown as minus symbols in Figure 4).

[0026] Figure 5 shows an embodiment 500 in which a contrast image is generated based at least in part on the method 400 described above. In this embodiment, an input tone image 502 is used to generate a scale space representation 504. For example, the scale space representation 504 is generated by applying the above formula, defined at least by the convolution of Y(x,y) and a Gaussian kernel g(x,y,t) with a standard deviation (t) of t=32, to the input tone image 502. Furthermore, as described above, the input tone image 502 may be converted to include luminance value pixels of the input image or a subset thereof, or otherwise modified.

[0027] In various embodiments, multiple scale-space representations with various t-values ​​are generated. For example, a computing device that generates an HDR image as described herein generates multiple scale-space representations by applying the above formula to an input tone image 502 using at least several different t-values. In one embodiment, the input tone image 502 is compared to the scale-space representation 504 to generate a contrast image 506. In one embodiment, the contrast image 506 includes positive and negative contrast values ​​for specific pixels of the image. In one embodiment, grayscale is used to represent positive and / or negative contrast in the contrast image 506. Furthermore, in this embodiment, the magnitude of the contrast is proportional to the saturation, where white represents a contrast of 0.

[0028] Figure 6 shows a method 600 in which a set of contrast curves is generated based at least partially on statistical data, according to one embodiment. Three contrast gain curves are generated as shown in Figure 6. In one embodiment, the three contrast gain curves are then modified by a set of control parameters 616, merged together, and then the final control parameter from the set of control parameters 616 is used to modify the merged curves. In various embodiments, an input image 602 (e.g., input image 102 as described above) is provided to a statistical analyzer 606 (e.g., statistical analyzer 106 as described above). Furthermore, as described above, the statistical analyzer 606 determines statistical data, such as mean shadow, mean highlight, and mean midtone, based at least partially on the input image 602. Although only three contrast curves are shown in Figure 6, any number of contrast gain curves can be determined based at least partially on individual parts of the input image. For example, the number of contrast gain curves can be varied at least partially on purpose for application to adapt to different levels of contrast.

[0029] Returning to FIG. 6, the contrast gain curves for various representative tonal regions are defined by corresponding equations. For example, the midtone contrast gain curve generation 608, the shadow contrast gain curve generation 610, and the highlight contrast gain curve generation 612 are defined by the following equations. Midtone contrast gain curve: m curve (y) = m gain × Gaussian(m, S m ) Highlight contrast gain curve generation: h curve (y) = h gain × Gaussian(h, S h ) Shadow contrast gain curve: s curve (y) = s gain × Gaussian(s, S s ) In the above equations, m, h, and s include the key statistical data described above, where these values represent the average values for the individual tonal regions of the input image 602 (e.g., shadow average value, highlight average value, and midtone average value). Further, S m , S h , and S s include sigma values representing the Gaussian curve, where it is adjusted by the statistical analyzer 606 based on the standard deviation (std defined above) according to the distribution of values within the individual tonal region.

[0030] As shown in FIG. 6, the set of control parameters 616 includes the midtone gain 618, the shadow gain 620, the highlight gain 622, and the strength 624. In various embodiments, these control parameters include a set of values used to modify the contrast gain curve. For example, the user defines the set of control parameters 616 based at least in part on a specific application, such as photographing, object / person detection, low-light images, or other uses. Returning to the above equations, m gain , h gain , and s gainThese represent the midtone gain 618, the highlight gain 622, and the shadow gain 620. In various embodiments, these values ​​are scaling factors that define the maximum size of the corresponding contrast gain curves (e.g., midtone-contrast gain curve generation 608, shadow-contrast gain curve generation 610, and highlight-contrast gain curve generation 612). In one example, the control parameter 616 is determined at least in part based on scene conditions from the input image 602.

[0031] In one embodiment, the maximum magnitude of the contrast gain curve may be defined by the ratio of the amount of expandable contrast over a sufficiently broad intrinsic contrast in a particular tonal region, where the sufficiently broad intrinsic contrast may be represented as a weighted combination (w) of the maximum possible contrast of the image and the limited contrast between neighboring tone regions. Furthermore, in some embodiments, the weight w is adaptively adjusted based at least in part on the statistical properties of the input image 602, such as the modality and skewness of the histogram of the input image 602. For example, an image showing a single-tone distribution and higher skewness in its histogram has a higher probability of both low and high tones coexisting in the same local region. In this example, w needs to be set lower so that the contrast estimate is closer to the full range of contrast (e.g., from highlight to shadow regions). In another example, if the image has multi-modality as shown in its histogram, w may be set higher than for a single tone so that there is improved separation in the tone distribution and a lower probability of different tones intersecting. In various embodiments, the set of control parameters 616 is adapted to achieve maximum contrast within a limited range, as shown below in Figure 7 as h_to_max and s_to_min. In one embodiment, w can be formalized by the constraint w∝(modality × 1 / skewness).

[0032] In other embodiments, the contrast gain curves are not limited to these equations, and any other method for determining the gain curves may be employed. For example, a method for generating a gain curve that includes a sufficient amount of representative tone regions and ensures that tone regions are distinguishable from at least one other tone region with a smooth transition on the boundary may be used according to this disclosure. In various embodiments, the midtone contrast gain curve generator 608, the shadow contrast gain curve generator 610, and the highlight contrast gain curve generator 612 are merged together after the corresponding control parameters (e.g., midtone gain 618, shadow gain 620, highlight gain 622) are applied (indicated by "×" in Figure 6). Merging the contrast gain curves 614 generates a combined contrast gain curve shown in Figure 7.

[0033] Figure 7 shows an exemplary graph 700 of a set of contrast gain curves determined at least partially based on the input image in one embodiment. Graph 700 shows the image histogram 708, the shadow contrast gain curve 702, the midtone contrast gain curve 704, the highlight contrast gain curve 706, and the combined contrast gain curve 710. In various embodiments, the shadow contrast gain curve 702, the midtone contrast gain curve 704, and the highlight contrast gain curve 706 are determined as described above with respect to Figure 6. In one embodiment, the combined contrast gain curve 710 is determined at least partially based on the following equation. G(y) = P0(P1 × s curve (y) + P2 × m curve (y) + P3 × s curve (y))

[0034] In the above equation, P0, P1, P2, and P3 define control parameters, such as those described above with respect to Figure 6. In one example, as described above, P1 is the shadow gain parameter, P2 is the midtone gain parameter, P3 is the highlight gain parameter, and P0 is the intensity parameter. As described above, in various embodiments, four contrast gain curves are described with respect to Figure 7, but the input image can be divided into any number of distinct regions, each with a corresponding contrast gain curve.

[0035] Figure 8 shows Method 800 in one embodiment, in which a set of contrast curves is applied to the scale-space representation of the input image to generate a contrast image. In various embodiments, the contrast gain curve is a function of the input tone value to produce the corresponding gain defined by the contrast gain curve. Furthermore, as shown in Figure 8, the contrast gain curves 802A and 802B are applied to the scale-space representations 804A and 804B in one embodiment to produce the contrast gain maps 810A and 810B. The contrast gain curves 802A and 802B are generated in one embodiment using Method 600 described above. Furthermore, the scale-space representations 804A and 804B are generated in various embodiments, at least in part, based on the convolution of Y(x,y) with a Gaussian kernel g(x,y,t) having a standard deviation (t), as described above.

[0036] In one embodiment, contrast gain maps 810A and 810B define tone values ​​at least partially based on input tone values ​​and corresponding gains defined by contrast gain curves 802A and 802B defined by the following equations. G M (x,y,t)=G(Y(x,y,t))

[0037] In the example shown in Figure 8, the contrast gain maps 810A and 810B are then applied to the corresponding contrast images 806A and 806B (indicated by "×" in Figure 8). In various embodiments, the contrast images 806A and 806B are generated using the method 400 described above. For example, contrast gain map G M (x,y,t) is applied to the contrast image C(x,y,t) pixels of the contrast image. In various embodiments, adjusted contrast images 812A and 812B are generated as a result of applying contrast gain maps 810A and 810B to the contrast images 806A and 806B. As described above, the adjusted contrast images 812A and 812B show tone values ​​generated by applying contrast gain curves 802A and 802B to input tone values ​​(for example, tone values ​​shown in scale space representations 804A and 804B).

[0038] In one embodiment, multiple contrast images (for example, adjusted contrast images 812A and 812B) are added together, as shown in Figure 8, and indicated by a "+". In one example, the addition of adjusted contrast images 812A and 812B is defined by the following formula:

number

[0039] In the above formula, C F This defines the final contrast image 814, which is generated by summing a set of adjusted contrast images (e.g., adjusted contrast images 812A and 812B) for at least a set of t-values.

[0040] Figure 9 shows an exemplary graph 900 of a global tone map applied to an image according to one embodiment. Graph 900 shows a global tone mapping curve 902 applied to values ​​in an input image in various embodiments. In one example, the global tone mapping curve is applied to the luminance values ​​of the input image. The process for determining and applying global tone mapping to an input image is defined in U.S. Patent Application No. 16 / 526,902, filed July 30, 2019, entitled "ENHANCED HIGH-DYNAMIC-RANGE IMAGING AND TONE MAPPING," which is incorporated by reference as if it were described in whole.

[0041] Figure 10 shows a method 1000 according to one embodiment, in which the output image 1020 is generated by combining at least the final contrast image 1014 with global tone mapping 1004. For example, contrast image C F (x,y) (for example, the final contrast image 814) is merged with the result of applying global tone mapping to the input image as described above (indicated by the "+" symbol in Figure 10). In one example, the input image includes the input image 102 described above with respect to Figure 1. In various embodiments, the output image 1020 shows the result of applying separate tone gain functions to individual tonal regions of the HDR image.

[0042] The following description includes numerous specific details to provide a more complete understanding of at least one embodiment. However, it will be apparent to those skilled in the art that the inventive concept can be implemented without one or more of these specific details.

[0043] Data center Figure 11 shows an exemplary data center 1100 according to at least one embodiment. In at least one embodiment, the data center 1100 includes, but is not limited to, a data center infrastructure layer 1110, a framework layer 1120, a software layer 1130, and an application layer 1140.

[0044] In at least one embodiment, as shown in Figure 11, the data center infrastructure layer 1110 may include a resource orchestrator 1112, grouped computing resources 1114, and node computing resources ("node CRs") 1116(1) to 1116(N), where "N" represents any positive integer. In at least one embodiment, nodes CR1116(1) to 1116(N) may include, but not limited to, any number of central processing units ("CPU") or other processors (including accelerators, field-programmable gate arrays ("FPGA"), graphics processors, etc.), 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, virtual machines ("VM"), power modules, and cooling modules. In at least one embodiment, one or more nodes CR from among nodes CR1116(1) to 1116(N) may be servers having one or more of the computing resources described above.

[0045] In at least one embodiment, the grouped computing resources 1114 may include separate groupings of node CRs housed in one or more racks (not shown), or many racks housed in a data center at various geographical locations (also not shown). A separate grouping of node CRs within the grouped computing resources 1114 may include grouped compute resources, network resources, memory resources, or storage resources that can be configured or allocated to support one or more workloads. In at least one embodiment, several node CRs, including CPUs or processors, may be grouped in one or more racks to provide compute resources to support one or more workloads. In at least one embodiment, one or more racks may also include any number of power modules, cooling modules, and network switches in any combination.

[0046] In at least one embodiment, the resource orchestrator 1112 may configure or otherwise control one or more nodes CR1116(1) to 1116(N) and / or a grouped computing resource 1114. In at least one embodiment, the resource orchestrator 1112 may include a software design infrastructure ("SDI") management entity for the data center 1100. In at least one embodiment, the resource orchestrator 1112 may include hardware, software, or any combination thereof.

[0047] In at least one embodiment, as shown in Figure 11, the framework layer 1120 includes, but is not limited to, a job scheduler 1132, a configuration manager 1134, a resource manager 1136, and a distributed file system 1138. In at least one embodiment, the framework layer 1120 may include a framework for supporting the software 1152 of the software layer 1130 and / or one or more applications 1142 of the application layer 1140. In at least one embodiment, the software 1152 or (one or more) applications 1142 may include web-based service software or applications, such as those provided by Amazon Web Services, Google Cloud, and Microsoft Azure, respectively. In at least one embodiment, the framework layer 1120 may be a type of free and open-source software web application framework, such as Apache Spark® ("Spark"), which can leverage the distributed file system 1138 for large-scale data processing (e.g., "big data"). In at least one embodiment, the job scheduler 1132 may include a Spark driver to facilitate scheduling of workloads supported by various layers of the data center 1100. In at least one embodiment, the configuration manager 1134 may be able to configure different layers, such as the software layer 1130 and the framework layer 1120, which includes Spark and a distributed file system 1138 to support large-scale data processing. In at least one embodiment, the resource manager 1136 may be able to manage clustered or grouped computing resources that are mapped or allocated to support the distributed file system 1138 and the job scheduler 1132.In at least one embodiment, clustered or grouped computing resources may include grouped computing resources 1114 in the data center infrastructure layer 1110. In at least one embodiment, a resource manager 1136 may work in conjunction with a resource orchestrator 1112 to manage these mapped or allocated computing resources.

[0048] In at least one embodiment, the software 1152 contained within the software layer 1130 may include software used by nodes CR1116(1) to 1116(N), grouped computing resources 1114, and / or at least a portion of the distributed file system 1138 of the framework layer 1120. One or more types of software may include, but are not limited to, internet web page search software, email virus scanning software, database software, and streaming video content software.

[0049] In at least one embodiment, the applications 1142 (one or more) contained within the application layer 1140 may include one or more types of applications used by nodes CR1116(1) to 1116(N), grouped computing resources 1114, and / or at least a portion of the distributed file system 1138 of the framework layer 1120. At least one or more types of applications may include, but are not limited to, CUDA applications.

[0050] In at least one embodiment, any of the configuration manager 1134, resource manager 1136, and resource orchestrator 1112 may implement any number and type of self-correcting actions based on any amount and type of data obtained in any technically feasible manner. In at least one embodiment, the self-correcting actions may relieve the data center operator of data center 1100 of the task of determining potentially faulty configurations and potentially avoiding underutilized and / or underperforming portions of the data center.

[0051] Computer-based systems The following diagrams illustrate exemplary computer-based systems that may be used to implement at least one embodiment, but are not limited to it.

[0052] Figure 12 shows a processing system 1200 according to at least one embodiment. In at least one embodiment, the processing system 1200 includes one or more processors 1202 and one or more graphics processors 1208, and may be a single-processor desktop system, a multi-processor workstation system, or a server system having a large number of processors 1202 or processor cores 1207. In at least one embodiment, the processing system 1200 is a processing platform embedded in a system-on-a-chip ("SoC") integrated circuit for use in a mobile device, handheld device, or embedded device.

[0053] In at least one embodiment, the processing system 1200 may include, or be incorporated into, a server-based gaming platform, game console, media console, mobile gaming console, handheld game console, or online game console. In at least one embodiment, the processing system 1200 is a mobile phone, smartphone, tablet computing device, or mobile internet device. In at least one embodiment, the processing system 1200 may also include, be coupled to, or be incorporated into wearable devices such as a smartwatch wearable device, smart eyewear device, augmented reality device, or virtual reality device. In at least one embodiment, the processing system 1200 is a television or set-top box device having one or more processors 1202 and a graphical interface produced by one or more graphics processors 1208.

[0054] In at least one embodiment, one or more processors 1202 each include one or more processor cores 1207 for processing instructions that, when executed, perform actions for system and user software. In at least one embodiment, each of the one or more processor cores 1207 is configured to process a particular instruction set 1209. In at least one embodiment, the instruction set 1209 may facilitate computing via complex instruction set computing ("CISC"), reduced instruction set computing ("RISC"), or very long instruction word ("VLIW"). In at least one embodiment, each processor core 1207 may process a different instruction set 1209, and the instruction set 1209 may include instructions to facilitate emulation of other instruction sets. In at least one embodiment, the processor core 1207 may also include other processing devices, such as a digital signal processor ("DSP").

[0055] In at least one embodiment, the processor 1202 includes a cache memory ("cache") 1204. In at least one embodiment, the processor 1202 may have a single internal cache or multiple levels of internal caches. In at least one embodiment, the cache memory is shared among the various components of the processor 1202. In at least one embodiment, the processor 1202 also uses an external cache (e.g., a Level 3 ("L3") cache or a Last Level Cache ("LLC")) (not shown), and the external cache may be shared among the processor cores 1207 using known cache coherency techniques. In at least one embodiment, additionally, a register file 1206 is included in the processor 1202, and the register file 1206 may contain different types of registers for storing different types of data (e.g., integer registers, floating-point registers, status registers, and instruction pointer registers). In at least one embodiment, the register file 1206 may contain general-purpose registers or other registers.

[0056] In at least one embodiment, one or more processors 1202 are coupled with one or more interface buses 1210 to transmit communication signals, such as addresses, data, or control signals, between the processors 1202 and other components in the processing system 1200. In at least one embodiment, the interface bus 1210 in one embodiment may be a processor bus, such as a version of the Direct Media Interface ("DMI") bus. In at least one embodiment, the interface bus 1210 is not limited to the DMI bus and may include one or more peripheral component interconnect buses (e.g., "PCI": Peripheral Component Interconnect, PCI Express ("PCIe")), memory buses, or other types of interface buses. In at least one embodiment, the (one or more) processors 1202 include an integrated memory controller 1216 and a platform controller hub 1230. In at least one embodiment, the memory controller 1216 facilitates communication between the memory device and other components of the processing system 1200, and the platform controller hub ("PCH") 1230 provides connectivity to I / O devices via a local input / output ("I / O") bus.

[0057] In at least one embodiment, the memory device 1220 may be a dynamic random access memory ("DRAM") device, a static random access memory ("SRAM") device, a flash memory device, a phase-change memory device, or any other memory device having performance suitable for acting as processor memory. In at least one embodiment, the memory device 1220 may act as system memory for the processing system 1200 to store data 1222 and instructions 1221 for use when one or more processors 1202 execute an application or process. In at least one embodiment, the memory controller 1216 may also be coupled to an optional external graphics processor 1212, which may communicate with one or more graphics processors 1208 in the processor 1202 to perform graphics and media operations. In at least one embodiment, the display device 1211 may be connected to one or more processors 1202. In at least one embodiment, the display device 1211 may include one or more internal display devices, such as those found in mobile electronic devices or laptop devices, or external display devices attached via a display interface (e.g., DisplayPort). In at least one embodiment, the display device 1211 may include a head-mounted display (HMD), such as a stereoscopic display device for use in virtual reality ("VR") or augmented reality ("AR") applications.

[0058] In at least one embodiment, the platform controller hub 1230 enables peripherals to connect to the memory device 1220 and processor 1202 via a high-speed I / O bus. In at least one embodiment, the I / O peripherals include, but are not limited to, an audio controller 1246, a network controller 1234, a firmware interface 1228, a wireless transceiver 1226, a touch sensor 1225, and a data storage device 1224 (e.g., a hard disk drive, flash memory, etc.). In at least one embodiment, the data storage device 1224 may be connected via a storage interface (e.g., SATA) or via a peripheral bus such as PCI or PCIe. In at least one embodiment, the touch sensor 1225 may include a touchscreen sensor, a pressure sensor, or a fingerprint sensor. In at least one embodiment, the wireless transceiver 1226 may 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, the firmware interface 1228 enables communication with system firmware and may be, for example, a Unified Extensible Firmware Interface (UEFI). In at least one embodiment, the network controller 1234 may enable network connectivity to a wired network. In at least one embodiment, a high-performance network controller (not shown) is coupled to the interface bus 1210. In at least one embodiment, the audio controller 1246 is a multi-channel high-definition audio controller.In at least one embodiment, the processing system 1200 includes an optional legacy I / O controller 1240 for connecting legacy devices (e.g., Personal System 2 ("PS / 2")) to the processing system 1200. In at least one embodiment, the platform controller hub 1230 may also connect to one or more Universal Serial Bus ("USB") controller 1242-connected input devices, such as a keyboard and mouse combination 1243, a camera 1244, or other USB input devices.

[0059] In at least one embodiment, instances of the memory controller 1216 and the platform controller hub 1230 may be integrated into a discreet external graphics processor, such as an external graphics processor 1212. In at least one embodiment, the platform controller hub 1230 and / or the memory controller 1216 may be external to one or more processors 1202. For example, in at least one embodiment, the processing system 1200 may include an external memory controller 1216 and a platform controller hub 1230, which may be configured as a memory controller hub and peripheral controller hub within a system chipset communicating with (one or more) processors 1202.

[0060] Figure 13 shows a computer system 1300 according to at least one embodiment. In at least one embodiment, the computer system 1300 may be a system, a SOC, or any combination of interconnected devices and components. In at least one embodiment, the computer system 1300 is formed together with a processor 1302 which may include an execution unit for executing instructions. In at least one embodiment, the computer system 1300 may include components such as the processor 1302 for employing an execution unit which includes logic for implementing algorithms for processing data. In at least one embodiment, the computer system 1300 may include a processor such as the PENTIUM® processor family, Xeon®, Itanium®, XScale®, and / or StrongARM®, Intel® Core®, or Intel® Nervana® microprocessors, available from Intel Corporation in Santa Clara, California, but other systems (including PCs with other microprocessors, engineering workstations, set-top boxes, etc.) may also be used. In at least one embodiment, the computer system 1300 may run a version of the WINDOWS® operating system available from Microsoft Corporation in Redmond, Washington, but other operating systems (e.g., UNIX® and Linux®), embedded software, and / or graphical user interfaces may also be used.

[0061] In at least one embodiment, the computer system 1300 may 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, the embedded application may include a microcontroller, a digital signal processor (DSP), a SoC, a network computer ("NetPC"), a set-top box, a network hub, a wide area network ("WAN") switch, or any other system capable of executing one or more instructions.

[0062] In at least one embodiment, the computer system 1300 may include, but is not limited to, a processor 1302, which may include, but is not limited to, one or more execution units 1308 that can be configured to run a Compute Unified Device Architecture ("CUDA") program (CUDA® is a registered trademark developed by NVIDIA Corporation in Santa Clara, California). In at least one embodiment, the CUDA program is at least part of a software application written in the CUDA programming language. In at least one embodiment, the computer system 1300 is a single-processor desktop or server system. In at least one embodiment, the computer system 1300 may be a multi-processor system. In at least one embodiment, the processor 1302 may include, but is not limited to, 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, the processor 1302 may be coupled to a processor bus 1310, and the processor bus 1310 may transmit data signals between the processor 1302 and other components in the computer system 1300.

[0063] In at least one embodiment, the processor 1302 may include, but is not limited to, a level 1 ("L1") internal cache memory ("cache") 1304. In at least one embodiment, the processor 1302 may have a single internal cache or multiple levels of internal caches. In at least one embodiment, the cache memory may reside outside the processor 1302. In at least one embodiment, the processor 1302 may also include a combination of both internal and external caches. In at least one embodiment, the register file 1306 may store different types of data in various registers, including, but is not limited to, integer registers, floating-point registers, status registers, and instruction pointer registers.

[0064] In at least one embodiment, but not limited to, an execution unit 1308 containing logic for performing integer and floating-point arithmetic may also be present in the processor 1302. The processor 1302 may also include a microcode ("u-code") read-only memory ("ROM") for storing microcode for several macro instructions. In at least one embodiment, the execution unit 1308 may include logic for handling a packed instruction set 1309. In at least one embodiment, by including the packed instruction set 1309, along with the associated circuit elements for executing the instructions, in the instruction set of the general-purpose processor 1302, arithmetic used by many multimedia applications can be performed using packed data in the general-purpose processor 1302. In at least one embodiment, many multimedia applications may be accelerated and run more efficiently by using the full width of the processor's data bus to perform arithmetic on packed data, which may eliminate the need to transfer smaller units of data across the processor's data bus to perform one or more arithmetic operations, one data element at a time.

[0065] In at least one embodiment, the execution unit 1308 may also be used in a microcontroller, embedded processor, graphics device, DSP, and other types of logic circuits. In at least one embodiment, the computer system 1300 may include, but is not limited to, memory 1320. In at least one embodiment, memory 1320 may be implemented as a DRAM device, SRAM device, flash memory device, or other memory device. Memory 1320 may store (one or more) instructions 1319 and / or data 1321, which are represented by data signals that can be executed by the processor 1302.

[0066] In at least one embodiment, a system logic chip may be coupled to a processor bus 1310 and memory 1320. In at least one embodiment, the system logic chip may include, but is not limited to, a memory controller hub ("MCH") 1316, and the processor 1302 may communicate with the MCH 1316 via the processor bus 1310. In at least one embodiment, the MCH 1316 may provide a high-bandwidth memory path 1318 to memory 1320 for instruction and data storage, as well as for the storage of graphics commands, data, and textures. In at least one embodiment, the MCH 1316 may direct data signals between the processor 1302, memory 1320, and other components in the computer system 1300, and bridge data signals between the processor bus 1310, memory 1320, and system I / O 1322. In at least one embodiment, the system logic chip may provide a graphics port for coupling to a graphics controller. In at least one embodiment, the MCH1316 may be coupled to memory 1320 through a high-bandwidth memory path 1318, and the graphics / video card 1312 may be coupled to the MCH1316 via an Accelerated Graphics Port ("AGP") interconnect 1314.

[0067] In at least one embodiment, the computer system 1300 may use a system I / O 1322, which is a proprietary hub interface bus for coupling the MCH 1316 to the I / O controller hub ("ICH") 1330. In at least one embodiment, the ICH 1330 may provide direct connectivity to several I / O devices via a local I / O bus. In at least one embodiment, the local I / O bus may include, but is not limited to, a high-speed I / O bus for connecting peripherals to memory 1320, the chipset, and the processor 1302. Examples may include, but are not limited to, an audio controller 1329, a firmware hub ("Flash BIOS") 1328, a wireless transceiver 1326, data storage 1324, a legacy I / O controller 1323 including a user input interface 1325 and a keyboard interface, a serial expansion port 1327 such as USB, and a network controller 1334. The data storage 1324 may comprise a hard disk drive, a floppy disk drive, a CD-ROM device, a flash memory device, or other mass storage device.

[0068] In at least one embodiment, Figure 13 shows a system including interconnected hardware devices or “chips.” In at least one embodiment, Figure 13 may show an exemplary SoC. In at least one embodiment, the devices shown in Figure 13 may be interconnected by proprietary interconnects, standard interconnects (e.g., PCIe), or any combination thereof. In at least one embodiment, one or more components of system 1300 are interconnected using a Compute Express Link ("CXL") interconnect.

[0069] Figure 14 shows System 1400 according to at least one embodiment. In at least one embodiment, System 1400 is an electronic device utilizing Processor 1410. In at least one embodiment, System 1400 may be, for example, a notebook, tower server, rack server, blade server, laptop, desktop, tablet, mobile device, telephone, embedded computer, or any other suitable electronic device, without limitation.

[0070] In at least one embodiment, system 1400 may include a processor 1410 communicably coupled to any number or type of preferred components, peripherals, modules, or devices, but not limited to these. In at least one embodiment, the processor 1410 is coupled using a bus or interface, such as an I2C 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 Advance Technology Attachment ("SATA") bus, a USB (versions 1, 2, or 3), or a Universal Asynchronous Receiver / Transmitter ("UART") bus. In at least one embodiment, Figure 14 shows a system including interconnected hardware devices or "chips." In at least one embodiment, Figure 14 may show an exemplary SoC. In at least one embodiment, the devices shown in Figure 14 may be interconnected by proprietary interconnects, standard interconnects (e.g., PCIe), or any combination thereof. In at least one embodiment, one or more components of Figure 14 are interconnected using a CXL interconnect.

[0071] In at least one embodiment, Figure 14 includes a display 1424, a touchscreen 1425, a touchpad 1430, a Near Field Communication ("NFC") unit 1445, a sensor hub 1440, a thermal sensor 1446, an Express Chipset ("EC") 1435, a Trusted Platform Module ("TPM") 1438, a BIOS / firmware / flash memory ("BIOS,FW flash") 1422, a DSP 1460, a Solid State Disk ("SSD") or Hard Disk Drive ("HDD") 1420, a Wireless Local Area Network ("WLAN") unit 1450, a Bluetooth unit 1452, and a Wireless Wide Area Network ("WWAN") unit. This may include a network (1456), a Global Positioning System ("GPS") (1455), a camera such as a USB 3.0 camera ("USB 3.0 camera") (1454), or a Low Power Double Data Rate ("LPDDR") memory unit ("LPDDR3") (1415) implemented, for example, in the LPDDR3 standard. Each of these components may be implemented in any preferred manner.

[0072] In at least one embodiment, other components may be communicatively coupled to the processor 1410 through the components described above. In at least one embodiment, the accelerometer 1441, the ambient light sensor ("ALS") 1442, the compass 1443, and the gyroscope 1444 may be communicatively coupled to the sensor hub 1440. In at least one embodiment, the thermal sensor 1439, the fan 1437, the keyboard 1436, and the touchpad 1430 may be communicatively coupled to the EC 1435. In at least one embodiment, the speaker 1463, the headphones 1464, and the microphone ("mic") 1465 may be communicatively coupled to the audio unit ("audio codec and class d amplifier") 1462, and the audio unit 1462 may be communicatively coupled to the DSP 1460. In at least one embodiment, the audio unit 1462 may include, for example, an audio coder / decoder ("codec") and a Class D amplifier. In at least one embodiment, a SIM card ("SIM") 1457 may be communicatively coupled to the WWAN unit 1456. In at least one embodiment, components such as the WLAN unit 1450 and the Bluetooth unit 1452, as well as the WWAN unit 1456, may be implemented in a Next Generation Form Factor ("NGFF").

[0073] Figure 15 shows an exemplary integrated circuit 1500 according to at least one embodiment. In at least one embodiment, the exemplary integrated circuit 1500 is a SoC that may be fabricated using one or more IP cores. In at least one embodiment, the integrated circuit 1500 includes one or more application processors 1505 (e.g., CPUs), at least one graphics processor 1510, and additionally, an image processor 1515 and / or a video processor 1520, any of which may be modular IP cores. In at least one embodiment, the integrated circuit 1500 includes peripherals or bus logic, including a USB controller 1525, a UART controller 1530, an SPI / SDIO controller 1535, and an I2S / I2C controller 1540. In at least one embodiment, the integrated circuit 1500 may include a display device 1545 coupled to one or more of the following: a high-definition multimedia interface ("HDMI®") controller 1550 and a mobile industry processor interface ("MIPI") display interface 1555. In at least one embodiment, storage may be provided by a flash memory subsystem 1560, which includes flash memory and a flash memory controller. In at least one embodiment, a memory interface may be provided via a memory controller 1565 for access to SDRAM or SRAM memory devices. In at least one embodiment, some integrated circuits may additionally include an embedded security engine 1570.

[0074] Figure 16 shows a computing system 1600 according to at least one embodiment. In at least one embodiment, the computing system 1600 includes a processing subsystem 1601 having one or more processors 1602 and system memory 1604 communicating via an interconnection path which may include a memory hub 1605. In at least one embodiment, the memory hub 1605 may be a separate component within a chipset component or may be integrated within one or more processors 1602. In at least one embodiment, the memory hub 1605 is coupled to an I / O subsystem 1611 via a communication link 1606. In at least one embodiment, the I / O subsystem 1611 includes an I / O hub 1607 which can enable the computing system 1600 to receive input from one or more input devices 1608. In at least one embodiment, the I / O hub 1607 can enable a display controller, which may be included in one or more processors 1602, to provide output to one or more display devices 1610A. In at least one embodiment, one or more display devices 1610A coupled with the I / O hub 1607 may include local, internal, or embedded display devices.

[0075] In at least one embodiment, the processing subsystem 1601 includes one or more parallel processors 1612 coupled to a memory hub 1605 via a bus or other communication link 1613. In at least one embodiment, the communication link 1613 may be one of any number of standards-based communication link technologies or protocols, such as PCIe, or it may be a vendor-specific communication interface or communication fabric. In at least one embodiment, one or more parallel processors 1612 may form a computation-focused parallel or vector processing system, which may include a large number of processing cores and / or processing clusters, such as a many-integrated-core processor. In at least one embodiment, one or more parallel processors 1612 may form a graphics processing subsystem, which may output pixels to one of one or more display devices 1610A coupled via an I / O hub 1607. In at least one embodiment, one or more parallel processors 1612 may also include a display controller and a display interface (not shown) for enabling direct connection to one or more display devices 1610B.

[0076] In at least one embodiment, the system storage unit 1614 can be connected to the I / O hub 1607 to provide storage functionality for the computing system 1600. In at least one embodiment, an I / O switch 1616 may be used to provide an interface mechanism for enabling connections between the I / O hub 1607 and other components such as a network adapter 1618 and / or a wireless network adapter 1619 that may be incorporated into the platform, as well as various other devices that may be added via one or more add-in devices 1620. In at least one embodiment, the network adapter 1618 may be an Ethernet adapter or another wired network adapter. In at least one embodiment, the wireless network adapter 1619 may include one or more other network devices, including Wi-Fi, Bluetooth, NFC, or one or more wireless radios.

[0077] In at least one embodiment, the computing system 1600 may include other components not explicitly shown, which may also be connected to the I / O hub 1607, including USB or other port connections, optical storage drives, video capture devices, etc. In at least one embodiment, the communication paths interconnecting the various components in Figure 16 may be implemented using any preferred protocol, such as a PCI-based protocol (e.g., PCIe), or other bus or point-to-point communication interfaces and / or protocols, or interconnection protocols, such as NVLink High Speed ​​Interconnection.

[0078] In at least one embodiment, one or more parallel processors 1612 incorporate circuit elements optimized for graphics and video processing, such as video output circuit elements, to constitute a graphics processing unit ("GPU"). In at least one embodiment, one or more parallel processors 1612 incorporate circuit elements optimized for general-purpose processing. In at least one embodiment, the components of the computing system 1600 may be integrated with one or more other system elements on a single integrated circuit. For example, in at least one embodiment, one or more parallel processors 1612, a memory hub 1605, one or more processors 1602, and an I / O hub 1607 may be incorporated into a SoC integrated circuit. In at least one embodiment, the components of the computing system 1600 may be incorporated into a single package to form a system-in-package ("SIP") configuration. In at least one embodiment, at least a portion of the components of the computing system 1600 may be incorporated into a multi-chip module ("MCM"), which can be interconnected with other multi-chip modules to form a modular computing system. In at least one embodiment, the I / O subsystem 1611 and the display device 1610B are omitted from the computing system 1600.

[0079] Processing system The following diagram illustrates an exemplary processing system that may be used to implement at least one embodiment, but is not limited to it.

[0080] Figure 17 shows an accelerated processing unit ("APU") 1700 according to at least one embodiment. In at least one embodiment, the APU 1700 is developed by AMD Corporation in Santa Clara, California. In at least one embodiment, the APU 1700 may be configured to run application programs, such as CUDA programs. In at least one embodiment, the APU 1700 includes, but is not limited to, a core complex 1710, a graphics complex 1740, a fabric 1760, an I / O interface 1770, a memory controller 1780, a display controller 1792, and a multimedia engine 1794. In at least one embodiment, the APU 1700 may include, but is not limited to, any number of core complexes 1710, any number of graphics complexes 1750, any number of display controllers 1792, and any number of multimedia engines 1794 in any combination. For illustrative purposes, multiple instances of similar objects are shown herein with a reference number identifying the object and, where necessary, a number in parentheses identifying the instance.

[0081] In at least one embodiment, the core complex 1710 is a CPU, the graphics complex 1740 is a GPU, and the APU 1700 is a processing unit that incorporates the 1710 and 1740 on a single chip, but is not limited to this. In at least one embodiment, some tasks may be assigned to the core complex 1710, and other tasks may be assigned to the graphics complex 1740. In at least one embodiment, the core complex 1710 is configured to run main control software related to the APU 1700, such as the operating system. In at least one embodiment, the core complex 1710 is the master processor of the APU 1700, controlling and coordinating the operation of the other processors. In at least one embodiment, the core complex 1710 issues commands that control the operation of the graphics complex 1740. In at least one embodiment, the core complex 1710 may be configured to run host executable code derived from CUDA source code, and the graphics complex 1740 may be configured to run device executable code derived from CUDA source code.

[0082] In at least one embodiment, the core complex 1710 includes, but is not limited to, cores 1720(1) to 1720(4) and an L3 cache 1730. In at least one embodiment, the core complex 1710 may include, but is not limited to, any number of cores 1720 and any number and type of caches in any combination. In at least one embodiment, the core 1720 is configured to execute instructions of a particular instruction set architecture ("ISA"). In at least one embodiment, each core 1720 is a CPU core.

[0083] In at least one embodiment, each core 1720 includes, but is not limited to, 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, the fetch / decode unit 1722 fetches instructions, decodes such instructions, generates micro-operations, and dispatches separate microinstructions to the integer execution engine 1724 and the floating-point execution engine 1726. In at least one embodiment, the fetch / decode unit 1722 can simultaneously dispatch one microinstruction to the integer execution engine 1724 and another microinstruction to the floating-point execution engine 1726. In at least one embodiment, the integer execution engine 1724 performs, but is not limited to, integer and memory operations. In at least one embodiment, the floating-point engine 1726 performs, but is not limited to, floating-point and vector operations. In at least one embodiment, the fetch-decoder unit 1722 dispatches microinstructions to a single execution engine that replaces both the integer execution engine 1724 and the floating-point execution engine 1726.

[0084] In at least one embodiment, each core 1720(i), where i is an integer representing a particular instance of core 1720, can access the L2 cache 1728(i) contained within core 1720(i). In at least one embodiment, each core 1720 contained within core complex 1710(j), where j is an integer representing a particular instance of core complex 1710, is connected to other cores 1720 contained within core complex 1710(j) via the L3 cache 1730(j) contained within core complex 1710(j). In at least one embodiment, a core 1720 contained within core complex 1710(j), where j is an integer representing a particular instance of core complex 1710, can access all of the L3 cache 1730(j) contained within core complex 1710(j). In at least one embodiment, the L3 cache 1730 may contain any number of slices, but is not limited to any number.

[0085] In at least one embodiment, the graphics complex 1740 may be configured to perform compute operations in a highly parallel manner. In at least one embodiment, the graphics complex 1740 is configured to perform graphics pipeline operations, such as drawing commands, pixel operations, geometric calculations, and other operations related to rendering an image to a display. In at least one embodiment, the graphics complex 1740 is configured to perform non-graphics operations. In at least one embodiment, the graphics complex 1740 is configured to perform both graphics-related and non-graphics operations.

[0086] In at least one embodiment, the graphics complex 1740 includes, but not limited to, any number of compute units 1750 and an L2 cache 1742. In at least one embodiment, the compute units 1750 share the L2 cache 1742. In at least one embodiment, the L2 cache 1742 is partitioned. In at least one embodiment, the graphics complex 1740 includes, but not limited to, any number of compute units 1750 and any number and type of cache (including zero). In at least one embodiment, the graphics complex 1740 includes, but not limited to, any amount of dedicated graphics hardware.

[0087] In at least one embodiment, each compute unit 1750 includes, but not limited to, any number of SIMD units 1752 and shared memory 1754. In at least one embodiment, each SIMD unit 1752 may implement a SIMD architecture and be configured to perform operations in parallel. In at least one embodiment, each compute unit 1750 may execute any number of thread blocks, but each thread block executes on a single compute unit 1750. In at least one embodiment, a thread block includes, but not limited to, any number of execution threads. In at least one embodiment, a workgroup is a thread block. In at least one embodiment, each SIMD unit 1752 executes a different warp. In at least one embodiment, a warp is a group of threads (e.g., 16 threads), where each thread in the warp belongs to a single thread block and is configured to process a different set of data based on a single set of instructions. In at least one embodiment, predication may be used to disable one or more threads in a warp. 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 warp. In at least one embodiment, different wavefronts in a thread block may be synchronized with each other and communicate via shared memory 1754.

[0088] In at least one embodiment, Fabric 1760 is a system interconnect that facilitates data and control transmissions across the core complex 1710, graphics complex 1740, I / O interface 1770, memory controller 1780, display controller 1792, and multimedia engine 1794. In at least one embodiment, APU 1700 may include, but not limited to, any number and type of system interconnects in addition to or instead of Fabric 1760, which facilitate data and control transmissions across any number and type of directly or indirectly linked components that may be inside or outside APU 1700. In at least one embodiment, I / O interface 1770 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 1770. In at least one embodiment, peripheral devices coupled to the I / O interface 1770 may include, but are not limited to, a keyboard, mouse, printer, scanner, joystick or other type of game controller, media recording device, external storage device, network interface card, and the like.

[0089] In at least one embodiment, the display controller AMD92 displays images on one or more display devices, such as liquid crystal display (LCD) devices. In at least one embodiment, the multimedia engine 240 includes, but is not limited to, any amount and type of circuit elements related to multimedia, such as video decoders, video encoders, and image signal processors. In at least one embodiment, the memory controller 1780 facilitates data transfer between the APU 1700 and the unified system memory 1790. In at least one embodiment, the core complex 1710 and the graphics complex 1740 share the unified system memory 1790.

[0090] In at least one embodiment, the APU 1700 implements a memory subsystem including any amount and type of memory controller 1780 and memory devices (e.g., shared memory 1754) which may be dedicated to one component or shared among multiple components. In at least one embodiment, the APU 1700 implements a cache subsystem including one or more cache memories (e.g., L2 cache 1828, L3 cache 1730, and L2 cache 1742), each of which may be private to any number of components (e.g., core 1720, core complex 1710, SIMD unit 1752, compute unit 1750, and graphics complex 1740) or shared among any number of components.

[0091] Figure 18 shows a CPU 1800 according to at least one embodiment. In at least one embodiment, the CPU 1800 is developed by AMD Corporation in Santa Clara, California. In at least one embodiment, the CPU 1800 may be configured to run application programs. In at least one embodiment, the CPU 1800 is configured to run main control software, such as an operating system. In at least one embodiment, the CPU 1800 issues commands to control the operation of an external GPU (not shown). In at least one embodiment, the CPU 1800 may be configured to run host executable code derived from CUDA source code, and the external GPU may be configured to run device executable code derived from such CUDA source code. In at least one embodiment, the CPU 1800 includes, but is not limited to, any number of core complexes 1810, a fabric 1860, an I / O interface 1870, and a memory controller 1880.

[0092] In at least one embodiment, the core complex 1810 includes, but is not limited to, cores 1820(1) to 1820(4) and an L3 cache 1830. In at least one embodiment, the core complex 1810 may include, but is not limited to, any number of cores 1820 and any number and type of caches in any combination. In at least one embodiment, the core 1820 is configured to execute instructions of a particular ISA. In at least one embodiment, each core 1820 is a CPU core.

[0093] In at least one embodiment, each core 1820 includes, but is not limited to, a fetch / decode unit 1822, an integer execution engine 1824, a floating-point execution engine 1826, and an L2 cache 1828. In at least one embodiment, the fetch / decode unit 1822 fetches instructions, decodes such instructions, generates microoperations, and dispatches separate microinstructions to the integer execution engine 1824 and the floating-point execution engine 1826. In at least one embodiment, the fetch / decode unit 1822 can simultaneously dispatch one microinstruction to the integer execution engine 1824 and another microinstruction to the floating-point execution engine 1826. In at least one embodiment, the integer execution engine 1824 performs, but is not limited to, integer and memory operations. In at least one embodiment, the floating-point engine 1826 performs, but is not limited to, floating-point and vector operations. In at least one embodiment, the fetch-decoder unit 1822 dispatches microinstructions to a single execution engine that replaces both the integer execution engine 1824 and the floating-point execution engine 1826.

[0094] In at least one embodiment, each core 1820(i), where i is an integer representing a particular instance of core 1820, can access the L2 cache 1828(i) contained within core 1820(i). In at least one embodiment, each core 1820 contained within core complex 1810(j), where j is an integer representing a particular instance of core complex 1810, is connected to other cores 1820 in core complex 1810(j) via the L3 cache 1830(j) contained within core complex 1810(j). In at least one embodiment, a core 1820 contained within core complex 1810(j), where j is an integer representing a particular instance of core complex 1810, can access all of the L3 cache 1830(j) contained within core complex 1810(j). In at least one embodiment, the L3 cache 1830 may contain any number of slices, but is not limited to any number.

[0095] In at least one embodiment, the fabric 1860 is a system interconnect that facilitates data and control transmissions across the core complexes 1810(1) to 1810(N) (where N is an integer greater than 0), the I / O interface 1870, and the memory controller 1880. In at least one embodiment, the CPU 1800 may include, but not limited to, any number and type of system interconnects in addition to or instead of the fabric 1860, which facilitates data and control transmissions across any number and type of directly or indirectly linked components that may be inside or outside the CPU 1800. In at least one embodiment, the I / O interface 1870 represents any number and type of I / O interfaces (e.g., PCI, PCI-X, PCIe, GBE, USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to the I / O interface 1870. In at least one embodiment, peripheral devices coupled to the I / O interface 1870 may include, but are not limited to, a display, keyboard, mouse, printer, scanner, joystick or other type of game controller, media recording device, external storage device, network interface card, and the like.

[0096] In at least one embodiment, the memory controller 1880 facilitates data transfer between the CPU 1800 and the system memory 1890. In at least one embodiment, the core complex 1810 and the graphics complex 1840 share the system memory 1890. In at least one embodiment, the CPU 1800 implements a memory subsystem including any number and type of memory controllers 1880 and memory devices, which may be dedicated to one component or shared among multiple components, but are not limited. In at least one embodiment, the CPU 1800 implements a cache subsystem including one or more cache memories (e.g., L2 cache 1828 and L3 cache 1830), each of which may be private to any number of components (e.g., core 1820 and core complex 1810) or shared among any number of components.

[0097] Figure 19 shows an exemplary accelerator-integrated slice 1990 according to at least one embodiment. As used herein, a “slice” comprises a designated portion of the processing resources of an accelerator-integrated circuit. In at least one embodiment, the accelerator-integrated circuit provides cache management, memory access, context management, and interrupt management services on behalf of multiple graphics processing engines contained within a graphics acceleration module. Each graphics processing engine may comprise a separate GPU. Alternatively, the graphics processing engine may comprise different types of graphics processing engines within the GPU, such as a graphics execution unit, a media processing engine (e.g., a video encoder / decoder), a sampler, and a Blit engine. In at least one embodiment, the graphics acceleration module may be a GPU with multiple graphics processing engines. In at least one embodiment, the graphics processing engines may be individual GPUs integrated on a common package, line card, or chip.

[0098] The application effective address space 1982 in system memory 1914 stores process elements 1983. In one embodiment, process elements 1983 are stored in response to a GPU call 1981 from an application 1980 running on processor 1907. Process elements 1983 contain the process state of the corresponding application 1980. The work descriptor ("WD") 1984 contained in process elements 1983 may be a single job requested by the application or may contain a pointer to a queue of jobs. In at least one embodiment, WD 1984 is a pointer to a job request queue in the application effective address space 1982.

[0099] The graphics acceleration module 1946 and / or individual graphics processing engines may be shared by all or a subset of processes in the system. In at least one embodiment, infrastructure may be included for setting process states and sending WD1984 to the graphics acceleration module 1946 to start jobs in a virtualized environment.

[0100] In at least one embodiment, the dedicated process programming model is implementation-specific. In this model, a single process owns the graphics acceleration module 1946 or an individual graphics processing engine. Since the graphics acceleration module 1946 is owned by a single process, the hypervisor initializes the accelerator integration circuitry for the owning partition, and when the graphics acceleration module 1946 is allocated, the operating system initializes the accelerator integration circuitry for the owning process.

[0101] During operation, the WD fetch unit 1991 in the accelerator integrated slice 1990 fetches the next WD 1984 containing instructions for the work to be performed by one or more graphics processing engines of the graphics acceleration module 1946. As shown, the data from WD 1984 is stored in register 1945 and can be used by the memory management unit ("MMU") 1939, the interrupt management circuit 1947, and / or the context management circuit 1948. For example, one embodiment of the MMU 1939 includes a segment / page walk circuit element for accessing the segment / page table 1986 in the OS virtual address space 1985. The interrupt management circuit 1947 may process an interrupt event ("INT") 1992 received from the graphics acceleration module 1946. When performing graphics operations, the effective address 1993 generated by the graphics processing engine is translated to a real address by the MMU 1939.

[0102] In one embodiment, the same set of registers 1945 may be duplicated for each graphics processing engine and / or graphics acceleration module 1946 and initialized by the hypervisor or operating system. Each of these duplicated registers may be included in the accelerator integration slice 1990. Exemplary registers that may be initialized by the hypervisor are shown in Table 1. [Table 1] Table 2 shows exemplary registers that can be initialized by the operating system. [Table 2]

[0103] In one embodiment, each WD1984 is specific to a particular graphics acceleration module 1946 and / or a particular graphics processing engine. The WD1984 contains all the information required by the graphics processing engine to perform its work, or the WD1984 This could be a pointer to a memory location set by the application that contains the command queue for the work to be completed.

[0104] Figures 20A and 20B show exemplary graphics processors according to at least one embodiment. In at least one embodiment, any of the exemplary graphics processors may be fabricated using one or more IP cores. In addition to those shown, at least one embodiment may include other logic and circuitry, including additional graphics processors / cores, peripheral interface controllers, or general-purpose processor cores. In at least one embodiment, the exemplary graphics processors are intended for use within a SoC.

[0105] Figure 20A shows an exemplary graphics processor 2010 of an SoC integrated circuit that may be fabricated using one or more IP cores according to at least one embodiment. Figure 20B shows an additional exemplary graphics processor 2040 of an SoC integrated circuit that may be fabricated using one or more IP cores according to at least one embodiment. In at least one embodiment, the graphics processor 2010 in Figure 20A is a low-power graphics processor core. In at least one embodiment, the graphics processor 2040 in Figure 20B is a higher-performance graphics processor core. In at least one embodiment, each of the graphics processors 2010 and 2040 may be a variation of the graphics processor 1510 in Figure 15.

[0106] In at least one embodiment, the graphics processor 2010 includes a vertex processor 2005 and one or more fragment processors 2015A-2015N (e.g., 2015A, 2015B, 2015C, 2015D-2015N-1, and 2015N). In at least one embodiment, the graphics processor 2010 can execute different shader programs via separate logic, thereby optimizing the vertex processor 2005 to perform operations for a vertex shader program, and the one or more fragment processors 2015A-2015N to perform fragment (e.g., pixel) shading operations for a fragment or pixel shader program. In at least one embodiment, the vertex processor 2005 performs the vertex processing stage of the 3D graphics pipeline and generates primitive and vertex data. In at least one embodiment, one or more fragment processors 2015A–2015N use primitive and vertex data generated by the vertex processor 2005 to create a frame buffer that is displayed on the display device. In at least one embodiment, one or more fragment processors 2015A–2015N are optimized to execute fragment shader programs such as those provided in the OpenGL API, and the OpenGL API may be used to perform similar operations to pixel shader programs such as those provided in the Direct 3D API.

[0107] In at least one embodiment, the graphics processor 2010 additionally includes one or more MMUs 2020A-2020B, (one or more) caches 2025A-2025B, and (one or more) circuit interconnects 2030A-2030B. In at least one embodiment, one or more MMUs 2020A-2020B provide virtual-physical address mappings for the graphics processor 2010, including vertex processors 2005 and / or (one or more) fragment processors 2015A-2015N, which may reference vertex or image / texture data stored in memory, in addition to vertex or image / texture data stored in one or more caches 2025A-2025B. In at least one embodiment, one or more MMUs 2020A-2020B may be synchronized with other MMUs in the system, including one or more MMUs associated with one or more application processors 1505, image processor 1515, and / or video processor 1520 in Figure 15, thereby allowing each processor 1505-1520 to participate in a shared or unified virtual memory system. In at least one embodiment, one or more circuit interconnects 2030A-2030B enable the graphics processor 2010 to interface with other IP cores in the SoC, either via the SoC's internal bus or via a direct connection.

[0108] In at least one embodiment, the graphics processor 2040 includes one or more MMUs 2020A to 2020B of the graphics processor 2010 in Figure 20A, caches 2025A to 2025B, and circuit interconnects 2030A to 2030B. In at least one embodiment, the graphics processor 2040 includes one or more shader cores 2055A-2055N (for example, 2055A, 2055B, 2055C, 2055D, 2055E, 2055F-2055N-1, and 2055N), and one or more shader cores 2055A-2055N provide 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 for implementing vertex shaders, fragment shaders, and / or compute shaders. In at least one embodiment, the number of shader cores can vary. In at least one embodiment, the graphics processor 2040 includes an intercore task manager 2045 that acts as a thread dispatcher for dispatching execution threads to one or more shader cores 2055A-2055N, and a tiling unit 2058 for accelerating tiling operations for tile-based rendering, where rendering operations for a scene are subdivided in image space, for example, to take advantage of local space coherence within the scene or to optimize the use of an internal cache.

[0109] Figure 21A shows a graphics core 2100 according to at least one embodiment. In at least one embodiment, the graphics core 2100 may be contained within the graphics processor 1510 in Figure 15. In at least one embodiment, the graphics core 2100 may be a unified shader core 2055A-2055N, as in Figure 20B. In at least one embodiment, the graphics core 2100 includes a shared instruction cache 2102, a texture unit 2118, and a cache / shared memory 2120, which are common to the execution resources within the graphics core 2100. In at least one embodiment, the graphics core 2100 may include multiple slices 2101A-2101N, or partitions for each core, and the graphics processor may include multiple instances of the graphics core 2100. Slices 2101A to 2101N may contain supporting logic including local instruction caches 2104A to 2104N, thread schedulers 2106A to 2106N, thread dispatchers 2108A to 2108N, and register sets 2110A to 2110N. In at least one embodiment, slices 2101A to 2101N may include a set of additional function units ("AFU") 2112A to 2112N, floating-point units ("FPU") 2114A to 2114N, integer arithmetic logic units ("ALU") 2116 to 2116N, address computational units ("ACU") 2113A to 2113N, double-precision floating-point units ("DPFPU") 2115A to 2115N, and matrix processing units ("MPU") 2117A to 2117N.

[0110] In at least one embodiment, the FPU2114A-2114N can perform single-precision (32-bit) and half-precision (16-bit) floating-point operations, and the DPFPU2115A-2115N can perform double-precision (64-bit) floating-point operations. In at least one embodiment, the ALU2116A-2116N can perform variable-precision integer operations with 8-bit, 16-bit, and 32-bit precision and may be configured for mixed-precision operations. In at least one embodiment, the MPU2117A-2117N may also be configured for mixed-precision matrix operations, including half-precision floating-point operations and 8-bit integer operations. In at least one embodiment, the MPU2117-2117N can perform various matrix operations to accelerate CUDA programs, including enabling support for accelerated general-purpose matrix-to-matrix multiplication ("GEMM"). In at least one embodiment, AFU2112A~2112N can perform additional logical operations not supported by the floating-point unit or integer unit, including trigonometric function operations (e.g., sine, cosine, etc.).

[0111] Figure 21B shows a general-purpose graphics processing unit ("GPGPU") 2130 according to at least one embodiment. In at least one embodiment, the GPGPU 2130 is highly parallel and suitable for deployment on a multi-chip module. In at least one embodiment, the GPGPU 2130 may be configured to allow highly parallel compute operations to be performed by an array of GPUs. In at least one embodiment, the GPGPU 2130 may be directly linked to other instances of the GPGPU 2130 to create a multi-GPU cluster to improve execution time for CUDA programs. In at least one embodiment, the GPGPU 2130 includes a host interface 2132 for enabling connectivity to a host processor. In at least one embodiment, the host interface 2132 is a PCIe interface. In at least one embodiment, the host interface 2132 may be a vendor-specific communication interface or communication fabric. In at least one embodiment, the GPGPU 2130 receives commands from the host processor and uses the global scheduler 2134 to distribute the execution threads associated with those commands across a set of compute clusters 2136A–2136H. In at least one embodiment, the compute clusters 2136A–2136H share a cache memory 2138. In at least one embodiment, the cache memory 2138 can act as a higher-level cache for the cache memory within the compute clusters 2136A–2136H.

[0112] In at least one embodiment, the GPGPU 2130 includes memory 2144A-2144B coupled to compute clusters 2136A-2136H via a set of memory controllers 2142A-2142B. In at least one embodiment, the memory 2144A-2144B may include various types of memory devices, including graphics random access memory such as DRAM or synchronous graphics random access memory ("SGRAM") including graphics double data rate ("GDDR") memory.

[0113] In at least one embodiment, compute clusters 2136A to 2136H each include a set of graphics cores, such as the graphics core 2100 in Figure 21A, and the set of graphics cores may include multiple types of integer and floating-point logic units capable of performing computational operations at varying precisions, including those suitable for computations related to CUDA programs. For example, in at least one embodiment, at least a subset of floating-point units in each of compute clusters 2136A to 2136H may be configured to perform 16-bit or 32-bit floating-point operations, and different subsets of floating-point units may be configured to perform 64-bit floating-point operations.

[0114] In at least one embodiment, multiple instances of GPGPU2130 may be configured to operate as a compute cluster. Compute clusters 2136A-2136H may implement any technically feasible communication techniques for synchronization and data exchange. In at least one embodiment, multiple instances of GPGPU2130 communicate via a host interface 2132. In at least one embodiment, GPGPU2130 includes an I / O hub 2139, which couples GPGPU2130 to a GPU link 2140 that enables direct connections to other instances of GPGPU2130. In at least one embodiment, the GPU link 2140 is coupled to a dedicated GPU-GPU bridge that enables communication and synchronization between multiple instances of GPGPU2130. In at least one embodiment, the GPU link 2140 is coupled to a high-speed interconnect for sending and receiving data to and from other GPGPU2130s or parallel processors. In at least one embodiment, multiple instances of the GPGPU 2130 are located on separate data processing systems and communicate via network devices accessible through the host interface 2132. In at least one embodiment, the GPU link 2140 may be configured to enable connectivity to a host processor in addition to, or as an alternative to, the host interface 2132. In at least one embodiment, the GPGPU 2130 may be configured to run CUDA programs.

[0115] Figure 22A shows a parallel processor 2200 according to at least one embodiment. In at least one embodiment, various components of the parallel processor 2200 may be implemented using one or more integrated circuit devices, such as a programmable processor, an application-specific integrated circuit ("ASIC"), or an FPGA.

[0116] In at least one embodiment, the parallel processor 2200 includes a parallel processing unit 2202. In at least one embodiment, the parallel processing unit 2202 includes an I / O unit 2204 that enables communication with other devices, including other instances of the parallel processing unit 2202. In at least one embodiment, the I / O unit 2204 may be directly connected to other devices. In at least one embodiment, the I / O unit 2204 connects to other devices via the use of a hub or switch interface, such as a memory hub 2205. In at least one embodiment, the connection between the memory hub 2205 and the I / O unit 2204 forms a communication link. In at least one embodiment, the I / O unit 2204 connects to a host interface 2206 and a memory crossbar 2216, the host interface 2206 receiving commands intended to perform processing operations and the memory crossbar 2216 receiving commands intended to perform memory operations.

[0117] In at least one embodiment, when the host interface 2206 receives a command buffer via the I / O unit 2204, the host interface 2206 can direct work operations to the front end 2208 to execute those commands. In at least one embodiment, the front end 2208 is coupled with a scheduler 2210, which is configured to distribute commands or other work items to the processing array 2212. In at least one embodiment, the scheduler 2210 ensures that the processing array 2212 is properly configured and enabled before tasks are distributed to the processing array 2212. In at least one embodiment, the scheduler 2210 is implemented via firmware logic running on a microcontroller. In at least one embodiment, the microcontroller-implemented scheduler 2210 can be configured to perform complex scheduling and work distribution operations at both coarse and fine granularity, enabling rapid preemption and context switching of threads running on the processing array 2212. In at least one embodiment, host software can prove workloads for scheduling on the processing array 2212 via one of a plurality of graphics processing doorbells. In at least one embodiment, the workload can then be automatically distributed across the processing array 2212 by scheduler 2210 logic within the microcontroller, which includes the scheduler 2210.

[0118] In at least one embodiment, the processing array 2212 may contain up to "N" clusters (e.g., cluster 2214A, cluster 2214B to cluster 2214N). In at least one embodiment, each cluster 2214A to 2214N of the processing array 2212 may execute a large number of concurrent threads. In at least one embodiment, the scheduler 2210 may allocate work to clusters 2214A to 2214N of the processing array 2212 using various scheduling and / or work distribution algorithms, which may vary depending on the workload arising for each type of program or computation. In at least one embodiment, scheduling may be handled dynamically by the scheduler 2210 or partially assisted by compiler logic during the compilation of program logic configured for execution by the processing array 2212. In at least one embodiment, different clusters 2214A to 2214N of the processing array 2212 may be allocated to process different types of programs or to perform different types of computations.

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

[0120] In at least one embodiment, the processing array 2212 is configured to perform parallel graphics processing operations. In at least one embodiment, the processing array 2212 may include, but is not limited to, additional logic to support the execution of such graphics processing operations, including texture sampling logic for performing texture operations, as well as tessellation logic and other vertex processing logic. In at least one embodiment, the processing array 2212 may be configured to execute graphics processing-related shader programs, such as vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. In at least one embodiment, the parallel processing unit 2202 may transfer data from system memory via the I / O unit 2204 for processing. In at least one embodiment, during processing, the transferred data may be stored in on-chip memory (e.g., parallel processor memory 2222) during processing and then written back to system memory.

[0121] In at least one embodiment, when the parallel processing unit 2202 is used to perform graphics processing, the scheduler 2210 may be configured to divide the processing workload into tasks of approximately equal size in order to better enable the distribution of graphics processing operations across multiple clusters 2214A to 2214N of the processing array 2212. In at least one embodiment, the parts of the processing array 2212 may be configured to perform different types of processing. For example, in at least one embodiment, to produce a rendered image for display, a first part may be configured to perform vertex shading and topology generation, a second part may be configured to perform tessellation and geometry shading, and a third part may be configured to perform pixel shading or other screen-space operations. In at least one embodiment, intermediate data produced by one or more of the clusters 2214A to 2214N may be stored in a buffer to enable the intermediate data to be transmitted between clusters 2214A to 2214N for further processing.

[0122] In at least one embodiment, the processing array 2212 may receive processing tasks to be executed via the scheduler 2210, which receives commands defining the processing tasks from the front end 2208. In at least one embodiment, the processing task may include an index of data to be processed, e.g., surface (patch) data, primitive data, vertex data, and / or pixel data, as well as state parameters and commands defining how the data should be processed (e.g., which program should be executed). In at least one embodiment, the scheduler 2210 may be configured to fetch the index corresponding to the task or to receive the index from the front end 2208. In at least one embodiment, the front end 2208 may be configured to ensure that the processing array 2212 is configured to a valid state before the workload specified by an incoming command buffer (e.g., a batch buffer, a push buffer, etc.) is started.

[0123] In at least one embodiment, each of one or more instances of the parallel processing unit 2202 can be coupled to a parallel processor memory 2222. In at least one embodiment, the parallel processor memory 2222 can be accessed via a memory crossbar 2216, which can receive memory requests from the processing array 2212 and the I / O unit 2204. In at least one embodiment, the memory crossbar 2216 can access the parallel processor memory 2222 via a memory interface 2218. In at least one embodiment, the memory interface 2218 can include a plurality of partition units (for example, partition unit 2220A, partition unit 2220B to partition unit 2220N), each of which can be coupled to a portion of the parallel processor memory 2222 (for example, a memory unit). In at least one embodiment, the number of partition units 2220A to 2220N is configured to be equal to the number of memory units, such that the first partition unit 2220A has a corresponding first memory unit 2224A, the second partition unit 2220B has a corresponding memory unit 2224B, and the nth partition unit 2220N has a corresponding nth memory unit 2224N. In at least one embodiment, the number of partition units 2220A to 2220N may not be equal to the number of memory devices.

[0124] In at least one embodiment, memory units 2224A-2224N may include various types of memory devices, including DRAM or graphics random access memory, such as SGRAM containing GDDR memory. In at least one embodiment, memory units 2224A-2224N may also include 3D stacked memory, including, but not limited to, high bandwidth memory ("HBM"). In at least one embodiment, to efficiently use the available bandwidth of parallel processor memory 2222, render targets such as frame buffers or texture maps may be stored across memory units 2224A-2224N, allowing partition units 2220A-2220N to write portions of each render target in parallel. In at least one embodiment, local instances of parallel processor memory 2222 may be excluded to favor a unified memory design that utilizes system memory in conjunction with local cache memory.

[0125] In at least one embodiment, any one of the clusters 2214A to 2214N of the processing array 2212 can process data that will be written to any of the memory units 2224A to 2224N in the parallel processor memory 2222. In at least one embodiment, the memory crossbar 2216 may be configured to forward the output of each cluster 2214A to 2214N to any partition unit 2220A to 2220N that can perform additional processing operations on the output, or to another cluster 2214A to 2214N. In at least one embodiment, each cluster 2214A to 2214N can communicate with the memory interface 2218 through the memory crossbar 2216 to read from or write to various external memory devices. In at least one embodiment, the memory crossbar 2216 has connections to a memory interface 2218 for communicating with the I / O unit 2204, as well as to a local instance of the parallel processor memory 2222, which allows processing units in different clusters 2214A-2214N to communicate with system memory or other memory that is not local to the parallel processing unit 2202. In at least one embodiment, the memory crossbar 2216 can use virtual channels to isolate traffic streams between clusters 2214A-2214N and partition units 2220A-2220N.

[0126] In at least one embodiment, multiple instances of the parallel processing unit 2202 may be provided on a single add-in card, or multiple add-in cards may be interconnected. In at least one embodiment, different instances of the parallel processing unit 2202 may be configured to interoperate even if the different instances have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences. For example, in at least one embodiment, some instances of the parallel processing unit 2202 may include higher-precision floating-point units than other instances. In at least one embodiment, a system incorporating one or more instances of the parallel processing unit 2202 or parallel processor 2200 may 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.

[0127] Figure 22B shows a processing cluster 2294 according to at least one embodiment. In at least one embodiment, processing cluster 2294 is contained within a parallel processing unit. In at least one embodiment, processing cluster 2294 is one of the processing clusters 2214A to 2214N in Figure 22. In at least one embodiment, processing cluster 2294 may be configured to run many threads in parallel, the term “thread” refers to a particular instance of a program running on a particular set of input data. In at least one embodiment, a single instruction, multiple data ("SIMD") instruction issuing technique is used to support the parallel execution of many threads without providing multiple independent instruction units. In at least one embodiment, a single instruction, multiple thread ("SIMT") technique is used to support the parallel execution of many threads with a common instruction unit configured to issue instructions to a set of processing engines within each processing cluster 2294.

[0128] In at least one embodiment, the operation of the processing cluster 2294 may be controlled via a pipeline manager 2232 that distributes processing tasks to SIMT parallel processors. In at least one embodiment, the pipeline manager 2232 receives instructions from the scheduler 2210 in Figure 22 and manages the execution of those instructions via the graphics multiprocessor 2234 and / or texture unit 2236. In at least one embodiment, the graphics multiprocessor 2234 is an exemplary instance of a SIMT parallel processor. However, in at least one embodiment, various types of SIMT parallel processors of different architectures may be included within the processing cluster 2294. In at least one embodiment, one or more instances of the graphics multiprocessor 2234 may be included within the processing cluster 2294. In at least one embodiment, the graphics multiprocessor 2234 may process data, and a data crossbar 2240 may be used to distribute the processed data to one of several possible destinations, including other shader units. In at least one embodiment, the pipeline manager 2232 can facilitate the distribution of processed data by specifying destinations for the processed data that will be distributed via the data crossbar 2240.

[0129] In at least one embodiment, each graphics multiprocessor 2234 within the processing cluster 2294 may contain an identical set of function execution logic (e.g., arithmetic logic units, load / store units ("LSU"), etc.). In at least one embodiment, the function execution logic may be configured in a pipelined manner in which a new instruction can be issued before the previous instruction is completed. In at least one embodiment, the function execution logic supports a variety of operations, including integer and floating-point arithmetic, comparison operations, Boolean operations, bit shifts, and the computation of various algebraic functions. In at least one embodiment, the same function unit hardware may be utilized to perform different operations, and any combination of function units may exist.

[0130] In at least one embodiment, instructions sent to processing cluster 2294 constitute a thread. In at least one embodiment, a set of threads running across a set of parallel processing engines is a thread group. In at least one embodiment, a thread group executes a program for different input data. In at least one embodiment, each thread in a thread group may be assigned to a different processing engine in the graphics multiprocessor 2234. In at least one embodiment, a thread group may contain fewer threads than the number of processing engines in the graphics multiprocessor 2234. In at least one embodiment, when a thread group contains fewer threads than the number of processing engines, one or more of the processing engines may be idle during the cycle in which the thread group is being processed. In at least one embodiment, a thread group may also contain more threads than the number of processing engines in the graphics multiprocessor 2234. In at least one embodiment, when a thread group contains more threads than the number of processing engines in the graphics multiprocessor 2234, processing may be carried out over consecutive clock cycles. In at least one embodiment, multiple thread groups may run simultaneously on the graphics multiprocessor 2234.

[0131] In at least one embodiment, the graphics multiprocessor 2234 includes internal cache memory for performing load and store operations. In at least one embodiment, the graphics multiprocessor 2234 may omit its internal cache and use cache memory in the processing cluster 2294 (e.g., L1 cache 2248). In at least one embodiment, each graphics multiprocessor 2234 also has access to a Level 2 ("L2") cache in a partition unit (e.g., partition units 2220A-2220N in Figure 22A), and these L2 caches may be shared among all processing clusters 2294 and used to transfer data between threads. In at least one embodiment, the graphics multiprocessor 2234 may also have access to off-chip global memory, which may include one or more of the local parallel processor memory and / or system memory. In at least one embodiment, any memory outside the parallel processing unit 2202 may be used as global memory. In at least one embodiment, the processing cluster 2294 includes multiple instances of a graphics multiprocessor 2234, which can share common instructions and data, and which can be stored in an L1 cache 2248.

[0132] In at least one embodiment, each processing cluster 2294 may include an MMU 2245 configured to map virtual addresses to physical addresses. In at least one embodiment, one or more instances of the MMU 2245 may reside within the memory interface 2218 in Figure 22. In at least one embodiment, the MMU 2245 includes a set of page table entries ("PTEs") used to map virtual addresses to physical addresses of tiles and optionally cache line indices. In at least one embodiment, the MMU 2245 may include an address translation lookaside buffer ("TLB") or cache, which may reside within the graphics multiprocessor 2234 or the L1 cache 2248 or the processing cluster 2294. In at least one embodiment, physical addresses are processed to distribute surface data access locality and enable efficient request interleaving between partition units. In at least one embodiment, a cache line index may be used to determine whether a request for a cache line is a hit or a miss.

[0133] In at least one embodiment, the processing cluster 2294 may be configured such that each graphics multiprocessor 2234 is coupled to a texture unit 2236 for performing 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 from an internal texture L1 cache (not shown) or from an L1 cache within the graphics multiprocessor 2234 and, if necessary, fetched from an L2 cache, local parallel processor memory, or system memory. In at least one embodiment, each graphics multiprocessor 2234 outputs the processed tasks to a data crossbar 2240 to provide the processed tasks to another processing cluster 2294 for further processing, or stores the processed tasks in an L2 cache, local parallel processor memory, or system memory via a memory crossbar 2216. In at least one embodiment, a pre-raster operation unit ("pre-ROP") 2242 is configured to receive data from a graphics multiprocessor 2234 and direct the data to a ROP unit, which may be located with partition units as described herein (for example, partition units 2220A-2220N in Figure 22). In at least one embodiment, the pre-ROP 2242 can perform optimizations for color blending, organize pixel color data, and perform address translation.

[0134] Figure 22C shows a graphics multiprocessor 2296 according to at least one embodiment. In at least one embodiment, the graphics multiprocessor 2296 is the graphics multiprocessor 2234 in Figure 22B. In at least one embodiment, the graphics multiprocessor 2296 is coupled with a pipeline manager 2232 of the processing cluster 2294. In at least one embodiment, the graphics multiprocessor 2296 has an execution pipeline that includes, but is not limited to, an instruction cache 2252, an instruction unit 2254, an address mapping unit 2256, a register file 2258, one or more GPGPU cores 2262, and one or more LSUs 2266. The GPGPU cores 2262 and LSUs 2266 are coupled with cache memory 2272 and shared memory 2270 via a memory and cache interconnect 2268.

[0135] In at least one embodiment, the instruction cache 2252 receives a stream of instructions to be executed from the pipeline manager 2232. In at least one embodiment, the instructions are cached in the instruction cache 2252 and dispatched for execution by the instruction unit 2254. In at least one embodiment, the instruction unit 2254 can dispatch the instructions as a thread group (e.g., a warp), where each thread in the thread group is assigned to a different execution unit within the GPGPU core 2262. In at least one embodiment, the instructions can access either the local, shared, or global address space by specifying an address within the unified address space. In at least one embodiment, the address mapping unit 2256 may be used to translate addresses in the unified address space to individual memory addresses that can be accessed by the LSU 2266.

[0136] In at least one embodiment, the register file 2258 provides a set of registers to the functional units of the graphics multiprocessor 2296. In at least one embodiment, the register file 2258 provides temporary storage for operands connected to the data paths of the functional units of the graphics multiprocessor 2296 (e.g., GPGPU core 2262, LSU 2266). In at least one embodiment, the register file 2258 is divided among the functional units such that each functional unit is allocated a dedicated portion of the register file 2258. In at least one embodiment, the register file 2258 is divided among different thread groups being executed by the graphics multiprocessor 2296.

[0137] In at least one embodiment, each GPGPU core 2262 may include an FPU and / or an integer ALU used to execute instructions of the graphics multiprocessor 2296. The GPGPU cores 2262 may have similar or different architectures. In at least one embodiment, a first part of the GPGPU core 2262 includes a single-precision FPU and an integer ALU, and a second part of the GPGPU core 2262 includes a double-precision FPU. In at least one embodiment, the FPU may implement the IEEE 754-2008 standard for floating-point arithmetic or enable variable-precision floating-point arithmetic. In at least one embodiment, the graphics multiprocessor 2296 may additionally include one or more fixed-function units or special-function units for performing specific functions such as rectangular copy operations or pixel blending operations. In at least one embodiment, one or more of the GPGPU cores 2262 may also include fixed or special-function logic.

[0138] In at least one embodiment, the GPGPU core 2262 includes SIMD logic capable of executing a single instruction for multiple data sets. In at least one embodiment, the GPGPU core 2262 can physically execute SIMD4, SIMD8, and SIMD16 instructions and logically execute SIMD1, SIMD2, and SIMD32 instructions. In at least one embodiment, SIMD instructions for the GPGPU core 2262 may be generated at compile time by the shader compiler or automatically generated when running a program written and compiled for a single program multiple data ("SPMD") or SIMT architecture. In at least one embodiment, multiple threads of a program configured for a SIMT execution model may be executed via a single SIMD instruction. For example, in at least one embodiment, eight SIMT threads performing the same or similar operation may be executed in parallel via a single SIMD8 logic unit.

[0139] In at least one embodiment, the memory and cache interconnect 2268 is an interconnect network connecting each functional unit of the graphics multiprocessor 2296 to the register file 2258 and shared memory 2270. In at least one embodiment, the memory and cache interconnect 2268 is a crossbar interconnect that enables the LSU 2266 to implement load and store operations between the shared memory 2270 and the register file 2258. In at least one embodiment, the register file 2258 can operate at the same frequency as the GPGPU core 2262, and therefore data transfer between the GPGPU core 2262 and the register file 2258 is very low latency. In at least one embodiment, the shared memory 2270 may be used to enable communication between threads running on functional units within the graphics multiprocessor 2296. In at least one embodiment, the cache memory 2272 may be used as a data cache to cache texture data communicated between functional units and texture unit 2236, for example. In at least one embodiment, shared memory 2270 may also be used as a cached managed program. In at least one embodiment, a thread running on the GPGPU core 2262 may programmatically store data in shared memory in addition to automatically cached data stored in cache memory 2272.

[0140] In at least one embodiment, a parallel processor or GPGPU as described herein is communicatively coupled to a host / processor core to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general-purpose GPU (GPGPU) functions. In at least one embodiment, the GPU may be communicatively coupled to a host processor / core via a bus or other interconnect (e.g., a high-speed interconnect such as PCIe or NVLink). In at least one embodiment, the GPU may be integrated as a core on the same package or chip and communicatively coupled to the core via a processor bus / interconnection located within the package or chip. In at least one embodiment, regardless of how the GPU is connected, the processor core may allocate work to the GPU in the form of a sequence of commands / instructions contained in the WD. In at least one embodiment, the GPU then uses dedicated circuit elements / logic to efficiently process these commands / instructions.

[0141] Figure 23 shows a graphics processor 2300 according to at least one embodiment. In at least one embodiment, the graphics processor 2300 includes a ring interconnect 2302, a pipeline front end 2304, a media engine 2337, and graphics cores 2380A to 2380N. In at least one embodiment, the ring interconnect 2302 connects the graphics processor 2300 to other graphics processors or other processing units including one or more general-purpose processor cores. In at least one embodiment, the graphics processor 2300 is one of many processors incorporated within a multicore processing system.

[0142] In at least one embodiment, the graphics processor 2300 receives batches of commands via a ring interconnect 2302. In at least one embodiment, incoming commands are interpreted by a command streamer 2303 in a pipeline front end 2304. In at least one embodiment, the graphics processor 2300 includes scalable execution logic for performing 3D geometry processing and media processing via one or more graphics cores 2380A-2380N. In at least one embodiment, for 3D geometry processing commands, the command streamer 2303 feeds the commands to a geometry pipeline 2336. In at least one embodiment, for at least some media processing commands, the command streamer 2303 feeds the commands to a video front end 2334, which then couples with a media engine 2337. In at least one embodiment, the media engine 2337 includes a Video Quality Engine ("VQE") 2330 for video and image post-processing and a multi-format encode / decode ("MFX") engine 2333 for providing hardware-accelerated media data encoding and decoding. In at least one embodiment, the geometry pipeline 2336 and the media engine 2337 each generate execution threads for thread execution resources provided by at least one graphics core 2380A.

[0143] In at least one embodiment, the graphics processor 2300 includes scalable thread execution resources featuring modular graphics cores 2380A-2380N (sometimes called core slices), each having multiple subcores 2350A-550N, 2360A-2360N, each having multiple subcores (sometimes called core sub-slices). In at least one embodiment, the graphics processor 2300 may have any number of graphics cores 2380A-2380N. In at least one embodiment, the graphics processor 2300 includes a graphics core 2380A having at least a first subcore 2350A and a second subcore 2360A. In at least one embodiment, the graphics processor 2300 is a low-power processor having a single subcore (e.g., subcore 2350A). In at least one embodiment, the graphics processor 2300 includes a plurality of graphics cores 2380A to 2380N, each including a first set of subcores 2350A to 2350N and a second set of subcores 2360A to 2360N. In at least one embodiment, each subcore in the first set of subcores 2350A to 2350N includes at least a first set of execution units ("EUs") 2352A to 2352N and media / texture samplers 2354A to 2354N. In at least one embodiment, each subcore in the second set of subcores 2360A to 2360N includes at least a second set of execution units 2362A to 2362N and samplers 2364A to 2364N. In at least one embodiment, each sub-core 2350A-2350N, 2360A-2360N shares a set of shared resources 2370A-2370N. In at least one embodiment, the shared resource 2370 includes shared cache memory and pixel operation logic.

[0144] Figure 24 shows a processor 2400 according to at least one embodiment. In at least one embodiment, the processor 2400 may include, but is not limited to, logic circuits for executing instructions. In at least one embodiment, the processor 2400 may execute instructions including x86 instructions, AMR instructions, special instructions for ASICs, etc. In at least one embodiment, the processor 2410 may include registers for storing packed data, such as 64-bit wide MMX registers in a microprocessor enabled by MMX™ technology from Intel Corporation, Santa Clara, California. In at least one embodiment, MMX registers available in both integer and floating-point formats may operate with packed data elements accompanied by SIMD and streaming SIMD extension ("SSE") instructions. In at least one embodiment, 128-bit wide XMM registers related to SSE2, SSE3, SSE4, AVX, or higher technologies (collectively referred to as "SSEx") may hold such packed data operands. In at least one embodiment, the processor 2410 may implement instructions to accelerate the CUDA program.

[0145] In at least one embodiment, the processor 2400 includes an in-order front-end ("front-end") 2401 for fetching instructions to be executed and preparing instructions to be used later in the processor pipeline. In at least one embodiment, the front-end 2401 may include several units. In at least one embodiment, an instruction prefetcher 2426 fetches instructions from memory and feeds the instructions to an instruction decoder 2428, which decodes or interprets the instructions. For example, in at least one embodiment, the instruction decoder 2428 decodes the received instruction into one or more operations called "microinstructions" or "microoperations" (also called "microops" or "uops") for execution. In at least one embodiment, the instruction decoder 2428 parses the instruction into opcodes and corresponding data and control fields that may be used by the microarchitecture to perform the operation. In at least one embodiment, the trace cache 2430 may assemble the decoded uops into a program-order sequence or trace in the uop queue 2434 for execution. In at least one embodiment, when the trace cache 2430 encounters a complex instruction, the microcode ROM 2432 provides the uops necessary to complete the operation.

[0146] In at least one embodiment, some instructions can be converted to a single micro-op, while others require several micro-ops to complete the entire operation. In at least one embodiment, if five or more micro-ops are required to complete the instruction, the instruction decoder 2428 may access the microcode ROM 2432 to execute the instruction. In at least one embodiment, an instruction may be decoded into a small number of micro-ops for processing in the instruction decoder 2428. In at least one embodiment, if several micro-ops are required to accomplish the operation, the instruction may be stored in the microcode ROM 2432. In at least one embodiment, the trace cache 2430 refers to the entry-point programmable logic array ("PLA") to determine the correct microinstruction pointer for reading the microcode sequence in order to complete one or more instructions from the microcode ROM 2432. In at least one embodiment, after the microcode ROM 2432 has finished sequencing the micro-ops for an instruction, the machine's front-end 2401 may resume fetching micro-ops from the trace cache 2430.

[0147] In at least one embodiment, the out-of-order execution engine ("out-of-order engine") 2403 may prepare instructions for execution. In at least one embodiment, the out-of-order execution logic has several buffers to smooth and reorder the flow of instructions in order to optimize performance as instructions move down the pipeline and are scheduled for execution. The out-of-order execution engine 2403 includes, but is not limited to, an allocator / register renamer 2440, a memory uop queue 2442, an integer / floating-point uop queue 2444, a memory scheduler 2446, a fast scheduler 2402, a slow / general-purpose floating-point scheduler ("slow / general-purpose FP (floating point) scheduler") 2404, and a simple floating-point scheduler ("simple FP scheduler") 2406. In at least one embodiment, the fast scheduler 2402, the slow / general-purpose floating-point scheduler 2404, and the simple floating-point scheduler 2406 are collectively referred to herein as "uop schedulers 2402, 2404, and 2406." The allocator / register renamer 2440 allocates the machine buffers and resources that each uop needs to run. In at least one embodiment, the allocator / register renamer 2440 renames logical registers upon entry into the register file. In at least one embodiment, the allocator / register renamer 2440 also allocates entries for each uop in one of two uop queues, namely the memory uop queue 2442 for memory operations and the integer / floating-point uop queue 2444 for non-memory operations, prior to the memory scheduler 2446 and the uop schedulers 2402, 2404, 2406. In at least one embodiment, the uop schedulers 2402, 2404, 2406 determine when a uop is ready to execute based on whether their dependent input register operand sources are prepared and the availability of execution resources required by the uop to complete their operations.In at least one embodiment, the high-speed scheduler 2402 of at least one embodiment may schedule every half of the main clock cycle, while the slow / general-purpose floating-point scheduler 2404 and the simple floating-point scheduler 2406 may schedule once per main processor clock cycle. In at least one embodiment, the uop schedulers 2402, 2404, and 2406 arbitrate dispatch ports to schedule uops for execution.

[0148] In at least one embodiment, the execution block 2411 includes, but is not limited to, an integer register file / bypass network 2408, a floating-point register file / bypass network ("FP register file / bypass network") 2410, address generation units ("AGUs") 2412 and 2414, high-speed ALUs 2416 and 2418, low-speed ALU 2420, a floating-point ALU ("FP") 2422, and a floating-point move unit ("FP move") 2424. In at least one embodiment, the integer register file / bypass network 2408 and the floating-point register file / bypass network 2410 are also referred to herein as "register files 2408, 2410". In at least one embodiment, AGU2412 and 2414, high-speed ALU2416 and 2418, low-speed ALU2420, floating-point ALU2422, and floating-point movement unit 2424 are also referred to herein as “execution units 2412, 2414, 2416, 2418, 2420, 2422, and 2424”. In at least one embodiment, the execution block may include, but is not limited to, any number and type of register files (including zero), bypass networks, address generation units, and execution units in any combination.

[0149] In at least one embodiment, register files 2408, 2410 may be located between the uop schedulers 2402, 2404, 2406 and the execution units 2412, 2414, 2416, 2418, 2420, 2422, and 2424. In at least one embodiment, the integer register file / bypass network 2408 performs integer arithmetic. In at least one embodiment, the floating-point register file / bypass network 2410 performs floating-point arithmetic. In at least one embodiment, each of the register files 2408, 2410 may include, but not limited to, a bypass network that can bypass or forward recently completed results that have not yet been written to the register file to a new dependent uop. In at least one embodiment, the register files 2408, 2410 may communicate data with each other. In at least one embodiment, the integer register file / bypass network 2408 may include, but not limited to, two separate register files: one register file for low-order 32-bit data and a second register file for high-order 32-bit data. In at least one embodiment, since floating-point instructions typically have operands that are 64 to 128 bits wide, the floating-point register file / bypass network 2410 may include, but not limited to, 128-bit wide entries.

[0150] In at least one embodiment, execution units 2412, 2414, 2416, 2418, 2420, 2422, and 2424 may execute instructions. In at least one embodiment, register files 2408 and 2410 store integer and floating-point data operand values ​​that microinstructions need to execute. In at least one embodiment, processor 2400 may include, but not limited to, any number and combination of execution units 2412, 2414, 2416, 2418, 2420, 2422, and 2424. In at least one embodiment, floating-point ALU 2422 and floating-point movement unit 2424 may perform floating-point, MMX, SIMD, AVX, and SSE, or other operations. In at least one embodiment, the floating-point ALU2422 may include, but is not limited to, 64-bit floating-point dividers for performing division, square root, and remainder micro-ops. In at least one embodiment, instructions involving floating-point values ​​may be handled by floating-point hardware. In at least one embodiment, ALU operations may be passed to the high-speed ALU2416, 2418. In at least one embodiment, the high-speed ALU2416, 2418 may perform high-speed operations with an effective latency of half a clock cycle. In at least one embodiment, the slow ALU2420 may include, but is not limited to, integer execution hardware for high-latency types of operations such as multipliers, shifts, flag logic, and branching, so that most complex integer operations are passed to the slow ALU2420. In at least one embodiment, memory load / store operations may be performed by the ALU2412, 2414. In at least one embodiment, the high-speed ALU2416, high-speed ALU2418, and low-speed ALU2420 may perform integer arithmetic with 64-bit data operands. In at least one embodiment, the high-speed ALU2416, high-speed ALU2418, and low-speed ALU2420 may be implemented to support various data bit sizes, including 16, 32, 128, 256, and so on. In at least one embodiment, the floating-point ALU2422 and floating-point movement unit 2424 may be implemented to support various operands having various bit widths.In at least one embodiment, the floating-point ALU 2422 and the floating-point movement unit 2424 may operate with 128-bit wide packed data operands in conjunction with SIMD and multimedia instructions.

[0151] In at least one embodiment, the uop schedulers 2402, 2404, and 2406 dispatch dependent operations before the parent load has finished executing. In at least one embodiment, since uops may be speculatively scheduled and executed in processor 2400, processor 2400 may also include logic for handling memory misses. In at least one embodiment, if a data load misses in the data cache, there may be ongoing dependent operations in the pipeline that have passed the scheduler with temporarily inaccurate data. In at least one embodiment, a replay mechanism tracks and redelivers instructions that use inaccurate data. In at least one embodiment, dependent operations may need to be replayed, allowing independent operations to complete. In at least one embodiment, the scheduler and replay mechanism of at least one embodiment of the processor may also be designed to capture instruction sequences for text string comparison operations.

[0152] In at least one embodiment, the term “register” may refer to an onboard processor storage location that can be used as part of an instruction to identify an operand. In at least one embodiment, a register may be something that can be accessed from outside the processor (from the programmer’s perspective). In at least one embodiment, a register may not be limited to a particular type of circuit. Rather, in at least one embodiment, a register may store data, provide data, and perform the functions described herein. In at least one embodiment, the registers described herein may be implemented by circuit elements within the processor using any number of different techniques, such as dedicated physical registers, dynamically allocated physical registers using register renaming, or a combination of dedicated and dynamically allocated physical registers. In at least one embodiment, an integer register stores 32-bit integer data. The register file in at least one embodiment also includes eight multimedia SIMD registers for packed data.

[0153] Figure 25 shows a processor 2500 according to at least one embodiment. In at least one embodiment, the processor 2500 includes, but is not limited to, one or more processor cores ("cores") 2502A to 2502N, an integrated memory controller 2514, and an integrated graphics processor 2508. In at least one embodiment, the processor 2500 may include additional cores up to an additional processor core 2502N, represented by a dashed box. In at least one embodiment, each of the processor cores 2502A to 2502N includes one or more internal cache units 2504A to 2504N. In at least one embodiment, each processor core also has access to one or more shared cache units 2506.

[0154] In at least one embodiment, the internal cache units 2504A–2504N and the shared cache unit 2506 represent a cache memory hierarchy within the processor 2500. In at least one embodiment, the cache memory units 2504A–2504N may include at least one level of instruction and data cache within each processor core, and one or more levels of shared intermediate level caches such as L2, L3, level 4 ("L4"), or other levels of cache, where the highest level cache prior to external memory is classified as LLC. In at least one embodiment, cache coherency logic maintains coherency among the various cache units 2506 and 2504A–2504N.

[0155] In at least one embodiment, the processor 2500 may also include a set of one or more bus controller units 2516 and a system agent core 2510. In at least one embodiment, one or more bus controller units 2516 manage a set of peripheral buses, such as one or more PCI or PCI Express buses. In at least one embodiment, the system agent core 2510 provides management functionality for various processor components. In at least one embodiment, the system agent core 2510 includes one or more integrated memory controllers 2514 for managing access to various external memory devices (not shown).

[0156] In at least one embodiment, one or more of the processor cores 2502A to 2502N include support for simultaneous multithreading. In at least one embodiment, the system agent core 2510 includes components for coordinating and operating the processor cores 2502A to 2502N during multithreaded processing. In at least one embodiment, the system agent core 2510 may additionally include a power control unit ("PCU"), the PCU including logic and components for regulating the power states of one or more of the processor cores 2502A to 2502N and the graphics processor 2508.

[0157] In at least one embodiment, the processor 2500 additionally includes a graphics processor 2508 for performing graphics processing operations. In at least one embodiment, the graphics processor 2508 is coupled with a system agent core 2510 which includes a shared cache unit 2506 and one or more integrated memory controllers 2514. In at least one embodiment, the system agent core 2510 also includes a display controller 2511 for driving graphics processor outputs to one or more coupled displays. In at least one embodiment, the display controller 2511 may also be a separate module coupled with the graphics processor 2508 via at least one interconnection, or it may be incorporated within the graphics processor 2508.

[0158] In at least one embodiment, a ring-based interconnect unit 2512 is used to connect the internal components of the processor 2500. In at least one embodiment, alternative interconnect units such as point-to-point interconnects, switching interconnects, or other techniques may be used. In at least one embodiment, the graphics processor 2508 is connected to the ring interconnect 2512 via an I / O link 2513.

[0159] In at least one embodiment, I / O link 2513 represents at least one of several types of I / O interconnects, including on-package I / O interconnects that facilitate communication between various processor components and high-performance embedded memory modules 2518, such as eDRAM modules. In at least one embodiment, each of the processor cores 2502A to 2502N and the graphics processor 2508 use the embedded memory module 2518 as a shared LLC.

[0160] In at least one embodiment, the processor cores 2502A to 2502N are homogeneous cores that execute a common instruction set architecture. In at least one embodiment, the processor cores 2502A to 2502N are heterogeneous in terms of ISA, where one or more of the processor cores 2502A to 2502N execute a common instruction set, and one or more other cores of the processor cores 2502A to 2502N execute a subset of the common instruction set or a different instruction set. In at least one embodiment, the processor cores 2502A to 2502N are heterogeneous in terms of microarchitecture, where one or more cores with relatively high power consumption are coupled with one or more cores with lower power consumption. In at least one embodiment, the processor 2500 may be implemented on one or more chips or as an SoC integrated circuit.

[0161] Figure 26 shows a graphics processor core 2600 according to at least one embodiment described. In at least one embodiment, the graphics processor core 2600 is contained within a graphics core array. In at least one embodiment, the graphics processor core 2600, sometimes referred to as a core slice, may be one or more graphics cores in a modular graphics processor. In at least one embodiment, the graphics processor core 2600 is an example of a single graphics core slice, and the graphics processor described herein may include multiple graphics core slices based on target power and performance envelopes. In at least one embodiment, each graphics core 2600 may include a fixed-function block 2630 coupled with a plurality of sub-cores 2601A to 2601F, also referred to as sub-slices, which include modular blocks of general-purpose and fixed-function logic.

[0162] In at least one embodiment, the fixed-function block 2630 includes a geometry / fixed-function pipeline 2636 which can be shared by all subcores in the graphics processor 2600, for example, in a lower-performance and / or lower-power graphics processor implementation. In at least one embodiment, the geometry / fixed-function pipeline 2636 includes a 3D fixed-function pipeline, a video front-end unit, a thread spawner and thread dispatcher, and a unified return buffer manager that manages a unified return buffer.

[0163] In at least one embodiment, the fixed-function block 2630 also includes a graphics SoC interface 2637, a graphics microcontroller 2638, and a media pipeline 2639. The graphics SoC interface 2637 provides an interface between the graphics core 2600 and other processor cores within the SoC integrated circuit. In at least one embodiment, the graphics microcontroller 2638 is a programmable sub-processor configurable to manage various functions of the graphics processor 2600, including thread dispatch, scheduling, and preemption. In at least one embodiment, the media pipeline 2639 includes logic to facilitate decoding, encoding, preprocessing, and / or postprocessing of multimedia data, including image and video data. In at least one embodiment, the media pipeline 2639 implements media operations via requests to compute logic or sampling logic within sub-cores 2601-2601F.

[0164] In at least one embodiment, the SoC interface 2637 enables the graphics core 2600 to communicate with a general-purpose application processor core (e.g., a CPU) and / or other components within the SoC, the other components within the SoC include memory hierarchy elements such as shared LLC memory, system RAM, and / or embedded on-chip or on-package DRAM. In at least one embodiment, the SoC interface 2637 can also enable communication with fixed-function devices within the SoC, such as a camera imaging pipeline, and enable and / or implement the use of a global memory atomic that may be shared between the graphics core 2600 and the CPU within the SoC. In at least one embodiment, the SoC interface 2637 can also implement power management controls for the graphics core 2600 and enable facilitating between the clock domain of the graphics core 2600 and other clock domains within the SoC. In at least one embodiment, the SoC interface 2637 enables the reception of command buffers from a command streamer and a global thread dispatcher configured to provide commands and instructions to each of the one or more graphics cores in the graphics processor. In at least one embodiment, commands and instructions can be dispatched to the media pipeline 2639 when media operations should be performed, or to the geometry and fixed-function pipelines (e.g., geometry and fixed-function pipeline 2636, geometry and fixed-function pipeline 2614) when graphics processing operations should be performed.

[0165] In at least one embodiment, the graphics microcontroller 2638 may be configured to perform various scheduling and management tasks for the graphics core 2600. In at least one embodiment, the graphics microcontroller 2638 may perform graphics and / or calculate workload scheduling for various graphics parallel engines in the execution unit (EU) arrays 2602A-2602F and 2604A-2604F within the sub-cores 2601A-2601F. In at least one embodiment, host software running on the CPU core of the SoC, including the graphics core 2600, may submit a workload to one of several graphics processor doorbells, which then invokes scheduling operations for the appropriate graphics engine. In at least one embodiment, the scheduling operation includes determining which workload should run next, submitting the workload to a command streamer, preempting existing workloads running on the engine, monitoring the progress of the workload, and notifying the host software when the workload is complete. In at least one embodiment, the graphics microcontroller 2638 can also facilitate a low-power or idle state for the graphics core 2600, providing the graphics core 2600 with the ability to save and restore registers within the graphics core 2600 during transitions to low-power states, independently of the operating system and / or graphics driver software on the system.

[0166] In at least one embodiment, the graphics core 2600 may have up to N modular subcores, more or less than the subcores 2601A-2601F shown. For each set of N subcores, in at least one embodiment, the graphics core 2600 may also include shared function logic 2610, shared and / or cache memory 2612, geometry / fixed function pipeline 2614, and additional fixed function logic 2616 for accelerating various graphics and calculating processing operations. In at least one embodiment, the shared function logic 2610 may include logic units (e.g., sampler, mathematical, and / or inter-thread communication logic) that can be shared by each of the N subcores in the graphics core 2600. The shared and / or cache memory 2612 may be LLC for the N subcores 2601A-2601F in the graphics core 2600 and may also function as shared memory accessible by multiple subcores. In at least one embodiment, the geometry / fixed function pipeline 2614 may be included in place of the geometry / fixed function pipeline 2636 within the fixed function block 2630 and may include the same or similar logical units.

[0167] In at least one embodiment, the graphics core 2600 includes an additional fixed-function logic 2616 which can include various fixed-function acceleration logic for use by the graphics core 2600. In at least one embodiment, the additional fixed-function logic 2616 includes an additional geometry pipeline for use in position-only shading. In position-only shading, there are at least two geometry pipelines: a full geometry pipeline in geometry / fixed-function pipeline 2616, 2636, and a cull pipeline, the cull pipeline being an additional geometry pipeline which may be contained within the additional fixed-function logic 2616. In at least one embodiment, the cull pipeline is a reduced version of the full geometry pipeline. In at least one embodiment, the full pipeline and the cull pipeline can run different instances of the application, each instance having a separate context. In at least one embodiment, position-only shading can hide long cull runs of truncated triangles, which allows shading to complete faster in some instances. For example, in at least one embodiment, the sorting pipeline fetches and shades the vertex position attributes without performing rasterization and rendering of pixels into a frame buffer, so that the sorting pipeline logic within the additional fixed-function logic 2616 can run the position shader in parallel with the main application and produce critical results faster overall than the full pipeline. In at least one embodiment, the sorting pipeline can use the generated critical results to calculate visibility information for all triangles, regardless of whether those triangles are sorted or not. In at least one embodiment, the full pipeline (sometimes called the replay pipeline in this instance) can consume the visibility information to skip the sorted triangles and shade only the visible triangles, which are then passed to the rasterization phase.

[0168] In at least one embodiment, additional fixed-function logic 2616 may also include general-purpose processing acceleration logic, such as fixed-function matrix multiplication logic, to accelerate CUDA programs.

[0169] In at least one embodiment, each graphics sub-core 2601A-2601F includes a set of execution resources which may be used to perform graphics operations, media operations, and compute operations in response to requests from the graphics pipeline, media pipeline, or shader program. In at least one embodiment, the graphics sub-cores 2601A-2601F include a plurality of EU arrays 2602A-2602F, 2604A-2604F, thread dispatch and inter-thread communication ("TD / IC") logic 2603A-2603F, 3D (e.g., texture) samplers 2605A-2605F, media samplers 2606A-2606F, shader processors 2607A-2607F, and shared local memory ("SLM") 2608A-2608F. EU arrays 2602A-2602F and 2604A-2604F each include multiple execution units, which are GPGPUs capable of performing floating-point and integer / fixed-point logical operations in graphics operations, media operations, or compute operations, including graphics, media, or compute shader programs. In at least one embodiment, TD / IC logic 2603A-2603F perform local thread dispatch and thread control operations for execution units within a subcore, facilitating communication between threads running on the execution units in the subcore. In at least one embodiment, 3-D samplers 2605A-2605F can read texture or other 3D graphics-related data into memory. In at least one embodiment, the 3D sampler can read texture data in different ways based on a configured sample state and texture format associated with a given texture. In at least one embodiment, media samplers 2606A to 2606F can perform similar reading operations based on the type and format associated with the media data.In at least one embodiment, each graphics sub-core 2601A-2601F may alternatively include a unified 3D and media sampler. In at least one embodiment, threads running on execution units within each of the sub-cores 2601A-2601F may utilize shared local memory 2608A-2608F within each sub-core to enable threads running within a thread group to run using a common pool of on-chip memory.

[0170] Figure 27 shows a parallel processing unit ("PPU") 2700 according to at least one embodiment. In at least one embodiment, the PPU 2700 consists of machine-readable code that, when executed by the PPU 2700, causes the PPU 2700 to perform some or all of the processes and techniques described herein. In at least one embodiment, the PPU 2700 is a multithreaded processor, which is implemented on one or more integrated circuit devices and utilizes multithreading as a latency-hiding technique designed to process computer-readable instructions (also called machine-readable instructions or simply instructions) in parallel on multiple threads. In at least one embodiment, a thread refers to a thread of execution and is an instantiation of a set of instructions configured to be executed by the PPU 2700. In at least one embodiment, the PPU2700 is a GPU configured to implement a graphics rendering pipeline for processing three-dimensional ("3D") graphics data to generate two-dimensional ("2D") image data for display on a display device such as an LCD device. In at least one embodiment, the PPU2700 is used to perform calculations such as linear algebra and machine learning operations. Figure 27 shows an exemplary parallel processor for illustrative purposes only and should be interpreted as a non-limiting example of a processor architecture that may be implemented in at least one embodiment.

[0171] In at least one embodiment, one or more PPUs 2700 are configured to accelerate high-performance computing ("HPC"), data center, and machine learning applications. In at least one embodiment, one or more PPUs 2700 are configured to accelerate CUDA programs. In at least one embodiment, the PPU 2700 includes, but is not limited to, an I / O unit 2706, a front-end unit 2710, a scheduler unit 2712, a work distribution unit 2714, a hub 2716, a crossbar ("X-bar") 2720, one or more general-purpose processing clusters ("GPC") 2718, and one or more partition units ("memory partition units") 2722. In at least one embodiment, the PPU 2700 is connected to a host processor or other PPU 2700 via one or more high-speed GPU interconnects ("GPU interconnects") 2708. In at least one embodiment, the PPU 2700 is connected to a host processor or other peripheral devices via a system bus or interconnect 2702. In at least one embodiment, the PPU 2700 is connected to local memory comprising one or more memory devices ("memory") 2704. In at least one embodiment, the memory device 2704 includes, but is not limited to, one or more dynamic random-access memory (DRAM) devices. In at least one embodiment, one or more DRAM devices are configured and / or configurable as a high-bandwidth memory ("HBM") subsystem in which multiple DRAM dies are stacked within each device.

[0172] In at least one embodiment, the high-speed GPU interconnect 2708 may refer to a wire-based multi-lane communication link, which is used by the system to scale and include one or more PPUs 2700 in combination with one or more CPUs, and supports cache coherence and CPU mastering between the PPUs 2700 and the CPUs. In at least one embodiment, data and / or commands are transmitted by the high-speed GPU interconnect 2708 through the hub 2716 to / from other units of the PPUs 2700, such as one or more copy engines, video encoders, video decoders, power management units, and other components that may not be explicitly shown in Figure 27.

[0173] In at least one embodiment, the I / O unit 2706 is configured to send and receive communications (e.g., commands, data) from a host processor (not shown in Figure 27) via the system bus 2702. In at least one embodiment, the I / O unit 2706 communicates with the host processor directly via the system bus 2702 or through one or more intermediate devices such as a memory bridge. In at least one embodiment, the I / O unit 2706 may communicate with one or more other processors, such as one or more of the PPU 2700, via the system bus 2702. In at least one embodiment, the I / O unit 2706 implements a PCIe interface for communication via the PCIe bus. In at least one embodiment, the I / O unit 2706 implements an interface for communicating with external devices.

[0174] In at least one embodiment, I / O unit 2706 decodes packets received via system bus 2702. In at least one embodiment, at least some packets represent commands configured to cause the PPU 2700 to perform various operations. In at least one embodiment, I / O unit 2706 transmits the decoded commands to various other units of the PPU 2700 specified by the commands. In at least one embodiment, the commands are transmitted to the front-end unit 2710 and / or to the hub 2716, or to one or more copy engines, video encoders, video decoders, power management units, etc. (not explicitly shown in Figure 27). In at least one embodiment, I / O unit 2706 is configured to route communication between and between various logical units of the PPU 2700.

[0175] In at least one embodiment, a program executed by the host processor encodes a command stream in a buffer that provides the workload to the PPU 2700 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 accessible (e.g., writable / readable) by both the host processor and the PPU 2700, and the host interface unit may be configured to access the buffer in system memory connected to the system bus 2702 via memory requests sent by the I / O unit 2706 over the system bus 2702. In at least one embodiment, the host processor writes a command stream to the buffer and then sends a pointer to the start of the command stream to the PPU 2700, thereby the front-end unit 2710 receives a pointer to one or more command streams, manages one or more command streams, reads commands from the command streams, and forwards the commands to various units of the PPU 2700.

[0176] In at least one embodiment, the front-end unit 2710 is coupled to a scheduler unit 2712 that configures various GPCs 2718 to handle tasks defined by one or more command streams. In at least one embodiment, the scheduler unit 2712 is configured to track state information relating to the various tasks managed by the scheduler unit 2712, which may indicate which of the GPCs 2718 a task is assigned to, whether the task is active or inactive, the priority level associated with the task, etc. In at least one embodiment, the scheduler unit 2712 manages the execution of multiple tasks on one or more of the GPCs 2718.

[0177] In at least one embodiment, the scheduler unit 2712 is coupled to a work distribution unit 2714 configured to dispatch tasks for execution on the GPC 2718. In at least one embodiment, the work distribution unit 2714 tracks the number of scheduled tasks received from the scheduler unit 2712, and the work distribution unit 2714 manages a pending task pool and an active task pool for each of the GPCs 2718. In at least one embodiment, the pending task pool comprises several slots (e.g., 32 slots) containing tasks assigned to be processed by a particular GPC 2718, and the active task pool comprises several slots (e.g., 4 slots) for tasks being actively processed by the GPCs 2718, so that when one of the GPCs 2718 completes the execution of a task, that task is removed from the active task pool for the GPC 2718, and one of the other tasks from the pending task pool is selected and scheduled for execution on the GPC 2718. In at least one embodiment, if an active task is idle on GPC2718, such as while waiting for data dependencies to be resolved, the active task is removed from GPC2718 and returned to the pending task pool, in which case another task is selected from the pending task pool and scheduled for execution on GPC2718.

[0178] In at least one embodiment, the work distribution unit 2714 communicates with one or more GPCs 2718 via the X-bar 2720. In at least one embodiment, the X-bar 2720 is an interconnection network that connects many units of the PPU 2700 to other units of the PPU 2700 and may be configured to connect the work distribution unit 2714 to a specific GPC 2718. In at least one embodiment, one or more other units of the PPU 2700 may also be connected to the X-bar 2720 via the hub 2716.

[0179] In at least one embodiment, tasks are managed by a scheduler unit 2712 and dispatched to one of the GPCs 2718 by a work distribution unit 2714. The GPC 2718 is configured to process tasks and generate results. In at least one embodiment, the results may be consumed by other tasks within the GPC 2718, routed to a different GPC 2718 via the X-bar 2720, or stored in memory 2704. In at least one embodiment, the results may be written to memory 2704 via a partition unit 2722, which implements a memory interface for reading and writing data to and from memory 2704. In at least one embodiment, the results may be sent to another PPU 2704 or CPU via a high-speed GPU interconnect 2708. In at least one embodiment, the PPU 2700 includes, but is not limited to, U partition units 2722 equal to the number of separate individual memory devices 2704 coupled to the PPU 2700.

[0180] In at least one embodiment, the host processor runs a driver kernel, which implements an application programming interface ("API") that enables one or more applications running on the host processor to schedule their operations for execution on the PPU2700. In at least one embodiment, multiple compute applications run concurrently on the PPU2700, and the PPU2700 provides isolation, quality of service ("QoS"), and independent address spaces for the multiple compute applications. In at least one embodiment, an application generates instructions (e.g., in the form of API calls) that cause the driver kernel to generate one or more tasks for execution on the PPU2700, and the driver kernel outputs the tasks to one or more streams being processed by the PPU2700. In at least one embodiment, each task comprises one or more groups of relational threads, which may be called warps. In at least one embodiment, a warp comprises multiple relational threads (e.g., 32 threads) that can run in parallel. In at least one embodiment, linked threads can refer to multiple threads that include instructions for performing tasks and exchange data through shared memory.

[0181] Figure 28 shows a GPC2800 according to at least one embodiment. In at least one embodiment, the GPC2800 is the GPC2718 in Figure 27. In at least one embodiment, each GPC2800 includes, but is not limited to, several hardware units for processing tasks, and each GPC2800 includes, but is not limited to, a pipeline manager 2802, a pre-raster operation unit ("PROP") 2804, a raster engine 2808, a work distribution crossbar ("WDX") 2816, an MMU 2818, one or more data processing clusters ("DPC") 2806, and any preferred combination of parts.

[0182] In at least one embodiment, the operation of the GPC2800 is controlled by a pipeline manager 2802. In at least one embodiment, the pipeline manager 2802 manages the configuration of one or more DPC2806s for handling tasks assigned to the GPC2800. In at least one embodiment, the pipeline manager 2802 configures at least one of the one or more DPC2806s to implement at least a portion of the graphics rendering pipeline. In at least one embodiment, the DPC2806 is configured to run a vertex shader program on a programmable streaming multiprocessor ("SM": streaming multiprocessor) 2814. In at least one embodiment, the pipeline manager 2802 is configured to route packets received from work distribution units to appropriate logical units within the GPC 2800, and in at least one embodiment, some packets may be routed to fixed-function hardware units and / or raster engine 2808 in PROP 2804, and other packets may be routed to DPC 2806 for processing by primitive engine 2812 or SM 2814. In at least one embodiment, the pipeline manager 2802 configures at least one of the DPC 2806 to implement a compute pipeline. In at least one embodiment, the pipeline manager 2802 configures at least one of the DPC 2806 to execute at least a portion of a CUDA program.

[0183] In at least one embodiment, the PROP unit 2804 is configured to route data generated by the raster engine 2808 and DPC 2806 to raster operation ("ROP") units in partition units, such as the memory partition unit 2722, which is described in more detail above in conjunction with Figure 27. In at least one embodiment, the PROP unit 2804 is configured to perform tasks such as performing optimization for color blending, organizing pixel data, and performing address translation. In at least one embodiment, the raster engine 2808 includes, but is not limited to, several fixed-function hardware units configured to perform various raster operations, and in at least one embodiment, the raster engine 2808 includes, but is not limited to, a setup engine, a coarse raster engine, a sorting engine, a clipping engine, a fine raster engine, a tile merging engine, and any preferred combination thereof. In at least one embodiment, the setup engine receives the transformed vertices and generates a plane equation related to the geometric primitives defined by the vertices, which is sent to a coarse raster engine to generate coverage information about the primitives (e.g., an x,y coverage mask for tiles), the output of the coarse raster engine is sent to a sorting engine to sort out fragments related to primitives that have fallen into the z test, and the output is sent to a clipping engine to clip fragments that are outside the view frustum. In at least one embodiment, the fragments that have passed clipping and sorting are passed to a fine raster engine to generate attributes about the pixel fragments based on the plane equation generated by the setup engine. In at least one embodiment, the output of the raster engine 2808 contains fragments to be processed by any preferred entity, such as a fragment shader implemented within the DPC2806.

[0184] In at least one embodiment, each DPC2806 contained within the GPC2800 includes, but is not limited to, an M-Pipe Controller ("MPC") 2810, a primitive engine 2812, one or more SM2814s, and any preferred combination thereof. In at least one embodiment, the MPC2810 controls the operation of the DPC2806 to route packets received from the pipeline manager 2802 to the appropriate unit in the DPC2806. In at least one embodiment, packets related to vertices may be routed to a primitive engine 2812 configured to fetch vertex attributes related to the vertices from memory, while packets related to shader programs may be sent to an SM2814.

[0185] In at least one embodiment, the SM2814 includes a programmable streaming processor configured to handle tasks represented by several threads, but not limited to these. In at least one embodiment, the SM2814 is multithreaded and configured to execute multiple threads (e.g., 32 threads) from a particular group of threads concurrently, implementing a SIMD architecture where each thread in a group of threads (e.g., a warp) is configured to process a different set of data based on the same set of instructions. In at least one embodiment, all threads in a group of threads execute the same instruction. In at least one embodiment, the SM2814 implements a SIMT architecture where each thread in a group of threads is configured to process a different set of data based on the same set of instructions, but individual threads in a group of threads are allowed to diverge during execution. In at least one embodiment, program counters, call stacks, and execution states are maintained for each warp to enable concurrent processing between warps and serial execution within warps when threads in a warp diverge. In another embodiment, a program counter, call stack, and execution state are maintained for each individual thread to enable equal concurrent processing among all threads, within warps, and between warps. In at least one embodiment, the execution state is maintained for each individual thread, and threads executing the same instruction may converge and execute in parallel for better efficiency. At least one embodiment of the SM2814 is described in further detail in conjunction with Figure 29.

[0186] In at least one embodiment, the MMU2818 provides an interface between the GPC2800 and a memory partition unit (for example, partition unit 2722 in Figure 27), and the MMU2818 provides virtual address-to-physical address translation, memory protection, and memory request arbitration. In at least one embodiment, the MMU2818 provides one or more translation lookaside buffers (TLBs) for performing virtual address-to-physical address translation in memory.

[0187] Figure 29 shows a streaming multiprocessor ("SM") 2900 according to at least one embodiment. In at least one embodiment, the SM2900 is the SM2814 in Figure 28. In at least one embodiment, the SM2900 includes, but is not limited to, an instruction cache 2902, one or more scheduler units 2904, a register file 2908, one or more processing cores ("cores") 2910, one or more special function units ("SFUs") 2912, one or more LSUs 2914, an interconnect network 2916, a shared memory / L1 cache 2918, and any preferred combination thereof. In at least one embodiment, a work distribution unit dispatches tasks for execution on a GPC of parallel processing units (PPUs), each task is assigned to a specific data processing cluster (DPC) within the GPC, and if the task is related to a shader program, the task is assigned to one of the SM2900s. In at least one embodiment, a scheduler unit 2904 receives tasks from the work distribution unit and manages instruction scheduling for one or more thread blocks assigned to the SM2900s. In at least one embodiment, the scheduler unit 2904 schedules thread blocks for execution as warps of parallel threads, each thread block is assigned at least one warp. In at least one embodiment, each warp executes a thread. In at least one embodiment, the scheduler unit 2904 manages multiple different thread blocks, allocates warps to different thread blocks, and then dispatches instructions from multiple different interlocking groups to various functional units (e.g., processing core 2910, SFU 2912, and LSU 2914) during each clock cycle.

[0188] In at least one embodiment, “dependency group” may refer to a programming model for organizing groups of communicating threads, and the programming model allows developers to express the granularity at which threads are communicating, enabling a richer and more efficient representation of parallel decomposition. In at least one embodiment, the dependent invocation API supports synchronization between thread blocks for the execution of parallel algorithms. In at least one embodiment, the API of the conventional programming model provides a single, simple construct for synchronizing dependent threads, namely a barrier across all threads in a thread block (e.g., the syncthreads() function). However, in at least one embodiment, the programmer may define groups of threads at a granularity smaller than a thread block and synchronize them within the defined group, enabling higher performance, design flexibility, and software reuse in the form of a functional interface across the collective group as a whole. In at least one embodiment, dependency groups allow programmers to explicitly define groups of threads at sub-block and multi-block granularity and perform collective actions such as synchronization of threads in the dependency group. In at least one embodiment, the sub-block granularity is as small as a single thread. In at least one embodiment, the programming model supports clean composition across software boundaries, thereby allowing libraries and utility functions to be safely synchronized within their local contexts without the need to make assumptions about convergence. In at least one embodiment, the interdependent group primitive enables new patterns of interdependent parallelism, including, but not limited to, producer-consumer parallelism, opportunistic parallelism, and global synchronization across the entire grid of thread blocks.

[0189] In at least one embodiment, the dispatch unit 2906 is configured to send instructions to one or more of the functional units, and the scheduler unit 2904 includes, but is not limited to, two dispatch units 2906 that enable two different instructions from the same warp to be dispatched during each clock cycle. In at least one embodiment, each scheduler unit 2904 includes a single dispatch unit 2906 or an additional dispatch unit 2906.

[0190] In at least one embodiment, each SM2900 includes a register file 2908 that provides a set of registers to the functional units of the SM2900, in at least one embodiment, but not limited to. In at least one embodiment, the register file 2908 is divided among the functional units such that each functional unit is allocated a dedicated portion of the register file 2908. In at least one embodiment, the register file 2908 is divided among different warps being executed by the SM2900, and the register file 2908 provides temporary storage for operands connected to the data paths of the functional units. In at least one embodiment, each SM2900 includes, but not limited to, a plurality of L processing cores 2910. In at least one embodiment, the SM2900 includes, but not limited to, a large number (e.g., 128 or more) individual processing cores 2910. In at least one embodiment, each processing core 2910 includes, but is not limited to, fully pipelining, single-precision, double-precision, and / or mixed-precision processing units, which include, but is not limited to, floating-point arithmetic logic units and integer arithmetic logic units. In at least one embodiment, the floating-point arithmetic logic units implement the IEEE 754-2008 standard for floating-point arithmetic. In at least one embodiment, the processing core 2910 includes, but is not limited to, 64 single-precision (32-bit) floating-point cores, 64 integer cores, 32 double-precision (64-bit) floating-point cores, and 8 tensor cores.

[0191] In at least one embodiment, the tensor core is configured to perform matrix operations. In at least one embodiment, one or more tensor cores are included in the processing core 2910. In at least one embodiment, the tensor core is configured to perform deep learning matrix arithmetic, such as convolution operations for neural network training and inference. In at least one embodiment, each tensor core operates on a 4x4 matrix and performs a matrix multiply and accumulate operation D = A × B + C, where A, B, C, and D are 4x4 matrices.

[0192] In at least one embodiment, the matrix multiplication inputs A and B are 16-bit floating-point matrices, and the addition matrices C and D are either 16-bit floating-point or 32-bit floating-point matrices. In at least one embodiment, the Tensor Core operates with 16-bit floating-point input data with a 32-bit floating-point sum. In at least one embodiment, the 16-bit floating-point multiplication uses 64 operations, resulting in a full-precision product, which is then added using a 32-bit floating-point addition with other intermediate products for a 4x4x4 matrix multiplication. In at least one embodiment, the Tensor Core is used to perform much larger 2D or even higher-dimensional matrix operations built from these smaller elements. In at least one embodiment, APIs such as the CUDA-C++ API expose special matrix load, matrix multiplicative sum, and matrix store operations to efficiently use the Tensor Core from CUDA-C++ programs. In at least one embodiment, at the CUDA level, the warp-level interface assumes a 16x16 matrix that spans all 32 threads of the warp.

[0193] In at least one embodiment, each SM2900 includes M SFU2912s that perform special functions (e.g., attribute evaluation, reciprocal square root, etc.). In at least one embodiment, the SFU2912 includes a tree traversal unit configured to traverse a hierarchical tree data structure (e.g.,). In at least one embodiment, the SFU2912 includes a texture unit configured to perform texture map filtering operations (e.g., texture map filtering). In at least one embodiment, the texture unit is configured to load a texture map (e.g., a 2D array of texels) from memory and sampled texture maps to produce sampled texture values ​​for use in a shader program executed by the SM2900. In at least one embodiment, the texture map is stored in shared memory / L1 cache 2918. In at least one embodiment, the texture unit implements texture operations such as filtering operations using mip maps (e.g., texture maps with different levels of detail). In at least one embodiment, each SM2900 includes, but is not limited to, two texture units.

[0194] In at least one embodiment, each SM2900 includes, but is not limited to, N LSUs2914 that implement load and store operations between the shared memory / L1 cache 2918 and the register file 2908. In at least one embodiment, each SM2900 includes, but is not limited to, an interconnection network 2916 that connects each of the functional units to the register file 2908 and connects the LSUs2914 to the register file 2908 and the shared memory / L1 cache 2918. In at least one embodiment, the interconnection network 2916 may be a crossbar that connects any of the functional units to any of the registers in the register file 2908 and connects the LSUs2914 to memory locations in the register file 2908 and the shared memory / L1 cache 2918.

[0195] In at least one embodiment, the shared memory / L1 cache 2918 is an array of on-chip memory that enables data storage and communication between the SM2900 and the primitive engine and between threads in the SM2900. In at least one embodiment, the shared memory / L1 cache 2918 has a storage capacity of 128KB, but is not limited to, and lies on the path from the SM2900 to the partition unit. In at least one embodiment, the shared memory / L1 cache 2918 is used to cache reads and writes. In at least one embodiment, one or more of the shared memory / L1 cache 2918, the L2 cache, and the memory are auxiliary stores.

[0196] In at least one embodiment, combining data caching and shared memory functionality into a single memory block provides improved performance for both types of memory access. In at least one embodiment, the capacity is used or made available as a cache by programs that do not use shared memory, such as when shared memory is configured to use half of the capacity and texture and load / store operations can use the remaining capacity. In at least one embodiment, integration within the shared memory / L1 cache 2918 allows the shared memory / L1 cache 2918 to function as a high-throughput tube for streaming data while simultaneously providing high-bandwidth and low-latency access to frequently reused data. In at least one embodiment, a simpler configuration may be used compared to graphics processing when configured for general-purpose parallel computing. In at least one embodiment, a fixed-function GPU is bypassed to create a much simpler programming model. In at least one embodiment and in the general-purpose parallel computing configuration, the work distribution unit directly allocates and distributes blocks of threads to the DPC. In at least one embodiment, a blocked thread runs the same program using a unique thread ID in the computation to ensure that each thread produces a unique result, executes the program using the SM2900, performs computations, communicates between threads using the shared memory / L1 cache 2918, and reads and writes global memory through the shared memory / L1 cache 2918 and the memory partition unit using the LSU2914. In at least one embodiment, when configured for general-purpose parallel computation, the SM2900 writes a command that the scheduler unit 2904 can use to start new work on the DPC.

[0197] In at least one embodiment, the PPU is included in or coupled to a desktop computer, laptop computer, tablet computer, server, supercomputer, smartphone (e.g., wireless handheld device), PDA, digital camera, vehicle, head-mounted display, handheld electronic device, etc. In at least one embodiment, the PPU is embodied on a single semiconductor substrate. In at least one embodiment, the PPU is included in an SoC together with one or more other devices such as additional PPUs, memory, RISC CPU, MMU, digital-to-analog converter ("DAC").

[0198] In at least one embodiment, the PPU may be contained on a graphics card that includes one or more memory devices. In at least one embodiment, the graphics card may be configured to interface with a PCIe slot on the motherboard of a desktop computer. In at least one embodiment, the PPU may be an integrated GPU ("iGPU") contained within the chipset of the motherboard.

[0199] Software constructs for general-purpose computing The following diagrams illustrate exemplary software constructs for implementing at least one embodiment, though they are not limiting.

[0200] Figure 30 shows a software stack for a programming platform according to at least one embodiment. In at least one embodiment, the programming platform is a platform for leveraging hardware on a computing system to accelerate computational tasks. In at least one embodiment, the programming platform may be accessible to software developers through libraries, compiler directives, and / or extensions to programming languages. In at least one embodiment, the programming platform may be, but is not limited to, CUDA, Radeon Open Compute Platform (“ROCm”), OpenCL (OpenCL® is developed by the Khronos group), SYCL, or Intel One API.

[0201] In at least one embodiment, the programming platform software stack 3000 provides an execution environment for application 3001. In at least one embodiment, application 3001 may include any computer software that can be launched on the software stack 3000. In at least one embodiment, application 3001 may 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.

[0202] In at least one embodiment, application 3001 and software stack 3000 may run on hardware 3007. In at least one embodiment, hardware 3007 may include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of compute devices that support programming platforms. In at least one embodiment, such as in the case of CUDA, software stack 3000 may be vendor-specific and compatible only with devices from a specific vendor (one or more). In at least one embodiment, such as in the case of OpenCL, software stack 3000 may be used with devices from different vendors. In at least one embodiment, hardware 3007 includes a host connected to another device that may be accessed to perform compute tasks via application programming interface ("API") calls. In at least one embodiment, the device in hardware 3007 may include, but is not limited to, a GPU, FPGA, AI engine, or other compute device (which may also include a CPU) and its memory, in contrast to the host in hardware 3007, which may include a CPU and its memory.

[0203] In at least one embodiment, the software stack 3000 of the programming platform includes, but is not limited to, several libraries 3003, a runtime 3005, and a device kernel driver 3006. In at least one embodiment, each of the libraries 3003 may include data and programming code that can be used by computer programs and utilized during software development. In at least one embodiment, the library 3003 may include, but is not limited to, prewritten code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. In at least one embodiment, the library 3003 includes functionality optimized for execution on one or more types of devices. In at least one embodiment, the library 3003 may include, but is not limited to, functionality for performing mathematical, deep learning, and / or other types of operations on a device. In at least one embodiment, the library 3003 may relate to a corresponding API 3002 that includes one or more APIs that expose functionality implemented in the library 3003.

[0204] In at least one embodiment, application 3001 is written as source code that is compiled into executable code, as will be described in more detail below in conjunction with Figures 35-37. In at least one embodiment, the executable code of application 3001 may run, at least partially, on an execution environment provided by the software stack 3000. In at least one embodiment, during the execution of application 3001, it may be possible to reach code that needs to run on the device, as opposed to the host. In at least one embodiment, in such a case, runtime 3005 may be called to load and start the code required on the device. In at least one embodiment, runtime 3005 may include any technically feasible runtime system capable of supporting the execution of application S01.

[0205] In at least one embodiment, runtime 3005 is implemented as one or more runtime libraries relating to the corresponding APIs, indicated as API 3004. In at least one embodiment, one or more such runtime libraries may include, but are not limited to, functions for memory management, execution control, device management, error handling, and / or synchronization. In at least one embodiment, memory management functions may include, but are not limited to, functions for allocating, deallocating, copying device memory, and transferring data between host memory and device memory. In at least one embodiment, execution control functions may include, but are not limited to, functions for invoking functions on a device (sometimes called "kernels" when the functions are global functions callable from the host) and setting attribute values ​​in buffers maintained by runtime libraries for given functions to be executed on the device.

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

[0207] In at least one embodiment, the device kernel driver 3006 is configured to facilitate communication with the underlying device. In at least one embodiment, the device kernel driver 3006 may provide low-level functionality on which APIs and / or other software, such as API 3004, rely. In at least one embodiment, the device kernel driver 3006 may be configured to compile intermediate representation ("IR") code into binary code at runtime. In at least one embodiment, for CUDA, the device kernel driver 3006 may compile non-hardware-specific parallel thread execution ("PTX") IR code into binary code for a specific target device at runtime (with caching of the compiled binary code), which is sometimes referred to as "finalizing" the code. In at least one embodiment, doing so may allow the finalized code to run on the target device, which may not have existed when the source code was first compiled into PTX code. Alternatively, in at least one embodiment, the device source code may be compiled into binary code offline without requiring the device kernel driver 3006 to compile the IR code at runtime.

[0208] Figure 31 shows a CUDA implementation of the software stack 3000 of Figure 30 according to at least one embodiment. In at least one embodiment, the CUDA software stack 3100, from which application 3101 can be invoked, includes a CUDA library 3103, a CUDA runtime 3105, a CUDA driver 3107, and a device kernel driver 3108. In at least one embodiment, the CUDA software stack 3100 runs on hardware 3109, which may include a GPU that supports CUDA and is developed by NVIDIA Corporation in Santa Clara, California.

[0209] In at least one embodiment, application 3101, CUDA runtime 3105, and device kernel driver 3108 may perform similar functionality to application 3001, runtime 3005, and device kernel driver 3006, respectively, as described above in conjunction with Figure 30. In at least one embodiment, CUDA driver 3107 includes a library (libcuda.so) that implements CUDA driver API 3106. In at least one embodiment, CUDA driver API 3106 may expose functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability, but not limited to these. In at least one embodiment, the CUDA driver API 3106 differs from the CUDA runtime API 3104 in that it simplifies device code management by providing implicit initialization, context management (similar to a process), and module management (similar to a dynamically loaded library). In at least one embodiment, in contrast to the high-level CUDA runtime API 3104, the CUDA driver API 3106 is a low-level API that provides finer-grained control of the device, particularly with respect to context and module loading. In at least one embodiment, the CUDA driver API 3106 may expose functionality for context management that is not exposed by the CUDA runtime API 3104. In at least one embodiment, the CUDA driver API 3106 is also language-independent and supports OpenCL, for example, in addition to the CUDA runtime API 3104. Furthermore, in at least one embodiment, the development library containing the CUDA runtime 3105 may be considered separate from the driver components containing the user-mode CUDA driver 3107 and the kernel-mode device driver 3108 (sometimes referred to as the “display” driver).

[0210] In at least one embodiment, the CUDA library 3103 may include, but is not limited to, mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries, which can be used by parallel computing applications such as application 3101. In at least one embodiment, the CUDA library 3103 may include mathematical libraries such as the cuBLAS library, which is an implementation of Basic Linear Algebra Subprograms ("BLAS") for performing linear algebra operations; the cuFFT library for calculating the Fast Fourier Transform ("FFT"); and the cuRAND library for generating random numbers. In at least one embodiment, the CUDA library 3103 may include deep learning libraries such as the cuDNN library of primitives for deep neural networks and the TensorRT platform for high-performance deep learning inference.

[0211] Figure 32 shows an ROCm implementation of the software stack 3000 of Figure 30 according to at least one embodiment. In at least one embodiment, the ROCm software stack 3200, from which application 3201 can be launched, includes a language runtime 3203, a system runtime 3205, a thunk 3207, and a ROCm kernel driver 3208. In at least one embodiment, the ROCm software stack 3200 runs on hardware 3209, which may include a GPU that supports ROCm and is developed by AMD Corporation in Santa Clara, California.

[0212] In at least one embodiment, application 3201 may perform similar functionality to application 3001 as described above in conjunction with Figure 30. In at least one embodiment, the language runtime 3203 and system runtime 3205 may further perform similar functionality to runtime 3005 as described above in conjunction with Figure 30. In at least one embodiment, the language runtime 3203 and system runtime 3205 differ in that system runtime 3205 is a language-independent runtime that implements the ROCr system runtime API 3204 and utilizes the Heterogeneous System Architecture ("HSA") runtime API. In at least one embodiment, the HSA runtime API is a thin user-mode API that exposes an interface for accessing and interacting with the AMD GPU, including, among other things, memory management, execution control via kernel-designed dispatch, error handling, system and agent information, and functions for runtime initialization and shutdown. In at least one embodiment, as opposed to the system runtime 3205, the language runtime 3203 is an implementation of a language-specific runtime API 3202 layered on top of the ROCr system runtime API 3204. In at least one embodiment, the language runtime API may include, but is not limited to, a Heterogeneous compute Interface for Portability ("HIP") language runtime API, a Heterogeneous Compute Compiler ("HCC") language runtime API, or an OpenCL API. The HIP language, in particular, is an extension of the C++ programming language with a functionally similar version of the CUDA mechanism, and in at least one embodiment, the HIP language runtime API includes, among other things, functions similar to those of the CUDA runtime API 3104 described above in conjunction with Figure 31, such as functions for memory management, execution control, device management, error handling, and synchronization.

[0213] In at least one embodiment, the thunk (ROCt) 3207 is an interface 3206 that may be used to interact with the underlying ROCm driver 3208. In at least one embodiment, the ROCm driver 3208 is a ROCk driver, which is a combination of the AMDGPU driver and the HSA kernel driver (amdkfd). In at least one embodiment, the AMDGPU driver is a device kernel driver for GPUs developed by AMD that implements similar functionality to the device kernel driver 3006 described above in conjunction with Figure 30. In at least one embodiment, the HSA kernel driver is a driver that allows different types of processors to share system resources more effectively through hardware features.

[0214] In at least one embodiment, various libraries (not shown) may be included in the ROCm software stack 3200 above the language runtime 3203, providing a functional similarity to the CUDA library 3103 described above, in conjunction with Figure 31. In at least one embodiment, the various libraries may include, but are not limited to, a hipBLAS library that implements functionality similar to that of CUDA cuBLAS, a rocFFT library for calculating FFTs similar to those of CUDA cuFFT, and other libraries for mathematics, deep learning, and / or other purposes.

[0215] Figure 33 shows an OpenCL implementation of the software stack 3000 of Figure 30 in at least one embodiment. In at least one embodiment, the OpenCL software stack 3300, from which application 3301 can be invoked, includes the OpenCL framework 3310, the OpenCL runtime 3306, and the driver 3307. In at least one embodiment, the OpenCL software stack 3300 runs on vendor-non-vendor specific hardware 3109. In at least one embodiment, since OpenCL is supported by devices developed by different vendors, a specific OpenCL driver may be required to interact with hardware from such vendors.

[0216] In at least one embodiment, application 3301, OpenCL runtime 3306, device kernel driver 3307, and hardware 3308 may perform similar functionality to application 3001, runtime 3005, device kernel driver 3006, and hardware 3007, respectively, as described above in conjunction with Figure 30. In at least one embodiment, application 3301 further includes an OpenCL kernel 3302 having code to be executed on the device.

[0217] In at least one embodiment, OpenCL defines a “platform” that enables a host to control devices connected to the host. In at least one embodiment, the OpenCL framework provides platform layer APIs and runtime APIs, shown as platform API 3303 and runtime API 3305. In at least one embodiment, runtime API 3305 uses contexts to manage kernel execution on devices. In at least one embodiment, each identified device may be associated with its own context, and runtime API 3305 may use its respective context to manage, among other things, command queues, program objects, and kernel objects, and share memory objects for that device. In at least one embodiment, platform API 3303 exposes functionality that allows device contexts to be used to select and initialize devices, submit work to devices via command queues, and enable data transfer to and from devices. In at least one embodiment, the OpenCL framework further provides a variety of built-in functions (not shown), among other things, including mathematical functions, relational functions, and image processing functions.

[0218] In at least one embodiment, compiler 3304 is also included in the OpenCL framework 3310. In at least one embodiment, source code may be compiled offline before the application is executed or online during the execution of the application. In contrast to CUDA and ROCm, in at least one embodiment, an OpenCL application may be compiled online by compiler 3304, which is included to represent any number of compilers that can be used to compile source code and / or IR code into binary code, such as Standard Portable Intermediate Representation ("SPIR-V") code. Alternatively, in at least one embodiment, an OpenCL application may be compiled offline before such application is executed.

[0219] Figure 34 shows software supported by a programming platform in at least one embodiment. In at least one embodiment, the programming platform 3404 is configured to support various programming models 3403, middleware and / or libraries 3402, and frameworks 3401 on which application 3400 may rely. In at least one embodiment, application 3400 may be an AI / ML application implemented using a deep learning framework such as MXNet, PyTorch, or TensorFlow, which may rely on libraries such as cuDNN, NVIDIA Collective Communications Library (“NCCL”), and / or NVIDIA Developer Data Loading Library (“DALI”) CUDA libraries to provide accelerated computing on the underlying hardware.

[0220] In at least one embodiment, the programming platform 3404 may be one of the CUDA, ROCm, or OpenCL platforms described above in conjunction with Figures 31, 32, and 33, respectively. In at least one embodiment, the programming platform 3404 supports a plurality of programming models 3403, which are abstractions of the underlying computing system that allow representation of algorithms and data structures. In at least one embodiment, the programming model 3403 may expose the features of the underlying hardware to improve performance. In at least one embodiment, the programming model 3403 may include, but is not limited to, CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism ("C++AMP"), Open Multi-Processing ("OpenMP"), Open Accelerators ("OpenACC"), and / or Vulcan Compute.

[0221] In at least one embodiment, the library and / or middleware 3402 provides an implementation of the abstraction of the programming model 3404. In at least one embodiment, such a library includes data and programming code that can be used by a computer program and utilized during software development. In at least one embodiment, such middleware includes software that provides services to the application in addition to the software available from the programming platform 3404. In at least one embodiment, the library and / or middleware 3402 may include, but is not limited to, cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Furthermore, in at least one embodiment, the library and / or middleware 3402 may include libraries of NCCL and ROCm Communication Collectives Library ("RCCL") providing communication routines for the GPU, the MIOpen library for deep learning acceleration, and / or the Eigen library for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0222] In at least one embodiment, the application framework 3401 depends on libraries and / or middleware 3402. In at least one embodiment, each of the application frameworks 3401 is a software framework used to implement a standard structure of application software. In at least one embodiment, returning to the AI / ML example described above, the AI / ML application may be implemented using a framework such as Caffe, Caffe2, TensorFlow, Keras, PyTorch, or the MxNet deep learning framework.

[0223] Figure 35 illustrates the compilation of code for execution on one of the programming platforms shown in Figures 30–33, according to at least one embodiment. In at least one embodiment, compiler 3501 receives source code 3500, which includes both host code and device code. In at least one embodiment, compiler 3501 is configured to convert source code 3500 into host executable code 3502 for execution on a host and device executable code 3503 for execution on a device. In at least one embodiment, source code 3500 may be compiled offline before application execution or online during application execution.

[0224] In at least one embodiment, source code 3500 may include code in any programming language supported by compiler 3501, such as C++, C, or Fortran. In at least one embodiment, source code 3500 may be contained in a single source file having a mixture of host code and device code, where the location of the device code is indicated. In at least one embodiment, the single source file may be a .cu file containing CUDA code or a .hip.cpp file containing HIP code. Alternatively, in at least one embodiment, source code 3500 may include multiple source code files rather than a single source file in which host code and device code are separated.

[0225] In at least one embodiment, compiler 3501 is configured to compile source code 3500 into host executable code 3502 for execution on a host and device executable code 3503 for execution on a device. In at least one embodiment, compiler 3501 performs operations including parsing source code 3500 into an abstract system tree (AST), performing optimizations, and generating executable code. In at least one embodiment in which source code 3500 comprises a single source file, compiler 3501 may separate device code from host code in such a single source file, as will be described in more detail below with respect to Figure 36, and compile the device code and host code into device executable code 3503 and host executable code 3502, respectively, and link the device executable code 3503 and host executable code 3502 to each other in a single file.

[0226] In at least one embodiment, the host executable code 3502 and the device executable code 3503 may be in any preferred format, such as binary code and / or IR code. In at least one embodiment, for CUDA, the host executable code 3502 may contain native object code, and the device executable code 3503 may contain code in the PTX intermediate representation. In at least one embodiment, for ROCm, both the host executable code 3502 and the device executable code 3503 may contain target binary code.

[0227] Figure 36 is a more detailed diagram of compiling code for execution on one of the programming platforms shown in Figures 30–33, according to at least one embodiment. In at least one embodiment, compiler 3601 is configured to receive source code 3600, compile source code 3600, and output executable file 3610. In at least one embodiment, source code 3600 is a single source file, such as a .cu file, a .hip.cpp file, or a file in another format, containing both host code and device code. In at least one embodiment, compiler 3601 may be, but is not limited to, an NVIDIA CUDA compiler ("NVCC": NVIDIA CUDA compiler) for compiling CUDA code in a .cu file, or an HCC compiler for compiling HIP code in a .hip.cpp file.

[0228] In at least one embodiment, compiler 3601 includes a compiler front-end 3602, a host compiler 3605, a device compiler 3606, and a linker 3609. In at least one embodiment, compiler front-end 3602 is configured to separate device code 3604 from host code 3603 in source code 3600. In at least one embodiment, device code 3604 is compiled by device compiler 3606 into device executable code 3608, which may include binary code or IR code as described. In at least one embodiment, separately, host code 3603 is compiled by host compiler 3605 into host executable code 3607. In at least one embodiment, for NVCC, the host compiler 3605 may be a general-purpose C / C++ compiler that outputs native object code, but is not limited to that which, while the device compiler 3606 may be a Low-Level Virtual Machine ("LLVM") based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code, but is not limited to that which. In at least one embodiment, for HCC, both the host compiler 3605 and the device compiler 3606 may be LLVM-based compilers that output target binary code, but is not limited to that which.

[0229] In at least one embodiment, after compiling source code 3600 into host executable code 3607 and device executable code 3608, linker 3609 links host executable code 3607 and device executable code 3608 to each other in executable file 3610. In at least one embodiment, native object code for the host and PTX or binary code for the device may be linked to each other in an Executable and Linkable Format ("ELF") file, which is a container format used to store the object code.

[0230] Figure 37 shows, in at least one embodiment, that the source code is translated before compiling the source code. In at least one embodiment, the source code 3700 is passed through a translation tool 3701, which translates the source code 3700 into translated source code 3702. In at least one embodiment, the compiler 3703 is used to compile the translated source code 3702 into host executable code 3704 and device executable code 3705 in a process similar to the compilation of the source code 3500 by the compiler 3501 into host executable code 3502 and device executable code 3503, as described above in conjunction with Figure 35.

[0231] In at least one embodiment, the translation performed by the translation tool 3701 is used to port the source code 3700 for execution in an environment different from the environment in which it was originally intended to run. In at least one embodiment, the translation tool 3701 may include, but is not limited to, a HIP translator used to "hipify" CUDA code targeting the CUDA platform into HIP code that can be compiled and executed on the ROCm platform. In at least one embodiment, the translation of the source code 3700 may include parsing the source code 3700 and converting calls to one or more APIs provided by one programming model (e.g., CUDA) to corresponding calls to one or more APIs provided by another programming model (e.g., HIP), as will be described in more detail below in conjunction with Figures 38A to 39. In at least one embodiment, returning to the example of hipify CUDA code, calls to the CUDA runtime API, CUDA driver API, and / or CUDA libraries can be converted to corresponding HIP API calls. In at least one embodiment, the automated translation performed by the translation tool 3701 is sometimes incomplete and may require additional manual effort to fully port the source code 3700.

[0232] Configuring GPUs for general-purpose computing The following diagram illustrates, but is not limited to, an exemplary architecture for compiling and executing compute source code, in at least one embodiment.

[0233] Figure 38A shows a system 38A00 configured to compile and execute CUDA source code 3810 using different types of processing units, according to at least one embodiment. In at least one embodiment, system 38A00 includes, but is not limited to, CUDA source code 3810, CUDA compiler 3850, host executable code 3870(1), host executable code 3870(2), CUDA device executable code 3884, CPU 3890, CUDA-enabled GPU 3894, GPU 3892, CUDA-to-HIP translation tool 3820, HIP source code 3830, HIP compiler driver 3840, HCC 3860, and HCC device executable code 3882.

[0234] In at least one embodiment, CUDA source code 3810 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, but is not limited to, a mechanism for defining device code and distinguishing device code from host code. In at least one embodiment, device code is source code that can be executed in parallel on a device after compilation. In at least one embodiment, the device may be a processor optimized for parallel instruction processing, such as a CUDA-enabled GPU 3890, GPU 38192, or another GPGPU. In at least one embodiment, host code is source code that can be executed on a host after compilation. In at least one embodiment, the host may be a processor optimized for sequential instruction processing, such as a CPU 3890.

[0235] In at least one embodiment, the CUDA source code 3810 includes, but is not limited to, any number of global functions 3812 (including zero), any number of device functions 3814 (including zero), any number of host functions 3816 (including zero), and any number of host / device functions 3818 (including zero). In at least one embodiment, the global functions 3812, device functions 3814, host functions 3816, and host / device functions 3818 may be mixed in the CUDA source code 3810. In at least one embodiment, each of the global functions 3812 is executable on a device and callable from a host. In at least one embodiment, one or more of the global functions 3812 can therefore act as an entry point to a device. In at least one embodiment, each of the global functions 3812 is a kernel. In at least one embodiment, and in a technique known as dynamic parallelism, one or more of the global functions 3812 define a kernel, which is executable on a device and callable from such a device. In at least one embodiment, the kernel is executed N times in parallel by N (where N is any positive integer) different threads on the device during execution.

[0236] In at least one embodiment, each of the device functions 3814 runs on a device and is callable only from such a device. In at least one embodiment, each of the host functions 3816 runs on a host and is callable only from such a host. In at least one embodiment, each of the host / device functions 3816 defines both a host version of the function that is run on a host and is callable only from such a host, and a device version of the function that is run on a device and is callable only from such a device.

[0237] In at least one embodiment, CUDA source code 3810 may also include any number of calls to any number of functions defined via CUDA runtime API 3802, without limiting it. In at least one embodiment, CUDA runtime API 3802 may include any number of functions that run on the host, such as allocating and deallocating device memory, transferring data between host memory and device memory, and managing a system with multiple devices, without limiting it. In at least one embodiment, CUDA source code 3810 may 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 for use by CUDA code. In at least one embodiment, a CUDA API may include, but is not limited to, CUDA runtime API 3802, CUDA driver APIs, APIs for any number of CUDA libraries, etc. In at least one embodiment, and with respect to the CUDA runtime API 3802, the CUDA driver API is a lower-level API but provides finer-grained control of the device. In at least one embodiment, examples of CUDA libraries include, but are not limited to, cuBLAS, cuFFT, cuRAND, cuDNN, etc.

[0238] In at least one embodiment, the CUDA compiler 3850 compiles input CUDA code (e.g., CUDA source code 3810) to generate host executable code 3870(1) and CUDA device executable code 3884. In at least one embodiment, the CUDA compiler 3850 is an NVCC. In at least one embodiment, the host executable code 3870(1) is a compiled version of the host code contained in the input source code, which is executable on the CPU 3890. In at least one embodiment, the CPU 3890 may be any processor optimized for sequential instruction processing.

[0239] In at least one embodiment, CUDA device executable code 3884 is a compiled version of the device code contained in the input source code, executable on a CUDA-enabled GPU 3894. In at least one embodiment, CUDA device executable code 3884 includes, but is not limited to, binary code. In at least one embodiment, CUDA device executable code 3884 includes, but is not limited to, IR code such as PTX code, which is further compiled at runtime by the device driver into binary code for a specific target device (e.g., a CUDA-enabled GPU 3894). In at least one embodiment, the CUDA-enabled GPU 3894 can be any processor optimized for parallel instruction processing and supporting CUDA. In at least one embodiment, the CUDA-enabled GPU 3894 is developed by NVIDIA Corporation in Santa Clara, California.

[0240] In at least one embodiment, the CUDA-to-HIP translation tool 3820 is configured to translate CUDA source code 3810 into functionally similar HIP source code 3830. In at least one embodiment, HIP source code 3830 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, but is not limited to, a functionally similar version of the CUDA mechanism for defining device code and distinguishing device code from host code. In at least one embodiment, the HIP programming language may include a subset of the functionality of the CUDA programming language. In at least one embodiment, for example, a HIP programming language includes (one or more) mechanisms for defining global functions 3812, but such a HIP programming language may not support dynamic parallelism, and therefore global functions 3812 defined in the HIP code may only be callable from the host.

[0241] In at least one embodiment, the HIP source code 3830 includes, without limitation, any number (including zero) of global functions 3812, any number (including zero) of device functions 3814, any number (including zero) of host functions 3816, and any number (including zero) of host / device functions 3818. In at least one embodiment, the HIP source code 3830 may also include any number of calls to any number of functions specified in the HIP runtime API 3832. In at least one embodiment, the HIP runtime API 3832 includes, without limitation, a functionally similar version of a subset of the functions included in the CUDA runtime API 3802. In at least one embodiment, the HIP source code 3830 may also include any number of calls to any number of functions specified in any number of other HIP APIs. In at least one embodiment, the HIP API can be any API designed for use with HIP code and / or ROCm. In at least one embodiment, the HIP API includes, without limitation, the HIP runtime API 3832, the HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, and the like.

[0242] In at least one embodiment, the CUDA-to-HIP translation tool 3820 converts each kernel call in the CUDA code from CUDA syntax to HIP syntax and converts any number of other CUDA calls in the 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 the CUDA API, and a HIP call is a call to a function specified in the HIP API. In at least one embodiment, the CUDA-to-HIP translation tool 3820 converts any number of calls to functions specified in the CUDA runtime API 3802 to any number of calls to functions specified in the HIP runtime API 3832.

[0243] In at least one embodiment, the CUDA-to-HIP translation tool 3820 is a tool known as hipify-Perl that performs a text-based translation process. In at least one embodiment, the CUDA-to-HIP translation tool 3820 is a tool known as hipify-clang, which performs a more complex and robust translation process that involves parsing CUDA code using clang (a compiler front-end) and then translating the resulting symbols. In at least one embodiment, properly converting CUDA code to HIP code may require modifications (e.g., manual edits) in addition to the modifications performed by the CUDA-to-HIP translation tool 3820.

[0244] In at least one embodiment, the HIP compiler driver 3840 is a front-end that determines the target device 3846 and then configures a compiler that is compatible with the target device 3846 to compile the HIP source code 3830. In at least one embodiment, the target device 3846 is a processor optimized for parallel instruction processing. In at least one embodiment, the HIP compiler driver 3840 may determine the target device 3846 in any technically feasible manner.

[0245] In at least one embodiment, if the target device 3846 is compatible with CUDA (for example, a CUDA-enabled GPU 3894), the HIP compiler driver 3840 generates a HIP / NVCC compile command 3842. In at least one embodiment, and as described in more detail in conjunction with Figure 38B, the HIP / NVCC compile command 3842 configures the CUDA compiler 3850 to compile the HIP source code 3830 using, but not limited to, HIP-to-CUDA translation headers and CUDA runtime libraries. In at least one embodiment, and in response to the HIP / NVCC compile command 3842, the CUDA compiler 3850 generates host executable code 3870(1) and CUDA device executable code 3884.

[0246] In at least one embodiment, if the target device 3846 is not CUDA compatible, the HIP compiler driver 3840 generates a HIP / HCC compile command 3844. In at least one embodiment, and as described in more detail in conjunction with Figure 38C, the HIP / HCC compile command 3844 configures HCC3860 to compile HIP source code 3830 using, but not limited to, HCC headers and HIP / HCC runtime libraries. In at least one embodiment, and in response to the HIP / HCC compile command 3844, HCC3860 generates host executable code 3870(2) and HCC device executable code 3882. In at least one embodiment, HCC device executable code 3882 is a compiled version of the device code contained in HIP source code 3830, which is executable on GPU 3892. In at least one embodiment, the GPU3892 could be any processor optimized for parallel instruction processing, incompatible with CUDA, and compatible with HCC. In at least one embodiment, the GPU3892 is developed by AMD Corporation in Santa Clara, California. In at least one embodiment, the GPU3892 is a non-CUDA-enabled GPU3892.

[0247] For illustrative purposes only, three different flows that may be implemented in at least one embodiment to compile CUDA source code 3810 for execution on CPU 3890 and different devices are illustrated in Figure 38A. In at least one embodiment, a direct CUDA flow compiles CUDA source code 3810 for execution on CPU 3890 and CUDA-enabled GPU 3894 without translating CUDA source code 3810 to HIP source code 3830. In at least one embodiment, an indirect CUDA flow translates CUDA source code 3810 to HIP source code 3830 and then compiles HIP source code 3830 for execution on CPU 3890 and CUDA-enabled GPU 3894. In at least one embodiment, the CUDA / HCC flow translates CUDA source code 3810 into HIP source code 3830, and then compiles HIP source code 3830 for execution on CPU 3890 and GPU 3892.

[0248] A direct CUDA flow, which can be implemented in at least one embodiment, is illustrated via a dashed line and a series of bubbles annotated A1-A3. In at least one embodiment, and as illustrated in the bubble annotated A1, the CUDA compiler 3850 receives CUDA source code 3810 and a CUDA compile command 3848 that configures the CUDA compiler 3850 to compile the CUDA source code 3810. In at least one embodiment, the CUDA source code 3810 used in the direct CUDA flow is written in a CUDA programming language based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In at least one embodiment, and in response to the CUDA compile command 3848, the CUDA compiler 3850 generates host executable code 3870(1) and CUDA device executable code 3884 (illustrated in the bubble annotated A2). In at least one embodiment, and as illustrated in the bubble annotated with A3, the host executable code 3870(1) and the CUDA device executable code 3884 may run on a CPU 3890 and a CUDA-enabled GPU 3894, respectively. In at least one embodiment, the CUDA device executable code 3884 includes, but is not limited to, binary code. In at least one embodiment, the CUDA device executable code 3884 includes, but is not limited to, PTX code which is further compiled at runtime into binary code for a particular target device.

[0249] An indirect CUDA flow that may be implemented in at least one embodiment is illustrated via a dotted line and a series of bubbles annotated as B1-B6. In at least one embodiment, as illustrated in the bubble annotated as B1, the CUDA-to-HIP translation tool 3820 receives CUDA source code 3810. In at least one embodiment, as illustrated in the bubble annotated as B2, the CUDA-to-HIP translation tool 3820 translates CUDA source code 3810 into HIP source code 3830. In at least one embodiment, as illustrated in the bubble annotated as B3, the HIP compiler driver 3840 receives HIP source code 3830 and determines that the target device 3846 is CUDA-enabled.

[0250] In at least one embodiment, and as illustrated in the bubble annotated as B4, the HIP compiler driver 3840 generates a HIP / NVCC compile command 3842 and sends both the HIP / NVCC compile command 3842 and the HIP source code 3830 to the CUDA compiler 3850. In at least one embodiment, and as described in more detail in conjunction with Figure 38B, the HIP / NVCC compile command 3842 configures the CUDA compiler 3850 to compile the HIP source code 3830 using, but not limited to, a HIP-to-CUDA translation header and CUDA runtime libraries. In at least one embodiment, and in response to the HIP / NVCC compile command 3842, the CUDA compiler 3850 generates host executable code 3870(1) and CUDA device executable code 3884 (illustrated in the bubble annotated as B5). In at least one embodiment, and as illustrated in the bubble annotated as B6, the host executable code 3870(1) and the CUDA device executable code 3884 may run on a CPU 3890 and a CUDA-enabled GPU 3894, respectively. In at least one embodiment, the CUDA device executable code 3884 includes, but is not limited to, binary code. In at least one embodiment, the CUDA device executable code 3884 includes, but is not limited to, PTX code which is further compiled at runtime into binary code for a specific target device.

[0251] A CUDA / HCC flow that may be implemented in at least one embodiment is illustrated via solid lines and a series of bubbles annotated C1-C6. In at least one embodiment, as illustrated in the bubble annotated C1, the CUDA-to-HIP translation tool 3820 receives CUDA source code 3810. In at least one embodiment, as illustrated in the bubble annotated C2, the CUDA-to-HIP translation tool 3820 translates CUDA source code 3810 into HIP source code 3830. In at least one embodiment, as illustrated in the bubble annotated C3, the HIP compiler driver 3840 receives HIP source code 3830 and determines that the target device 3846 is not CUDA-enabled.

[0252] In at least one embodiment, the HIP compiler driver 3840 generates a HIP / HCC compile command 3844 and sends both the HIP / HCC compile command 3844 and the HIP source code 3830 to the HCC 3860 (illustrated by a bubble annotated with C4). In at least one embodiment, and as described in more detail in conjunction with Figure 38C, the HIP / HCC compile command 3844 configures the HCC 3860 to compile the HIP source code 3830 using, but not limited to, the HCC header and the HIP / HCC runtime library. In at least one embodiment, and in response to the HIP / HCC compile command 3844, the HCC 3860 generates the host executable code 3870(2) and the HCC device executable code 3882 (illustrated by a bubble annotated with C5). In at least one embodiment, and as illustrated in the bubble annotated with C6, the host executable code 3870(2) and the HCC device executable code 3882 may be executed on the CPU 3890 and the GPU 3892, respectively.

[0253] In at least one embodiment, after the CUDA source code 3810 has been translated into HIP source code 3830, the HIP compiler driver 3840 may then be used to generate executable code for either the CUDA-enabled GPU 3894 or GPU 3892 without re-running the CUDA-to-HIP translation tool 3820. In at least one embodiment, the CUDA-to-HIP translation tool 3820 translates the CUDA source code 3810 into HIP source code 3830, which is then stored in memory. In at least one embodiment, the HIP compiler driver 3840 then configures the HCC 3860 to generate host executable code 3870(2) and HCC device executable code 3882 based on the HIP source code 3830. In at least one embodiment, the HIP compiler driver 3840 then configures the CUDA compiler 3850 to generate host executable code 3870(1) and CUDA device executable code 3884 based on the stored HIP source code 3830.

[0254] Figure 38B shows a system 3804 configured to compile and execute the CUDA source code 3810 of Figure 38A using a CPU 3890 and a CUDA-enabled GPU 3894, according to at least one embodiment. In at least one embodiment, system 3804 includes, but is not limited to, the CUDA source code 3810, a CUDA-to-HIP translation tool 3820, HIP source code 3830, a HIP compiler driver 3840, a CUDA compiler 3850, host executable code 3870(1), CUDA device executable code 3884, a CPU 3890, and a CUDA-enabled GPU 3894.

[0255] In at least one embodiment, and as previously described herein in conjunction with Figure 38A, the CUDA source code 3810 includes, but is not limited to, any number of global functions 3812 (including zero), any number of device functions 3814 (including zero), any number of host functions 3816 (including zero), and any number of host / device functions 3818 (including zero). In at least one embodiment, the CUDA source code 3810 also includes, but is not limited to, any number of calls to any number of functions specified in any number of CUDA APIs.

[0256] In at least one embodiment, the CUDA-to-HIP translation tool 3820 translates CUDA source code 3810 into HIP source code 3830. In at least one embodiment, the CUDA-to-HIP translation tool 3820 converts each kernel call in the CUDA source code 3810 from CUDA syntax to HIP syntax, and converts any number of other CUDA calls in the CUDA source code 3810 into any number of other functionally similar HIP calls.

[0257] In at least one embodiment, the HIP compiler driver 3840 determines that the target device 3846 is CUDA-enabled and generates a HIP / NVCC compile command 3842. In at least one embodiment, the HIP compiler driver 3840 then configures the CUDA compiler 3850 via the HIP / NVCC compile command 3842 to compile the HIP source code 3830. In at least one embodiment, as part of configuring the CUDA compiler 3850, the HIP compiler driver 3840 provides access to a HIP-to-CUDA translation header 3852. In at least one embodiment, the HIP-to-CUDA translation header 3852 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, the CUDA compiler 3850 uses a HIP-to-CUDA translation header 3852, in conjunction with a CUDA runtime library 3854 corresponding to the CUDA runtime API 3802, to generate host executable code 3870(1) and CUDA device executable code 3884. In at least one embodiment, the host executable code 3870(1) and CUDA device executable code 3884 may then be executed on a CPU 3890 and a CUDA-enabled GPU 3894, respectively. In at least one embodiment, the CUDA device executable code 3884 includes, but is not limited to, binary code. In at least one embodiment, the CUDA device executable code 3884 includes, but is not limited to, PTX code, which is further compiled at runtime into binary code for a specific target device.

[0258] Figure 38C shows a system 3806 configured to compile and run the CUDA source code 3810 of Figure 38A using a CPU 3890 and a non-CUDA-enabled GPU 3892, according to at least one embodiment. In at least one embodiment, system 3806 includes, but is not limited to, the CUDA source code 3810, a CUDA-to-HIP translation tool 3820, HIP source code 3830, a HIP compiler driver 3840, an HCC 3860, host executable code 3870(2), HCC device executable code 3882, a CPU 3890, and a GPU 3892.

[0259] In at least one embodiment, and as previously described herein in conjunction with Figure 38A, the CUDA source code 3810 includes, but is not limited to, any number of global functions 3812 (including zero), any number of device functions 3814 (including zero), any number of host functions 3816 (including zero), and any number of host / device functions 3818 (including zero). In at least one embodiment, the CUDA source code 3810 also includes, but is not limited to, any number of calls to any number of functions specified in any number of CUDA APIs.

[0260] In at least one embodiment, the CUDA-to-HIP translation tool 3820 translates CUDA source code 3810 into HIP source code 3830. In at least one embodiment, the CUDA-to-HIP translation tool 3820 converts each kernel call in the CUDA source code 3810 from CUDA syntax to HIP syntax, and converts any number of other CUDA calls in the source code 3810 into any number of other functionally similar HIP calls.

[0261] In at least one embodiment, the HIP compiler driver 3840 then determines that the target device 3846 is not CUDA-enabled and generates a HIP / HCC compile command 3844. In at least one embodiment, the HIP compiler driver 3840 then configures HCC3860 to execute the HIP / HCC compile command 3844 to compile the HIP source code 3830. In at least one embodiment, the HIP / HCC compile command 3844 configures HCC3860 to use the HIP / HCC runtime library 3858 and HCC header 3856 to generate host executable code 3870(2) and HCC device executable code 3882, but is not limited to. In at least one embodiment, the HIP / HCC runtime library 3858 corresponds to the HIP runtime API 3832. In at least one embodiment, the HCC header 3856 includes, but is not limited to, any number and type of interoperability mechanisms for HIP and HCC. In at least one embodiment, the host executable code 3870(2) and the HCC device executable code 3882 may be executed on the CPU 3890 and the GPU 3892, respectively.

[0262] Figure 39 shows an exemplary kernel translated by the CUDA-to-HIP translation tool 3820 of Figure 38C, according to at least one embodiment. In at least one embodiment, the CUDA source code 3810 divides an overall problem, which a given kernel is designed to solve, into relatively coarse subproblems that can be solved independently using thread blocks. In at least one embodiment, each thread block contains any number of threads, but is not limited. In at least one embodiment, each subproblem is divided into relatively fine pieces that can be solved in parallel and in conjunction by threads within a thread block. In at least one embodiment, threads within a thread block can be coordinated by sharing data through shared memory and by synchronizing their execution to coordinate memory access.

[0263] In at least one embodiment, the CUDA source code 3810 organizes the thread blocks associated with a given kernel into a one-dimensional grid, two-dimensional grid, or three-dimensional grid of thread blocks. In at least one embodiment, each thread block includes any number of threads, although not limited thereto, and the grid includes any number of thread blocks, although not limited thereto.

[0264] In at least one embodiment, the kernel is a function in device code defined using the "__global__" declaration specifier. In at least one embodiment, the dimensions of the grid that executes the kernel for a given kernel call and associated stream are specified using the CUDA kernel launch syntax 3910. In at least one embodiment, the CUDA kernel launch syntax 3910 is specified as "KernelName<<<GridSize,BlockSize,SharedMemorySize,Stream>>>(KernelArguments);". In at least one embodiment, the execution configuration syntax is the "<<<...>>>" construct inserted between the kernel name ("KernelName") and the list of kernel arguments enclosed in parentheses ("KernelArguments"). In at least one embodiment, the CUDA kernel launch syntax 3910 includes, although not limited thereto, the CUDA launch function syntax instead of the execution configuration syntax.

[0265] In at least one embodiment, "GridSize" is of type dim3 and specifies the dimensions and size of the grid. In at least one embodiment, type dim3 is a CUDA-defined structure containing, but not limited to, unsigned integers x, y, and z. In at least one embodiment, if z is not specified, z defaults to 1. In at least one embodiment, if y is not specified, y defaults to 1. In at least one embodiment, the number of thread blocks in the grid is equal to the product of GridSize.x, GridSize.y, and GridSize.z. In at least one embodiment, "BlockSize" is of type dim3 and specifies the dimensions and size of each thread block. In at least one embodiment, the number of threads per thread block is equal to the product of BlockSize.x, BlockSize.y, and BlockSize.z. In at least one embodiment, each thread executing the kernel is given a unique thread ID accessible within the kernel through an intrinsic variable (e.g., "threadIdx").

[0266] In at least one embodiment, and with respect to CUDA kernel boot syntax 3910, “SharedMemorySize” is an optional argument specifying the number of bytes in shared memory dynamically allocated per thread block for a given kernel call, in addition to statically allocated memory. In at least one embodiment, and with respect to CUDA kernel boot syntax 3910, SharedMemorySize defaults to 0. In at least one embodiment, and with respect to CUDA kernel boot syntax 3910, “Stream” is an optional argument specifying the associated stream, and defaults to 0 to specify a default stream. In at least one embodiment, a stream is a sequence of commands to be executed in order (possibly issued by different host threads). In at least one embodiment, different streams may execute commands out of order or simultaneously with each other.

[0267] In at least one embodiment, CUDA source code 3810 includes, but is not limited to, a kernel definition and a main function for an exemplary kernel "MatAdd". In at least one embodiment, the main function is host code that runs on the host and includes, but is not limited to, kernel calls that cause the kernel MatAdd to run on the device. In at least one embodiment, and as shown, the kernel MatAdd adds two matrices A and B of size N×N, where N is a positive integer, and stores the result in matrix C. In at least one embodiment, the main function defines the threadsPerBlock variable as 16×16 and the numBlocks variable as N / 16×N / 16. In at least one embodiment, the main function then makes the kernel call "MatAdd<<<numBlocks,threadsPerBlock> Specify >>(A,B,C);. In at least one embodiment, and according to the CUDA kernel boot syntax 3910, the kernel MatAdd is executed using a grid of thread blocks having dimensions N / 16 × N / 16, where each thread block has dimensions 16 × 16. In at least one embodiment, each thread block contains 256 threads, and the grid is made up of enough blocks to have one thread per matrix element, and each thread in such a grid executes the kernel MatAdd to perform one pairwise addition.

[0268] In at least one embodiment, while translating CUDA source code 3810 to HIP source code 3830, the CUDA-to-HIP translation tool 3820 translates each kernel call in the CUDA source code 3810 from the CUDA kernel launch syntax 3910 to the HIP kernel launch syntax 3920, and converts any number of other CUDA calls in the source code 3810 to any number of other functionally similar HIP calls. In at least one embodiment, the HIP kernel launch syntax 3920 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 boot syntax 3920 as it does in CUDA kernel boot syntax 3910 (as previously described herein). In at least one embodiment, the arguments SharedMemorySize and Stream are required in HIP kernel boot syntax 3920 but optional in CUDA kernel boot syntax 3910.

[0269] In at least one embodiment, a portion of the HIP source code 3830 shown in Figure 39 is identical to a portion of the CUDA source code 3810 shown in Figure 39, except for the kernel call that causes the kernel MatAdd to execute on the device. In at least one embodiment, the kernel MatAdd is defined in the HIP source code 3830 using the same "__global__" declaration specifier that the kernel MatAdd is defined in the CUDA source code 3810. In at least one embodiment, the kernel call in the HIP source code 3830 is "hipLaunchKernelGGL(MatAdd,numBlocks,threadsPerBlock,0,0,A,B,C);", while the corresponding kernel call in the CUDA source code 3810 is "MatAdd<<<numBlocks,threadsPerBlock> >>(A,B,C);"

[0270] Figure 40 shows in more detail the CUDA-less GPU 3892 of Figure 38C according to at least one embodiment. In at least one embodiment, the GPU 3892 is developed by AMD Corporation in Santa Clara. In at least one embodiment, the GPU 3892 may be configured to perform compute operations in a highly parallel manner. In at least one embodiment, the GPU 3892 is configured to perform graphics pipeline operations, such as drawing commands, pixel operations, geometric calculations, and other operations related to rendering images to a display. In at least one embodiment, the GPU 3892 is configured to perform non-graphics operations. In at least one embodiment, the GPU 3892 is configured to perform both graphics-related and non-graphics operations. In at least one embodiment, the GPU 3892 may be configured to execute device code contained in the HIP source code 3830.

[0271] In at least one embodiment, the GPU 3892 includes, but not limited to, any number of programmable processing units 4020, a command processor 4010, an L2 cache 4022, a memory controller 4070, a DMA engine 4080(1), a system memory controller 4082, a DMA engine 4080(2), and a GPU controller 4084. In at least one embodiment, each programmable processing unit 4020 includes, but not limited to, a workload manager 4030 and any number of compute units 4040. In at least one embodiment, the command processor 4010 reads commands from one or more command queues (not shown) and distributes the commands to the workload manager 4030. In at least one embodiment, for each programmable processing unit 4020, the associated workload manager 4030 distributes the work to the compute units 4040 contained within the programmable processing unit 4020. In at least one embodiment, each compute unit 4040 may execute any number of thread blocks, but each thread block executes on a single compute unit 4040. In at least one embodiment, a workgroup is a thread block.

[0272] In at least one embodiment, each compute unit 4040 includes, but not limited to, any number of SIMD units 4050 and shared memory 4060. In at least one embodiment, each SIMD unit 4050 may implement a SIMD architecture and be configured to perform operations in parallel. In at least one embodiment, each SIMD unit 4050 includes, but not limited to, a vector ALU 4052 and a vector register file 4054. In at least one embodiment, each SIMD unit 4050 performs a different warp. In at least one embodiment, a warp is a group of threads (e.g., 16 threads), where each thread in the warp belongs to a single thread block and is configured to process a different set of data based on a single set of instructions. In at least one embodiment, predication may be used to disable one or more threads in a warp. 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, the wavefront is a warp. In at least one embodiment, different wavefronts in a thread block may be synchronized with each other and communicate via shared memory 4060.

[0273] In at least one embodiment, the programmable processing unit 4020 is referred to as the “shader engine”. In at least one embodiment, each programmable processing unit 4020 includes, but is not limited to, any amount of dedicated graphics hardware in addition to the compute unit 4040. In at least one embodiment, each programmable processing unit 4020 includes, but is not limited to, any number of geometry processors (including zero), any number of rasterizers (including zero), any number of render backends (including zero), a workload manager 4030, and any number of compute units 4040.

[0274] In at least one embodiment, compute units 4040 share L2 cache 4022. In at least one embodiment, L2 cache 4022 is partitioned. In at least one embodiment, GPU memory 4090 is accessible by all compute units 4040 in GPU 3892. In at least one embodiment, memory controllers 4070 and system memory controller 4082 facilitate data transfer between GPU 3892 and a host, and DMA engine 4080(1) enables asynchronous memory transfer between GPU 3892 and such a host. In at least one embodiment, memory controllers 4070 and GPU controller 4084 facilitate data transfer between GPU 3892 and other GPU 3892s, and DMA engine 4080(2) enables asynchronous memory transfer between GPU 3892 and other GPU 3892s.

[0275] In at least one embodiment, the GPU3892 includes any amount and type of system interconnects that facilitate data and control transmissions across any number and type of directly or indirectly linked components, which may be internal to or external to the GPU3892. In at least one embodiment, the GPU3892 includes any number and type of I / O interfaces (e.g., PCIe) that are coupled to any number and type of peripheral devices, which may not be limited. In at least one embodiment, the GPU3892 may include any number of display engines (including zero) and any number of multimedia engines (including zero), which may not be limited. In at least one embodiment, the GPU3892 implements a memory subsystem that includes any amount and type of memory controllers (e.g., memory controller 4070 and system memory controller 4082) and memory devices (e.g., shared memory 4060) that may be dedicated to one component or shared among multiple components. In at least one embodiment, the GPU3892 implements a cache subsystem including, but not limited to, one or more cache memories (e.g., L2 cache 4022), each of which may be private to any number of components (e.g., SIMD unit 4050, compute unit 4040, and programmable processing unit 4020) or shared among any number of components.

[0276] Figure 41 shows how threads of an exemplary CUDA grid 4120 are mapped to different compute units 4040 in Figure 40, according to at least one embodiment. In at least one embodiment, and for illustrative purposes only, the grid 4120 has a GridSize of BX × BY × 1 and a BlockSize of TX × TY × 1. In at least one embodiment, the grid 4120 thus includes, but not limited to, (BX * BY) thread blocks 4130, and each thread block 4130 includes, but not limited to, (TX * TY) threads 4140. The threads 4140 are illustrated in Figure 41 as squiggly arrows.

[0277] In at least one embodiment, grid 4120 is mapped to a programmable processing unit 4020(1) including, but not limited to, compute units 4040(1) to 4040(C). In at least one embodiment, and as shown, (BJ*BY) thread blocks 4130 are mapped to compute unit 4040(1), and the remaining thread blocks 4130 are mapped to compute unit 4040(2). In at least one embodiment, each thread block 4130 may contain, but not limited to, any number of warps, each warp being mapped to a different SIMD unit 4050 in Figure 40.

[0278] In at least one embodiment, warps in a given thread block 4130 can synchronize with each other and communicate through a shared memory 4060 contained in the associated compute unit 4040. For example, and in at least one embodiment, warps in thread block 4130(BJ,1) can synchronize with each other and communicate through a shared memory 4060(1). For example, and in at least one embodiment, warps in thread block 4130(BJ+1,1) can synchronize with each other and communicate through a shared memory 4060(2).

[0279] Figure 42 shows how to migrate existing CUDA code to Data Parallel C++ code using at least one embodiment. Data Parallel C++ (DPC++) can refer to an open, standards-based alternative to a single-architecture, proprietary language, which allows developers to reuse code across hardware targets (CPUs as well as accelerators such as GPUs and FPGAs) and to implement custom adjustments for specific accelerators. DPC++ uses similar and / or identical C and C++ constructs that adhere to ISO C++, which developers may be familiar with. DPC++ incorporates the standard SYCL from the Khronos Group to support data parallelism and heterogeneous programming. SYCL refers to a cross-platform abstraction layer based on the underlying concepts, portability, and efficiency of OpenCL, which allows code for heterogeneous processors to be written in a “single-source” style using standard C++. SYCL enables single-source development where C++ template functions contain both host and device code, allowing for the construction of complex algorithms that utilize OpenCL acceleration, and then enabling their reuse across their entire source code for different types of data.

[0280] 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 a DPC++ application that can be deployed across a variety of hardware targets, and DPC++ compatibility tools may be used to migrate a CUDA application into a DPC++ multi-platform program. In at least one embodiment, a DPC++ base toolkit includes a DPC++ compiler for deploying applications across a variety of hardware targets, DPC++ libraries for increasing productivity and performance across CPUs, GPUs, and FPGAs, DPC++ compatibility tools for migrating CUDA applications into multi-platform applications, and any preferred combination thereof.

[0281] In at least one embodiment, the DPC++ programming model is used for one or more aspects related to simply programming CPUs and accelerators by using modern C++ features to represent parallel processing using a programming language called Data Parallel C++. The DPC++ programming language is used for code reuse for hosts (e.g., CPUs) and accelerators (e.g., GPUs or FPGAs), using a single source language, and execution and memory dependencies can be clearly communicated. Mappings within DPC++ code can be used to migrate applications to run on the hardware or set of hardware devices that best accelerate the workload. Even on platforms without available accelerators, the host may be available to simplify the development and debugging of device code.

[0282] In at least one embodiment, CUDA source code 4200 is provided as input to a DPC++ compatibility tool 4202 to generate human-readable DPC++ 4204. In at least one embodiment, human-readable DPC++ 4204 includes inline comments generated by the DPC++ compatibility tool 4202, which guide the developer on how and / or where to modify the DPC++ code to complete the coding and adjustments to the desired performance 4206, thereby generating DPC++ source code 4208.

[0283] In at least one embodiment, CUDA source code 4200 is a collection of human-readable source code in the CUDA programming language, or includes such a collection. In at least one embodiment, CUDA source code 4200 is human-readable source 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, but is not limited to, a mechanism for defining device code and distinguishing device code from host code. In at least one embodiment, device code is source code that, after compilation, is executable on a device (e.g., a GPU or FPGA) and may run on one or more processor cores of the device, or may include a more parallelizable workflow. In at least one embodiment, the device may be a processor optimized for parallel instruction processing, such as a CUDA-enabled GPU, a GPU, or another GPGPU. In at least one embodiment, host code is source code that, after compilation, is executable on a host. In at least one embodiment, some or all of the host code and device code may be executed in parallel across the CPU and GPU / FPGA. In at least one embodiment, the host is a processor optimized for sequential instruction processing, such as a CPU. The CUDA source code 4200 described with respect to Figure 42 may conform to the CUDA source code described elsewhere in this specification.

[0284] In at least one embodiment, DPC++ Compatibility Tool 4202 refers to an executable tool, program, application, or any other suitable type of tool used to facilitate the migration of CUDA source code 4200 to DPC++ source code 4208. In at least one embodiment, DPC++ Compatibility Tool 4202 is a command-line based code migration tool available as part of a DPC++ toolkit used to port existing CUDA source to DPC++. In at least one embodiment, DPC++ Compatibility Tool 4202 converts some or all of the source code of a CUDA application from CUDA to DPC++ and generates a resulting file, called Human-Readable DPC++4204, which is at least partially written in DPC++. In at least one embodiment, Human-Readable DPC++4204 includes comments generated by DPC++ Compatibility Tool 4202 to indicate where user intervention may be required. In at least one embodiment, user intervention is required when CUDA source code 4200 calls a CUDA API that does not have a similar DPC++ API; other instances where user intervention is required will be described in more detail later.

[0285] In at least one embodiment, the workflow for migrating CUDA source code 4200 (e.g., an application or a portion thereof) includes creating one or more compile database files, migrating CUDA to DPC++ using the DPC++ compatibility tool 4002, completing and validating the migration to thereby generate DPC++ source code 4208, and compiling the DPC++ source code 4208 using a DPC++ compiler to generate a DPC++ application. In at least one embodiment, the compatibility tool provides a utility that intercepts the commands used when a Makefile is executed and stores them in a compile database file. In at least one embodiment, the file is stored in JSON format. In at least one embodiment, the intercept-built command converts Makefile commands into DPC-compatible commands.

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

[0287] In at least one embodiment, the DPC++ Compatibility Tool 4202 migrates CUDA code written in CUDA (e.g., an application) to DPC++ by generating DPC++ whenever possible. In at least one embodiment, the DPC++ Compatibility Tool 4202 is available as part of a toolkit. In at least one embodiment, the DPC++ toolkit includes the intercept-build tool. In at least one embodiment, the intercept-built tool creates a compilation database that captures compilation commands to migrate CUDA files. In at least one embodiment, the compilation database generated by the intercept-built tool is used by the DPC++ Compatibility Tool 4202 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, the DPC++ Compatibility Tool 4202 generates human-readable DPC++ 4204, which may be DPC++ code that, when generated by the DPC++ Compatibility Tool 4202, may not compile with the DPC++ compiler and may require additional plumbing to identify the parts of the code that were not migrated correctly, potentially involving manual intervention by the developer. In at least one embodiment, the DPC++ Compatibility Tool 4202 provides hints or tools embedded in the code to help the developer manually migrate additional code that may not be automatically migrated. In at least one embodiment, migration is a one-time activity for a source file, project, or application.

[0288] In at least one embodiment, the DPC++ compatibility tool 42002 is capable of successfully migrating all parts of CUDA code to DPC++, with only optional steps for manually verifying and adjusting the performance of the generated DPC++ source code. In at least one embodiment, the DPC++ compatibility tool 4202 directly generates DPC++ source code 4208 that is compiled by a DPC++ compiler without requiring or utilizing human intervention to modify the DPC++ code generated by the DPC++ compatibility tool 4202. In at least one embodiment, the DPC++ compatibility tool generates compilable DPC++ code that can be optionally adjusted by the developer for performance, readability, maintainability, various other considerations, or any combination thereof.

[0289] In at least one embodiment, one or more CUDA source files are migrated to DPC++ source files using the DPC++ compatibility tool 4202, at least partially. In at least one embodiment, the CUDA source code includes one or more header files, which may include CUDA header files. In at least one embodiment, the CUDA source files are<cuda.h> Header files and can be used to print text.<stdio.h> This includes a header file. In at least one embodiment, a portion of the vector addition kernel CUDA source file may be written as follows, or relating to the following: #include<cuda.h> #include<stdio.h> #define VECTOR_SIZE 256 [] global__ void VectorAddKernel(float* A,float* B,float* C) { A[threadIdx.x]=threadIdx.x+1.0f; B[threadIdx.x]=threadIdx.x+1.0f; C[threadIdx.x]=A[threadIdx.x]+B[threadIdx.x]; } int main() { float *d_A,*d_B,*d_C; cudaMalloc(&d_A,VECTOR_SIZE*sizeof(float)); cudaMalloc(&d_B,VECTOR_SIZE*sizeof(float)); cudaMalloc(&d_C,VECTOR_SIZE*sizeof(float)); VectorAddKernel<<<1,VECTOR_SIZE>>>(d_A,d_B,d_C); float Result[VECTOR_SIZE]={}; cudaMemcpy(Result,d_C,VECTOR_SIZE*sizeof(float),cudaMemcpyDeviceToHost); cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); for(int i=0; i<VECTOR_SIZE; i++ { if(i%16==0){ printf(“\n”); } printf(“%f”,Result[i]); } return 0; }

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

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

[0292] In at least one embodiment, and with respect to the CUDA source files presented above, memory allocation is migrated. In at least one embodiment, cudaMalloc() is migrated to malloc_device(), a unified shared memory SYCL call to which the device and context are passed, relying on the SYCL concept, including platform, device, context, and queues. In at least one embodiment, the SYCL platform may have multiple devices (e.g., host and GPU devices), each device may have multiple queues to which jobs can be submitted, each device may have a context, and a context may have multiple devices and manage shared memory objects.

[0293] In at least one embodiment, and with respect to the CUDA source file presented above, the main() function calls or calls VectorAddKernel() to add two vectors A and B to each other and store the result in vector C. In at least one embodiment, the CUDA code for calling VectorAddKernel() is replaced by DPC++ code for submitting the kernel to a command queue for execution. In at least one embodiment, the command group handler cgh passes the data, synchronization, and computation to be submitted to the queue, and parallel_for is called for the number of global elements and the number of work items in the work group where VectorAddKernel() is called.

[0294] In at least one embodiment, and with respect to the CUDA source files presented above, CUDA calls to copy device memory and then free memory for vectors A, B, and C are migrated to corresponding DPC++ calls. In at least one embodiment, C++ code (e.g., standard ISO C++ code for printing a vector of floating-point variables) is migrated as is without modification by the DPC++ compatibility tool 4202. In at least one embodiment, the DPC++ compatibility tool 4202 modifies the CUDA API for memory setup and / or host calls to run the kernel on an accelerating device. In at least one embodiment, and with respect to the CUDA source files presented above, the corresponding human-readable DPC++ 4204 (e.g., one that can be compiled) is written as follows or relates to the following: #include<CL / sycl.hpp> #include<dpct / dpct.hpp> #define VECTOR_SIZE 256 void VectorAddKernel(float* A,float* B,float* C,sycl::nd_item <3> item_ct1) { A[item_ct1.get_local_id(2)]=item_ct1.get_local_id(2)+1.0f; B[item_ct1.get_local_id(2)]=item_ct1.get_local_id(2)+1.0f; C[item_ct1.get_local_id(2)]= A[item_ct1.get_local_id(2)]+B[item_ct1.get_local_id(2)]; } int main() { float *d_A,*d_B,*d_C; d_A=(float *)sycl::malloc_device(VECTOR_SIZE * sizeof(float), dpct::get_current_device(), dpct::get_default_context()); d_B=(float *)sycl::malloc_device(VECTOR_SIZE * sizeof(float), dpct::get_current_device(), dpct::get_default_context()); d_C=(float *)sycl::malloc_device(VECTOR_SIZE * sizeof(float), dpct::get_current_device(), dpct::get_default_context()); dpct::get_default_queue_wait().submit([&](sycl::handler&cgh){ cgh.parallel_for( sycl::nd_range<3>( sycl::range<3>(1,1,1) * sycl::range<3>(1,1,VECTOR_SIZE) * sycl::range<3>(1,1,VECTOR_SIZE)), [=](sycl::nd_items<3> item_ct1){ VectorAddKernel(d_A,d_B,d_C,item_ct1); }); }); float Result[VECTOR_SIZE]={}; dpct::get_default_queue_wait() .memcpy(Result,d_C,VECTOR_SIZE * sizeof(float)) .wait(); sycl::free(d_A,dpct::get_default_context()); sycl::free(d_B,dpct::get_default_context()); sycl::free(d_C,dpct::get_default_context()); for(int i=0; i<VECTOR_SIZE; i++ { if(i%16==0){ printf(“\n”); } printf(“%f”,Result[i]); } return 0; }

[0295] In at least one embodiment, human-readable DPC++ 4204 refers to the output generated by the DPC++ compatibility tool 4202, which may be optimized in one or another manner. In at least one embodiment, human-readable DPC++ 4204 generated by the DPC++ compatibility tool 4202 may be manually edited by a developer after migration to make it more maintainable, for performance, or for other considerations. In at least one embodiment, DPC++ code generated by the DPC++ compatibility tool 42002, such as the disclosed DPC++, may 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, the DPC++ code generated above uses a three-dimensional nd_range, which may be refactored to use only a single dimension, thereby reducing memory usage. In at least one embodiment, a developer may manually edit the DPC++ code generated by the DPC++ compatibility tool 4202 and replace the use of unified shared memory with accessors. In at least one embodiment, the DPC++ compatibility tool 4202 has an option to change how it migrates CUDA code to DPC++ code. In at least one embodiment, the DPC++ compatibility tool 4202 is redundant because it uses a generic template for migrating CUDA code to DPC++ code that works for a large number of cases.

[0296] In at least one embodiment, the CUDA to DPC++ migration workflow includes steps of preparing for migration using an intercept-build script, performing the migration of a CUDA project to DPC++ using the DPC++ compatibility tool 4202, manually reviewing and editing the migrated source files for completion and validity, and compiling the final DPC++ code to generate a DPC++ application. In at least one embodiment, manual review of the DPC++ source code may be required in one or more scenarios, but not limited to, that the migrated API does not return error codes (CUDA code can return error codes, which can then be consumed by the application, whereas SYCL uses exceptions to report errors and therefore does not use error codes to surface errors), that CUDA compute capability-dependent logic is not supported by DPC++, and that statements may not be deleted. In at least one embodiment, scenarios in which DPC++ code requires manual intervention may include, but are not limited to, error code logic being replaced with (*,0) code or commented out, an equivalent DPC++ API not being available, CUDA compute capability-dependent logic, hardware-dependent APIs (clock()), missing features, unsupported APIs, execution time measurement logic, dealing with built-in vector type conflicts, and migration of cuBLAS APIs.

[0297] Other variations are within the scope of this disclosure. Thus, the disclosed techniques can be modified and constructed in various ways, some exemplary embodiments of which are shown in the drawings and described in detail above. However, this disclosure is not intended to limit itself to any particular one or more disclosed forms, but rather to encompass all modifications, alternative constructions, and equivalents that fall within the spirit and scope of the disclosure, as defined in the appended claims.

[0298] In the context describing the disclosed embodiments (particularly in the context of the following claims), the terms “a,” “an,” and “the,” and similar demonstrative pronouns, should be interpreted as encompassing both singular and plural, and not as definitions of terms, unless otherwise stated herein or clearly refuted by the context. The terms “comprising,” “having,” “including,” and “containing” should be interpreted as open-ended terms (meaning “including, but not limited to,”) unless otherwise stated. The term “connected,” when unmodified and referring to a physical connection, should be interpreted as being partially or completely contained, attached, or joined to one another, even if there is something intervening. The detailing of ranges of values ​​herein is merely intended to serve as a concise way of individually referring to each distinct value that falls within a range, unless otherwise stated herein and unless each distinct value is incorporated into the specification as if it were individually detailed herein. The use of the terms “set” (for example, “set of items”) or “subset” should be interpreted as a non-empty set comprising one or more members, unless otherwise stated or denied by the context. Furthermore, unless otherwise stated or denied by the context, the term “subset” of a corresponding set does not necessarily refer to a strict subset of the corresponding set, and a subset and a corresponding set can be equivalent.

[0299] Combinations such as “at least one of A, B, and C” or “at least one of A, B, and C” are understood in contexts generally used to indicate that an item, term, etc., can be either A, B, or C, or any non-empty subset of the set of A, B, and C, unless otherwise specifically stated or explicitly denied by the context. For example, in a descriptive example of a set having three members, the combinations “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 combinations do not imply as a whole that some embodiments require the presence of each of A, B, and C. Furthermore, unless otherwise stated or negated by the context, the term "plurality" refers to the state of being multiple (for example, "a plurality of items" refers to multiple items). The number of items that are plural is at least two, but can be more when explicitly or by contextual indication. Furthermore, unless otherwise stated or clarified by the context, the phrase "based on" means "at least partially based on," and does not mean "based solely on."

[0300] The operation of the processes described herein may be carried out in any preferred order unless otherwise stated herein or explicitly refuted by the context. In at least one embodiment, a process such as the processes described herein (or variations and / or combinations thereof) is carried out under the control of one or more computer systems consisting of executable instructions and is implemented as code (e.g., executable instructions, one or more computer programs, or one or more applications) that is executed collectively on one or more processors, by hardware, or by a combination thereof. In at least one embodiment, the code is stored in a computer-readable storage medium, for example, in the form of a computer program comprising multiple instructions executable by one or more processors. In at least one embodiment, the computer-readable storage medium is a non-temporary computer-readable storage medium that excludes temporary signals (e.g., transient electrical or electromagnetic transmissions that propagate) but includes non-temporary data storage circuit elements (e.g., buffers, caches, and queues) in transceivers for temporary signals. In at least one embodiment, code (e.g., executable code or source code) is stored in a set of one or more non-temporary computer-readable storage media which, when executed by one or more processors of a computer system (i.e., as a result of execution), causes the computer system to perform the operations described herein (or has other memory for storing executable instructions). In at least one embodiment, the set of non-temporary computer-readable storage media comprises a plurality of non-temporary computer-readable storage media which one or more of the individual non-temporary storage media of the plurality of non-temporary computer-readable storage media do not contain all of the code, but the plurality of non-temporary computer-readable storage media collectively contain all of the code.In at least one embodiment, executable instructions are executed such that different instructions are executed by different processors, for example, a non-temporary computer-readable storage medium stores the instructions, the main central processing unit ("CPU") executes some of the instructions, and the graphics processing unit ("GPU") executes others. In at least one embodiment, different components of a computer system have separate processors, and different processors execute different subsets of instructions.

[0301] Accordingly, in at least one embodiment, the computer system is configured to implement one or more services that individually or collectively perform the operations of the processes described herein, and such a computer system consists of applicable hardware and / or software that enables the performance of the operations. Furthermore, a computer system implementing at least one embodiment of the present disclosure is a single device, and in another embodiment, a distributed computer system comprising multiple devices operating in different ways so that the distributed computer system performs the operations described herein and the single device does not perform all of the operations.

[0302] Any use of any examples or illustrative language provided herein (e.g., "such as") is intended solely to further illustrate the embodiments of this disclosure and, unless otherwise asserted, does not limit the scope of this disclosure. Nothing in this specification should be construed as indicating any unclaimed element as essential to the practice of this disclosure.

[0303] All references cited herein, including publications, patent applications, and patents, are incorporated herein by reference to the same extent as if they were included herein in their entirety, provided that each reference is clearly indicated individually to be incorporated by reference.

[0304] In the specification and claims, the terms “joined” and “connected” may be used together with their derivatives. It should be understood that these terms may not be intended to be synonymous with one another. Rather, in certain instances, “connected” or “joined” may be used to indicate that two or more elements are in direct or indirect physical or electrical contact with one another. “Joined” may also mean that two or more elements are not in direct contact with one another, but still interact or communicate with one another.

[0305] Unless otherwise specified, throughout this specification, terms such as “processing,” “computing,” “calculating,” or “determining” refer to the actions and / or processes of a computer or computing system or similar electronic computing device that manipulate and / or transform data, represented as electronic or other physical quantities, in the registers and / or memory of a computing system into other data, similarly represented as physical quantities in the memory, registers, or other such information storage, transmission, or display device of a computing system.

[0306] Similarly, the term “processor” may refer to any device or part of a device that processes electronic data from registers and / or memory and converts that electronic data into other electronic data that can be stored in registers and / or memory. In non-limiting examples, “processor” may be a CPU or a GPU. A “computing platform” may comprise one or more processors. As used herein, “software” processes may include software and / or hardware entities that perform work over time, such as tasks, threads, and intelligent agents. Each process may also refer to multiple processes for executing instructions serially or in parallel, continuously or intermittently. The terms “system” and “method” are used interchangeably herein only if one or more methods can be embodied by a system, and a method can be considered a system.

[0307] This specification may refer to acquiring, obtaining, receiving, or inputting analog or digital data into subsystems, computer systems, or computer-implemented machines. The process of acquiring, obtaining, receiving, or inputting analog or digital data can be implemented in various ways, such as by receiving data as parameters to function calls or calls to application programming interfaces. In some implementations, the process of acquiring, obtaining, receiving, or inputting analog or digital data can be implemented by transferring data via serial or parallel interfaces. In other implementations, the process of acquiring, obtaining, receiving, or inputting analog or digital data can be implemented by transferring data via a computer network from a providing entity to a receiving entity. This specification may also refer to providing, outputting, transmitting, sending out, or presenting analog or digital data. In various examples, the process of providing, outputting, transmitting, sending out, or presenting analog or digital data can be implemented by transferring data as input or output parameters to function calls, application programming interfaces, or parameters to inter-process communication mechanisms.

[0308] While the above description outlines exemplary implementations of the techniques described, other architectures may be used to implement the described functionality and are intended to be within the scope of this disclosure. Furthermore, while specific distributions of responsibility are defined for illustrative purposes, various functions and responsibilities may be distributed and divided in different ways depending on the context.

[0309] Furthermore, while the subject matter is described in language specific to structural features and / or methodological actions, it should be understood that the subject matter claimed in the attached claims is not necessarily limited to the described specific features or actions. Rather, the specific features and actions are disclosed as exemplary forms that implement the claims.

Claims

1. A method performed by one or more processors, A step of determining a first set of values, wherein the first set of values ​​represents one or more tonal ranges corresponding to an input image, The steps include generating a contrast image based at least partially on the difference between the input image and a second image generated by applying at least a smoothing factor, A step of generating a set of contrast gain curves corresponding to one or more tonal ranges based at least partially on the first set of values ​​and the second set of values, A step of generating an output image by applying a contrast gain map to the contrast image, wherein the contrast gain map is determined for a set of pixels in the input image by applying a first contrast gain curve from the set of contrast gain curves to a first scale space representation of the input image. Methods that include...

2. The method according to claim 1, wherein the second image further includes a scale-space representation of the input image.

3. The method according to claim 1, wherein the smoothing factor is a Gaussian kernel applied to the input image.

4. The method according to claim 1, further comprising the step of applying a global tone map to the input image in order to generate a modified input image.

5. The method according to claim 4, wherein the one or more ranges include a set of ranges that are represented as a histogram.

6. The method according to claim 5, further comprising the step of generating the histogram based at least partially on the modified input image.

7. The method according to claim 1, further comprising the step of applying a global tone map to the input image.

8. The method according to claim 1, wherein the second set of values ​​includes a set of sigma values ​​representing one or more Gaussian curves.

9. The method according to claim 1, wherein the one or more tonal ranges further include at least one of a shadow range, a highlight range, or a midtone range.

10. The method according to claim 9, wherein the first set of values ​​further includes a set of mean values ​​of the one or more tonal ranges.

11. The method according to claim 10, wherein the shadow range further includes a first set of pixels of the input image that fall below a first threshold.

12. The method according to claim 11, wherein the first average of the set of averages further includes the average of the first set of pixels.

13. The method according to claim 12, wherein the highlight range further includes a second set of pixels of the input image that exceed a second threshold.

14. The method according to claim 13, wherein the intermediate tone range further includes a third set of pixels of the input image, wherein the intermediate tone range is above the first threshold and below the second threshold.

15. The method according to claim 14, wherein the first threshold value represents the proportion of pixels in the input image.

16. The method according to claim 1, wherein the contrast gain curve among the set of contrast gain curves is a function of the input tone value for producing the corresponding gain.

17. The method according to claim 1, wherein the one or more tonal ranges further include a tone range defined at least partially on groups of pixels classified by a threshold pixel value, wherein the threshold pixel value is determined at least partially on a percentage of pixels in the group of pixels.

18. One or more processors, Memory for storing instructions and A system comprising, wherein the instruction, as a result of being executed by one or more processors, Obtaining image data representing the input image, Determining one or more tone ranges and a first set of values ​​based at least partially on the aforementioned image data, wherein a first value in the first set of values ​​corresponds to a first tone range among the one or more tone ranges. A contrast image is generated based at least partially on a set of scale space representations of the input image and a set of luminance values ​​of the input image. To generate a first contrast gain function for the first tone range, At a minimum, an output image is generated by applying a first contrast gain map to the contrast image, wherein the first contrast gain map is generated at least partially based on the first contrast gain function. A system that enables this to happen.

19. The system according to claim 18, wherein the image data is a histogram of the input image.

20. The system according to claim 18, wherein the image data is high dynamic range (HDR) image data.

21. The system according to claim 18, further comprising instructions that cause the system to generate a second contrast gain function for a second tone range of a set of tone ranges, as a result of the memory being executed by the one or more processors.

22. The system according to claim 21, wherein the memory further includes instructions causing the system to generate a second contrast gain map based at least partially on the second contrast gain function, as a result of being executed by the one or more processors.

23. The system according to claim 22, further comprising an instruction causing the system to generate the output image by applying at least the first contrast gain map and the second contrast gain map, as a result of the instruction causing the system to generate the output image by applying at least the first contrast gain map and the second contrast gain map, being executed by one or more processors.

24. The system according to claim 18, wherein the one or more tone ranges further include at least one of a shadow range, a highlight range, or a midtone range.

25. The system according to claim 18, wherein the image data further includes statistical image data representing an input image.

26. The system according to claim 25, wherein the statistical image data further includes a histogram.

Images