Application programming interface for indicating objects intersected by light rays

Abstract

The invention relates to an application programming interface for indicating objects intersected by rays. Apparatuses, systems, and techniques for indicating one or more objects intersected by one or more rays are disclosed. In at least one embodiment, one or more circuitry is to execute an application programming interface (API) such that one or more objects intersecting by one or more rays are indicated to a user.

Description

TECHNICAL FIELD

[0001] At least one embodiment relates to processing resources for executing one or more programs written for a parallel computing platform and application programming interface. For example, at least one embodiment relates to executing an application programming interface (API) to indicate an object intersected by a ray according to various novel techniques described herein. BACKGROUND

[0002] Performing operations using rays can consume a large amount of time, power, or computing resources. The amount of time, power, or computing resources can be improved. BRIEF DESCRIPTION OF DRAWINGS

[0003] Figure 1 is a block diagram illustrating a system, in accordance with at least one embodiment;

[0004] Figure 2 is a block diagram illustrating a computing environment, in accordance with at least one embodiment;

[0005] Figure 3 illustrates an API for performing object intersection operations, in accordance with at least one embodiment;

[0006] Figure 4 illustrates a flow diagram of techniques for performing object intersection operations, in accordance with at least one embodiment;

[0007] Figure 5 is a flow diagram of techniques for performing object intersection operations, in accordance with at least one embodiment;

[0008] Figure 6 is a block diagram for illustrating an example of a processor, in accordance with at least one embodiment;

[0009] Figure 7 is a block diagram for illustrating a driver and / or runtime environment, in accordance with at least one embodiment;

[0010] Figure 8 illustrates an example data center, in accordance with at least one embodiment;

[0011] Figure 9 illustrates a processing system, in accordance with at least one embodiment;

[0012] Figure 10 illustrates a computer system, in accordance with at least one embodiment;

[0013] Figure 11 illustrates a system, in accordance with at least one embodiment;

[0014] Figure 12 illustrates an example integrated circuit, in accordance with at least one embodiment;

[0015] Figure 13 A computing system is shown in accordance with at least one embodiment;

[0016] Figure 14 An APU is shown in accordance with at least one embodiment;

[0017] Figure 15 A CPU is shown in accordance with at least one embodiment;

[0018] Figure 16 An exemplary accelerator integration slice is shown in accordance with at least one embodiment;

[0019] Figures 17A-17B An exemplary graphics processor is shown in accordance with at least one embodiment;

[0020] Figure 18A A graphics core is shown in accordance with at least one embodiment;

[0021] Figure 18B A GPGPU is shown in accordance with at least one embodiment;

[0022] Figure 19A A parallel processor is shown in accordance with at least one embodiment;

[0023] Figure 19B A processing cluster is shown in accordance with at least one embodiment;

[0024] Figure 19C A graphics multiprocessor is shown in accordance with at least one embodiment;

[0025] Figure 20 A graphics processor is shown in accordance with at least one embodiment;

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

[0027] Figure 22 A processor is shown in accordance with at least one embodiment;

[0028] Figure 23 A graphics processor core is shown in accordance with at least one embodiment;

[0029] Figure 24 A PPU is shown in accordance with at least one embodiment;

[0030] Figure 25 A GPC is shown in accordance with at least one embodiment;

[0031] Figure 26 A streaming multiprocessor is shown in accordance with at least one embodiment;

[0032] Figure 27 A software stack of a programming platform according to at least one embodiment is shown;

[0033] Figure 28 The illustration shows an embodiment according to at least one of the embodiments. Figure 27 The CUDA implementation of the software stack;

[0034] Figure 29 The illustration shows an embodiment according to at least one of the embodiments. Figure 27 The ROCm implementation of the software stack;

[0035] Figure 30 The illustration shows an embodiment according to at least one of the embodiments. Figure 27 The OpenCL implementation of the software stack;

[0036] Figure 31 Software supported by a programming platform according to at least one embodiment is shown;

[0037] Figure 32 The illustration shows an embodiment of at least one of the following: Figures 27-30 Compiled code executed on the programming platform;

[0038] Figure 33 The illustration shows an embodiment of at least one of the following: Figures 27-30 More detailed compiled code executed on the programming platform;

[0039] Figure 34 This illustrates the transformation of source code before compilation, according to at least one embodiment;

[0040] Figure 35A A system configured to compile and execute CUDA source code using different types of processing units, according to at least one embodiment, is shown;

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

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

[0043] Figure 36 The diagram illustrates a method according to at least one embodiment. Figure 35C An example kernel converted by the CUDA to HIP conversion tool;

[0044] Figure 37 A more detailed description is provided according to at least one embodiment.Figure 35C GPUs without CUDA enabled; and

[0045] Figure 38 This illustrates how threads of an exemplary CUDA grid, according to at least one embodiment, are mapped to... Figure 37 Different computational units;

[0046] Figure 39 This illustrates how to migrate existing CUDA code to data-parallel C++ code according to at least one embodiment;

[0047] Figure 40 Components of a system for accessing a large language model according to at least one embodiment are shown. Detailed Implementation

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

[0049] Figure 1 This is a block diagram illustrating a system 100 according to at least one embodiment. In at least one embodiment, system 100 includes a computing device 102 comprising a central processing unit (CPU) 104 and a parallel processing unit (PPU) 106 (e.g., an accelerator, such as a graphics processing unit (GPU), a field-programmable gate array (FPGA), an application-specific integrated circuit (ASIC), and / or some other suitable device). In at least one embodiment, computing device 102 includes one or more other components (not shown for clarity), such as main memory used by CPU 104, storage devices, one or more networking components, one or more additional CPUs, and / or one or more additional PPUs. In at least one embodiment, API 108 provides at least one function accessible to application 110. In at least one embodiment, compiler 112 translates requests received via API 108 into instructions executable on PPU 106 (e.g., instructions as part of the instruction set architecture of PPU 106). In at least one embodiment, a user (e.g., computer program code, such as a kernel running on PPU 106) provides one or more inputs to one or more APIs and / or instructions to API 108.

[0050] In at least one embodiment, PPU 106 includes processors 114 and 116. In at least one embodiment, one or more of processors 114 and 116 are streaming multi-processors (SMs) or some other suitable type of processor. In at least one embodiment, a streaming multi-processor (SM) is referred to as a compute unit, execution unit, sub-slice, shader engine, single instruction multiple data (SIMD) processor, SIMD unit, and / or some other suitable terminology. In at least one embodiment, although two processors are illustrated, it is to be understood that PPU 106 can include different numbers of processors. In at least one embodiment, PPU 106 includes memory 118. In at least one embodiment, memory 118 includes global memory, shared memory, and / or some other suitable type of memory. In at least one embodiment, shared memory is referred to as local memory. In at least one embodiment, although one memory is illustrated, it is to be understood that PPU 106 can include different numbers of memories. In at least one embodiment, a set of threads executed by processors 114 and / or processors 116 can access memory 118.

[0051] In at least one embodiment, PPU 106 includes ray tracing (RT) hardware (H / W) 120. In at least one embodiment, ray tracing hardware is referred to as a ray tracing (RT) core, a tree traversal unit (TTU), a ray tracing unit (RTU), a ray accelerator, and / or some other suitable name. In at least one embodiment, RT hardware 120 includes one or more RT cores and / or one or more TTUs.

[0052] In at least one embodiment, API 108 includes functions and / or instructions for generating and / or encoding tree data structure 122. In at least one embodiment, tree data structure 122 is a bounding volume hierarchy (BVH), an n-ary tree, or some other suitable tree data structure. In at least one embodiment, tree data structure 122 is a hierarchical spatial data structure. In at least one embodiment, API 108 includes functions for generating more than one type of tree data structure. In at least one embodiment, API 108 includes functions and / or instructions for generating and / or encoding lookup information 124. In at least one embodiment, lookup information 124 is information that specifies a ray. In at least one embodiment, lookup information 124 corresponds to an object and / or information to be found in another data structure. In at least one embodiment, a user (e.g., computer program code executed by CPU 104 and / or PPU 106) need not generate tree data structure 122 and / or lookup information 124 using API 108.

[0053] In at least one embodiment, API 108 includes functions and / or instructions to indicate to a user one or more objects intersected by one or more rays. In at least one embodiment, functions and / or instructions of API 108 to indicate to a user one or more objects intersected by one or more rays use ray tracing hardware 120 to identify one or more intersected objects based at least in part on tree data structure 122 and lookup information 124. In at least one embodiment, functions to indicate one or more objects intersected by one or more rays indicate these objects based at least in part on storing one or more identifiers of intersected objects in object information 126. In at least one embodiment, object information 126 includes a data structure including one or more identifiers of corresponding one or more intersected objects. In at least one embodiment, intersected objects can be triangles. In at least one embodiment, intersected objects can be triangles in a mesh. In at least one embodiment, tree data structure 122 includes triangles as leaf nodes, and object intersection information 126 indicates triangles that intersect these leaf nodes with lookup information 124 (e.g., rays).

[0054] In at least one embodiment, intersecting objects can not be graphics, but rather information that a programmer causes to be represented. In at least one embodiment, a set of elements to be queried is mapped to different geometric representations in one or more dimensions (e.g., two or three dimensions), and then a corresponding tree data structure and lookup information (e.g., a ray built with parameters so that it can be used in a manner equivalent to querying a particular element and / or number) is constructed for use by ray tracing H / W 120 to perform one or more aspects of API 108 to generate object information 126. In at least one embodiment, a set of elements is represented as points on a line, and small triangles are constructed around these points, which are then represented in a tree data structure, and a lookup operation is performed for a particular element in this set of elements using a ray by looking up one or more intersecting triangles. In at least one embodiment, triangles are not constructed around these points, data used to represent these points is contained in a tree data structure, and API 108 directly uses ray tracing H / W 120 to traverse the tree data structure (e.g., perform one or more intersection operations of lookup information (e.g., a ray) with point objects). In at least one embodiment, directly using ray tracing hardware 120 refers to using ray tracing hardware 120 from other code executed by an accelerator (e.g., GPU), where results from ray tracing hardware 120 are able to be accessed for use by that code. In at least one embodiment, this provides a technical advantage by enabling faster, more efficient use of ray tracing hardware results compared to using a separate ray tracing pipeline and / or kernel to perform object intersection operations. In at least one embodiment, a tree data structure includes data that can represent an arbitrary set of elements (also referred to as objects) that API 108 can query based on lookup information (e.g., a ray) using ray tracing hardware 120. In at least one embodiment, a set of objects and / or elements to be queried is contained in an array that is mapped to a tree data structure, and API 108 directly uses ray tracing hardware 120 to traverse the tree data structure (e.g., perform one or more intersection operations) and / or return an indication of intersecting objects and / or elements.

[0055] In at least one embodiment, a user (e.g., one or more threads of a kernel running on PPU 106) uses one or more functions and / or instructions of API 108 to initiate generation and / or encoding of object information 126. In at least one embodiment, a kernel is part of application 110. In at least one embodiment, processor 114, processor 116, and / or ray tracing hardware 120 execute threads of a kernel. In at least one embodiment, it is useful and provides a technical advantage to generate object information 126 using a kernel that is to use object information 126, as object information 126 can be used more quickly compared to generating object information 126 by a separate kernel.

[0056] Figure 2 is a block diagram illustrating a computing environment 200, in accordance with at least one embodiment. In at least one embodiment, computer system 202 includes a processor 204, a memory 206, and a set of graphics processing units (GPUs) 208. In at least one embodiment, computer system 202 includes Figure 1 one or more components of system 100. In at least one embodiment, GPU set 208 includes GPU 210 and GPU 212. In at least one embodiment, GPU set 208 includes a different number of GPUs (e.g., fewer or more than two GPUs). In at least one embodiment, GPU 210 includes GPU memory 214 and GPU 212 includes GPU memory 216. In at least one embodiment, GPU memory 214 and / or GPU memory 216 includes more than one level and / or type of memory (e.g., global memory accessible to an entire GPU, memory accessible to a subset of processors on a GPU, cache memory accessible to a single processor on a GPU, shared memory accessible to a particular thread group). In at least one embodiment, GPU memory 214 includes global memory 218 and shared memory 220. In at least one embodiment, GPU memory 214 includes a different number of shared memories (e.g., more than one shared memory).

[0057] In at least one embodiment, GPU 210 includes ray tracing hardware (H / W) 222. In at least one embodiment, ray tracing H / W 222 is Figure 1 ray tracing H / W 120. In at least one embodiment, ray tracing H / W 222 includes one or more ray tracing cores. In at least one embodiment, ray tracing H / W 222 accelerates one or more types of ray tracing operations. In at least one embodiment, ray tracing H / W 222 can be used to accelerate tree traversal operations. In at least one embodiment, ray tracing H / W 222 can be used to accelerate object intersection operations.

[0058] In at least one embodiment, GPU 210 includes one or more processors 224. In at least one embodiment, one or more processors 224 include Figure 1processor 114 and / or processor 116. In at least one embodiment, GPU 210 includes one or more caches, which are not shown for clarity. In at least one embodiment, GPU 212 includes one or more processors, one or more ray-tracing H / W components, and / or one or more memories, which are not shown for clarity. In at least one embodiment, different numbers of processors (e.g., more than one processor 204) and / or different numbers of memories (e.g., more than one memory 206) are included in computer system 202. In at least one embodiment, processor 204 is a central processing unit (CPU). In at least one embodiment, computer system 202 includes one or more other components (not shown for clarity) (e.g., a network interface card, a persistent storage device, one or more input devices, one or more output devices, and / or one or more other suitable components).

[0059] In at least one embodiment, processor 204 is a single-core processor. In at least one embodiment, processor 204 is a multi-core processor. In at least one embodiment, processor 204 is an element of a processing system (e.g., processing system 900 described herein). In at least one embodiment, processor 204 is an element of a computer system (e.g., computer system 1000 described herein). In at least one embodiment, processor 204 is an element of a system (e.g., system 1100 described herein). In at least one embodiment, processor 204 is an element of a computing system (e.g., computing system 1300 described herein). In at least one embodiment, processor 204 is an element of a computing unit (e.g., computing unit 3740 described herein). In at least one embodiment, processor 204 is some other processor shown and / or described herein.

[0060] In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are graphics processor 1710 described herein. In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are graphics processor 1740 described herein. In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are graphics multiprocessor 1934 described herein. In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are graphics processor 2000 described herein. In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are graphics processor 2208 described herein. In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are GPU 3592 described herein. In at least one embodiment, one or more GPUs (e.g., GPU 210) in GPU set 208 are some other GPU shown and / or described herein. In at least one embodiment, computer system 202 includes one or more accelerators (e.g., one or more parallel processing units (PPUs), FPGAs, ASICS, and / or other suitable accelerators) in place of or in addition to GPUs in GPU set 208.

[0061] In at least one embodiment, computer system 202 includes API set 226. In at least one embodiment, when one or more APIs are said to perform an action or technical effect, one or more hardware components of a computer system running the API(s) (e.g., CPUs, GPUs, ray-tracing H / W, and / or other hardware components) perform that action or technical effect. In at least one embodiment, API set 226 is a set of APIs for GPUs in GPU set 208. In at least one embodiment, one or more operations described with respect to GPUs in GPU set 208 and / or APIs in API set 226 are performed by one or more accelerators (not shown for clarity) that are not GPUs. In at least one embodiment, API set 226 is referred to as an API that includes a number of callable functions (e.g., a driver API). In at least one embodiment, API set 226 is implemented in the form of a dynamic library. In at least one embodiment, API set 226 is a handle-based, imperative API. In at least one embodiment, API set 226 is a parallel processing framework API (e.g., a Compute Unified Device Architecture (CUDA) driver API, a Heterogeneous Compute Portability Interface (HIP) API, or some other API).

[0062] In at least one embodiment, one or more APIs in API set 226 are high-level APIs (e.g., accessed using a high-level programming language such as C++, Python, Java, Fortran, C, or some other suitable language). In at least one embodiment, one or more APIs in API set 226 are low-level APIs (e.g., accessed using instructions of a programming framework such as CUDA PTX instructions or some other suitable intermediate representation that can be compiled to machine-level binary representations for specific hardware architectures). In at least one embodiment, one or more APIs in API set 226 can also be implemented as instructions such as PTX, assembly, x86, GPU instruction set architectures (ISAs), machine-level, or some other suitable type of instruction. In at least one embodiment, API set 226 is a set of APIs for a programming platform. In at least one embodiment, a programming platform can be, but is not limited to, CUDA, Radeon Open Compute Platform (“ROCm”), OpenCL (OpenCL TM developed by the Khronos Group), SYCL, or Intel OneAPI. In at least one embodiment, although certain aspects of APIs and / or techniques to identify object intersections, generate trees, and generate lookup parameters are discussed in relation to CUDA, including CUDA APIs and / or CUDA kernels, it should be understood that ROCm, OpenCL, SYCL, OneAPI, and / or any other suitable APIs and / or kernels can be used. In at least one embodiment, one or more APIs in API set 226 are accessed at least in part by including a header file in one or more code portions that define one or more functions of the one or more APIs. In at least one embodiment, one or more APIs in API set are functions (e.g., defined in a function library).

[0063] In at least one embodiment, API set 226 includes an object intersection API 228. In at least one embodiment, object intersection API 228 is used to indicate to a user one or more objects intersected by one or more rays of light. In at least one embodiment, object intersection API 228 is callable and / or used by a kernel executed by GPU. In at least one embodiment, object intersection API 228 includes and / or uses one or more instructions to identify one or more intersected objects using ray tracing hardware 222. In at least one embodiment, object intersection API 228 is referred to as a space query API, a tracing API, a ray tracing tracing API, and / or some other suitable name. In at least one embodiment, object intersection API 228 is based at least in part on a tree data structure (e.g., Figure 1Tree data structure 124) and rays (e.g., Figure 1 The search information 126) indicates the intersecting objects. In at least one embodiment, the object intersection API 228 provides the technical effect of enabling users to directly access the ray tracing core from within a kernel (e.g., a CUDA kernel) executed by an accelerator (e.g., a GPU). In at least one embodiment, the object intersection API 228 provides the technical advantage of enabling users to directly access the ray tracing core without a pipeline (e.g., a ray tracing pipeline executed outside of other kernels executed by an accelerator) and / or without using regular graphics APIs. In at least one embodiment, the object intersection API 228 provides the technical advantage of enabling users to directly use the ray tracing core for a range of non-ray tracing applications (e.g., spatial queries). In at least one embodiment, directly using the ray tracing core refers to using the ray tracing core in other code executed by an accelerator (e.g., GPU 210), where results from the ray tracing core are accessible for use by that code. In at least one embodiment, the object intersection API 228 provides the technical effect of enabling users to interleave other computations with accelerated spatial queries using the ray tracing core. In at least one embodiment, API 228 provides a technical advantage by enabling users to directly utilize ray tracing hardware to accelerate search and / or lookup operations that can be modeled and / or converted into tree traversal and / or intersection operations. In at least one embodiment, the object intersection API 228 provides a technical advantage by enabling other code to accelerate query and lookup operations using tree data structures. In at least one embodiment, the object intersection API 228 provides a technical advantage by using a ray tracing core to accelerate spatial queries, binary search, lookup operations, and / or other suitable applications and / or use cases. In at least one embodiment, the object intersection API 228 provides a technical effect of more effectively utilizing GPU hardware resources by enabling GPU-executed code to use ray tracing hardware to accelerate query and lookup operations, which would otherwise be idle or underutilized while other code is executing. In at least one embodiment, the object intersection API 228 provides a technical effect by enabling ray tracing hardware to provide the results of accelerated object intersection and / or tree traversal operations to other code executing in the same kernel, which is faster and more efficient than having to use a separate pipeline and / or kernel to access the ray tracing hardware.

[0064] In at least one embodiment, API set 226 includes a spanning tree API 230. In at least one embodiment, the spanning tree API 230 spans a tree data structure (e.g., Figure 1tree data structure 122). In at least one embodiment, a tree data structure includes data describing one or more geometric objects and / or parameters (e.g., triangles, points, distances, areas, volumes, and / or other suitable geometric objects). In at least one embodiment, a tree data structure includes data representing some other type of information. In at least one embodiment, a tree data structure includes data that can be used to represent an arbitrary set of elements. In at least one embodiment, generate tree API 230 generates a tree data structure based at least in part on converting non-geometric information into a geometric representation. In at least one embodiment, generate tree API 230 stores a generated tree data structure in GPU memory 214.

[0065] In at least one embodiment, API set 226 includes a generate lookup parameter API 232. In at least one embodiment, generate lookup parameter API 232 generates a lookup parameter that specifies a ray (e.g., a ray 126 of FIG. 1). Figure 1 In at least one embodiment, generate lookup parameter API 232 generates a lookup parameter corresponding to information to be looked up based at least in part on converting non-geometric information to be looked up into a geometric representation (e.g., by encoding a search for non-geometric information as a ray). In at least one embodiment, generate lookup parameter API 232 stores a generated lookup parameter in GPU memory 214. In at least one embodiment, object intersection API 228 is used to indicate one or more objects intersected by one or more rays based at least in part on tree data structures generated by API 230 and / or rays generated by lookup parameter API 232. In at least one embodiment, API set 226 includes one or more other APIs, such as an API to initialize ray tracing hardware, an API to de-initialize ray tracing hardware, and / or other suitable APIs.

[0066] In at least one embodiment, compiler 234 converts requests received via an API of API set 226 into instructions that can be executed on and / or by GPU 210 (e.g., generates instructions that are part of an instruction set architecture of GPU 210). In at least one embodiment, generated instructions are stored as code 236 (e.g., as part of a kernel) that is copied to one or more GPUs of GPU set 208 (e.g., GPU 210) for execution. In at least one embodiment, one or more threads use one or more APIs of API set 226 and can pass one or more arguments to an API of API set. In at least one embodiment, API set 226 includes one or more APIs that are usable by code implemented at a higher level (e.g., a C++-style implementation), are usable by code implemented at an intermediate level (e.g., as PTX-style instructions), and / or are usable by code implemented at a lower level (e.g., as instructions of a GPU ISA).

[0067] In at least one embodiment, computer system 202 includes node set 238. In at least one embodiment, node set 238 includes node 240, node 242, and node 244. In at least one embodiment, node set 238 includes different numbers of nodes. In at least one embodiment, nodes of node set 238 include one or more GPUs. In at least one embodiment, kernel information (e.g., based at least in part on code 236) is copied to one or more GPUs included in one or more nodes of node set 238. In at least one embodiment, one or more components and / or aspects of computer system 202 and / or node set 238 are implemented by one or more hardware components, one or more software components, one or more circuits, dedicated hardware (e.g., fixed function circuits), and / or any other suitable type of hardware, software, or combinations of them. In at least one embodiment, with respect to Figure 2 one or more aspects shown or described are with respect to Figure 1 one or more aspects shown or described are with respect to Figure 1 one or more aspects shown or described. In at least one embodiment, API set 226 is included in API 108 of Figure 1 GPU 210 is Figure 1 PPU 106 of Figure 1 CPU 104 of Figure 1 compiler 112 of In at least one embodiment, API set 226 includes one or more other APIs, which are not shown for purposes of clarity.

[0068] In at least one embodiment, as used in any implementation described herein, unless the context explicitly specifies otherwise or explicitly to the contrary, the terms such as “module” and nominalized verbs (e.g., compiler and / or other terms) each refer to software logic, firmware logic, hardware logic, and / or any combination of circuitry configured to provide the functionality described herein. In at least one embodiment, software may be embodied as a software package, code, and / or instruction set or instructions, and “hardware” as used in any implementation described herein may include, for example, hardwired circuitry, programmable circuitry, state machine circuitry, fixed-function circuitry, execution unit circuitry, and / or firmware storing instructions executed by programmable circuitry, either individually or in any combination. In at least one embodiment, a module may be embodied collectively or individually as a circuitry system that forms part of a larger system, such as an integrated circuit (IC), a system-on-a-chip (SoC), etc.

[0069] Figure 3 API 300 for performing object intersection operations according to at least one embodiment is shown. In at least one embodiment, ray tracing hardware (e.g., Figure 1 Ray tracing H / W 120 and / or ray tracing H / W 222) are used to perform one or more aspects of API 300. In at least one embodiment, API 300 will use Figure 2 The object is intersected by API 228 to perform the operation. In at least one embodiment, API 300 includes one or more parameters. In at least one embodiment, the parameters of API 300 include one or more parameters for indicating a tree data structure, one or more lookup parameters, start parameters, output parameters, and status indicator parameters. In at least one embodiment, API 300 uses a different number of parameters and / or one or more other types of parameters.

[0070] In at least one embodiment, the object intersection API response 302 is used to return an operation status. In at least one embodiment, the operation status returns an unsigned integer. In at least one embodiment, a returned operation status of 0 indicates that an API call to API 300 returned with no errors. In at least one embodiment, a returned operation status of 1 indicates that one or more parameters passed to API 300 were incorrect. In at least one embodiment, a returned operation status of 2 indicates that a call to API 300 failed due to a lack of underlying resources. In at least one embodiment, a returned operation status of 3 indicates that an internal error occurred during initialization. In at least one embodiment, a returned operation status of 4 indicates that an operating system (OS) error occurred. In at least one embodiment, a returned operation status of 5 indicates that an error existed in another operation (e.g., a CUDA operation) as part of a call to API 300. In at least one embodiment, a returned operation status of 7 indicates that an error existed in an underlying driver (e.g., an underlying ray tracing (RT) driver). In at least one embodiment, a returned operation status of 0x7fffffff indicates that an unknown error occurred. In at least one embodiment, a returned operation status value can correspond to additional and / or different success or error states than those specified above.

[0071] In at least one embodiment, aspects of API 300 and API response 302 can be described in conjunction with the following pseudo code:

[0072] __device__ unsigned int spatialQuery(void *bvh,

[0073] Ray *ray,

[0074] unsigned int traversalState,

[0075] QueryResult *out,

[0076] unsigned int *inoutState) ;

[0077] wherein the API 300 is described for a function named spatialQuery, the parameter void *bvh corresponds to the tree data structure parameter of the API 300, Ray *ray corresponds to the search information parameter of the API 300, unsigned int traversalState corresponds to the start parameter of the API 300, QueryResult *out corresponds to the output parameter of the API 300, unsigned int *inoutState corresponds to the state indicator of the API 300, __device__ unsigned int indicates that spatialQuery is executed on a device (e.g., a GPU) and returns an unsigned int (e.g., an operation status returned by the object intersection API 302). In at least one embodiment, the API 300 modifies the output parameter QueryResult *out to include an identifier of an intersected object, and the output parameter is accessible to a user (e.g., computer program code using a kernel of the API 300). In at least one embodiment, the modified output parameter is referred to as a result, an output, and / or some other suitable term.

[0078] In at least one embodiment, a data structure of the output parameter of the API 300 can be described for the following pseudo code:

[0079] typedef struct QueryResult_st{

[0080] unsigned int ID;

[0081] float T;

[0082] float U;

[0083] float V;

[0084] unsigned int validHit;

[0085] }QueryResult;

[0086] where unsigned int ID is an identifier of the intersecting object, and unsigned int validHit indicates whether the object intersected. In at least one embodiment, a first value of validHit (e.g., a value of 1) indicates that the object intersected, and a second value of validHit (e.g., a value of 0) indicates that no object intersected. In at least one embodiment, float T, float U, and float V are not present and / or are not used. In at least one embodiment, when used, T, U, and V more specifically indicate a location (e.g., a particular coordinate on the intersecting object) at which the object with the object ID intersects the ray. In at least one embodiment, the value of the ID parameter of the output parameter data structure is used to indicate to a user the object that intersected the lookup parameter (e.g., the ray). In at least one embodiment, some other suitable data structure is used to specify the output parameter.

[0087] In at least one embodiment, the ray data structure of the lookup information parameter of API 300 can be described with respect to the following pseudo code:

[0088] typedef struct Ray_st {

[0089] float originX;

[0090] float originY;

[0091] float originZ;

[0092] float dirX;

[0093] float dirY;

[0094] float dirZ;

[0095] float tmin;

[0096] float tmax;

[0097] } Ray;

[0098] where originX, originY, and originZ specify the origin of the ray, dirX, dirY, and dirZ specify the direction of the ray, and tmin and tmax specify the segment between tmin and tmax to be tested for intersection. In at least one embodiment, some other suitable data structure is used to specify the lookup information and / or the ray.

[0099] In at least one embodiment, a call to an object intersection API (e.g., object intersection API 300) can be described in conjunction with the following pseudo code:

[0100] result = spatialQuery(bvh, ray, CUTTU_TRAVERSAL_START, &out, state); where bvh represents a tree data structure, ray represents a query parameter, CUTTU_TRAVERSAL_START represents that the call to the object intersection API is starting, not continuing, &out represents an output parameter data structure passed by reference, and state represents a state indicator value. In at least one embodiment, passing CUTTU_TRAVERSAL_CONTINUE instead of CUTTU_TRAVERSAL_START indicates that the call to the object intersection API is continuing from a previous state. In at least one embodiment, CUTTU_TRAVERSAL_START and CUTU_TRAVERSAL_CONTINUE correspond to specific numerical values (e.g., values of 0 and 1, respectively) that can be passed to the object intersection API. In at least one embodiment, the object intersection API does not include a start parameter, and if the state indicator parameter is 0 or NULL, the API infers that it will start a new traversal. In at least one embodiment, the result returned by a call to the API is an operation status value (e.g., the object intersection API returns an operation status of 302). In at least one embodiment, the object intersection API includes a different number of parameters and / or different names. In at least one embodiment, a call to the object intersection API can be described in relation to the following pseudocode: trace(bvh, ray, &out); or trace(bvh, ray, &out, NULL); where these calls are to a function call with an empty return value (e.g., the object intersection API return 302 does not exist). In at least one embodiment, the object intersection API 300 is used to indicate a single intersecting object (e.g., the first object that intersects with a ray), and does not use and / or does not have a start parameter and a state indicator parameter. In at least one embodiment, the object intersection API 300 updates a state variable (e.g., by increasing a value of the state variable), such that the updated state variable can be used as an argument to a subsequent call to the object intersection API 300 and as a state indicator parameter. In at least one embodiment, for a subsequent call to the object intersection API 300 that continues using the updated state, a ray used as a query parameter remains unchanged in direction, but can change in length and / or have different tmin / tmax values as described above. In at least one embodiment, a user iteratively uses the object intersection API 300 in a loop until the object intersection API 300 indicates that no intersecting object was found (e.g., by using a specific value of a validHit variable in the output parameter data structure).In at least one embodiment, the object intersection API 300 returns (e.g., by modifying one or more output data structures) all intersecting objects (e.g., between tmin and tmax along a ray) without requiring iterative calls using different values ​​of state indicators (e.g., state indicator parameters).

[0101] Figure 4 This is a flowchart of a technique 400 for performing object intersection operations according to at least one embodiment. In at least one embodiment, one or more aspects of technique 400 are described by... Figure 1 and / or Figure 2 The implementation includes one or more aspects shown or described herein (e.g., CPU 104, PPU 106, API 108, ray tracing hardware 120, processor 204, processor 224, and / or ray tracing hardware 222) and / or with respect to other components, technologies, and / or aspects shown or described herein. In at least one embodiment, technology 400 includes the implementation of Figure 1 API 108 Figure 2 Object Intersection API 228 and / or Figure 3 One or more aspects of API 300.

[0102] In at least one embodiment, at block 402, technique 400 includes obtaining an object intersection request. In at least one embodiment, obtaining the object intersection request is done via an API (e.g., Figure 1 API 108 Figure 2 Object Intersection API 228 and / or Figure 3 The API 300 is executed. In at least one embodiment, obtaining an object intersection request includes receiving and / or obtaining a lookup request, a spatial query, an RT core-accelerated query, or some other suitable request and / or function call. In at least one embodiment, obtaining an object intersection request includes receiving an object intersection request from software code (e.g., a kernel) executed by the GPU.

[0103] In at least one embodiment, at block 404, technique 400 includes providing intersection object information. In at least one embodiment, providing intersection object information includes using ray tracing hardware (e.g., Figure 1 Ray tracing hardware 120 and / or Figure 2 The ray tracing hardware 222) performs accelerated tree traversal and / or object intersection operations. In at least one embodiment, providing intersection object information includes storing one or more identifiers of the corresponding intersecting one or more objects (e.g., as...). Figure 1 The object information 122 and / or is stored in a specific data structure and / or location, for example, using Figure 3 (The output parameters specified by API 300).

[0104] In at least one embodiment, at block 406, technique 400 includes performing additional actions. In at least one embodiment, performing additional actions includes returning to block 402 to obtain another object intersection request. In at least one embodiment, performing additional actions includes returning an operation status code (e.g., Figure 3 The object intersection API returns a 302 operation status value.

[0105] Figure 5 This is a flowchart of a technique 500 for performing object intersection operations according to at least one embodiment. In at least one embodiment, one or more aspects of technique 500 are described by... Figure 1 and / or Figure 2 The implementation includes one or more aspects shown or described (e.g., CPU 104, PPU 106, API 108, ray tracing hardware 120, processor 204, processor 224 and / or ray tracing hardware 222) and / or one or more components, technologies and / or other aspects shown or described herein with respect to other figures. In at least one embodiment, technology 500 includes the implementation of Figure 1 API 108 Figure 2 Object Intersection API 228 and / or Figure 3 One or more aspects of API 300.

[0106] In at least one embodiment, at block 502, technique 500 includes obtaining an object intersection request. In at least one embodiment, obtaining the object intersection request is performed via an API, such as... Figure 1 API 108 Figure 2 Object Intersection API 228 and / or Figure 3 API 300. In at least one embodiment, obtaining an object intersection request includes receiving and / or obtaining a lookup request, a spatial query, an RT core-accelerated query, or some other suitable request and / or function call. In at least one embodiment, obtaining an object intersection request includes receiving an object intersection request from software code (e.g., a kernel) executed by the GPU. In at least one embodiment, obtaining an object intersection request includes receiving one or more arguments and / or parameters (e.g., ...). Figure 3 The object intersection API 300 provides a tree data structure, lookup information, start parameters, output parameters, and / or status indicators.

[0107] In at least one embodiment, at decision block 504, technique 500 includes: determining whether the object intersection request is a valid request. In at least one embodiment, at block 506, technique 500 includes: returning an error code if the object intersection request is not a valid request. In at least one embodiment, the error code may be an error code indicating that the object intersection API returned a 302 operation status, or some other suitable error code.

[0108] In at least one embodiment, at block 508, technique 500 includes: identifying a tree data structure if the object intersection request is a valid request. In at least one embodiment, identifying the tree data structure includes: identifying the tree data structure based on an indication of the tree data structure received using the object intersection request obtained at block 502. In at least one embodiment, the tree data structure is... Figure 3 The tree data structure parameters of API 300 intersect the objects. In at least one embodiment, the tree data structure is... Figure 1 The tree data structure 122 and / or can be derived from... Figure 1 API 108 generation and / or generation Figure 2 Tree API 230.

[0109] In at least one embodiment, at block 510, technique 500 includes: identifying lookup information. In at least one embodiment, the lookup information is information specifying a ray. In at least one embodiment, identifying the lookup information includes identifying the lookup information based on an indication received using the lookup information obtained at block 502 for an object intersection request. In at least one embodiment, the lookup information is... Figure 3 The lookup information for the object intersection API 300. In at least one embodiment, the lookup information is... Figure 1 The search information 124 and / or can be obtained from Figure 1 API 108 generation and / or generation Figure 2 The lookup parameter is API 232.

[0110] In at least one embodiment, at block 512, technique 500 includes: identifying an output data structure. In at least one embodiment, identifying the output data structure includes identifying the output data structure based on an indication of the output data structure received using the object intersection request obtained at block 502. In at least one embodiment, the output data structure is... Figure 3 The output parameters of the Object Intersection API 300. In at least one embodiment, the output data structure is some other data structure, and / or a data structure generated by the Object Intersection API instead of the data structure passed as a parameter, and / or a data structure generated by the user (e.g., other program code that calls the Object Intersection API).

[0111] In at least one embodiment, at determination block 514, technique 500 includes: determining whether an object intersection request is used for a new start. In at least one embodiment, determining whether an object intersection request is used for a new start is based on the start parameters received using the object intersection request obtained at block 502 (e.g., Figure 3 The start parameter of the object intersection API 300). In at least one embodiment, at block 516, technique 500 includes: identifying state information if the object intersection request is not for a new start (e.g., continuing from a previous request to find other object intersections). In at least one embodiment, identifying the state information includes: indicating, based on the state information received using the object intersection request obtained at block 502 (e.g., ... Figure 3 The object intersection API 300 status information parameters are used to identify status information. In at least one embodiment, technique 500 does not include: determining whether the object intersection request is for a new start or identifying status information (e.g., when technique 500 is used to find a single intersecting object, or when technique 500 automatically provides all intersecting objects without using the status indicator passed through the API).

[0112] In at least one embodiment, at decision block 518, technique 500 includes: determining whether an intersection exists. In at least one embodiment, determining whether an intersection exists is performed by ray tracing hardware (e.g., Figure 1 Ray tracing hardware 120 and / or Figure 2 The ray tracing hardware 222 performs this operation. In at least one embodiment, determining whether an intersection exists includes traversing a tree data structure (e.g., the tree data structure identified at block 508). In at least one embodiment, determining whether an intersection exists includes determining whether a ray (e.g., lookup information constructed as a ray identified at block 510) intersects with an object in a data structure (e.g., the tree data structure identified at block 508 and / or traversed at decision block 518). In at least one embodiment, if the object intersection request is for a continuation request rather than a new start, determining whether an intersection exists is based at least in part on state information received at block 516 (e.g., by finding the next intersecting object corresponding to the state information, and / or by changing one or more parameters of the lookup information based on the state information before using the modified lookup information to identify the intersecting object).

[0113] In at least one embodiment, if there is an intersection at decision block 518, then technique 500 includes generating object information at block 520. In at least one embodiment, generating object information includes updating an output data structure identified at block 512 to include an identifier and / or other information for the intersecting object. In at least one embodiment, at block 522, technique 500 includes storing the object information. In at least one embodiment, storing the object information includes storing the object information generated at block 520. In at least one embodiment, storing the object information includes storing the object information in an output data structure (e.g., the output data structure identified at block 512). In at least one embodiment, the stored object information is Figure 1 object information 126. In at least one embodiment, storing the object information includes storing an indication of an intersection in an output data structure (e.g., the output data structure identified at block 512). In at least one embodiment, the indication of an intersection is a particular value stored in an output parameter data structure of the object intersection API 300 (e.g., a particular value stored in a validHit variable, such as 1). Figure 3

[0114] In at least one embodiment, at block 524, technique 500 includes returning a success code. In at least one embodiment, the success code indicates that the object intersection request has been successfully executed and that the information in the output data structure is available for use. In at least one embodiment, the success code is a particular operation status value (e.g., a particular operation status value returned by the object intersection API 300, indicating that the API has been successfully executed). Figure 3 In at least one embodiment, the success code indicates that the object intersection API completed successfully, but does not indicate whether there was an object intersection or intersection object information, which is indicated in an output parameter data structure. In at least one embodiment, rather than returning information about an intersecting object and whether there was a valid hit in an output data structure passed by reference, technique 500 returns this information as a response to the function call.

[0115] In at least one embodiment, if there is no intersection at decision block 518, then technique 500 includes indicating that no intersection was found at block 526. In at least one embodiment, indicating that no intersection was found includes storing a not found information. In at least one embodiment, storing the not found information includes storing an indication of a miss in an output data structure (e.g., the output data structure identified at block 512). In at least one embodiment, the indication of a miss is a particular value stored in an output parameter data structure of the object intersection API 300 (e.g., a particular value stored in a validHit variable, such as 0). Figure 3 In at least one embodiment, the success code indicates that the object intersection API completed successfully, but does not indicate whether there was an object intersection or intersection object information, which is indicated in an output parameter data structure. In at least one embodiment, rather than returning information about an intersecting object and whether there was a valid hit in an output data structure passed by reference, technique 500 returns this information as a response to the function call.

[0116] ​In at least one embodiment, at block 528, techniques 500 include performing other actions. In at least one embodiment, performing other actions includes returning to block 502 to obtain another object intersection request. In at least one embodiment, performing other actions includes setting and / or updating state information. In at least one embodiment, setting and / or updating state information includes incrementing a state variable, resetting a state variable to zero, or setting a state variable to NULL.

[0117] In at least one embodiment, Figure 4 techniques 400 and / or Figure 5 one or more aspects of techniques 500 include executing an application programming interface (API) to cause one or more objects intersected by one or more rays to be indicated to a user. In at least one embodiment, the API (e.g., an object intersection API) is configured to cause one or more objects intersected by one or more rays to be indicated to a user based at least in part on a set of input parameters including a tree data structure (e.g., a tree data structure 122) indicative of one or more objects. Figure 2 object intersection API 228 and / or Figure 3 object intersection API 300) is configured to cause one or more objects to be indicated using one or more ray tracing cores based at least in part on a set of input parameters including a tree data structure (e.g., a tree data structure 122) indicative of one or more objects. In at least one embodiment, the API is configured to cause one or more objects to be indicated based at least in part on causing one or more ray tracing cores to identify one or more objects intersected by one or more rays. In at least one embodiment, the API is configured to cause identifiers of intersected objects in the one or more objects to be included in an output data structure (e.g., an output parameter of an object intersection API 300) indicated to a user. In at least one embodiment, the API is configured to cause one or more objects to be indicated using one or more ray tracing cores of a GPU (e.g., a ray tracing H / W 120 and / or a ray tracing H / W 222) based at least in part on a set of input parameters including ray information and a tree data structure including the one or more objects. In at least one embodiment, a non-transitory computer-readable medium stores a set of instructions which, if executed by one or more processors, cause the one or more processors to perform one or more aspects of techniques 400 and / or techniques 500. Figure 1 Figure 3 Figure 1 Figure 2 Figure 4 Figure 5

[0118] Figure 6 is a block diagram illustrating an example of a processor 600, in accordance with at least one embodiment. In at least one embodiment, processor 602 performs one or more processes (such as processes described herein) to identify intersected objects to a user. In at least one embodiment, processor 602 performs aspects of techniques 400 and / or techniques 500 described herein. Figure 1 ​​​​​​CPU 104, PPU 106, and / or API 108 of FIG. 1, Figure 2 processor 204, GPU 210, GPU 212, one or more APIs of API set 226, Figure 3 API 300 of FIG. 3, Figure 4 techniques 400 of FIG. 4, Figure 5 techniques 500 of FIG. 5, and / or Figure 7 one or more aspects described with respect to one or more APIs 710. In at least one embodiment, processor 602 performs one or more processes such as those described in conjunction with Figure 1 ray tracing H / W 120 of FIG. 1, and / or Figure 2 ray tracing H / W 222 of FIG. 2. In at least one embodiment, processor 602 performs one or more processes such as those described in conjunction with Figures 1-5

[0119] In at least one embodiment, processor 602 includes one or more processors such as those described in conjunction with one or more of FIGS. 17 through Figure 38 In at least one embodiment, processor 602 is any suitable processing unit and / or combination of processing units such as one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof. In at least one embodiment, processor 602 includes a tree generation module 604, a lookup parameter generation module 606, and an object intersection module 608. In at least one embodiment, tree generation module 604, lookup parameter generation module 606, and / or object intersection module 608 are part of processor 602 and / or one or more other processors. In at least one embodiment, tree generation module 604, lookup parameter generation module 606, and / or object intersection module 608 are distributed between multiple processors that communicate over buses, networks, by writing to shared memory, and / or any suitable communication process such as those described herein.

[0120] ​In at least one embodiment, as used in any embodiment described herein, a module refers to any combination of software, firmware, hardware, and / or circuitry configured to provide the functionality described herein. In at least one embodiment, software can be embodied as a software package, code and / or instruction set or instructions, and “hardware”, as used in any embodiment described herein, can include, for example, singly or in any combination, hardwired circuitry, programmable circuitry, state machine circuitry, fixed function circuitry, execution unit circuitry, and / or firmware that stores instructions executed by programmable circuitry. In at least one embodiment, a module can be embodied by a combination of all of the foregoing. In at least one embodiment, a module can also include one or more processing units, and / or one or more processing units executing one or more software programs or firmware routines.

[0121] In at least one embodiment, a build tree module 604 is a module that builds a tree data structure (e.g., a tree data structure 122 of FIG. 1 and / or a tree data structure of an object intersection API 300 of FIG. 1). In at least one embodiment, a build tree module 604 performs one or more processes (e.g., processes described herein) by including or otherwise encoding instructions that result in or are otherwise usable to perform the processes (e.g., by processor 602). In at least one embodiment, a build tree module 604 performs one or more aspects illustrated or described in relation to an API 108 of FIG. 1 and / or a build tree API 230 of FIG. 1. Figure 1 Figure 3 In at least one embodiment, a build tree module 604 is a module that builds a tree data structure (e.g., a tree data structure 122 of FIG. 1 and / or a tree data structure of an object intersection API 300 of FIG. 1). In at least one embodiment, a build tree module 604 performs one or more processes (e.g., processes described herein) by including or otherwise encoding instructions that result in or are otherwise usable to perform the processes (e.g., by processor 602). In at least one embodiment, a build tree module 604 performs one or more aspects illustrated or described in relation to an API 108 of FIG. 1 and / or a build tree API 230 of FIG. 1. Figure 1 Figure 2 In at least one embodiment, a build tree module 604 is a module that builds a tree data structure (e.g., a tree data structure 122 of FIG. 1 and / or a tree data structure of an object intersection API 300 of FIG. 1). In at least one embodiment, a build tree module 604 performs one or more processes (e.g., processes described herein) by including or otherwise encoding instructions that result in or are otherwise usable to perform the processes (e.g., by processor 602). In at least one embodiment, a build tree module 604 performs one or more aspects illustrated or described in relation to an API 108 of FIG. 1 and / or a build tree API 230 of FIG. 1.

[0122] In at least one embodiment, a build lookup parameter module 606 is a module that builds lookup information (e.g., information specifying a ray (e.g., lookup information 124 of FIG. 1 and / or lookup information of an object intersection API 300 of FIG. 1). In at least one embodiment, a build lookup parameter module 606 performs one or more processes (e.g., processes described herein) by including or otherwise encoding instructions that result in or are otherwise usable to perform the processes (e.g., by processor 602). In at least one embodiment, a build lookup parameter module 606 performs one or more aspects illustrated or described in relation to an API 108 of FIG. 1 and / or a build lookup parameter API 232 of FIG. 1. Figure 1 Figure 3 In at least one embodiment, a build tree module 604 is a module that builds a tree data structure (e.g., a tree data structure 122 of FIG. 1 and / or a tree data structure of an object intersection API 300 of FIG. 1). In at least one embodiment, a build tree module 604 performs one or more processes (e.g., processes described herein) by including or otherwise encoding instructions that result in or are otherwise usable to perform the processes (e.g., by processor 602). In at least one embodiment, a build tree module 604 performs one or more aspects illustrated or described in relation to an API 108 of FIG. 1 and / or a build tree API 230 of FIG. 1. Figure 1 Figure 2 In at least one embodiment, a build tree module 604 is a module that builds a tree data structure (e.g., a tree data structure 122 of FIG. 1 and / or a tree data structure of an object intersection API 300 of FIG. 1). In at least one embodiment, a build tree module 604 performs one or more processes (e.g., processes described herein) by including or otherwise encoding instructions that result in or are otherwise usable to perform the processes (e.g., by processor 602). In at least one embodiment, a build tree module 604 performs one or more aspects illustrated or described in relation to an API 108 of FIG. 1 and / or a build tree API 230 of FIG. 1. ​​​​

[0123] In at least one embodiment, the object intersection module 608 is a module that indicates to a user one or more objects intersected by one or more rays. In at least one embodiment, the object intersection generation module 608 performs one or more processes by at least including or otherwise encoding instructions, such as those described herein, which cause the execution of the one or more processes or are otherwise available for execution of the one or more processes (e.g., executed by processor 602). In at least one embodiment, the object intersection module 608 performs actions related to... Figure 1 API 108 Figure 2 Object Intersection API 228 Figure 3 API 300 Figure 4 Technology 400 and / or Figure 5 The technology 500 shows or describes one or more aspects.

[0124] In at least one embodiment, the processor includes one or more circuits for executing an application programming interface (API) to indicate one or more objects intersecting with one or more rays of light to a user. In at least one embodiment, the API (e.g., Figure 2 Object intersection API 228 and / or Figure 3 The object intersection API 300 is used to enable objects to intersect, at least in part, based on tree data structures (e.g., Figure 1 Tree data structures 122 and / or Figure 3 The intersecting objects in API 300 (a tree data structure) indicate one or more objects. In at least one embodiment, the API is used to enable the use of a ray tracing kernel (e.g., Figure 1 Ray tracing H / W 120 and / or Figure 2 Ray tracing (H / W 222) in the API indicates one or more objects. In at least one embodiment, the API is used to make it at least partially GPU-based (e.g., Figure 2 The software executed by the GPU 210 in the GPU indicates one or more objects. In at least one embodiment, the one or more objects include one or more triangles. In at least one embodiment, the API is used to make the input parameters at least in part based on a set of input parameters (e.g., Figure 3 The API (300) uses one or more parameters to indicate one or more objects, the input parameters including rays in one or more rays and a data structure including the one or more objects. In at least one embodiment, the API is used to indicate one or more objects using one or more ray tracing cores of the GPU, at least in part based on a tree data structure including the one or more objects.

[0125] In at least one embodiment, the system includes one or more processors for executing an application programming interface (API) to indicate to a user one or more objects intersected by one or more rays of light. In at least one embodiment, the API (e.g., Figure 2 Object intersection API 228 and / or Figure 3 The Object Intersection API 300 is used to enable the use of one or more ray tracing cores (e.g., Figure 1 Ray tracing H / W 120 and / or Figure 2 In at least one embodiment, the API is configured to indicate one or more objects based at least in part on a set of input parameters including a tree data structure. In at least one embodiment, the API is configured to indicate one or more objects based at least in part on one or more software function calls executed by a GPU including one or more ray tracing cores. In at least one embodiment, one or more objects comprise one or more triangles in a triangle mesh. In at least one embodiment, the API is configured to indicate one or more objects using one or more ray tracing cores based at least in part on a set of input parameters including rays from one or more rays and a data structure including one or more objects.

[0126] In at least one embodiment, Figure 6 The processor 602 can execute Figure 2 Object Intersection API 228 Figure 3 Object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name), Figure 4 Technology 400 and / or Figure 5 Technology 500. In at least one embodiment, Figure 6 The processor 602 can execute Figure 2 The tree generation API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name). In at least one embodiment, Figure 6 The processor 602 can execute Figure 2 The API 232 for generating lookup parameters (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name). In at least one embodiment, Figure 1 The CPU 104 and / or PPU 106 can execute Figure 2 Object Intersection API 228Figure 3 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name) of the techniques 400 of the Figure 4 techniques 400 of the Figure 5 techniques 500 of the. In at least one embodiment, Figure 1 CPU 104 and / or PPU 106 can perform a generate tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of the techniques 400 of the Figure 2 CPU 104 and / or PPU 106 can perform a generate tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of the techniques 400 of the Figure 1 CPU 104 and / or PPU 106 can perform a generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of the techniques 400 of the Figure 2 CPU 104 and / or PPU 106 can perform a generate tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of the techniques 400 of the Figure 2 object intersection API 228 of the techniques 400 of the Figure 2 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name) of the techniques 400 of the Figure 3 techniques 400 of the Figure 4 techniques 500 of the. In at least one embodiment, Figure 5 CPU 104 and / or PPU 106 can perform a generate tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of the techniques 400 of the Figure 2 CPU 104 and / or PPU 106 can perform a generate tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of the techniques 400 of the Figure 2 CPU 104 and / or PPU 106 can perform a generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of the techniques 400 of the Figure 2 CPU 104 and / or PPU 106 can perform a generate tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of the techniques 400 of the Figure 2 CPU 104 and / or PPU 106 can perform a generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of the techniques 400 of the

[0127] Figure 7is a block diagram illustrating a driver and / or runtime environment 700 including one or more libraries to provide one or more application programming interfaces (APIs), in accordance with at least one embodiment. In at least one embodiment, software program 702 is a software module. In at least one embodiment, software program 702 includes one or more software modules including, but not limited to, at least those described in conjunction with Figure 6 software modules described. In at least one embodiment, software modules are further described, non-exclusively, in conjunction with Figure 6 one or more APIs 710 are sets of software instructions that, if executed by a set of software instructions, cause one or more processors to perform one or more computational operations. In at least one embodiment, execution of software instructions includes execution of software instructions. In at least one embodiment, one or more aspects of one or more software modules shown or described in conjunction with Figure 7 software program 702 (e.g., instructions and / or code for calling functions, when executed by a processor), and are partially included in one or more APIs 710 and / or functions 712 (e.g., instructions and / or code for implementing called functions, when executed by a processor).

[0128] In at least one embodiment, one or more APIs 710 are distributed or otherwise provided as part of one or more libraries 706, drivers and / or runtime 704, and / or any other software, non-transitory computer-readable instructions, and / or executable code grouping further described herein. In at least one embodiment, one or more APIs 710 perform one or more computational operations in response to a call by software program 702. In at least one embodiment, software program 702 is a collection of software code, commands, instructions, or other sequences of text to instruct a computing device to perform one or more computational operations and / or call one or more other sets of instructions (e.g., APIs 710 or API functions 712) to perform. In at least one embodiment, functionality provided by one or more APIs 710 includes software functions 712, such as can be used to accelerate one or more portions of software program 702 using one or more parallel processing units (PPUs) (e.g., graphics processing units (GPUs)). In at least one embodiment, APIs 710 and / or API functions 712 include APIs and / or functions to perform one or more aspects of APIs 108 of Figure 1 API set 226 of Figure 2 APIs, Figure 3 API 300 of Figure 4 techniques 400 of, and / or Figure 5 techniques 500 of.

[0129] In at least one embodiment, API 710 is a hardware interface for one or more circuits for performing one or more computational operations. In at least one embodiment, the one or more software APIs 710 described herein are implemented as one or more circuits for performing the operations described herein. Figures 1-6 One or more techniques described herein. In at least one embodiment, one or more software programs 702 include instructions that, if executed, cause one or more hardware devices and / or circuits to perform the functions described herein. Figures 1-6 One or more techniques are described. In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 provided by the driver and / or runtime 704 to manage virtual memory and / or allocate or otherwise reserve one or more memory blocks 714 for one or more PPUs (e.g., GPUs). In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 provided by the driver and / or runtime 704 to manage virtual memory and / or allocate or otherwise reserve memory blocks, and / or perform one or more operations on one or more caches.

[0130] In at least one embodiment, software program 702 (e.g., user-implemented software program (e.g., ...) Figure 1 Application 110 utilizes one or more application programming interfaces (APIs) 710 to perform various computational operations, such as memory reservation, tree traversal, object intersection operations, tensor graph encoding, tensor copying, matrix multiplication, arithmetic operations, or any computational operation performed by a parallel processing unit (PPU), such as a graphics processing unit (GPU) and GPU-specific acceleration features, such as ray tracing features, as further described herein. In at least one embodiment, one or more APIs 710 provide a set of callable functions 712, referred to herein as APIs, API functions, and / or functions, which individually perform one or more computational operations, such as computational operations related to parallel computing. In at least one embodiment, the processor uses an API that includes a spanning tree function 716. In at least one embodiment, the spanning tree function 716 performs... Figure 1 API 108 Figure 2 Spanning Tree API 230 and / or Figure 6 The processor utilizes one or more aspects of the spanning tree module 604. In at least one embodiment, the processor uses an API including a lookup parameter generation function 718. In at least one embodiment, the lookup parameter generation function 718 is executed. Figure 1 API 108 Figure 2 Generate lookup parameters API 232 and / orFigure 6 one or more aspects of the generating lookup parameters module 606 of FIG. 7. In at least one embodiment, the processor uses an API including an object intersection function 720. In at least one embodiment, the object intersection function 720 performs Figure 1 the API 108 of FIG. 1, Figure 2 the object intersection API 228 of FIG. 2, Figure 3 the API 300 of FIG. 3, Figure 4 the techniques 400 of FIG. 4, Figure 5 the techniques 500 of FIG. 5, and / or Figure 6 one or more aspects of the object intersection module 608 of FIG. 6.

[0131] In at least one embodiment, one or more software programs 702 interact with or otherwise communicate with one or more APIs 710 to perform one or more computational operations using one or more PPUs (e.g., GPUs). In at least one embodiment, one or more computational operations using one or more PPUs include at least one or more sets of computational operations that are at least partially performed by the one or more PPUs to accelerate. In at least one embodiment, one or more software programs 702 interact with one or more APIs 710 to facilitate parallel computing using a remote or local interface.

[0132] In at least one embodiment, an interface is a software instruction that, if executed, provides access to one or more functions 712 provided by one or more APIs 710. In at least one embodiment, a software program 702 uses a local interface when a software developer compiles the software program 702 in conjunction with one or more libraries 706 that include or otherwise provide access to one or more APIs 710. In at least one embodiment, a software program 702 is statically compiled in conjunction with a pre-compiled library 706 that includes instructions to perform one or more APIs 710 or source code that is not compiled. In at least one embodiment, a software program 702 is dynamically compiled and the software program utilizes a linker to link to one or more pre-compiled libraries 706 that include one or more APIs 710.

[0133] In at least one embodiment, software programs 702 use remote interfaces when a software program utilizes or otherwise communicates with libraries 706 including one or more APIs 710 over a network or other remote communication medium. In at least one embodiment, one or more libraries 706 including one or more APIs 710 are to be executed by a remote computing service (e.g., a computing resource service provider). In another embodiment, one or more libraries 706 including one or more APIs 710 are to be executed by any other computing host that provides said one or more APIs 710 to one or more software programs 702.

[0134] In at least one embodiment, a processor executing or using one or more software programs 702 calls upon, uses, performs, or otherwise implements one or more APIs 710 to allocate and otherwise manage memory to be used by said software programs 702. In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 to allocate and otherwise manage memory to be used by one or more portions of said software programs 702 to use one or more PPUs (e.g., GPUs or any other accelerator or processor described further herein) for acceleration.

[0135] In at least one embodiment, an API in one or more APIs 710 is an API for facilitating parallel computing. In at least one embodiment, one or more APIs 710 include any other API described further herein. In at least one embodiment, one or more APIs 710 are provided by a driver and / or runtime 704. In at least one embodiment, an API in one or more APIs 710 is provided by a CUDA user mode driver. In at least one embodiment, an API in one or more APIs 710 is provided by a CUDA runtime. In at least one embodiment, driver 704 is data values and software instructions that, if executed, perform or otherwise facilitate operations of one or more functions 712 of APIs 710 during loading and execution of one or more portions of software programs 702. In at least one embodiment, driver and / or runtime 704 is data values and software instructions that, if executed, perform or otherwise facilitate operations of one or more functions 712 of APIs 710 during execution of software programs 702. In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 implemented or otherwise provided by driver and / or runtime 704 to perform combined arithmetic operations by one or more software programs 702 during execution by one or more PPUs (e.g., GPUs).

[0136] Data Center

[0137] The following figure illustrates, without limitation, an exemplary data center system that can be used to implement at least one embodiment. In at least one embodiment, one or more data center components in the figure can implement... Figures 1-7 One or more aspects of one or more embodiments described herein. In at least one embodiment, one or more data center components include Figure 1 The computing device 102 includes one or more components (e.g., CPU 104, PPU 106, and / or API 108). In at least one embodiment, one or more data center components include... Figure 2 One or more components of the computer system 202 (e.g., processor 204, GPU 210, API set 226, and / or node set 238). In at least one embodiment, one or more data center components perform... Figure 3 One or more aspects of API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name), Figure 4 Technology 400 and / or Figure 5 Technology 500. In at least one embodiment, one or more data center components include... Figure 6 Processor 600 and / or about Figure 6 One or more modules are shown or described. In at least one embodiment, one or more data center components include and / or perform... Figure 7 One or more aspects of the environment 700 (e.g., driver and / or runtime 704, one or more APIs 710 and / or one or more functions 712).

[0138] Figure 8 An example data center 800 according to at least one embodiment is shown. In at least one embodiment, the data center 800 includes, but is not limited to, a data center infrastructure layer 810, a framework layer 820, a software layer 830, and an application layer 840.

[0139] In at least one embodiment, such as Figure 8As shown, the data center infrastructure layer 810 can include a resource orchestrator 812, grouped computing resources 814, and node computing resources (“node C.R.s”) 816(1)-816(N), where “N” represents any whole, positive integer. In at least one embodiment, node C.R.s 816(1)-816(N) can include, but are not limited to, any number of central processing units (“CPUs” or “processors”), including accelerators, field programmable gate arrays (“FPGAs”), data processing units (DPUs) in network devices, graphics processors, etc., memory devices such as dynamic read only memories, storage devices such as solid state drives or disk drives, network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power modules, and cooling modules, etc. In at least one embodiment, one or more node C.R.s of node C.R.s 816(1)-816(N) can be a server having one or more of above-described computing resources.

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

[0141] In at least one embodiment, resource orchestrator 812 can configure or otherwise control one or more node C.R.s 816(1)-816(N) and / or grouped computing resources 814. In at least one embodiment, resource orchestrator 812 can include a software design infrastructure (“SDI”) management entity for data center 800. In at least one embodiment, resource orchestrator 812 can include hardware, software, or some combination thereof. In at least one embodiment, Figure 8 In at least one embodiment, one or more node C.R.s 816(1)-816(N) and / or grouped computing resources 814 can perform Figure 2 an object intersection API 228 in Figure 3 an object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name) in Figure 4Technology 400 and / or Figure 5 The technology 500. In at least one embodiment, Figure 8 One or more nodes CR816(1)-816(N) and / or grouped computing resources 814 can perform Figure 2 The Spanning Tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API names). In at least one embodiment, Figure 8 One or more nodes CR816(1)-816(N) and / or grouped computing resources 814 can perform Figure 2 Generate lookup parameters in API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name).

[0142] In at least one embodiment, such as Figure 8 As shown, the framework layer 820 includes, but is not limited to, a job scheduler 832, a configuration manager 834, a resource manager 836, and a distributed file system 838. In at least one embodiment, the framework layer 820 may include a framework of software 852 supporting the software layer 830 and / or one or more applications 842 of the application layer 840. In at least one embodiment, the software 852 or application 842 may respectively include web-based service software or applications, such as services or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. In at least one embodiment, the framework layer 820 may be, but is not limited to, a free and open-source software web application framework, such as Apache Spark, which can utilize the distributed file system 838 for large-scale data processing (e.g., "big data"). TM(“Spark”). In at least one embodiment, a job scheduler 832 can include a Spark driver to facilitate scheduling of workloads supported by various layers of the datacenter 800. In at least one embodiment, a configuration manager 834 can be capable of configuring different layers, such as software layer 830 and framework layer 820 including Spark and a distributed file system 838 for supporting large-scale data processing. In at least one embodiment, a resource manager 836 can be capable of managing clustered or grouped computing resources mapped to or allocated for supporting distributed file system 838 and job scheduler 832. In at least one embodiment, clustered or grouped computing resources can include grouped computing resources 814 on datacenter infrastructure layer 810. In at least one embodiment, resource manager 836 can coordinate with resource orchestrator 812 to manage these mapped or allocated computing resources.

[0143] In at least one embodiment, software 852 included in software layer 830 can include software used by at least portions of node C.R.s 816(1)-816(N), grouped computing resources 814, and / or distributed file system 838 of framework layer 820. One or more types of software can include, but are not limited to, Internet web page search software, email virus scanning software, database software, and streaming video content software.

[0144] In at least one embodiment, one or more application programs 842 included in application layer 840 can include one or more types of application programs used by at least portions of node C.R.s 816(1)-816(N), grouped computing resources 814, and / or distributed file system 838 of framework layer 820. One or more types of application programs can include, but are not limited to, CUDA application programs.

[0145] In at least one embodiment, any of configuration manager 834, resource manager 836, and resource orchestrator 812 can implement any number and type of self-modifying actions based on any number and type of data acquired in any technically feasible fashion. In at least one embodiment, self-modifying actions can mitigate datacenter operators of datacenter 800 making possibly poor configuration decisions and can avoid underutilized and / or poorly performing portions of a datacenter.

[0146] Computer-based system

[0147] The following figures set forth, without limitation, example computer-based systems that can be used to implement at least one embodiment. In at least one embodiment, one or more computer-based systems shown in the following figures can implement aspects of the systems and techniques described with respect to Figures 1-7one or more aspects of one or more embodiments in the Figure 1 one or more components of computing device 102 (e.g., CPU 104, PPU 106, and / or API 108). In at least one embodiment, one or more computer-based systems include Figure 2 one or more components of computer system 202 (e.g., processor 204, GPU 210, API set 226, and / or one or more components of node set 238). In at least one embodiment, one or more computer-based systems perform Figure 3 one or more aspects of API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names), Figure 4 techniques 400 in Figure 5 techniques 500. In at least one embodiment, one or more computer-based systems include Figure 6 processor 600 in Figure 6 shown or described with respect to Figure 7 one or more aspects of environment 700 (e.g., one or more of drivers and / or runtime 704, one or more APIs 710, and / or one or more of functions 712).

[0148] Figure 9 A processing system 900, in accordance with at least one embodiment, is shown. In at least one embodiment, system 900 includes one or more processor(s) 902 and one or more graphics processor(s) 908, and can be a single processor desktop system, a multiprocessor workstation system, or a server system having many processor cores 907. In at least one embodiment, processor(s) 902 and / or graphics processor(s) 908 can be integrated into a SoC. In at least one embodiment, processing system 900 is a server system, interconnected to one or more other systems through a network 910.

[0149] In at least one embodiment, processing system 900 can include or be incorporated into a server-based gaming platform, including a game console, a gaming console including a game and media console, a mobile gaming console, a handheld game console, or an online game console. In at least one embodiment, processing system 900 is a mobile phone, a smart phone, a tablet device, or a mobile internet device. In at least one embodiment, processing system 900 can also include or be incorporated into a wearable device such as a smart watch wearable device, a smart eyewear device, an augmented reality device, or a virtual reality device. In at least one embodiment, processing system 900 is a television or set-top box device having one or more processors 902 and a graphical interface generated by one or more graphics processors 908. In at least one embodiment, Figure 9 One or more processors 902 and / or one or more graphics processors 908 of Figure 2 Object Intersection API 228 of Figure 3 Object Intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name) of Figure 4 Techniques 400 of Figure 5 Techniques 500 of In at least one embodiment, Figure 9 One or more processors 902 and / or one or more graphics processors 908 of Figure 2 Generation Tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API name) of Figure 9 One or more processors 902 and / or one or more graphics processors 908 of Figure 2 Generation Lookup Parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of

[0150] In at least one embodiment, one or more processors 902 each include one or more processor cores 907 to process instructions which, when executed, perform operations for system and user software. In at least one embodiment, each of the one or more processor cores 907 is configured to process a specific instruction set 909. In at least one embodiment, instruction set 909 can facilitate Complex Instruction Set Computing (CISC), Reduced Instruction Set Computing (RISC), or computing via a Very Long Instruction Word (VLIW). In at least one embodiment, multiple processor cores 907 can each process a different instruction set 909, which can include instructions to facilitate emulation of other instruction sets. In at least one embodiment, processor core 907 can also include other processing devices, such as a digital signal processor (DSP).

[0151] In at least one embodiment, processor 902 includes cache memory 904. In at least one embodiment, processor 902 can have single-level or multi-level internal cache. In at least one embodiment, cache memory is shared among various components of processor 902. In at least one embodiment, processor 902 also uses an external cache (e.g., a Level 3 (L3) cache or last level cache (LLC)) (not shown), which can be shared among processor cores 907 using known cache coherency techniques. In at least one embodiment, additionally included in processor 902 is a register file 906, which can include different types of registers to store different types of data (e.g., integer registers, floating point registers, status registers, and instruction pointer registers). In at least one embodiment, register file 906 can include a general register file or other registers.

[0152] In at least one embodiment, one or more processors 902 are coupled with one or more interface buses 910 for passing control information and receiving data signals between processor 902 and other components of system 900. In at least one embodiment, interface bus 910 can be versions of a processor bus, such as a direct media interface (DMI) bus. In at least one embodiment, interface bus 910 is not limited to DMI bus, and can include one or more peripheral component interconnect buses (e.g., a PCI, PCI Express), a memory bus, or other types of interface buses. In at least one embodiment, processor 902 includes integrated memory controller 916 and platform controller hub 930. In at least one embodiment, memory controller 916 facilitates communication between memory devices and other components of processing system 900, while platform controller hub 930 provides connections to input / output (I / O) devices to local I / O bus. In at least one embodiment, one or more peripheral component interconnect buses include PCIe Gen 5, which provides an interface for processors.

[0153] In at least one embodiment, storage devices 920 can be dynamic random access memory (DRAM) devices, static random access memory (SRAM) devices, flash memory devices, phase-change memory devices, or some other memory devices having suitable performance to serve as external memory devices for processor 902. In at least one embodiment, storage devices 920 can be used as system memory for processing system 900, to store data 922 and instructions 921 for use when one or more processors 902 executes an application or process. In at least one embodiment, memory controller 916 also couples with an optional external graphics processor 912, which can communicate with one or more graphics processors 908 in processor 902 to perform graphics and media operations. In at least one embodiment, display device 911 can be connected to processor 902. In at least one embodiment, display device 911 can include one or more internal display devices, such as for a mobile electronic device or a head-mounted display (HMD) for a virtual reality (VR) application or an augmented reality (AR) application, for example. In at least one embodiment, display device 911 can include an external display device, such as for a display port, DisplayPort®, HDMI®, etc.

[0154] In at least one embodiment, platform controller hub 930 enables peripherals to connect to storage device 920 and processor 902 via a high-speed I / O bus. In at least one embodiment, I / O peripherals include, but are not limited to, audio controller 946, network controller 934, firmware interface 928, wireless transceiver 926, touch sensors 925, data storage device 924 (e.g., hard disk drive, flash memory, etc.). In at least one embodiment, data storage device 924 can connect to the storage interface via a storage interface bus, e.g., SATA or via a peripheral bus, e.g., a peripheral component interconnect bus (such as PCI, PCI Express). In at least one embodiment, touch sensors 925 can include touch screen sensors, pressure sensors, or fingerprint sensors. In at least one embodiment, wireless transceiver 926 can be a Wi-Fi transceiver, a Bluetooth transceiver, or a mobile network transceiver such as a 3G, 4G, or Long Term Evolution (LTE) transceiver. In at least one embodiment, firmware interface 928 enables communication with system firmware, and can be, e.g., a unified extensible firmware interface (UEFI) according to at least one embodiment. In at least one embodiment, network controller 934 can enable network connectivity to a wired network. In at least one embodiment, a high-performance network controller (not shown) is coupled

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

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

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

[0158] In at least one embodiment, computer system 1000 can include, without limitation, a processor 1002 that can include, without limitation, one or more execution units 1008 that can be configured to employ a compute unified device architecture (“CUDA”) (registered trademark) to perform algorithms for process data. In at least one embodiment, computer system 1000 can include a processor such as a a CUDA program. In at least one embodiment, a CUDA program is at least a portion of a software application written in the CUDA programming language. In at least one embodiment, computer system 1000 is a single processor desktop or server system. In at least one embodiment, computer system 1000 can be a multiprocessor system. In at least one embodiment, processor 1002 can include, without limitation, a CISC microprocessor, a RISC microprocessor, a VLIW microprocessor, a processor implementing a combination of instruction sets, or any other processor device, such as a digital signal processor. In at least one embodiment, processor 1002 can be coupled to a processor bus 1010 that can transmit data signals between processor 1002 and other components in computer system 1000.

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

[0160] In at least one embodiment, execution unit 1008, including, without limitation, logic to perform integer and floating point operations, also resides in processor 1002. Processor 1002 can also include microcode (“ucode”) read-only memory (“ROM”) that stores microcode for certain macro instructions. In at least one embodiment, execution unit 1008 can include logic to handle a packed instruction set 1009. In at least one embodiment, by including the packed instruction set 1009 in the instruction set of a general-purpose processor 1002, along with associated circuitry to execute the instructions, the general-purpose processor 1002 can be used to perform the operations on packed data that many multimedia applications use. In at least one embodiment, by using the full width of a processor’s data bus to transfer data to and from the processor’s registers, many multimedia applications can be accelerated as compared to integer-only processors lacking a packed data execution facility. In at least one embodiment, without requiring the transfer of smaller units of data to a processor for the performance of a single operation on a single data element at a time, many multimedia applications run efficiently.

[0161] In at least one embodiment, execution unit 1008 can also be used in microcontrollers, embedded processors, graphics devices, DSPs, and other types of logic circuits. In at least one embodiment, computer system 1000 can include, but not limited to, memory 1020. In at least one embodiment, memory 1020 can be implemented using DRAM, SRAM, Flash memories, or other storage devices. Memory 1020 can store instruction(s) 1019 and / or data 1021 represented by data signals that can be executed by processor 1002.

[0162] In at least one embodiment, system logic chip can be coupled to processor bus 1010 and memory 1020. In at least one embodiment, system logic chip can include, without limitation, a memory controller hub (“MCH”) 1016, and processor 1002 can communicate with MCH 1016 via processor bus 1010. In at least one embodiment, MCH 1016 can provide a high bandwidth memory path 1018 to memory 1020 for instruction and data storage and for storage of graphics commands, data, and textures. In at least one embodiment, MCH 1016 can direct data signals between processor 1002, memory 1020, and other components in computer system 1000, and can bridge data signals between processor bus 1010, memory 1020, and system I / O 1022.

[0163] In at least one embodiment, system logic chip can provide a graphics port for coupling to a graphics controller. In at least one embodiment, MCH 1016 can be coupled to memory 1020 through high bandwidth memory path 1018, and a graphics / video card 1012 can be coupled to MCH 1016 through an Accelerated Graphics Port (“AGP”) interconnect 1014.

[0164] In at least one embodiment, computer system 1000 can use system I / O 1022 as a proprietary hub interface bus to couple MCH 1016 to I / O controller hub (“ICH”) 1030. In at least one embodiment, ICH 1030 can provide direct connection to some I / O devices via local I / O bus. In at least one embodiment, local I / O bus can include, without limitation, a high-speed I / O bus for connecting peripheral devices to memory 1020, chipset, and processor 1002. Examples can include, without limitation, audio controller 1029, firmware hub (“Flash BIOS”) 1028, wireless transceiver 1026, data storage 1024, legacy I / O controller 1023 containing user input 1025 and keyboard interface, serial expansion port 1027 (e.g., USB), and network controller 1034. Data storage 1024 can include a hard disk drive, a floppy disk drive, a CD-ROM device, a flash memory device, or other mass storage device.

[0165] In at least one embodiment, Figure 10 A system including interconnected hardware devices or “chips” is shown. In at least one embodiment, Figure 10 An exemplary SoC can be shown. In at least one embodiment, Figure 10 Devices shown in FIG. 1 can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, one or more components of system 1000 are interconnected using Compute Express Link (CXL) interconnects. In at least one embodiment, Figure 10 Processor 1002 and / or graphics / video card 1012 of FIG. 1 can perform Figure 2 Object intersection API 228 of FIG. 2, Figure 3 Object intersection API 300 of FIG. 3 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name), Figure 4 Techniques 400 of FIG. 4, and / or Figure 5 Techniques 500 of FIG. 5. In at least one embodiment, Figure 10 Processor 1002 and / or graphics / video card 1012 of FIG. 1 can perform Figure 2 Generation tree API 230 of FIG. 2 (e.g., generateTree(), buildTree(), and / or other suitable API name). In at least one embodiment, Figure 10 Processor 1002 and / or graphics / video card 1012 of FIG. 1 can perform Figure 2generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name).

[0166] Figure 11 A system 1100, in accordance with at least one embodiment, is shown. In at least one embodiment, system 1100 is an electronic device that utilizes a processor 1110. In at least one embodiment, system 1100 can be, for example and without limitation, a notebook, a tower server, a rack server, a blade server, an edge device communicatively coupled to one or more local or cloud service providers, a laptop, a desktop, a tablet, a mobile device, a phone, an embedded computer, or any other suitable electronic device. In at least one embodiment, system 1100 includes processor 1110 coupled to memory 1120 and input / output (I / O) components 1130. Figure 11 Processor 1110 of system 1100 can execute instructions Figure 2 Object Intersection API 228 of system 1100, Figure 3 Object Intersection API 300 of system 1100 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name), Figure 4 Techniques 400 of system 1100, and / or Figure 5 Techniques 500 of system 1100. In at least one embodiment, Figure 11 Processor 1110 of system 1100 can execute Figure 2 Generation Tree API 230 of system 1100 (e.g., generateTree(), buildTree(), and / or other suitable API name). In at least one embodiment, Figure 11 Processor 1110 of system 1100 can execute Figure 2 Generation Lookup Parameter API 232 of system 1100 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API name).

[0167] In at least one embodiment, system 1100 can include, without limitation, processor 1110 communicatively coupled to any suitable number or kind of components, peripherals, modules, or devices. In at least one embodiment, processor 1110 is coupled using a bus or interface, such as an I 2 C bus, a System Management Bus (“SMBus”), a Low Pin Count (LPC) bus, a Serial Peripheral Interface (“SPI”), a High Definition Audio (“HDA”) bus, a Serial Advanced Technology Attachment (“SATA”) bus, a USB (versions 1, 2, 3), or a Universal Asynchronous Receiver / Transmitter (“UART”) bus. In at least one embodiment, Figure 11Systems can be shown that include interconnected hardware devices or “chips.” In at least one embodiment, Figure 11 An exemplary SoC can be shown. In at least one embodiment, Figure 11 Devices shown in FIG. 6 can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, Figure 11 One or more components of FIG. 6 are interconnected using compute express link (CXL) interconnects.

[0168] In at least one embodiment, Figure 11 may include a display 1124, a touch screen 1125, a touch pad 1130, a near field communication unit (“NFC”) 1145, a sensor hub 1140, a thermal sensor 1146, an express chipset (“EC”) 1135, a trusted platform module (“TPM”) 1138, a BIOS / firmware / flash memory (“BIOS, FW Flash”) 1122, a DSP 1160, a solid state disk (“SSD”) or hard disk drive (“HDD”) 1120, a wireless local area network unit (“WLAN”) 1150, a Bluetooth unit 1152, a wireless wide area network unit (“WWAN”) 1156, a global positioning system (GPS) 1155, a camera (“USB 3.0 camera”) 1154 (e.g., USB 3.0 camera), or a low power double data rate (“LPDDR”) memory unit (“LPDDR3”) 1115 implemented in, e.g., LPDDR3 standard. These components can each be implemented in any suitable manner.

[0169] In at least one embodiment, other components can be communicatively coupled to processor 1110 by components discussed above. In at least one embodiment, an accelerometer 1141, an ambient light sensor (“ALS”) 1142, a compass 1143, and a gyroscope 1144 can be communicatively coupled to a sensor hub 1140. In at least one embodiment, a thermal sensor 1139, a fan 1137, a keyboard 1136, and a touchpad 1130 can be communicatively coupled to an EC 1135. In at least one embodiment, a speaker 1163, a headphone 1164, and a microphone (“mic”) 1165 can be communicatively coupled to an audio unit (“audio codec and class D amplifier”) 1162, which in turn can be communicatively coupled to a DSP 1160. In at least one embodiment, audio unit 1162 can include, for example and without limitation, an audio coder / decoder (“codec”) and a class D amplifier. In at least one embodiment, a SIM card (“SIM”) 1157 can be communicatively coupled to a WWAN unit 1156. In at least one embodiment, components such as WLAN unit 1150 and Bluetooth unit 1152, as well as WWAN unit 1156, can be implemented in a next generation form factor (NGFF).

[0170] Figure 12 An exemplary integrated circuit 1200 is shown in accordance with at least one embodiment. In at least one embodiment, exemplary integrated circuit 1200 is a SoC, which can be fabricated using one or more IP cores. In at least one embodiment, integrated circuit 1200 includes one or more application processors 1205 (e.g., CPUs, DPUs), at least one graphics processor 1210, and can additionally include an image processor 1215 and / or a video processor 1220, any of which can be a modular IP core. In at least one embodiment, integrated circuit 1200 includes peripheral or bus logic including USB controllers 1225, UART controllers 1230, SPI / SDIO controllers 1235, and I 2 S / I 2 C controllers 1240. In at least one embodiment, integrated circuit 1200 can include a display device 1245 coupled to one or more of a high-definition multimedia interface (HDMI) controller 1250 and a mobile industry processor interface (MIPI) display interface 1255. In at least one embodiment, storage can be provided by flash memory subsystem 1260, including flash memory and a flash memory controller. In at least one embodiment, a memory interface can be provided via a memory controller 1265 for access to SDRAM or SRAM memory devices. In at least one embodiment, some integrated circuits also include an embedded Secure Engine 1270. In at least one embodiment, Figure 12integrated circuit 1200, application processor 1205, and / or graphics processor 1210 in FIG. 12B can perform techniques 400 in FIG. 4 and / or Figure 2 object intersection API 228 in FIG. 23, Figure 3 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name) in FIG. 24, Figure 4 techniques 400 in FIG. 4 and / or Figure 5 techniques 500 in FIG. 5. In at least one embodiment, Figure 12 integrated circuit 1200, application processor 1205, and / or graphics processor 1210 in FIG. 12B can perform generation tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) in FIG. 23. In at least one embodiment, Figure 2 integrated circuit 1200, application processor 1205, and / or graphics processor 1210 in FIG. 12B can perform generation lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) in FIG. 23. Figure 12 Figure 2

[0171] Figure 13 A computing system 1300 is shown, in accordance with at least one embodiment. In at least one embodiment, computing system 1300 includes a processing subsystem 1301 having one or more processor(s) 1302 and system memory 1304, communicating via an interconnection path 1305 that can include a memory hub 1305. In at least one embodiment, memory hub 1305 can be a separate component coupled with one or more processors 1302 via interconnection path 1305. In at least one embodiment, memory hub 1305 can be integrated into one or more processors 1302. In at least one embodiment, memory hub 1305 couples with system memory 1304 through communication link 1306. In at least one embodiment, I / O subsystem 1311 includes I / O hub 1307 that can enable connectivity to one or more input devices 1308 and one or more output devices 1310. In at least one embodiment, I / O hub 1307 can include a display controller to couple with one or more display devices 1310A, which can be integrated with or external to computing system 1300. In at least one embodiment, one or more input devices 1308 can include a keyboard, a mouse, a trackpad, a trackball, a touchscreen, a microphone, and / or some other suitable input device.

[0172] ​​In at least one embodiment, processing subsystem 1301 includes one or more parallel processor(s) 1312 coupled to memory hub 1305 via bus or other communication link 1313. In at least one embodiment, communication link 1313 can be one of many based on well-known bus or Figure 13 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 2 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 3 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 4 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 5 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 13 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 2 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 13 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of Figure 2 In at least one embodiment, one or more processor(s) 1300 and / or one or more parallel processor(s) 1312 of

[0173] In at least one embodiment, system storage unit 1314 can connect to I / O hub 1307 to provide storage mechanisms for computing system 1300. In at least one embodiment, I / O switches 1316 can be used to provide interface mechanisms to enable connections between I / O hub 1307 and other components such as network adapter 1318 and / or wireless network adapter 1319 that can be integrated into a platform, as well as various other devices that can be added via one or more add-in devices 1320. In at least one embodiment, network adapter 1318 can be an Ethernet adapter or another wired

[0174] In at least one embodiment, computing system 1300 can include other components not explicitly shown, including USB or other port connections, optical storage drives, video capture devices, etc., that can also be connected to I / O hub 1307. In at least one embodiment, communication paths interconnecting various components in Figures 1-7 Communication paths interconnecting various components in computing system 1300 can use any suitable protocols, such as PCI (Peripheral Component Interconnect) based protocols (e.g., PCI Express), or other bus or point-to-point communication interfaces and / or protocols (e.g., NVLink high-speed interconnect, or interconnect protocols).

[0175] In at least one embodiment, one or more parallel processors 1312 include circuitry optimized for graphics and video processing (including, for example, video output circuitry) and constitute a graphics processing unit (GPU). In at least one embodiment, one or more parallel processors 1312 include circuitry optimized for general-purpose processing. In at least one embodiment, components of the computing system 1300 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 1312, memory hub 1305, processor 1302, and I / O hub 1307 may be integrated into a system-on-a-chip (SoC) integrated circuit. In at least one embodiment, components of the computing system 1300 may be integrated 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 1300 may be integrated into a multi-chip module (MCM) that can interconnect with other MCMs to a modular computing system. In at least one embodiment, the I / O subsystem 1311 and display device 1310B are omitted from the computing system 1300. In at least one embodiment, one or more parallel processors 1312 include one or more tensor memory accelerator (TMA) units that can transfer data blocks between global memory and shared memory. In at least one embodiment, one or more processors use or access one or more TMAs to perform bidirectional copy operations, such as from global memory to shared memory and vice versa.

[0176] Processing system

[0177] The following figures illustrate, in a non-limiting manner, exemplary processing systems that can be used to implement at least one embodiment. In at least one embodiment, one or more processing systems in the following figures can implement [the following description of processing systems]. Figure 1 One or more aspects of one or more embodiments described. In at least one embodiment, one or more processing systems include Figure 2 The computing device 102 includes one or more components (e.g., CPU 104, PPU 106, and / or API 108). In at least one embodiment, one or more processing systems include... Figure 3 One or more components of the computer system 202 (e.g., one or more components of the processor 204, GPU 210, API set 226, and / or node set 238). In at least one embodiment, one or more processing systems execute... Figure 4 One or more aspects of API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name),Figure 5 techniques 400 and / or Figure 6 techniques 500. In at least one embodiment, one or more processing systems include Figure 6 processor 600 and / or one or more modules shown or described with respect to Figure 7 In at least one embodiment, one or more processing systems include and / or execute one or more aspects of environment 700 (e.g., drivers and / or runtime 704, one or more APIs 710, and / or one or more functions 712). Figure 14

[0178] Figure 14 An accelerated processing unit (“APU”) 1400, in accordance with at least one embodiment, is shown. In at least one embodiment, APU 1400 is developed by AMD Corporation of Santa Clara, CA. In at least one embodiment, APU 1400 can be configured to execute application programs, such as CUDA programs. In at least one embodiment, APU 1400 includes, without limitation, core complex 1410, graphics complex 1440, fabric 1460, I / O interface 1470, memory controllers 1480, display controllers 1492, and multimedia engine 1494. In at least one embodiment, APU 1400 can include, without limitation, any combination of any number of core complexes 1410, any number of graphics complexes 1450, any number of display controllers 1492, and any number of multimedia engines 1494. For purposes of illustration, multiple instances of like objects are denoted with reference numerals including a like reference designator, in which case the number in the reference designator identifies the desired instance.

[0179] In at least one embodiment, core complex 1410 is a CPU, graphics complex 1440 is a GPU, and APU 1400 is a processing unit that integrates, without limitation, 1410 and 1440 onto a single chip. In at least one embodiment, some tasks can be assigned to core complex 1410, while other tasks can be assigned to graphics complex 1440. In at least one embodiment, core complex 1410 is configured to execute host control software associated with APU 1400, such as an operating system. In at least one embodiment, core complex 1410 is a master processor of APU 1400 that controls and coordinates operations of other processors. In at least one embodiment, core complex 1410 issues commands that control operations of graphics complex 1440. In at least one embodiment, core complex 1410 can be configured to execute host executable code derived from CUDA source code, and graphics complex 1440 can be configured to execute device executable code derived from CUDA source code. In at least one embodiment,​Figure 2 APU 1400, graphics complex 1440, and / or core complex 1410 can perform Figure 3 object intersection API 228 of Figure 4 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name) of Figure 5 techniques 400 of Figure 14 techniques 500 of In at least one embodiment, Figure 2 APU 1400, graphics complex 1440, and / or core complex 1410 can perform Figure 14 generation tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API name) of Figure 2 APU 1400, graphics complex 1440, and / or core complex 1410 can perform Figure 15 generation lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of

[0180] In at least one embodiment, core complex 1410 includes, without limitation, cores 1420(1)-1420(4) and L3 cache 1430. In at least one embodiment, core complex 1410 can include, without limitation, any number of cores 1420 and any combination of any number and type of caches. In at least one embodiment, cores 1420 are configured to execute instructions of a particular instruction set architecture (“ISA”). In at least one embodiment, each core 1420 is a CPU core. In at least one embodiment, cores 1420 are referred to as compute or arithmetic units.

[0181] In at least one embodiment, each core 1420 includes, without limitation, a fetch / decode unit 1422, an integer execution engine 1424, a floating point execution engine 1426, and an L2 cache 1428. In at least one embodiment, fetch / decode unit 1422 fetches instructions, decodes such instructions, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1424 and floating point execution engine 1426. In at least one embodiment, fetch / decode unit 1422 can simultaneously dispatch one micro-instruction to integer execution engine 1424 and another micro-instruction to floating point execution engine 1426. In at least one embodiment, integer execution engine 1424 executes not limited to integer and memory operations. In at least one embodiment, floating point engine 1426 executes not limited to floating point and vector operations. In at least one embodiment, fetch-decode unit 1422 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1424 and floating point execution engine 1426.

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

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

[0184] In at least one embodiment, graphics processing engine 1430 includes, without limitation, any number of graphics processing clusters 1440 and L2 cache 1442. In at least one embodiment, graphics processing clusters 1440 share L2 cache 1442. In at least one embodiment, L2 cache 1442 is partitioned, with each graphics processing cluster 1440 having access to partitioned L2 cache 1442. In at least one embodiment, graphics processing engine 1430 includes, without limitation, any number of graphics processing clusters 1440 and any number (including zero) and type of cache. In at least one embodiment, graphics processing engine 1430 includes, without limitation, any number of specialized graphics hardware.

[0185] In at least one embodiment, each graphics processing cluster 1440 includes, without limitation, any number of SIMD units 1452 and shared memory 1454. In at least one embodiment, each SIMD unit 1452 implements a SIMD architecture and is configured to execute operations in parallel. In at least one embodiment, each graphics processing cluster 1440 can execute any number of thread blocks, but each thread block executes on a single graphics processing cluster 1440. In at least one embodiment, a thread block includes, without limitation, any number of execution threads. In at least one embodiment, a workgroup is a thread block. In at least one embodiment, each SIMD unit 1452 executes a different thread warp. In at least one embodiment, a thread warp is a group of threads (e.g., 16 threads), where each thread in a thread warp belongs to a single thread block and is configured to process a different set of data based on a single instruction set. In at least one embodiment, one or more threads in a thread warp can be disabled using predication. In at least one embodiment, a lane is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a thread warp. In at least one embodiment, different wavefronts in a thread block can be synchronized together and communicate via shared memory 1454. In at least one embodiment, each graphics processing cluster 1440 includes one or more thread block clusters, where a thread block cluster can enable programmatic control of locality at a granularity larger than a single thread block of a single streaming multi-processor (SM). In at least one embodiment, a thread block cluster (also referred to as a“cluster”) allows multiple thread blocks running across multiple streaming multi-processors to synchronize and cooperatively fetch, exchange, or otherwise use data.

[0186] In at least one embodiment, fabric 1460 is a system interconnect that facilitates data and control transmissions across core complex 1410, graphics complex 1440, I / O interface 1470, memory controllers 1480, display controller 1492, and multimedia engine 1494. In at least one embodiment, APU 1400 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1460 that facilitate data and control transmissions across any number and type of directly or indirectly linked components that can be internal or external to APU 1400. In at least one embodiment, I / O interface 1470 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 1470. In at least one embodiment, peripheral devices coupled to I / O interface 1470 can include, without limitation, a keyboard, a mouse, a printer, a scanner, a joystick or other type of game controller, a media recording device, an external storage device, a network interface card, etc.

[0187] In at least one embodiment, display controller 1492 displays images on one or more display devices, such as liquid crystal display (“LCD”) devices. In at least one embodiment, multimedia engine 1494 includes, without limitation, any number and type of multimedia-related circuitry, such as a video decoder, a video encoder, an image signal processor, etc. In at least one embodiment, memory controllers 1480 facilitate data transfers between APU 1400 and unified system memory 1490. In at least one embodiment, core complex 1410 and graphics complex 1440 share unified system memory 1490.

[0188] In at least one embodiment, APU 1400 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1480 and memory devices (e.g., shared memory 1454) that can be dedicated to one component or shared among multiple components. In at least one embodiment, APU 1400 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1528, L3 cache 1430, and L2 cache 1442), each of which can be private to a component or shared among any number of components (e.g., core 1420, core complex 1410, SIMD unit 1452, compute unit 1450, and graphics complex 1440).

[0189] Figure 15A CPU 1500, according to at least one embodiment, is shown. In at least one embodiment, CPU 1500 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, CPU 1500 can be configured to execute application programs. In at least one embodiment, CPU 1500 is configured to execute host control software, such as an operating system. In at least one embodiment, CPU 1500 issues commands that control the operation of an external GPU (not shown). In at least one embodiment, CPU 1500 can be configured to execute host executable code derived from CUDA source code, and an external GPU can be configured to execute device executable code derived from such CUDA source code. In at least one embodiment, CPU 1500 includes, without limitation, any number of core complexes 1510, fabric 1560, I / O interfaces 1570, and memory controllers 1580.

[0190] In at least one embodiment, core complex 1510 includes, without limitation, cores 1520(1)-1520(4) and L3 cache 1530. In at least one embodiment, core complex 1510 can include, without limitation, any number of cores 1520 and any combination and type of caches. In at least one embodiment, cores 1520 are configured to execute instructions of a particular ISA. In at least one embodiment, each core 1520 is a CPU core. In at least one embodiment, Figure 2 CPU 1500 and / or core complex 1510 can execute Figure 3 object intersection API 228 of Figure 4 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name) of Figure 5 techniques 400 of Figure 15 techniques 500. In at least one embodiment, Figure 2 CPU 1500 and / or core complex 1510 can execute Figure 15 tree generation API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of Figure 2 CPU 1500 and / or core complex 1510 can execute Figure 16 generation lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of

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

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

[0193] In at least one embodiment, fabric 1560 is a system interconnect that facilitates data and control transmissions across core complexes 1510(1)-1510(N) (where N is an integer greater than zero), I / O interface 1570, and memory controllers 1580. In at least one embodiment, CPU 1500 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1560 that facilitate data and control transmissions across any number and type of directly or indirectly linked components that can be internal or external to CPU 1500. In at least one embodiment, I / O interface 1570 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 I / O interface 1570. In at least one embodiment, peripheral devices coupled to I / O interface 1570 can include, without limitation, a display, a keyboard, a mouse, a printer, a scanner, a joystick or other types of game controller, a media recording device, an external storage device, a network interface card, etc.

[0194] In at least one embodiment, memory controllers 1580 facilitate data transfers between CPU 1500 and system memory 1590. In at least one embodiment, core complexes 1510 and graphics complex 1540 share system memory 1590. In at least one embodiment, CPU 1500 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1580 and memory devices that can be dedicated to one component or shared among multiple components. In at least one embodiment, CPU 1500 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1528 and L3 cache 1530), each of which can be private to a component or shared among any number of components (e.g., cores 1520 and core complexes 1510).

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

[0196] Application effective address space 1682 within system memory 1614 stores process elements 1683. In one embodiment, process elements 1683 are stored in response to GPU invocations 1681 from applications 1680 executing on processor 1607. Process elements 1683 contain processing state for a respective application 1680. A work descriptor (WD) 1684 contained in process element 1683 can be a single job requested by an application or can contain pointers to a queue of jobs. In at least one embodiment, WD 1684 is a pointer to a job request queue in application effective address space 1682.

[0197] Graphics acceleration module 1646 and / or individual graphics processing engines can be shared by all or a subset of processes in a system. In at least one embodiment, a graphics acceleration module 1646 can include infrastructure for establishing processing state and sending WDs 1684 to graphics acceleration module 1646 to start jobs in a virtualized environment. In at least one embodiment, Figure 2 Processor 1607 and / or graphics acceleration module 1646 can perform Figure 3 Object intersection API 228 of Figure 4 Object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names) of Figure 5 Techniques 400 of Figure 16 Techniques 500 of In at least one embodiment, Figure 2 Processor 1607 and / or graphics acceleration module 1646 can perform Figure 16 Generation tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API names) of In at least one embodiment, Figure 2The processors 1607 and / or graphics acceleration module 1646 can execute Process and thread identification The generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name).

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

[0199] In operation, a WD fetch unit 1691 in accelerator integration slice 1690 fetches a next WD 1684 including an indication of work to be completed by one or more graphics processing engines of graphics acceleration module 1646. Data from WD 1684 can be stored in registers 1645 used by memory management unit (MMU) 1639, interrupt management circuit 1647, and / or environment management circuit 1648, as shown. For example, one embodiment of MMU 1639 includes segment / page walk circuitry to access segment / page tables 1686 within an OS virtual address space 1685. Interrupt management circuit 1647 can handle interrupt events (INT) 1692 received from graphics acceleration module 1646. Effective addresses 1693 produced by graphics processing engines, when executing graphics operations, are translated to real addresses by MMU 1639.

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

[0201] Table 1 - Hypervisor Initialized Registers

[0202]

[0203]

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

[0205] Table 2 - Operating System Initialized Registers

[0206] 1 Effective address (EA) environment save / restore pointer 2 Virtual address (VA) accelerator utilization record pointer 3 Virtual address (VA) storage segment table pointer 4 Authority mask 5 Work descriptor 6 Figures 17A-17B

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

[0208] Figure 17A An exemplary graphics processor according to at least one embodiment is shown. In at least one embodiment, any of the exemplary graphics processors can be fabricated as one or more IP cores. In addition to the graphics processors illustrated, in at least one embodiment, other logic and circuits can be included, including additional graphics processors / cores, peripheral interface controllers, or general purpose processors. It will be appreciated that the graphics processor is one embodiment in the wide field of graphics processors.

[0209] Figure 17B An exemplary graphics processor 1710 of a SoC integrated circuit, according to at least one embodiment, is shown, which can be fabricated using one or more IP cores. Figure 17A An additional exemplary graphics processor 1740 of a SoC integrated circuit, according to at least one embodiment, is shown, which can be fabricated using one or more IP cores. In at least one embodiment, Figure 17B The graphics processor 1710 of FIG. 17 is a low power graphics processor core. In at least one embodiment, Figure 12 The graphics processor 1740 of FIG. 17 is a higher performance graphics processor core. In at least one embodiment, each graphics processor 1710, 1740 can be Figure 12 Variations of the graphics processor 1210 of FIG. 12.

[0210] In at least one embodiment, graphics processor 1710 includes a vertex processor 1705 and one or more fragment processor(s) 1715A-1715N (e.g., 1715A, 1715B, 1715C, 1715D, through 1715N-1, and 1715N). In at least one embodiment, graphics processor 1710 can execute different shader programs via separate logic for vertex processing and / or for fragment / pixel processing. In at least one embodiment, vertex processor 1705 is optimized to execute operations on vertex data that frequently process complex, high-precision calculations for geometry transformations, motion blur, and dynamic shading. In at least one embodiment, one or more fragment processor(s) 1715A-1715N are optimized to execute more

[0211] In at least one embodiment, graphics processor 1710 additionally includes one or more MMU(s) 1720A-1720B, cache memory 1725A-1725B, and circuit interconnect 1730A-1730B. In at least one embodiment, one or more MMU(s) 1720A-1720B provide for virtual to physical address mapping for graphics processor 1710, including for vertex processor 1705 and / or for fragment processor(s) 1715A-1715N, which can reference vertex or image / texture data stored in memory, in addition to vertex or image / texture data stored in one or more cache(s) 1725A-1725B. In at least one embodiment, one or more MMU(s) 1720A-1720B can be synchronized with other MMUs within Figure 17A In at least one embodiment, one or more application processor(s) 1205, image processor(s) 1215, and / or video processor(s) 1220 are associated with one or more MMUs such that each processor 1205-1220 can participate in a shared or unified virtual memory system. In at least one embodiment, one or more circuit interconnects 1730A-1730B enable graphics processor 1710 to interface with other IP cores within a SoC, either via an internal bus, or via a direct connection.

[0212] In at least one embodiment, graphics processor 1740 includes Figure 17Aone or more MMUs 1720A-1720B, caches 1725A-1725B, and circuit interconnect 1730A-1730B of graphics processor 1710. In at least one embodiment, graphics processor 1740 includes one or more shader core(s) 1755A-1755N (e.g., 1755A, 1755B, 1755C, 1755D, 1755E, 1755F, through 1755N-1, and 1755N), which provides unified shader core architecture in which a single core or type or core can execute all types of programmable shader code, including shader program code to implement vertex shaders, fragment shaders, and / or compute shaders. In at least one embodiment, a plurality of shader cores can be varied. In at least one embodiment, graphics processor 1740 includes an inter-core task manager 1745, which acts as a thread dispatcher to dispatch execution threads to one or more shader cores 1755A-1755N and a tiling unit 1758 to accelerate tiling operations for tile-based rendering in which render operations for a scene are subdivided in image space, for example, to utilize spatial coherence within a scene or to optimize use of internal caches; in at least one embodiment, Figure 17B graphics processor 1710 and / or Figure 2 graphics processor 1740 can perform Figure 3 object intersection API 228 of Figure 4 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name) of Figure 5 techniques 400 of Figure 17A techniques 500 of In at least one embodiment, Figure 17B graphics processor 1710 and / or Figure 2 graphics processor 1740 can perform Figure 17A generation tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) of Figure 17B graphics processor 1710 and / or Figure 2 graphics processor 1740 can perform Figure 18A generation lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of

[0213] Figure 12A graphics core 1800 according to at least one embodiment is shown. In at least one embodiment, the graphics core 1800 may include... Figure 17B The graphics processor 1210 is located within it. In at least one embodiment, the graphics core 1800 may be... Figure 18B The graphics core 1800 uses a unified shader core 1755A-1755N. In at least one embodiment, the graphics core 1800 includes a shared instruction cache 1802, texture units 1818, and cache / shared memory 1820, which are common to execution resources within the graphics core 1800. In at least one embodiment, the graphics core 1800 may include multiple slices 1801A-1801N or partitions of each core, and the graphics processor may include multiple instances of the graphics core 1800. Slices 1801A-1801N may include supporting logic, including local instruction caches 1804A-1804N, thread schedulers 1806A-1806N, thread dispatchers 1808A-1808N, and a set of registers 1810A-1810N. In at least one embodiment, slices 1801A-1801N may include a set of additional functional units (AFU) 1812A-1812N, floating-point units (FPU) 1814A-1814N, integer arithmetic logic units (ALU) 1816A-1816N, address calculation units (ACU) 1813A-1813N, double-precision floating-point units (DPFPU) 1815A-1815N, and matrix processing units (MPU) 1817A-1817N. In at least one embodiment, graphics core 1800 is referred to as a computing unit or arithmetic unit.

[0214] In one embodiment, the FPU 1814A-1814N can perform single-precision (32-bit) and half-precision (16-bit) floating-point operations, while the DPFPU 1815A-1815N can perform double-precision (64-bit) floating-point operations. In at least one embodiment, the ALU 1816A-1816N can perform variable-precision integer operations with 8-bit, 16-bit, and 32-bit precision, and can be configured for mixed-precision operations. In at least one embodiment, the MPU 1817A-1817N can also be configured for mixed-precision matrix operations, including half-precision floating-point operations and 8-bit integer operations. In at least one embodiment, the MPU 1817A-1817N can perform various matrix operations to accelerate CUDA programs, including enabling accelerated Generalized Matrix-to-Matrix Multiplication (GEMM). In at least one embodiment, the AFU 1812A-1812N can perform additional logical operations not supported by floating-point or integer units, including trigonometric operations (e.g., Sine, Cosine, etc.).

[0215] Figure 18AA general-purpose graphics processing unit (GPGPU) 1830 is illustrated in at least one embodiment. In at least one embodiment, the GPGPU 1830 is highly parallel and suitable for deployment on a multi-chip module. In at least one embodiment, the GPGPU 1830 can be configured to enable highly parallel computational operations to be performed by a GPU array. In at least one embodiment, the GPGPU 1830 can be directly linked to other instances of the GPGPU 1830 to create a multi-GPU cluster to improve execution time for CUDA programs. In at least one embodiment, the GPGPU 1830 includes a host interface 1832 for connection to a host processor. In at least one embodiment, the host interface 1832 is a PCIe interface. In at least one embodiment, the host interface 1832 can be a vendor-specific communication interface or communication structure. In at least one embodiment, the GPGPU 1830 receives commands from the host processor and uses a global scheduler 1834 to assign execution threads associated with those commands to a set of compute clusters 1836A-1836H. In at least one embodiment, computing clusters 1836A-1836H share cache memory 1838. In at least one embodiment, cache memory 1838 can be used as an advanced cache of cache memory within computing clusters 1836A-1836H.

[0216] In at least one embodiment, the GPGPU 1830 includes memories 1844A-1844B coupled to the computing cluster 1836A-1836H via a set of memory controllers 1842A-1842B. In at least one embodiment, memories 1844A-1844B may include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory. In at least one embodiment, Figure 18B Graphics core 1800 and / or Figure 2 The GPGPU 1830 can execute Figure 3 Object Intersection API 228 Figure 4 Object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name), Figure 5 Technology 400 and / or Figure 18A Technology 500. In at least one embodiment, Figure 18B Graphics core 1800 and / or Figure 2 The GPGPU 1830 can execute Figure 18Aof the spanning tree API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name). In at least one embodiment, Figure 18B of the graphics core 1800 and / or Figure 2 of the GPGPU 1830 can execute Figure 18A of the generate lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name).

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

[0218] In at least one embodiment, multiple instances of the GPGPU 1830 can be configured to operate as computing clusters. Computing clusters 1836A-1836H can implement any technically feasible communication technology for synchronization and data exchange. In at least one embodiment, multiple instances of the GPGPU 1830 communicate over host interface 1832. In at least one embodiment, GPGPU 1830 includes an I / O hub 1839 that couples the GPGPU 1830 with a GPU link 1840 that enables a direct connection to other instances of GPGPU 1830. In at least one embodiment, GPU link 1840 is coupled to a dedicated GPU-to-GPU bridge that enables communication and synchronization between multiple instances of GPGPU 1830. In at least one embodiment, GPU link 1840 is coupled with a high-speed interconnect to transmit and receive data to other GPGPUs or parallel processors. In at least one embodiment, multiple instances of GPGPU 1830 are located in separate data processing systems and communicate over a network device that is accessible via host interface 1832. In at least one embodiment, GPU link 1840 can be configured to connect to a host processor, in addition to or in lieu of host interface 1832. In at least one embodiment, GPGPU 1830 can be configured to execute CUDA programs.

[0219] ​A parallel processor 1900 is shown according to at least one embodiment. In at least one embodiment, various components of parallel processor 1900 can be implemented using one or more integrated circuits, for example, programmable processors, application specific integrated circuits (ASICs), or FPGAs.

[0220] In at least one embodiment, parallel processor 1900 includes a parallel processing unit 1902. In at least one embodiment, parallel processing unit 1902 includes an I / O unit 1904 that enables communication with other devices, including other instances of parallel processing unit 1902. In at least one embodiment, I / O unit 1904 can be directly connected to the other devices. In at least one embodiment, I / O unit 1904 connects with other devices using a hub or switch interface, for example, memory hub 1905. In at least one embodiment, connections between memory hub 1905 and I / O unit 1904 form a communication link. In at least one embodiment, I / O unit 1904 connects with a host interface 1906 and a memory crossbar switch 1916, where host interface 1906 receives commands directed to the processing operations and memory crossbar switch 1916 receives commands directed to the memory operations.

[0221] In at least one embodiment, when host interface 1906 receives a command buffer via I / O unit 1904, host interface 1906 can direct work operations to execute those commands to front end 1908. In at least one embodiment, front end 1908 couples with a scheduler 1910, which is configured to assign commands or other work items to processing array 1912. In at least one embodiment, scheduler 1910 ensures that processing array 1912 is properly configured and in an active state before assigning tasks to processing array 1912 of processing array 1912. In at least one embodiment, scheduler 1910 is implemented by firmware logic executing on a microcontroller. In at least one embodiment, microcontroller- implemented scheduler 1910 is configurable to perform complex scheduling and work distribution operations at a coarse and fine grain level to achieve a quick preemption and context switch for threads executing on processing array 1912. In at least one embodiment, host software can prove a work load for scheduling on processing array 1912 through one of a number of graphics processing doorbells. In at least one embodiment, work load can then be automatically distributed on processing array 1912 by scheduler 1910 logic within microcontroller that includes scheduler 1910.

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

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

[0224] In at least one embodiment, processing array 1912 is configured to perform parallel graphics processing operations. In at least one embodiment, processing array 1912 can include additional logic to support performance of such graphics processing operations, including but not limited to texture sampling logic to perform texture operations, tessellation logic, and other vertex processing logic. In at least one embodiment, processing array 1912 can be configured to execute graphics processing related shader programs, such as, but not limited to, vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. In at least one embodiment, parallel processing unit 1902 can transfer data from system memory for processing via I / O unit 1904. In at least one embodiment, during processing, results can be stored to on-chip memory (e.g., parallel processor memory 1922) for write back to system memory.

[0225] In at least one embodiment, when parallel processing unit 1902 is used to perform graphics processing, scheduler 1910 can be configured to divide incoming workloads into tasks of approximately equal size to better enable distribution of graphics processing operations across multiple clusters 1914A-1914N of processing array 1912. In at least one embodiment, portions of processing array 1912 can be configured to perform different types of processing. For example, in at least one embodiment, a first portion can be configured to

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

[0227] In at least one embodiment, each of one or more instances of parallel processing unit 1902 can be coupled to a parallel processor memory 1922. In at least one embodiment, parallel processor memory 1922 can be accessed by the processing array 1912, as well as the I / O unit 1904, via a memory crossbar 1916. In at least one embodiment, memory crossbar 1916 can be used to transfer data between memory elements and the processing array 1912, I / O unit 1904, and / or other components of the parallel processing unit 1902. In at least one embodiment, a memory controller 1918 can be included in the parallel processing unit 1902 to perform memory access operations. In at least one embodiment, memory controller 1918 can be a dedicated controller, or can be integrated within another component (e.g., an I / O unit 1904).

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

[0229] In at least one embodiment, any of clusters 1914A-1914N of processing array 1912 can process data that is to be written into any of memory units 1924A-1924N within parallel processor memory 1922. In at least one embodiment, memory crossbar 1916 can be configured to transmit outputs of each cluster 1914A-1914N to any partition unit 1920A-1920N or another cluster 1914A-1914N, which can perform other processing operations on the outputs. In at least one embodiment, each cluster 1914A-1914N can communicate with memory interface 1918 through memory crossbar 1916 to read from or write to various external memory devices. In at least one embodiment, memory crossbar 1916 has a connection to memory interface 1918 to communicate with I / O unit 1904 and a local instance of parallel processor memory 1922, to enable processing units within different processing clusters 1914A-1914N to communicate with system memory or other memories not local to the parallel processing units 1902. In at least one embodiment, memory crossbar 1916 can use virtual channels to separate traffic streams between clusters 1914A-1914N and partition units 1920A-1920N.

[0230] In at least one embodiment, multiple instances of parallel processing unit 1902 can be provided on a single add-in card, or multiple add-in cards can be interconnected. In at least one embodiment, different instances of parallel processing unit 1902 can be configured to operate in coordination with each other to enable single program multi-processing (SPMP). In at least one embodiment, SPMP enables a single program to utilize the resources of multiple parallel processing units 1902 by running all of the program’s threads on the different parallel processing units 1902. For example, in at least one embodiment, some instances of parallel processing unit 1902 can include a higher precision floating point unit relative to other instances. In at least one embodiment, systems incorporating one or more instances of parallel processing unit 1902 or parallel processor 1900 can be implemented in a variety of configurations and form factors, including but not limited to desktop, laptop, or hand-held personal computers, servers, workstations, game consoles, and / or embedded systems.

[0231] FIG. 19B A processing cluster 1994 is shown according to at least one embodiment. In at least one embodiment, processing cluster 1994 is included in a parallel processing unit. In at least one embodiment, processing cluster 1994 is a FIG. 19Aone of the processing clusters 1914A-1914N. In at least one embodiment, processing cluster 1994 can be configured to execute many threads in parallel, where the term “thread” refers to an instance of a particular program executed by a particular group of one or more processing clusters. In at least one embodiment, Single Instruction Multiple Data (SIMD) instruction issue techniques are used to support parallel execution of a large number of threads with no or negligible program overhead. In at least one embodiment, Single Instruction Multiple Thread (SIMT) techniques are used to support parallel execution of a large number of generally synchronous threads using a common instruction unit configured to issue instructions to a group of processing engines within each processing cluster 1994.

[0232] In at least one embodiment, operation of processing cluster 1994 can be controlled by a pipeline manager 1932 that allocates processing tasks to SIMT parallel processor. In at least one embodiment, pipeline manager 1932 receives instructions from scheduler 1910 of a graphics processing pipeline, and manages execution of those instructions via graphics multiprocessor 1934 and / or texture unit 1936. In at least one embodiment, graphics multiprocessor 1934 is an exemplary instance of a SIMT parallel processor. However, in at least one embodiment, various types of SIMT parallel processors of differing architectures can be included within processing cluster 1994. In at least one embodiment, one or more instances of graphics multiprocessor 1934 can be included within processing cluster 1994. In at least one embodiment, graphics multiprocessor 1934 can process data, and a data crossbar 1940 can be used to distribute processed data to one of a number of possible destinations, including other shader units. In at least one embodiment, pipeline manager 1932 can facilitate distribution by specifying destinations for processed data as a function of its origin. FIG. 19A

[0233] In at least one embodiment, each graphics multiprocessor 1934 within processing cluster 1994 can include an identical set of functional execution logic (e.g., arithmetic logic units, load store units (LSUs), etc.). In at least one embodiment, functional execution logic can be configured in a pipelined manner in which new instructions can be issued before previous instructions are complete. In at least one embodiment, functional execution logic supports a variety of operations including integer and floating point arithmetic, comparison operations, Boolean operations, shift operations, and the like. In at least one embodiment, same functional-unit hardware can be leveraged to perform different operations using different control signals. Any combination of

[0234] ​In at least one embodiment, instructions delivered to processing cluster 1994 form a thread. In at least one embodiment, a set of threads executing across a set of parallel processing engines forms a thread group. In at least one embodiment, a thread group executes a program on different input data. In at least one embodiment, each thread within a thread group can be assigned to a different processing engine within graphics multiprocessor 1934. In at least one embodiment, a thread group can include fewer threads than are present in a plurality of processing engines. In at least one embodiment, when a thread group includes fewer threads than the number of processing engines present in graphics multiprocessor 1934, one or more of the processing engines can be idle during the execution of the threads in the thread group. In at least one embodiment, a thread group can also include more threads than are present in a plurality of processing engines. In at least one embodiment, when a thread group includes more threads than the number of processing engines present in graphics multiprocessor 1934, multiple threads can be executed concurrently on the processing engines present in graphics multiprocessor 1934. In at least one embodiment, a plurality of thread groups can be executed concurrently by graphics multiprocessor 1934.

[0235] In at least one embodiment, graphics multiprocessor 1934 includes internal cache memory, to perform load and store operations. In at least one embodiment, graphics multiprocessor 1934 can bypass internal cache and use cache memory within processing cluster 1994 (e.g., Ll cache 1948). In at least one embodiment, each graphics multiprocessor 1934 can also have access to L2 Cache within a partition unit (e.g., partition units 1920A-1920N) that is shared among all processing clusters 1994 and can be used to transfer data between threads. In at least one embodiment, graphics multiprocessor 1934 can also have access to off-chip global memory, which can include one or more of local parallel processor memory and / or system memory. In at least one embodiment, any memory external to parallel processor 1902 can be used as global memory. In at least one embodiment, processing cluster 1994 includes multiple instances of graphics multiprocessor 1934 that share shared cache memory, which can be stored in Ll cache 1948, which can be used to store shared data. FIG. 19A

[0236] In at least one embodiment, one or more instances of MMU 1945 can reside within graphics multiprocessor 1934. In at least one embodiment, MMU 1945 maps virtual addresses into physical addresses as part of processing cluster 1994's access to cache memory and / or system memory. In at least one embodiment, MMU 1945 provides translation lookaside buffer (TLB) functionality, as is known in the art. In at least one embodiment, processing cluster 1994 includes hardware logic to implement a TLB for address translation and page permission checking. FIG. 19A ​The memory interface 1918 is located within the MMU. In at least one embodiment, the MMU 1945 includes a set of page table entries (PTEs) for mapping virtual addresses to physical addresses of tiles (more information about tiles is discussed) and optionally to cache line indices. In at least one embodiment, the MMU 1945 may include an address translation back buffer (TLB) or a cache that may reside within the graphics multiprocessor 1934, L1 cache 1948, or processing cluster 1994. In at least one embodiment, physical addresses are processed to allocate surface data access locality for efficient request interleaving between partition units. In at least one embodiment, cache line indices may be used to determine whether a request for a cache line is a hit or a miss.

[0237] In at least one embodiment, the processing cluster 1994 can be configured such that each graphics multiprocessor 1934 is coupled to a texture unit 1936 to perform texture mapping operations, which may involve determining texture sample locations, reading texture data, and filtering texture data. In at least one embodiment, texture data is read as needed from an internal texture L1 cache (not shown) or from an L1 cache within the graphics multiprocessor 1934, and texture data is also retrieved from an L2 cache, local parallel processor memory, or system memory. In at least one embodiment, each graphics multiprocessor 1934 outputs a processed task to a data crossbar switch 1940 to provide the processed task to another processing cluster 1994 for further processing or to store the processed task in an L2 cache, local parallel processor memory, or system memory via a memory crossbar switch 1916. In at least one embodiment, a pre-raster operation unit (preROP) 1942 is configured to receive data from the graphics multiprocessor 1934 and direct the data to a ROP unit, which may be associated with a partitioning unit (e.g., [missing information]). FIG. 19A The PreROP1942 unit is located together with the partition units 1920A-1920N. In at least one embodiment, the PreROP1942 unit can perform optimizations for color blending, organize pixel color data, and perform address translation.

[0238] FIG. 19C A graphics multiprocessor 1996 according to at least one embodiment is illustrated. In at least one embodiment, the graphics multiprocessor 1996 is... FIG. 19Bgraphics processor 1934. In at least one embodiment, graphics processor 1996 is coupled with pipeline manager 1932 of processing cluster 1994. In at least one embodiment, graphics processor 1996 has a thread execution pipeline that includes, without limitation, an instruction cache 1952, an instruction unit 1954, an address mapping unit 1956, a register file 1958, one or more GPGPU cores 1962, and one or more LSU’s 1966. GPGPU cores 1962 and LSUs 1966 are coupled with cache memory 1972 and shared memory 1970 via memory and cache interconnect 1968.

[0239] In at least one embodiment, instruction cache 1952 receives a stream of instructions 1950 to execute from pipeline manager 1932. In at least one embodiment, instructions are cached in instruction cache 1952 and dispatched for execution by instruction unit 1954. In one embodiment, instruction unit 1954 can dispatch instructions to the threads of a thread group, with each thread in the thread group assigned to a different execution lane within GPGPU core 1962. In at least one embodiment, instructions can access any of the local, shared, or global address spaces by specifying an address in a unified address space. In at least one embodiment, address mapping unit 1956 can be used to convert an address in the unified address space into an address that can be accessed by LSU 1966.

[0240] In at least one embodiment, register file 1958 provides a set of registers for functional units of graphics processor 1996. In at least one embodiment, register file 1958 provides temporary storage for operands of the data

[0241] In at least one embodiment, GPGPU cores 1962 can each include FPUs and / or ALUs for executing instructions for graphics processing. GPGPU cores 1962 can be similar to each other in architecture or can differ from each other in architecture. In at least one embodiment, a first portion of GPGPU cores 1962 includes single precision FPUs and integer ALUs, while a second portion of GPGPU cores 1962 includes double precision FPUs. In at least one embodiment, FPUs can implement IEEE 754-2008 standard for floating point arithmetic or enable variable precision floating point arithmetic. In at least one embodiment, graphics processor 1900 can additionally include one or more fixed function or special purpose logic units to perform specific functions such as copy rectangle or pixel blending operations. In at least one embodiment one or more of GPGPU cores 1962 can also include fixed or special function logic.

[0242] In at least one embodiment, GPGPU cores 1962 include SIMD logic capable of

[0243] In at least one embodiment, memory and cache interconnect 1968 is an interconnect network that connects each functional unit of graphics multiprocessor 1996 to register file 1958 and shared memory 1970. In at least one embodiment, memory and cache interconnect 1968 is a crossbar interconnect that allows LSUs 1966 to effect load and store operations between shared memory 1970 and register file 1958. In at least one embodiment, register file 1958 can operate at same frequency as GPGPU cores 1962, making the latency in transferring data between GPGPU cores 1962 and register file 1958 very low. In at least one embodiment, shared memory 1970 can be used to enable FIG. 19A parallel processor 1900 of FIG. 1, FIG. 19B processing cluster 1994 of FIG. 1, FIG. 19B graphics multiprocessor 1934 of FIG. 1, and / or FIG. 19C graphics multiprocessor 1996 of FIG. 1 can perform FIG. 2 object intersection API 228 of FIG. 1, FIG. 3 object intersection API 300 of FIG. 1 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name), FIG. 4 techniques 400 of FIG. 1, and / or FIG. 5 techniques 500 of FIG. 1. In at least one embodiment, FIG. 19A parallel processor 1900 of FIG. 1, FIG. 19B processing cluster 1994 of FIG. 1, FIG. 19B graphics multiprocessor 1934 of FIG. 1, and / or FIG. 19C graphics multiprocessor 1996 of FIG. 1 can perform FIG. 2 generation tree API 230 of FIG. 1 (e.g., generateTree(), buildTree(), and / or some other suitable API name). In at least one embodiment, FIG. 19A parallel processor 1900 of FIG. 1, FIG. 19B processing cluster 1994 of FIG. 1, FIG. 19B graphics multiprocessor 1934 of FIG. 1, and / or FIG. 19Cgraphics multiprocessor 1996 in FIG. 1 can execute FIG. 2 generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names) in FIG. 1.

[0244] In at least one embodiment, 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, GPU can be communicatively coupled to host processor / cores by a bus or other interconnect (e.g., a high speed

[0245] FIG. 20 A graphics processor 2000 according to at least one embodiment is shown. In at least one embodiment, graphics processor 2000 includes a ring interconnect 2002, a front-end 2004, a media engine 2037, and graphics cores 2080A-2080N. In at least one embodiment, ring interconnect 2002 couples graphics processor 2000 to other processing units, including other graphics processors or one or more general-purpose processor cores. In at least one embodiment, graphics processor 2000 is one of a number of processors integrated within a multi-core processing system. In at least one embodiment, FIG. 20 graphics processor 2000 of FIG. 2 can execute FIG. 2 object intersection API 228 of FIG. 2, FIG. 3 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names) of FIG. 3, FIG. 4 techniques 400 of FIG. 4, and / or FIG. 5 techniques 500 of FIG. 5. In at least one embodiment, FIG. 20 graphics processor 2000 of FIG. 2 can execute FIG. 2 tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API names) of FIG. 3. In at least one embodiment, FIG. 20graphics processor 2000 can execute FIG. 2 generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names).

[0246] In at least one embodiment, graphics processor 2000 receives batches of commands via ring interconnect 2002. In at least one embodiment, incoming commands are interpreted by command streamer 2003 in pipeline front-end 2004. In at least one embodiment, graphics processor 2000 includes scalable execution logic to perform 3D geometry processing and media processing via graphics core(s) 2080A-2080N. In at least one embodiment, for 3D geometry processing commands, command streamer 2003 supplies commands to geometry pipeline 2036. In at least one embodiment, for at least some media processing commands, command streamer 2003 supplies commands to video front-end 2034, which couples with a media engine 2037. In at least one embodiment, media engine 2037 includes a video quality engine (VQE) 2030 to perform video and image post processing and a multi-format encode / decode (MFX) 2033 engine to

[0247] In at least one embodiment, graphics processor 2000 includes a scalable thread execution resource featuring modular graphics core 2080A-2080N (sometimes referred to as core slices), each having multiple sub-cores 2050A-2050N, 2060A-2060N (sometimes referred to as core sub-slices). In at least one embodiment, graphics processor 2000 can have any number of graphics cores 2080A-2080N. In at least one embodiment, graphics processor 2000 includes graphics core 2080A having at least first sub-core 2050A and second sub-core 2060A. In at least one embodiment, graphics processor 2000 is a low power processor with a single sub-core (e.g., 2050A). In at least one embodiment, graphics processor 2000 includes multiple graphics cores 2080A-2080N each including a set of first sub-cores 2050A-2050N and a set of second sub-cores 2060A-2060N. In at least one embodiment, each of first sub-cores 2050A-2050N includes at least a first set of execution units (EUs) 2052A-2052N and a media / texture sampler 2054A-2054N. In at least one embodiment, each of second sub-cores 2060A-2060N includes at least a second set of execution units (EUs) 2062A-2062N and a sampler 2064A-2064N. In at least one embodiment, each sub-core 2050A-2050N, 2060A-2060N shares a set of shared resources 2070A-2070N. In at least one embodiment, shared resources include shared cache memory and pixel operation logic.

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

[0249] In at least one embodiment, processor 2100 includes an in-order front end (“front end”) 2101 to fetch instructions to be executed and prepare instructions to be used later in processor pipeline. In at least one embodiment, front end 2101 can include several units. In at least one embodiment, instruction prefetcher 2126 fetches instructions from memory and feeds instructions to instruction decoder 2128 which, in turn, decodes or interprets instructions. For example, in at least one embodiment, instruction decoder 2128 decodes a received instruction into one or more operations called “micro-instructions” or “micro-operations” (also called “micro ops” or “uops”) for execution. In at least one embodiment, instruction decoder 2128 parses the instruction into an operation code and corresponding data and control fields, which can be used by micro-architecture to perform the operation. In at least one embodiment, trace cache 2130 can assemble decoded uops into program ordered sequences or traces in uop queue 2134 for execution. In at least one embodiment, when trace cache 2130 encounters complex instructions, microcode ROM 2132 provides uops needed to complete an operation. In at least one embodiment, FIG. 21 Processor 2100 of at least one embodiment can perform FIG. 2 Object Intersection API 228 of at least one embodiment, FIG. 3 Object Intersection API 300 of at least one embodiment (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name), FIG. 4 Techniques 400 of at least one embodiment, and / or FIG. 5 Techniques 500 of at least one embodiment. In at least one embodiment, FIG. 21 Processor 2100 of at least one embodiment can perform FIG. 2 Generate Tree API 230 of at least one embodiment (e.g., generateTree(), buildTree(), and / or other suitable API name). In at least one embodiment, FIG. 21 Processor 2100 of at least one embodiment can perform FIG. 2 Generate Lookup Parameters API 232 of at least one embodiment (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API name).

[0250] In at least one embodiment, some instructions can be converted into a single micro- operation, while others can require several micro-operations to complete the full operation. In at least one embodiment, if more than four micro-instructions are needed to complete an instruction, then instruction decoder 2128 can access microcode ROM 2132 to execute the instruction. In at least one embodiment, an instruction can be decoded into a small number of micro-instructions to handle at instruction decoder 2128. In at least one embodiment, if multiple micro-instructions are needed to complete an operation, then an instruction can be stored in microcode ROM 2132. In at least one embodiment, a trace cache 2130 references an entry point programmable logic array (“PLA”) to determine a correct micro-instruction pointer for reading a microcode sequence from microcode ROM 2132 to complete one or more instructions, in accordance with at least one embodiment. In at least one embodiment, after microcode ROM 2132 completes sequencing of micro-operations for an instruction, a front end 2101 of a machine can resume fetching micro-operations from trace cache 2130.

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

[0252] In at least one embodiment, execution block 2111 includes, without limitation, an integer register file / bypass network 2108, a floating point register file / bypass network (“FP register file / bypass network”) 2110, address generation units (“AGUs”) 2112 and 2114, fast ALUs 2116 and 2118, a slow ALU 2120, a floating point ALU (“FP”) 2122, and a floating point move unit (“FP move”) 2124. In at least one embodiment, integer register file / bypass network 2108 and floating point register file / bypass network 2110 are also referred to herein as “register files 2108, 2110.” In at least one embodiment, AGUs 2112 and 2114, fast ALUs 2116 and 2118, slow ALU 2120, floating point ALU 2122, and floating point move unit 2124 are also referred to herein as “execution units 2112, 2114, 2116, 2118, 2120, 2122, and 2124.” In at least one embodiment, execution block can include, without limitation, any number (including zero) and type of register files, bypass networks, address generation units, and execution units (in any combination).

[0253] In at least one embodiment, register files 2108, 2110 can be arranged between microinstruction schedulers 2102, 2104, 2106 and execution units 2112, 2114, 2116, 2118, 2120, 2122, and 2124. In at least one embodiment, integer register file / bypass network 2108 performs integer operations. In at least one embodiment, floating point register file / bypass network 2110 performs floating point operations. In at least one embodiment, each of register files 2108, 2110 can include, without limitation, a bypass network that can bypass or forward a just-completed result that has not yet been written into a register file to a new dependee. In at least one embodiment, register files 2108, 2110 can communicate data with each other. In at least one embodiment, integer register file / bypass network 2108 can include, without limitation, two separate register files, one for lower 32 bits of data and a second for upper 32 bits of data. In at least one embodiment, floating point register file / bypass network 2110 can include, without limitation, 128 bit wide entries, as floating point instructions typically have operands that are 64 to 128 bits wide.

[0254] In at least one embodiment, execution units 2112, 2114, 2116, 2118, 2120, 2122, 2124 can execute instructions. In at least one embodiment, register files 2108, 2110 store integer and floating point data operand values upon which microinstructions require execution. In at least one embodiment, processor 2100 can include, without limitation, any number of execution units 2112, 2114, 2116, 2118, 2120, 2122, 2124 and combinations thereof. In at least one embodiment, floating point ALU 2122 and floating point move unit 2124 can execute floating point, MMX, SIMD, AVX and SSE, or other operations, including specialized machine learning instructions. In at least one embodiment, floating point ALU 2122 can include, without limitation, a 64 bit by 64 bit floating point divider to execute divide, square root, and remainder micro-ops. In at least one embodiment, instructions for floating point

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

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

[0257] FIG. 22 A processor 2200 according to at least one embodiment is shown. In at least one embodiment, processor 2200 includes, without limitation, one or more processor cores (cores) 2202A-2202N, an integrated memory controller 2214, and an integrated graphics processor 2208. In at least one embodiment, processor 2200 can include additional cores up to and including the additional processor cores 2202N represented by the dashed box. In at least one embodiment, each processor core 2202A-2202N includes one or more internal cache units 2204A-2204N. In at least one embodiment, each processor core can also access one or more shared cache units 2206. In at least one embodiment, one or more processor cores 220A-2202N are referred to as one or more compute units or arithmetic units.

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

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

[0260] In at least one embodiment, one or more processor cores 2202A-2202N include support to run in multiple threads simultaneously. In at least one embodiment, system agent core 2210 includes components for coordination and operation of processor cores 2202A-2202N during multi-threaded processing. In at least one embodiment, system agent core 2210 can additionally include a power control unit (PCU), including logic and components to regulate one or more power states of processor cores 2202A-2202N and graphics processor 2208.

[0261] In at least one embodiment, processor 2200 additionally includes a graphics processor 2208 to perform graphics processing operations. In at least one embodiment, graphics processor 2208 couples with shared cache unit 2206 and system agent core 2210, including one or more integrated memory controllers 2214. In at least one embodiment, system agent core 2210 also includes a display controller 2211 to drive graphics processor output to one or more coupled displays. In at least one embodiment, display controller 2211 can also be a separate module coupled with graphics processor 2208 via at least one interconnect, or can be integrated within graphics processor 2208. In at least one embodiment, FIG. 22The processor 2200 and / or graphics processor 2208 can execute FIG. 2 Object Intersection API 228 FIG. 3 Object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names), FIG. 4 Technology 400 and / or FIG. 5 Technology 500. In at least one embodiment, FIG. 22 The processor 2200 and / or graphics processor 2208 can execute FIG. 2 The tree generation API 230 (e.g., generateTree(), buildTree(), and / or other suitable API names). In at least one embodiment, FIG. 22 The processor 2200 and / or graphics processor 2208 can execute FIG. 2 Generate lookup parameters using API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name).

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

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

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

[0265] FIG. 23 A graphics processor core 2300 according to at least one described embodiment is shown. In at least one embodiment, graphics processor core 2300 is included within a graphics core array. In at least one embodiment, graphics processor core 2300 (sometimes called a core slice) can be one or more graphics cores within a modular graphics processor. In at least one embodiment, graphics processor core 2300 is an example of one graphics core slice, and a graphics processor described herein can include multiple graphics core slices based on target power and performance envelope. In at least one embodiment, each graphics core 2300 can include fixed function and programmable processing

[0266] In at least one embodiment, fixed function block 2330 includes a geometry / fixed function pipeline 2336, for example, that can be shared by all sub-cores in graphics processor 2300, for example, in lower performance and / or lower power graphics processor implementations. In at least one embodiment, geometry / fixed function pipeline 2336 includes a 3D fixed function pipeline, a video front-end unit, a thread generator and thread dispatcher, and a unified return buffer manager that manages a unified return buffer.

[0267] In at least one embodiment, fixed function block 2330 also includes a graphics SoC interface 2337, a graphics microcontroller 2338, and a media pipeline 2339. Graphics SoC interface 2337 provides an interface between graphics core 2300 and other processor cores in a SoC integrated circuit. In at least one embodiment, graphics microcontroller 2338 is a programmable sub-processor that is configurable for managing graphics processing engine 2300. In at least one embodiment, graphics microcontroller 2338 includes programmable logic to manage various functions of graphics processing engine 2300, including thread dispatch, scheduling, and pre-emption.

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

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

[0270] In at least one embodiment, graphics core 2300 can have more or less than the illustrated sub-cores 2301A-2301F, up to N modular sub-cores. For each set of N sub-cores, in at least one embodiment, graphics core 2300 can also include shared function logic 2310, shared and / or cache memory 2312, geometry / fixed function pipeline 2314, and additional fixed function logic 2316 to accelerate various graphics and compute processing tasks. In at least one embodiment, shared function logic 2310 can include logic units (e.g., samplers, math, and / or inter-thread communication logic) that can be shared by each N sub-core within graphics core 2300. Shared and / or cache memory 2312 can be an LLC for N sub-cores 2301A-2301F within graphics core 2300, and can also serve as shared memory accessible by multiple sub-cores. In at least one embodiment, geometry / fixed function pipeline 2314 can be included in place of geometry / fixed function pipeline 2336 within fixed function block 2330, and can include the same or similar logic units.

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

[0272] In at least one embodiment, additional fixed function logic 2316 can also include general purpose processing acceleration logic, such as fixed function matrix multiplication logic, to implement a reduced CUAD program. In at least one embodiment, FIG. 23 In at least one embodiment, graphics core 2300 of FIG. 21 can perform FIG. 2 object intersection API 228 of FIG. 22, FIG. 3 object intersection API 300 of FIG. 23 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name), FIG. 4 techniques 400 of FIG. 24, and / or FIG. 5 techniques 500 of FIG. 25. In at least one embodiment, FIG. 23 In at least one embodiment, graphics core 2300 of FIG. 21 can perform FIG. 2 generate tree API 230 of FIG. 26 (e.g., generateTree(), buildTree(), and / or other suitable API name). In at least one embodiment, FIG. 23 In at least one embodiment, graphics core 2300 of FIG. 21 can perform FIG. 2generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names).

[0273] In at least one embodiment, each graphics sub-core 2301A-2301F includes a set of execution resources that can be used to perform graphics, media, and compute operations in response to requests by graphics pipeline, media pipeline, or shader programs. In at least one embodiment, graphics sub-cores 2301A-2301F include multiple arrays of EU’s 2302A-2302F, 2304A-2304F, thread dispatch and inter-thread communication (TD / IC) logic 2303A-2303F, 3D (e.g., texture) samplers 2305A-2305F, media samplers 2306A-2306F, shader processors 2307A-2307F, and shared local memory (SLM) 2308A-2308F. Arrays of EU’s 2302A-2302F, 2304A-2304F each include multiple execution units (EUs), which are GPGPU cores capable of performing floating-point and integer / fixed-point logic operations including graphics, media, and compute shader programs. In at least one embodiment, TD / IC logic 2303A-2303F performs local thread dispatch and thread control operations for execution units within a sub-core and facilitate

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

[0275] In at least one embodiment, PPU(s) 2400 are configured to accelerate high- performance computing (“HPC”), datacenter, and machine learning applications. In at least one embodiment, PPU(s) 2400 are configured to accelerate CUDA programs. In at least one embodiment, PPU 2400 includes, without limitation, I / O unit 2406, front-end unit 2410, scheduler unit 2412, work distribution unit 2414, hub 2416, crossbar (“Xbar”) 2420, one or more general processing clusters (“GPCs”) 2418, and one or more partition units (“memory partition units”) 2422. In at least one embodiment, PPU(s) 2400 connect to a host processor or other PPU(s) 2400 by one or more high-speed GPU interconnects 2408. In at least one embodiment, PPU(s) 2400 connect to host processors or other peripherals by a system bus or interconnect 2402. In at least one embodiment, PPU(s) 2400 connect to a local memory comprising one or more memory devices (“memory”) 2404. In at least one embodiment, memory devices 2404 include, without limitation, 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 high-bandwidth memory (“HBM”) subsystems with multiple DRAM dies stacked

[0276] In at least one embodiment, high-speed GPU interconnect 2408 can refer to a link-based parallel computer bus that systems use to scale, and includes one or more PPUs 2400 in conjunction with one or more CPUs (“CPU(s)”), supporting cache coherency between PPUs 2400 and CPUs, and CPU mastering. In at least one embodiment, high-speed GPU interconnect 2408 transports data and / or commands through hub 2416 to other units of PPU(s) 2400, such as one or more copy engines, video encoders, video decoders, power management units, and / or other components not explicitly shown in FIG. 2A. FIG. 24

[0277] In at least one embodiment, I / O unit 2406 is configured to facilitate communication between PPU(s) 2400 and a host processor (not shown), other PPU(s) 2400, and / or one or more FIG. 24 ​The I / O units 2406 send and receive communications (e.g., commands, data) to and from the system bus 2402. In at least one embodiment, the I / O units 2406 communicate directly with the host processor(s) via the system bus 2402 or through one or more intermediate devices such as a memory hub. In at least one embodiment, the I / O units 2406 can communicate with one or more other processors, such as one or more PPUs 2400, via the system bus 2402. In at least one embodiment, the I / O units 2406 implement a PCIe interface for communications over a PCIe bus. In at least one embodiment, the I / O units 2406 implement interfaces for communicating with external devices.

[0278] In at least one embodiment, the I / O units 2406 decode packets received via the system bus 2402. In at least one embodiment, at least some packets represent commands configured to cause the PPU 2400 to perform various operations. In at least one embodiment, the I / O units 2406 send decoded commands to various other units of the PPU 2400 as designated by the commands. In at least one embodiment, commands are sent to the front-end unit 2410 and / or to the hub 2416 or other units of the PPU 2400 such as one or more copy engines, a video encoder, a video decoder, a power management unit, etc. FIG. 24 In at least one embodiment, the I / O units 2406 are configured to route communications between various logical units of the PPU 2400 (not explicitly shown). In at least one embodiment, the I / O units 2406 are configured to provide a faster, more efficient way to route data through the PPU 2400 than data could be routed through the system bus 2402.

[0279] In at least one embodiment, a program executed by the host processor encodes a command stream in a buffer that provides a workload to the PPU 2400 for processing. In at least one embodiment, the workload includes instructions and data to be processed by those instructions. In at least one embodiment, the buffer is a region in memory that is accessible (e.g., read / write) by both the host processor and the PPU 2400 - the host interface unit can be configured to access memory requests that are transmitted via the I / O units 2406 over the system bus 2402 to the buffer in system memory connected to the system bus 2402. In at least one embodiment, the host processor writes the command stream to the buffer and then sends a pointer to the start of the command stream to the PPU 2400 so that the front-end unit 2410 receives the pointer to the one or more command streams and manages the one or more command streams, reading commands from the command stream and forwarding the commands to various units of the PPU 2400.

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

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

[0282] In at least one embodiment, work distribution unit 2414 communicates with one or more GPCs 2418 via XBar 2420. In at least one embodiment, XBar 2420 is an interconnect network that couples many units of PPU 2400 to other units of PPU 2400 and can be configured to couple work distribution unit 2414 to a particular GPC 2418. In at least one embodiment, other units of one or more PPU(s) 2400 can also be connected to XBar 2420 via hub 2416.

[0283] In at least one embodiment, tasks are managed by a scheduler unit 2412 and dispatched to one of GPCs 2418 by a work distribution unit 2414. GPCs 2418 are configured to process tasks and generate results. In at least one embodiment, results can be consumed by other tasks within GPC 2418, routed to different GPCs 2418 via XBar 2420 or stored in memory 2404. In at least one embodiment, results can be written to memory 2404 via a partition unit 2422, which implements a memory interface for reading and writing data to memory 2404. In at least one embodiment, results can be transmitted to another PPU 2400 or CPU via high-speed GPU interconnect 2408. In at least one embodiment, PPU 2400 includes, without limitation, U partition units 2422 equal to a number of separate and distinct memory devices 2404 coupled to PPU 2400.

[0284] In at least one embodiment, a host processor executes a driver core that implements an application programming interface (API) that enables one or more applications executing on a host processor to schedule operations to be performed on PPU 2400. In one embodiment, multiple compute applications are executed simultaneously by PPU 2400 and PPU 2400 provides isolation, quality of service (“QoS”), and independent address spaces for multiple compute applications. In at least one embodiment, an application generates instructions (e.g., in the form of API calls) that cause a driver core to generate one or more tasks for execution by PPU 2400 and driver core outputs tasks to one or more streams processed by PPU 2400. In at least one embodiment, each task includes one or more related thread groups, which can be referred to as warps. In at least one embodiment, a warp includes multiple related threads (e.g., 32 threads) that can be executed in parallel. In at least one embodiment, cooperating threads can refer to multiple threads including instructions for performing a task and exchanging data via shared memory. In at least one embodiment, FIG. 24 PPU 2400 can perform FIG. 2 an object intersection API 228 of FIG. 3 an object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name) of FIG. 4 techniques 400 of FIG. 5 techniques 500 of. In at least one embodiment, FIG. 24 PPU 2400 can perform FIG. 2of the tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API names). In at least one embodiment, FIG. 24 PPU 2400 can perform FIG. 2 of the generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names).

[0285] FIG. 25 A GPC 2500, according to at least one embodiment, is shown. In at least one embodiment, GPC 2500 is a FIG. 24 GPC 2418. In at least one embodiment, each GPC 2500 includes, without limitation, a plurality of hardware units for processing tasks and each GPC 2500 includes, without limitation, a pipeline manager 2502, a pre-raster operations unit (“PROP”) 2504, a raster engine 2508, a work distribution crossbar (“WDX”) 2516, a memory management unit (“MMU”) 2518, one or more Data Processing Clusters (“DPCs”) 2506, and any suitable combination of the components.

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

[0287] In at least one embodiment, PROP unit 2504 is configured to route data generated by raster engine 2508 and DPC 2506 to a raster operations (“ROP”) unit in a partition unit, such as described above in connection with FIG. 23. FIG. 24 In at least one embodiment, PROP unit 2504 is configured to perform optimizations for color blending, organize pixel data, perform address translations, and / or the like. In at least one embodiment, raster engine 2508 includes, without limitation, a number of fixed function hardware units configured to perform various raster operations, and in at least one embodiment, raster engine 2508 includes, without limitation, a setup engine, a coarse raster engine, a cull engine, a clip engine, a fine raster engine, a tile aggregate engine, and any suitable combination thereof. In at least one embodiment, setup engine receives transformed vertices and generates plane equations associated with geometric primitives defined by the vertices; plane equations are passed to coarse raster engine to generate coverage information (e.g., x, y coverage masks for tiles) for the primitive; output from coarse raster engine is passed to cull engine where fragments associated with primitives that fail a z-test are culled, and to clip engine where fragments that are outside of a viewing frustum are clipped. In at least one embodiment, clipped and culled fragments are passed to fine raster engine to generate attributes of pixel fragments based on plane equations generated by setup engine. In at least one embodiment, output from raster engine 2508 includes fragments to be processed by any suitable entity, such as by a fragment shader implemented within DPC 2506.

[0288] In at least one embodiment, each DPC 2506 included in GPC 2500 includes, without limitation, a M-Pipe Controller (“MPC”) 2510; a primitive engine 2512; one or more SMs 2514; and any suitable combination thereof. In at least one embodiment, MPC 2510 controls operation of DPC 2506, routing received packets from pipeline manager 2502 to appropriate units in DPC 2506. In at least one embodiment, packets associated with vertices are routed to primitive engine 2512, which is configured to fetch vertex attributes associated with the vertices from memory; in contrast, packets associated with a shader program can be transmitted to SM 2514.

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

[0290] In at least one embodiment, the MMU 2518 is integrated with the GPC 2500 and memory partitioning unit (e.g., FIG. 24 The MMU 2518 provides an interface between partition units 2422 and provides virtual address to physical address translation, memory protection, and memory request arbitration. In at least one embodiment, the MMU 2518 provides one or more Translation Backing Buffers (“TLBs”) for performing virtual address to physical address translation in memory. In at least one embodiment, FIG. 25 GPC 2500 in the middle can execute FIG. 2 Object Intersection API 228 FIG. 3 Object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names), FIG. 4 Technology 400 and / or FIG. 5 The technology 500. In at least one embodiment, FIG. 25 GPC 2500 in the middle can execute FIG. 2generateTree(), buildTree(), and / or other suitable API names) in the graphics driver 220. In at least one embodiment, FIG. 25 The GPC 2500 in the graphics driver 220 can perform FIG. 2 generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names) in the graphics driver 220.

[0291] FIG. 26 A streaming multi-processor (“SM”) 2600, according to at least one embodiment, is shown. In at least one embodiment, the SM 2600 is a FIG. 25 SM 2514 of FIG. 25A. In at least one embodiment, the SM 2600 includes, without limitation, an instruction cache 2602; one or more scheduler units 2604; a register file 2608; one or more processing cores (“cores”) 2610; one or more special-functionality units (“SFUs”) 2612; one or more load / store units (“LSUs”) 2614; an interconnect network 2616; shared memory / level one (“Ll”) cache 2618; and any suitable combination thereof. In at least one embodiment, a work distribution unit dispatches tasks to the general processing clusters (“GPCs”) of parallel processing units (“PPUs”), and each task is assigned a particular data processing cluster (“DPC”) within a GPC, and if the task is associated with a shader program, then the task is assigned to one of the SMs 2600. In at least one embodiment, the scheduler unit 2604 receives tasks from the work distribution unit and manages instructions scheduler for one or more thread blocks assigned to the SM 2600. In at least one embodiment, the scheduler unit 2604 schedules thread blocks for execution as thread warps of parallel threads, with each thread block being assigned at least one thread warp. In at least one embodiment, each thread warp executes a thread. In at least one embodiment, the scheduler unit 2604 manages a plurality of different thread blocks, assigning thread warps to different thread blocks, and then dispatching instructions from different ones of the plurality of different cooperative groups to various functional units (e.g., processing cores 2610, SFUs 2612, and LSUs 2614) during each clock cycle. In at least one embodiment, the SM 2600 contains one or more clusters of thread blocks, where a cluster of thread blocks can enable programming control over locality at a greater granularity than a single thread block of a single streaming multi-processor (SM). In at least one embodiment, a cluster of thread blocks (also referred to as a “cluster”) allows multiple thread blocks running concurrently across a streaming multi-processor to synchronize and cooperatively fetch, exchange, or otherwise use data.

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

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

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

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

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

[0297] In at least one embodiment, each SM 2600 includes, without limitation, M SFUs 2612 to perform special functions (e.g., certain mathematical functions). In at least one embodiment, SFU 2612 includes, without limitation, a tree traversal unit configured to traverse a hierarchical tree data structure. In at least one embodiment, SFU 2612 includes, without limitation, a texture unit configured to perform texture mapping filtering operations. In at least one embodiment, texture unit is configured to load a texture map (e.g., a 2D array of texture pixels) from memory and sample the texture map to produce sampled texture values for use by a shader program executed by SM 2600. In at least one embodiment, texture maps are stored in shared memory / L1 cache 2618. In at least one embodiment, texture unit uses mip-maps (e.g., different levels of detail for a texture map) to perform texture operations such as filtering operations. In at least one embodiment, each SM 2600 includes, without limitation, two texture units.

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

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

[0300] In at least one embodiment, combining data cache and shared memory functionality into a single memory block provides improved performance for both types of memory accesses. In at least one embodiment, capacity is used by programs that do not use shared memory or use it as a cache, e.g., if shared memory is configured to use half the capacity, then textures and load / store operations can use remaining capacity. According to at least one embodiment, integration within shared memory / L1 cache 2618 enables shared memory / L1 cache 2618 to be used as a high-throughput pipe for streaming data while also providing high bandwidth and low latency access to frequently reused data. In at least one embodiment, when configured for general purpose parallel computation, a simpler configuration can be used compared to graphics processing. In at least one embodiment, fixed function GPU is bypassed, creating a more straightforward programming model. In at least one embodiment, in a general purpose parallel computation configuration, work assignment unit assigns and distributes blocks of threads directly to DPCs. In at least one embodiment, threads in a block execute the same program, use unique thread IDs in a computation to ensure that each thread generates unique results, use SM 2600 to execute the program and perform the computation, use shared memory / L1 cache 2618 to communicate between threads, and use LSUs 2614 to read and write to global memory through shared memory / L1 cache 2618 and memory partition unit. In at least one embodiment, when configured for general purpose parallel computation, SM 2600 writes to scheduler unit 2604 commands that can be used to launch new work on DPCs. In at least one embodiment, SM 2600 includes one or more distributed shared memories (or distributed shared memories) that support direct SM to SM operations, such as cross-SM shared memory block loads, stores, and atomics.

[0301] In at least one embodiment, SM 2600 includes one or more asynchronous execution functions including a tensor memory accelerator (TMA) unit that can transfer data blocks between global memory and shared memory. In at least one embodiment, one or more processor cores use or access one or more TMAs to perform bi-directional copy operations, e.g., from global memory to shared memory, and vice versa. In at least one embodiment, SM 2600 includes one or more TMAs for asynchronous copying between thread blocks in a cluster. In at least one embodiment, SM 2600 includes one or more asynchronous transactional barriers for performing atomic data movement and synchronization. In at least one embodiment, SM 2600 includes a tensor core transformer engine containing software and one or more cores for accelerating training and inference of transformer models. In at least one embodiment, one or more processor cores executing one or more tensor core transformer engines manage and dynamically select between FP8 and 16-bit computing by re-converting and scaling between FP8 and 16-bit at each layer of one or more neural networks. In at least one embodiment, FIG. 26 SM 2600 of can perform FIG. 2 object intersection API 228 of FIG. 3 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name) of FIG. 4 techniques 400 of and / or FIG. 5 techniques 500 of In at least one embodiment, FIG. 26 SM 2600 of can perform FIG. 2 generation tree API 230 (e.g., generateTree(), buildTree(), and / or other suitable API name) of In at least one embodiment, FIG. 26 SM 2600 of can perform FIG. 2 generation lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) of

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

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

[0304] Software constructs for general purpose computing

[0305] The following figures illustrate, without limitation, example software constructs for implementing at least one embodiment.

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

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

[0308] ​In at least one embodiment, application 2701 and software stack 2700 run on hardware 2707. In at least one embodiment, hardware 2707 can include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices that support a programming platform. In at least one embodiment, software stack 2700 can be vendor specific, and only compatible with devices from a particular vendor, e.g., with CUDA. In at least one embodiment, software stack 2700 can be used with devices from different vendors, e.g., with OpenCL. In at least one embodiment, hardware 2707 includes a host connected to one or more devices that can be accessed via an application programming interface (API) call to perform a computational task. In at least one embodiment, a device within hardware 2707 can include, without limitation, a GPU, FPGA, AI engine, or other computing device (but also can include a CPU) and its memory, as opposed to a host within hardware 2707 which can include, without limitation, a CPU (but also can include a computing device) and its memory.

[0309] In at least one embodiment, software stack 2700 of a programming platform includes, without limitation, a plurality of libraries 2703, a runtime 2705, and a device kernel driver 2706. In at least one embodiment, each of libraries 2703 can include data and programming code that can be used by a computer program and utilized during software development. In at least one embodiment, libraries 2703 can include, without limitation, pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. In at least one embodiment, libraries 2703 include functions optimized for execution on one or more types of devices. In at least one embodiment, libraries 2703 can include, without limitation, functions for performing mathematical, deep learning, and / or other types of operations on a device. In at least one embodiment, libraries 2703 are associated with a corresponding API 2702, which can include one or more APIs that expose functions implemented in libraries 2703. In at least one embodiment, a processor (e.g., CPU, GPU) executes, calls, or otherwise uses one or more APIs to prioritize kernels. For example, a first kernel (e.g., a parent kernel) can launch a second kernel (e.g., a child kernel), and a processor can use the second kernel to launch an additional kernel (e.g., a grandchild kernel) independent of the first kernel. In at least one embodiment, a processor executes or calls an API from memory to perform to support dynamic stream prioritization (e.g., updating a priority when performing operations using a stream). For example, when a processor executes the API, it allows a programmer to copy a stream priority from one stream to one or more other streams.

[0310] In at least one embodiment, software stack 2700 includes an API to support dynamic stream priority (e.g., updating priority while using a stream to perform an operation), which allows a programmer to set a priority of a stream at any time after a stream is created. In at least one embodiment, software stack 2700 includes an API to support dynamic stream priority (e.g., updating priority while using a stream to perform an operation), which allows a programmer to get a current priority of a stream, where priority is one of a plurality of attributes of a stream. In at least one embodiment, software stack 2700 includes an API to support dynamic stream priority (e.g., updating priority while using a stream to perform an operation), which allows a programmer to get a current priority of a stream as a single attribute. In at least one embodiment, software stack 2700 includes an API to support dynamic stream priority (e.g., updating priority while using a stream to perform an operation), which allows a programmer to launch a kernel to perform an operation on a stream with a set priority that can be different from a stream priority. In at least one embodiment, software stack 2700 includes an API to indicate whether an object (e.g., a thread synchronization object such as a barrier) tracks whether all data movement operations of a set of threads running on a GPU have a specified status after a specified period of time, where a specified status can be a status indicating that data has been moved and is available for use, with an expected parity value as an input to the API to specify.

[0311] In at least one embodiment, software stack 2700 includes one or more APIs to update a kernel. In at least one embodiment, a processor executes an API or calls an API to be executed from memory to update an existing API to support context-independent kernels, which allows a programmer to add a kernel node to a graph without a graphics context so that a graphics context can be dynamically associated with the kernel at runtime. In at least one embodiment, software stack 2700 includes one or more APIs to allow a programmer to get a kernel identifier and a graphics context from a kernel node as separate parameters so that parameters are obtained from a kernel and a context-independent kernel. In at least one embodiment, software stack 2700 includes one or more APIs to launch a task graph (e.g., a task graph) and execute one or more task graphs (e.g., containing one or more programs) using a parallel processor (e.g., one or more graphics processing units).

[0312] In at least one embodiment, software stack 2700 includes one or more APIs for associating one or more instructions with one or more memory ordering operations (e.g., fence or membar operations). In at least one embodiment, instructions are associated with one or more domains such that memory ordering operations are performed in association with one or more particular domains without interfering with instructions of other domains. In addition, an API is included for indicating that a thread has reached (e.g., reached a thread synchronization barrier) or completed a certain phase of work related to an asynchronous data movement operation on a GPU. In at least one embodiment, software stack 2700 includes one or more APIs for allowing a programmer to manually indicate an expected transaction count when a thread completes a certain phase of work, which is used to update an object that tracks whether all data movement operations for a group of threads have completed.

[0313] In at least one embodiment, application 2701 is written as source code that is compiled into executable code, as discussed in more detail below in conjunction with FIG. 27B. FIG. 32-34 In at least one embodiment, executable code of application 2701 can run at least partially on an execution environment provided by software stack 2700. In at least one embodiment, during execution of application 2701, code can be derived that needs to run on a device (as opposed to a host). In such cases, in at least one embodiment, runtime 2705 can be invoked to load and launch the necessary code on a device. In at least one embodiment, runtime 2705 can include any technically feasible runtime system capable of supporting execution of application 2701.

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

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

[0316] In at least one embodiment, one or more processors disclosed in “Processing Systems” can execute, access, or otherwise use software stack 2700. For example, APU 1400, CPU 1500, FIG. 17A-17B In at least one embodiment, exemplary graphics processor, general purpose graphics processing unit (“GPGPU”) 1830, parallel processor 1900, processing cluster 1994, graphics multiprocessor 1934, graphics multiprocessor 1996, graphics processor 2000, processor 2100, processor 2200, parallel processing unit (“PPU”) 2400, GPC 2500, and / or streaming multiprocessor (“SM”) 2600 can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 27 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 27 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 2 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 3 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 4 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 5 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 27 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 2 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 27 In at least one embodiment, hardware 2707 of “Graphics Processors” can execute, use, call, or otherwise implement one or more APIs included in software stack 2700, for example by accessing memory. In at least one embodiment, FIG. 2generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names).

[0317] In at least one embodiment, device kernel driver 2706 is configured to facilitate communication with underlying devices. In at least one embodiment, device kernel driver 2706 can provide low-level functions relied upon by APIs such as API 2704 and / or other software. In at least one embodiment, device kernel driver 2706 can be configured to compile intermediate representation (“IR”) code into binary code at runtime. In at least one embodiment, for CUDA, device kernel driver 2706 can compile non-hardware-specific parallel thread execution (“PTX”) IR code into binary code for a particular target device at runtime (cache compiled binary code), which is sometimes also referred to as “final” code. In at least one embodiment, doing so can allow final code to run on a target device that can not have existed when source code was originally compiled into PTX code. Alternatively, in at least one embodiment, device source code can be compiled into binary code offline without requiring device kernel driver 2706 to compile IR code at runtime.

[0318] FIG. 28 FIG. 27 illustrates a CUDA implementation of software stack 2700, in accordance with at least one embodiment FIG. 27 In at least one embodiment, CUDA software stack 2800 on which application 2801 can be launched includes CUDA libraries 2803, CUDA runtime 2805, CUDA driver 2807, and device kernel driver 2808. In at least one embodiment, CUDA software stack 2800 executes on hardware 2809, which can include a CUDA-enabled GPU developed by NVIDIA Corporation of Santa Clara, CA.

[0319] In at least one embodiment, application 2801, CUDA runtime 2805, and device kernel driver 2808 can perform similar functions to application 2701, runtime 2705, and device kernel driver 2706, respectively, described above in conjunction with FIG. 27CUDA Driver 2807, which in at least one embodiment, includes a library (libcuda.so) that implements a CUDA Driver API 2806. In at least one embodiment, similar to the CUDA Runtime API 2804 implemented by the CUDA Runtime Library (cudart), the CUDA Driver API 2806 can expose, without limitation, functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability, etc. In at least one embodiment, the CUDA Driver API 2806 differs from the CUDA Runtime API 2804 in that the CUDA Runtime API 2804 simplifies device code management by providing implicit initialization, context (similar to a process) management, and module (similar to a dynamically loaded library) management. In contrast to the high-level CUDA Runtime API 2804, in at least one embodiment, the CUDA Driver API 2806 is a low-level API that provides more granular control of the device, particularly with respect to contexts and module loading. In at least one embodiment, the CUDA Driver API 2806 can expose functions for context management that are not exposed by the CUDA Runtime API 2804. In at least one embodiment, the CUDA Driver API 2806 is also language agnostic and supports, for example, OpenCL in addition to the CUDA Runtime API 2804. Further, in at least one embodiment, the development libraries, including the CUDA Runtime 2805, can be considered separate from the driver components, including the user-mode CUDA Driver 2807 and the kernel-mode device driver 2808 (sometimes also referred to as a “display” driver). In at least one embodiment, FIG. 28 Hardware 2809 of NVIDIA® Nsight™ 3D Renderer® can perform FIG. 2 Object Intersection API 228 of NVIDIA® Nsight™ 3D Renderer®, FIG. 3 Object Intersection API 300 of NVIDIA® Nsight™ 3D Renderer® (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names), FIG. 4 Techniques 400 of NVIDIA® Nsight™ 3D Renderer® and / or FIG. 5 Techniques 500 of NVIDIA® Nsight™ 3D Renderer®. In at least one embodiment, FIG. 28 Hardware 2809 of NVIDIA® Nsight™ 3D Renderer® can perform FIG. 2 Generation Tree API 230 of NVIDIA® Nsight™ 3D Renderer® (e.g., generateTree(), buildTree(), and / or other suitable API names). In at least one embodiment, FIG. 28 Hardware 2809 of NVIDIA® Nsight™ 3D Renderer® can perform FIG. 2generate lookup parameters API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names).

[0320] In at least one embodiment, CUDA libraries 2803 can include, without limitation, math libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries that can be utilized by parallel computing applications (e.g., application 2801). In at least one embodiment, CUDA libraries 2803 can include math libraries such as a cuBLAS library, which is an implementation of Basic Linear Algebra Subprograms (“BLAS”) for performing linear algebra operations; a cuFFT library for computing Fast Fourier Transforms (“FFTs”), and a cuRAND library for generating random numbers, etc. In at least one embodiment, CUDA libraries 2803 can include deep learning libraries such as a cuDNN library for primitives of deep neural networks and a TensorRT platform for high-performance deep learning inference, etc.

[0321] FIG. 29 FIG. 29 illustrates a ROCm implementation of software stack 2700, in accordance with at least one embodiment. FIG. 27 In at least one embodiment, ROCm software stack 2900 on which application 2901 can be launched includes language runtime 2903, system runtime 2905, thunk 2907, and ROCm kernel driver 2908. In at least one embodiment, ROCm software stack 2900 executes on hardware 2909, which can include a GPU that supports ROCm, which was developed by AMD Corporation of Santa Clara, CA.

[0322] In at least one embodiment, application 2901 can perform similar functionality as application 2701 discussed above in connection with FIG. 27 In at least one embodiment, language runtime 2903 and system runtime 2905 can perform similar functionality as language runtime 2703 and system runtime 2705 discussed above in connection with FIG. 27similar functions discussed above in connection with the CUDA runtime API 2804, such as functions for memory management, execution control, device management, error handling, and synchronization, etc. FIG. 28 similar functions discussed above in connection with the CUDA runtime API 2804, such as functions for memory management, execution control, device management, error handling, and synchronization, etc.

[0323] In at least one embodiment, the thunk (ROCt) 2907 is an interface 2906 that can be used to interface with an underlying ROCm driver 2908. In at least one embodiment, the ROCm driver 2908 is a ROCk driver, which is a combination of an AMDGPU driver and an HSA kernel driver (amdkfd). In at least one embodiment, the AMDGPU driver is a device kernel driver for GPUs developed by AMD that performs similar functions discussed above in connection with the device kernel driver 2706. In at least one embodiment, the HSA kernel driver is a driver that allows different types of processors to share system resources more efficiently via hardware features. FIG. 27 similar functions discussed above in connection with the device kernel driver 2706. In at least one embodiment, the HSA kernel driver is a driver that allows different types of processors to share system resources more efficiently via hardware features.

[0324] In at least one embodiment, various libraries (not shown) can be included in the ROCm software stack 2900 above the language runtime 2903 and provide similar functions discussed above in connection with the CUDA runtime API 2804, such as functions for memory management, execution control, device management, error handling, and synchronization, etc. FIG. 28The CUDA libraries 2803 discussed above have similar functionality. In at least one embodiment, various libraries can include, but are not limited to, math, deep learning, and / or other libraries such as the hipBLAS library that implements functions similar to CUDA cuBLAS, the rocFFT library similar to CUDA cuFFT for computing FFTs, and the like. In at least one embodiment, FIG. 29 The hardware 2909 of FIG. 2 the object intersection API 228 of FIG. 3 the object intersection API 300 of (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API names), FIG. 4 the techniques 400 of and / or FIG. 5 the techniques 500 of In at least one embodiment, FIG. 29 The hardware 2909 of FIG. 2 the generation tree API 230 of (e.g., generateTree(), buildTree(), and / or other suitable API names). In at least one embodiment, FIG. 29 The hardware 2909 of FIG. 2 the generation lookup parameters API 232 of (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API names).

[0325] FIG. 30 FIG. 20 illustrates an OpenCL implementation of the software stack 2700 of FIG. 27 In at least one embodiment, the OpenCL software stack 3000 on which the application 3001 can launch includes an OpenCL framework 3010, an OpenCL runtime 3006, and a driver 3007. In at least one embodiment, the OpenCL software stack 3000 executes on hardware 2809 that is not vendor-specific. In at least one embodiment, because devices developed by different vendors support OpenCL, specific OpenCL drivers can be required to interoperate with hardware from such vendors.

[0326] In at least one embodiment, the application 3001, the OpenCL runtime 3006, the device kernel driver 3007, and the hardware 3008 can perform functions similar to those described above in connection with the FIG. 27The application 2701, runtime 2705, device kernel driver 2706, and hardware 2707 discussed above have similar functionality. In at least one embodiment, the application 3001 also includes an OpenCL kernel 3002 with code to be executed on a device. In at least one embodiment, FIG. 30 The hardware 3008 of FIG. 2 The object intersection API 228 of FIG. 3 The object intersection API 300 of (e.g., spatialQuery(), trace(), curtxTrace(), and / or other suitable API name), FIG. 4 The techniques 400 of FIG. 5 The techniques 500 of In at least one embodiment, FIG. 30 The hardware 3008 of FIG. 2 The tree generation API 230 of (e.g., generateTree(), buildTree(), and / or other suitable API name). In at least one embodiment, FIG. 30 The hardware 3008 of FIG. 2 The lookup parameter generation API 232 of (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or other suitable API name).

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

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

[0329] FIG. 31 Software supported by a programming platform according to at least one embodiment is illustrated. In at least one embodiment, the programming platform 3104 is configured to support various programming models 3103, middleware and / or libraries 3102, and frameworks 3101 that the application 3100 may depend on. In at least one embodiment, the application 3100 may be an AI / ML application implemented using, for example, a deep learning framework (e.g., MXNet, PyTorch, or TensorFlow), which may depend on libraries such as cuDNN, the NVIDIA Collective Communications Library (“NCCL”), and / or the NVIDIA Developer Data Loading Library (“DALI”) CUDA library to provide accelerated computation on the underlying hardware.

[0330] In at least one embodiment, the programming platform 3104 can be a combination of the above-described components. FIG. 28 , FIG. 29 and FIG. 30 One of the described CUDA, ROCm, or OpenCL platforms. In at least one embodiment, the programming platform 3104 supports multiple programming models 3103, which are abstractions of the underlying computing system that allow for the expression of algorithms and data structures. In at least one embodiment, the programming model 3103 may expose features of the underlying hardware to improve performance. In at least one embodiment, the programming model 3103 may include, but is not limited to, CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++AMP”), Open Multiprocessing (“OpenMP”), Open Accelerator (“OpenACC”), and / or Vulcan Compute.

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

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

[0333] FIG. 32 FIG. 23 illustrates compiled code to execute on one of FIG. 27-30 In at least one embodiment, compiler 3201 receives source code 3200, which includes both host code as well as device code. In at least one embodiment, compiler 3201 is configured to convert source code 3200 into host executable code 3202 for execution on a host and device executable code 3203 for execution on a device. In at least one embodiment, source code 3200 can be compiled offline prior to execution of an application, or online during execution of an application. In at least one embodiment, compiler 3201 includes or has access to one or more libraries to identify sequences of API calls to perform a single fused API, where a single fused API is a combined API of two or more APIs.

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

[0335] In at least one embodiment, compiler 3201 is configured to compile source code 3200 into host executable code 3202 for execution on a host and device executable code 3203 for execution on a device. In at least one embodiment, compiler 3201 performs operations including parsing source code 3200 into an abstract syntax tree (AST), performing optimizations, and generating executable code. In at least one embodiment where source code 3200 includes a single-source file, compiler 3201 can separate device code from host code in such single-source file, compile device code and host code separately into device executable code 3203 and host executable code 3202, and link device executable code 3203 and host executable code 3202 together in a single file, as discussed below with respect to FIG. 4. FIG. 33 More detailed discussion.

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

[0337] FIG. 33 is compiled code to execute on a host in accordance with at least one embodiment FIG. 27-30more detail. In at least one embodiment, compiler 3301 is configured to receive source code 3300, compile source code 3300, and output executable 3310. In at least one embodiment, source code 3300 is a single source file, such as a.cu file, a.hip.cpp file, or other format of file, that includes both host code and device code. In at least one embodiment, compiler 3301 can be, without limitation, an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code in a.cu file, or an HCC compiler for compiling HIP code in a.hip.cpp file.

[0338] In at least one embodiment, compiler 3301 includes a compiler front end 3302, a host compiler 3305, a device compiler 3306, and a linker 3309. In at least one embodiment, compiler front end 3302 is configured to separate device code 3304 from host code 3303 in source code 3300. In at least one embodiment, device code 3304 is compiled by device compiler 3306 into device executable code 3308, which can include binary code or IR code, as described. In at least one embodiment, host code 3303 is separately compiled by host compiler 3305 into host executable code 3307. In at least one embodiment, for NVCC, host compiler 3305 can be, without limitation, a general C / C++ compiler that outputs native object code, while device compiler 3306 can be, without limitation, a low-level virtual machine (“LLVM”) based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code. In at least one embodiment, for HCC, both host compiler 3305 and device compiler 3306 can be, without limitation, LLVM based compilers that output target binary code.

[0339] In at least one embodiment, after source code 3300 is compiled into host executable code 3307 and device executable code 3308, linker 3309 links host and device executable code 3307 and 3308 together in executable 3310. In at least one embodiment, native object code for host and PTX or binary code for device can be linked together in an Executable and Linkable Format (“ELF”) file, which is a container format for storing object code.

[0340] FIG. 34Conversion of source code prior to compilation is shown in accordance with at least one embodiment. In at least one embodiment, source code 3400 is passed through a conversion tool 3401 that converts source code 3400 to converted source code 3402. In at least one embodiment, a compiler 3403 is used to compile converted source code 3402 to host executable 3404 and device executable 3305, a process similar to that of compiler 3201 compiling source code 3200 to host executable 3202 and device executable 3203, as discussed above in conjunction with FIG. 3. In at least one embodiment, compiler 3403 is a different compiler than compiler 3201, and / or is configured differently than compiler 3201. In at least one embodiment, compiler 3403 is a CUDA compiler, and source code 3400 is CUDA source code. In at least one embodiment, compiler 3403 is a HIP compiler, and source code 3400 is HIP source code. FIG. 32

[0341] In at least one embodiment, the conversion performed by conversion tool 3401 is used to port source code 3400 to execute in a different environment than originally intended. In at least one embodiment, conversion tool 3401 can include, without limitation, a HIP Transpiler, which is used to “hipify” CUDA code intended for a CUDA platform to HIP code that can be compiled and executed on a ROCm platform. In at least one embodiment, conversion of source code 3400 can include parsing source code 3400 and converting calls to APIs provided by one programming model (e.g., CUDA) to corresponding calls to APIs provided by another programming model (e.g., HIP), as discussed in more detail below in conjunction with FIG. 34. Returning to the example of porting CUDA code, in at least one embodiment, calls to CUDA runtime APIs, CUDA driver APIs, and / or CUDA libraries can be converted to corresponding HIP API calls. In at least one embodiment, automatic conversion performed by conversion tool 3401 can sometimes be incomplete, requiring additional human effort to fully port source code 3400. FIG. 35A FIG. 36

[0342] Configuring a GPU for general-purpose computing

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

[0344] FIG. 35A ​​​A system 3500 configured to compile and execute CUDA source code 3510 using different types of processing units is shown, in accordance with at least one embodiment. In at least one embodiment, system 3500 includes, without limitation, CUDA source code 3510, a CUDA compiler 3550, host executable code 3570(1), host executable code 3570(2), CUDA device executable code 3584, a CPU 3590, a CUDA-enabled GPU 3594, a GPU 3592, a CUDA to HIP translation tool 3520, a HIP source code 3530, a HIP compiler driver 3540, an HCC 3560, and HCC device executable code 3582.

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

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

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

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

[0349] In at least one embodiment, CUDA compiler 3550 compiles input CUDA code (e.g., CUDA source code 3510) to generate host executable code 3570(1) and CUDA device executable code 3584. In at least one embodiment, CUDA compiler 3550 is NVCC. In at least one embodiment, host executable code 3570(1) is a compiled version of host code included in input source code that is executable on CPU 3590. In at least one embodiment, CPU 3590 can be any processor optimized for sequential instruction processing.

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

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

[0352] In at least one embodiment, HIP source code 3530 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, HIP source code 3530 can also include any number of calls to any number of functions specified in a HIP runtime API 3532. In one embodiment, HIP runtime API 3532 includes, without limitation, functionally similar versions of a subset of functions included in CUDA runtime API 3502. In at least one embodiment, HIP source code 3530 can also include any number of calls to any number of functions specified in any number of other HIP APIs. In at least one embodiment, a HIP API can be any API designed for use with HIP code and / or ROCm. In at least one embodiment, a HIP API includes, without limitation, HIP runtime API 3532, a HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, and the like.

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

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

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

[0356] In at least one embodiment, if target device 3546 is compatible with CUDA (e.g., a CUDA-enabled GPU 3594), then HIP compiler driver 3540 generates HIP / NVCC compilation commands 3542. In at least one embodiment and in conjunction with FIG. 35B In more detail, HIP / NVCC compilation commands 3542 configure CUDA compiler 3550 to compile HIP source code 3530 using, without limitation, a HIP-to-CUDA translation header and a CUDA runtime library. In at least one embodiment and in response to HIP / NVCC compilation commands 3542, CUDA compiler 3550 generates host executable code 3570(1) and CUDA device executable code 3584.

[0357] In at least one embodiment, if target device 3546 is not compatible with CUDA, then HIP compiler driver 3540 generates HIP / HCC compilation commands 3544. In at least one embodiment and in conjunction with FIG. 35C In more detail, HIP / HCC compilation commands 3544 configure HCC 3560 to compile HIP source code 3530 using a HCC header and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3544, HCC 3560 generates host executable code 3570(2) and HCC device executable code 3582. In at least one embodiment, HCC device executable code 3582 is a compiled version of device code contained in HIP source code 3530 that is executable on GPU 3592. In at least one embodiment, GPU 3592 can be any processor optimized for parallel instruction processing that is not compatible with CUDA and is compatible with HCC. In at least one embodiment, GPU 3592 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, GPU 3592 is a GPU 3592 that is not CUDA-enabled.

[0358] For illustrative purposes only, in FIG. 35AFIG. 35 illustrates three different flows that can be implemented to compile CUDA source code 3510 for execution on CPUs 3590 and different devices, in at least one embodiment. In the direct CUDA flow, CUDA source code 3510 is compiled for execution on CPUs 3590 and CUDA-enabled GPUs 3594 without converting CUDA source code 3510 to HIP source code 3530, in at least one embodiment. In the indirect CUDA flow, CUDA source code 3510 is converted to HIP source code 3530, which is then compiled for execution on CPUs 3590 and CUDA-enabled GPUs 3594, in at least one embodiment. In the CUDA / HCC flow, CUDA source code 3510 is converted to HIP source code 3530, which is then compiled for execution on CPUs 3590 and GPUs 3592, in at least one embodiment. FIG. 35A-35C CPU 3590, GPU 3594, and / or GPU 3592 in FIG. 35 can execute FIG. 2 object intersection API 228 in FIG. 30, FIG. 3 object intersection API 300 (e.g., spatialQuery(), trace(), curtxTrace(), and / or some other suitable API name) in FIG. 31, FIG. 4 techniques 400 in FIG. 32, and / or FIG. 5 techniques 500 in FIG. 33. In at least one embodiment, FIG. 35A to FIG. 35C CPU 3590, GPU 3594, and / or GPU 3592 in FIG. 35 can execute FIG. 2 tree generation API 230 (e.g., generateTree(), buildTree(), and / or some other suitable API name) in FIG. 34. In at least one embodiment, FIG. 35A to FIG. 35C CPU 3590, GPU 3594, and / or GPU 3592 in FIG. 35 can execute FIG. 2 ray generation lookup parameter API 232 (e.g., generateRay(), generateLookupParams(), generateQueryParams(), and / or some other suitable API name) in FIG. 35.

[0359] The direct CUDA flow that can be implemented in at least one embodiment can be depicted by the dashed line and series of bubble annotations Al-A3. In at least one embodiment, and as shown by bubble annotation Al, CUDA compiler 3550 receives CUDA source code 3510 and a CUDA compile command 3548 that configures CUDA compiler 3550 to compile CUDA source code 3510. In at least one embodiment, the CUDA source code 3510 used in the direct CUDA flow is written in the CUDA programming language, which is based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In at least one embodiment, and in response to CUDA compile command 3548, CUDA compiler 3550 generates host executable code 3570(1) and CUDA device executable code 3584 (represented by bubble annotation A2). In at least one embodiment and as shown by bubble annotation A3, host executable code 3570(1) and CUDA device executable code 3584 can be executed on CPU 3590 and CUDA-enabled GPU 3594, respectively. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0360] The indirect CUDA flow that can be implemented in at least one embodiment can be described by the dashed line and series of bubble annotations Bl-B6. In at least one embodiment and as shown by bubble annotation Bl, CUDA-to-HIP translation tool 3520 receives CUDA source code 3510. In at least one embodiment and as shown by bubble annotation B2, CUDA-to-HIP translation tool 3520 translates CUDA source code 3510 into HIP source code 3530. In at least one embodiment and as shown by bubble annotation B3, HIP compiler driver 3540 receives HIP source code 3530 and determines whether a target device 3546 is CUDA-enabled.

[0361] In at least one embodiment and as shown by bubble annotation B4, HIP compiler driver 3540 generates a HIP / NVCC compile command 3542 and sends both HIP / NVCC compile command 3542 and HIP source code 3530 to CUDA compiler 3550. In at least one embodiment and as shown by bubble annotation B5, CUDA compiler 3550 receives HIP / NVCC compile command 3542 and configures CUDA compiler 3550 to compile HIP source code 3530. In at least one embodiment and as shown by bubble annotation B6, CUDA compiler 3550 generates host executable code 3570(2) and CUDA device executable code 3584. FIG. 35BIn more detail, the HIP / NVCC compilation command 3542 configures the CUDA compiler 3550 to compile the HIP source code 3530 using, without limitation, a HIP-to-CUDA translation header and a CUDA runtime library. In at least one embodiment and in response to the HIP / NVCC compilation command 3542, the CUDA compiler 3550 generates host executable code 3570(1) and CUDA device executable code 3584 (denoted with bubble annotation B5). In at least one embodiment and as shown with bubble annotation B6, the host executable code 3570(1) and the CUDA device executable code 3584 can execute on a CPU 3590 and a CUDA-enabled GPU 3594, respectively. In at least one embodiment, the CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, the CUDA device executable code 3584 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0362] A CUDA / HCC flow that can be implemented in at least one embodiment can be described by the solid line and series of bubble annotations CI - C6. In at least one embodiment and as shown with bubble annotation CI, the CUDA-to-HIP translation tool 3520 receives the CUDA source code 3510. In at least one embodiment and as shown with bubble annotation C2, the CUDA-to-HIP translation tool 3520 translates the CUDA source code 3510 into the HIP source code 3530. In at least one embodiment and as shown with bubble annotation C3, the HIP compiler driver 3540 receives the HIP source code 3530 and determines that the target device 3546 is not CUDA-enabled.

[0363] In at least one embodiment, the HIP compiler driver 3540 generates a HIP / HCC compilation command 3544 and sends both the HIP / HCC compilation command 3564 and the HIP source code 3530 to the HCC 3560 (denoted with bubble annotation C4). In at least one embodiment and as described in connection with bubble annotation C5, the HCC 3560 receives the HIP / HCC compilation command 3564 and the HIP source code 3530. FIG. 35C In more detail, the HIP / HCC compilation command 3564 configures the HCC 3560 to compile the HIP source code 3530 using, without limitation, an HCC header and a HIP / HCC runtime library. In at least one embodiment and in response to the HIP / HCC compilation command 3544, the HCC 3560 generates host executable code 3570(2) and HCC device executable code 3582 (denoted with bubble annotation C5). In at least one embodiment and as shown with bubble annotation C6, the host executable code 3570(2) and the HCC device executable code 3582 can execute on a CPU 3590 and a GPU 3592, respectively.

[0364] In at least one embodiment, after CUDA source code 3510 is converted to HIP source code 3530, a HIP compiler driver 3540 can then be used to generate executable code for CUDA-enabled GPU 3594 or GPU 3592 without CUDA being re-implemented as a HIP conversion tool 3520. In at least one embodiment, CUDA to HIP conversion tool 3520 converts CUDA source code 3510 to HIP source code 3530, which is then stored in memory. In at least one embodiment, HIP compiler driver 3540 then configures HCC 3560 to generate host executable code 3570(2) and HCC device executable code 3582 based on HIP source code 3530. In at least one embodiment, HIP compiler driver 3540 then configures CUDA compiler 3550 to generate host executable code 3570(l) and CUDA device executable code 3584 based on stored HIP source code 3530.

[0365] FIG. 35B FIG. 34 illustrates a system 3504 configured to compile and execute CUDA source code 3510 using CPU 3590 and CUDA-enabled GPU 3594, in accordance with at least one embodiment. In at least one embodiment, system 3504 includes, without limitation, CUDA source code 3510, CUDA to HIP conversion tool 3520, HIP source code 3530, HIP compiler driver 3540, CUDA compiler 3550, host executable code 3570(l), CUDA device executable code 3584, CPU 3590, and CUDA-enabled GPU 3594. FIG. 35A

[0366] In at least one embodiment and as previously described herein in conjunction with FIG. 35A FIG. 33, CUDA source code 3510 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, CUDA source code 3510 also includes, without limitation, any number of calls to any number of functions specified in any number of CUDA APIs.

[0367] ​In at least one embodiment, CUDA to HIP translation tool 3520 translates CUDA source code 3510 into HIP source code 3530. In at least one embodiment, CUDA to HIP translation tool 3520 translates each kernel call in CUDA source code 3510 from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in CUDA source code 3510 to any number of other functionally similar HIP calls.

[0368] In at least one embodiment, HIP compiler driver 3540 determines that target device 3546 is CUDA-enabled, and generates a HIP / NVCC compilation command 3542. In at least one embodiment, HIP compiler driver 3540 then configures CUDA compiler 3550 via HIP / NVCC compilation command 3542 to compile HIP source code 3530. In at least one embodiment, as part of configuring CUDA compiler 3550, HIP compiler driver 3540 provides access to a HIP to CUDA translation header 3552. In at least one embodiment, HIP to CUDA translation header 3552 translates any number of mechanisms (e.g., functions) specified in any number of HIP APIs to any number of mechanisms specified in any number of CUDA APIs. In at least one embodiment, CUDA compiler 3550 uses HIP to CUDA translation header 3552 in conjunction with a CUDA runtime library 3554 corresponding to CUDA runtime API 3502 to generate host executable code 3570(1) and CUDA device executable code 3584. In at least one embodiment, host executable code 3570(1) and CUDA device executable code 3584 can then be executed on CPU 3590 and CUDA-enabled GPU 3594, respectively. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0369] FIG. 35C A system 3506 configured to compile and execute HIP source code 3530 using CPU 3590 and non-CUDA-enabled GPU 3592 is shown, in accordance with at least one embodiment. FIG. 35ACUDA source code 3510. In at least one embodiment, system 3506 includes, without limitation, CUDA source code 3510, CUDA to HIP translation tool 3520, HIP source code 3530, HIP compiler driver 3540, HCC 3560, host executable code 3570(2), HCC device executable code 3582, CPU 3590, and GPU 3592.

[0370] In at least one embodiment, and as described previously herein in conjunction with FIG. 35A CUDA source code 3510 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, CUDA source code 3510 also includes, without limitation, any number of calls to any number of functions specified in any number of CUDA APIs.

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

[0372] In at least one embodiment, HIP compiler driver 3540 then determines that target device 3546 is not CUDA-enabled, and generates HIP / HCC compilation commands 3544. In at least one embodiment, HIP compiler driver 3540 then configures HCC 3560 to execute HIP / HCC compilation commands 3544, thereby compiling HIP source code 3530. In at least one embodiment, HIP / HCC compilation commands 3544 configure HCC 3560 to use, without limitation, HIP / HCC runtime library 3558 and HCC headers 3556 to generate host executable code 3570(2) and HCC device executable code 3582. In at least one embodiment, HIP / HCC runtime library 3558 corresponds to HIP runtime API 3532. In at least one embodiment, HCC headers 3556 include, without limitation, any number and type of interoperability mechanisms for HIP and HCC. In at least one embodiment, host executable code 3570(2) and HCC device executable code 3582 can be executed on CPU 3590 and GPU 3592, respectively.

[0373] FIG. 36 The diagram illustrates a method according to at least one embodiment. FIG. 35C An exemplary kernel is converted by the CUDA to HIP conversion tool 3520. In at least one embodiment, the CUDA source code 3510 divides the overall problem that 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 includes, but is not limited to, any number of threads. In at least one embodiment, each subproblem is divided into relatively small pieces that can be solved in parallel by the threads within the thread block. In at least one embodiment, threads within a thread block can cooperate by sharing data through shared memory and by coordinating memory accesses through synchronized execution.

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

[0375] In at least one embodiment, the kernel is a function in the device code defined using the "__global__" declaration specifier. In at least one embodiment, CUDA kernel startup syntax 3610 is used to specify the size of the mesh for executing the kernel for a given kernel call and the associated flow. In at least one embodiment, CUDA kernel startup sy...

Claims

1. A processor, comprising: One or more circuits for implementing an application programming interface (API) to indicate to the user one or more objects that are intersected by one or more rays.

2. The processor of claim 1, wherein the API is configured to indicate the one or more objects at least in part based on a tree data structure.

3. The processor of claim 1, wherein the API is configured to enable the use of a ray tracing core to indicate the one or more objects.

4. The processor of claim 1, wherein the API is configured to instruct the one or more objects based at least in part on software executed by a graphics processing unit (GPU).

5. The processor of claim 1, wherein the one or more objects comprise one or more triangles.

6. The processor of claim 1, wherein the API is configured to indicate the one or more objects at least in part based on a set of input parameters including the rays in the one or more rays and a data structure including the one or more objects.

7. The processor of claim 1, wherein the API is configured to enable the use of one or more ray tracing cores of a graphics processing unit (GPU) to indicate the one or more objects at least in part based on a tree data structure comprising the one or more objects.

8. A system comprising: One or more processors are used to execute an application programming interface (API) to indicate to the user one or more objects that are intersected by one or more rays.

9. The system of claim 8, wherein the API is configured to enable the use of one or more ray tracing cores to indicate the one or more objects.

10. The system of claim 8, wherein the API is configured to indicate the one or more objects based at least in part on a set of input parameters including a tree data structure.

11. The system of claim 8, wherein the API is configured to cause the use of the one or more ray tracing cores to indicate the one or more objects based at least in part on one or more software function calls executed by a graphics processing unit (GPU) including one or more ray tracing cores.

12. The system of claim 8, wherein the one or more objects comprise one or more triangles in a triangular mesh.

13. The system of claim 8, wherein the API is configured to enable the use of one or more ray tracing cores to indicate the one or more objects based at least in part on a set of input parameters including the rays in the one or more rays and a data structure including the one or more objects.

14. A method comprising: Execute the application programming interface (API) to indicate to the user one or more objects that are intersected by one or more rays.

15. The method of claim 14, wherein the API is configured to enable the use of the ray tracing core to indicate the one or more objects based at least in part on a set of input parameters including an indication of a tree data structure.

16. The method of claim 14, wherein the API is configured to indicate the one or more objects based at least in part on function calls of software executed by a graphics processing unit (GPU).

17. The method of claim 14, wherein the API is configured to indicate the one or more objects at least in part based on enabling one or more ray tracing cores to identify the one or more objects intersected by the one or more rays.

18. The method of claim 14, wherein the API is configured to include the identifiers of intersecting objects among the one or more objects in an output data structure indicated by the user.

19. The method of claim 14, wherein the API is configured to enable the use of one or more ray tracing cores of a graphics processing unit (GPU) to indicate the one or more objects based at least in part on a set of input parameters including ray information and a tree data structure including the one or more objects.

20. A non-transitory computer-readable medium having a set of instructions stored thereon, which, when executed by one or more processors, cause the one or more processors to perform at least the method according to claim 14.

Images