Memory replication

Abstract

The present disclosure relates to memory replication. Apparatuses, systems, and techniques for performing one or more memory replication operations via a single application programming interface are disclosed. In at least one embodiment, the memory copy operations are performed with a single set of boot and / or shutdown operations therebetween. In at least one embodiment, a processor includes one or more circuits to execute an application programming interface (API), the information is copied from the two or more first storage locations to the two or more second storage locations based, at least in part, on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations.

Description

Memory copying Technical Field

[0001] At least one embodiment relates to batch processing of memory copying application programming interface (API) operations. At least one embodiment relates to eliminating the need for multiple individual memory copying API operations to prevent startup and / or shutdown operations from being repeatedly performed. At least one embodiment relates to executing the application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations. Background Technology

[0002] The memory copy API is executed against memory stored in a non-contiguous memory selection, or executed to copy to a non-contiguous memory selection. This type of API repeatedly performs start and stop operations each time the API is executed. Methods for performing multiple memory copy operations can be improved. Attached Figure Description

[0003] Figure 1 illustrates an example bulk memory copy operation according to at least one embodiment;

[0004] Figure 2 illustrates example input requirements for a bulk memory copy operation according to at least one embodiment;

[0005] Figure 3 illustrates example variations of the bulk memory copy operation according to at least one embodiment;

[0006] Figure 4 illustrates an example bulk memory copy operation using a single engine according to at least one embodiment;

[0007] Figure 5 illustrates an example memory copy operation performed over a network according to at least one embodiment;

[0008] Figure 6 illustrates an example bulk memory copying process according to at least one embodiment;

[0009] Figure 7 illustrates an example processor according to at least one embodiment;

[0010] Figure 8 illustrates example API calls for performing a bulk memory copy operation according to at least one embodiment;

[0011] Figure 9 illustrates an example data center system according to at least one embodiment;

[0012] Figure 10 illustrates a system-on-a-chip (SOC) according to at least one embodiment;

[0013] Figure 11A illustrates a parallel processor according to at least one embodiment;

[0014] Figure 11B illustrates a processing cluster according to at least one embodiment;

[0015] Figure 11C illustrates a graphics multiprocessor according to at least one embodiment;

[0016] Figure 12 illustrates an accelerator processor according to at least one embodiment;

[0017] Figure 13A illustrates a central processing unit according to at least one embodiment;

[0018] Figure 13B shows the core of the central processing unit in Figure 13A according to at least one embodiment;

[0019] Figure 14 illustrates another accelerator processor according to at least one embodiment;

[0020] Figure 15 illustrates a neuromorphic processor according to at least one embodiment;

[0021] Figure 16 illustrates a supercomputer according to at least one embodiment;

[0022] Figure 17 illustrates another accelerator processor according to at least one embodiment;

[0023] Figure 18 illustrates another processor according to at least one embodiment;

[0024] Figure 19 illustrates another accelerator processor according to at least one embodiment;

[0025] Figure 20 illustrates a tensor processing unit according to at least one embodiment;

[0026] Figure 21 illustrates a RISC-V compatible processor according to at least one embodiment;

[0027] Figures 22A and 22B illustrate a language processing unit according to at least one embodiment;

[0028] Figure 23 illustrates the software stack of a programming platform according to at least one embodiment;

[0029] Figure 24 illustrates software supported by a programming platform according to at least one embodiment;

[0030] Figure 25 illustrates compiled code for execution on the programming platform of Figure 24 according to at least one embodiment;

[0031] Figure 26 illustrates an example of an autonomous vehicle and its system architecture according to at least one embodiment;

[0032] Figure 27A illustrates the inference and / or training logic according to at least one embodiment;

[0033] Figure 27B illustrates the inference and / or training logic according to at least one embodiment; and

[0034] Figure 27C illustrates the training and deployment of a neural network according to at least one embodiment. Detailed Implementation

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

[0036] In at least one embodiment, the system and method implemented according to this disclosure are used to execute an application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations.

[0037] In at least one embodiment, a memory copy operation between two non-directly communicating memory groups requires the execution of an application programming interface (API) that instructs the copying of memory from one location to another. In at least one embodiment, if the memory locations are not contiguous (e.g., there are memory address gaps in the source or destination memory groups that will not be copied), the API must be executed once for each distinct, contiguous memory group. In at least one embodiment, this means that if the memory to be copied will be copied to multiple source and / or destination locations, and / or the memory to be copied does not include some memory embedded in a larger memory selection, the API must be executed many times, each time requiring start-up and shutdown operations (e.g., reserving and / or releasing semaphores and / or memory buffers, and / or any other required resources). In at least one embodiment, this can lead to latency and / or inefficiency due to the excessive number of operations performed.

[0038] In at least one embodiment, such memory copying functionality can be batch-processed, allowing the execution of a single API, and thus also allowing the execution of a single set of start and stop operations to complete two or more information transfers between storage locations (e.g., memory addresses). In at least one embodiment, such batch-processing APIs will require input parameters for each memory copying operation (e.g., the size of the information to be copied, its position in a first storage location, its position in a second storage location, an indication of the order in which the information copying will be performed, and / or other required information). In at least one embodiment, a processor (e.g., processor 702 in FIG. 7) will then parse the parameters to determine which engines (e.g., engine 408 in FIG. 4) need to be used to perform the memory transfer, and / or make any other determinations to then execute the batch memory copying process (e.g., process 600 in FIG. 6).

[0039] Various techniques have been described in the preceding and following sections. For ease of explanation, specific configurations and details have been outlined to provide a comprehensive understanding of the possible ways to implement these techniques. However, it will be apparent that the techniques described below can be practiced in different configurations without specific details. Furthermore, well-known features may be omitted or simplified to avoid obscuring the described techniques.

[0040] In at least one embodiment, as used in any implementation described herein, unless the context explicitly states otherwise or to the contrary, terms such as “module” and nominalized verbs refer to any combination of software logic, firmware logic, hardware logic, and / or 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, while “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 circuitry forming part of a larger system, such as an integrated circuit (IC), a system-on-a-chip (SoC), etc.

[0041] In at least one embodiment, the system (e.g., Example 100, Example 200, Example 300, Example 400, Example 500, Process 600, Example 700, and / or Example 800) includes a collection of one or more hardware and / or software computing resources with instructions that, when executed, perform one or more communication processes, such as those described herein. In at least one embodiment, Example 100, Example 200, Example 300, Example 400, Example 500, Process 600, Example 700, and / or Example 800 includes one or more software programs executable on computer hardware, one or more application programs executable on computer hardware, and / or variations thereof. In at least one embodiment, one or more of the processes of Example 100, Example 200, Example 300, Example 400, Example 500, Process 600, Example 700 and / or Example 800 are executed by any suitable processing system or unit (e.g., graphics processing unit (GPU), general-purpose GPU (GPGPU), parallel processing unit (PPU), central processing unit (CPU), data processing unit (DPU) (described below)) in any suitable manner, including sequentially, in parallel and / or variations thereof. In at least one embodiment, Examples 100, 200, 300, 400, 500, 600, 700, and / or 800 use machine learning training frameworks, such as PYTORCH, TENSORFLOW, BOOST, CAFFE, MICROSOFT COGNITIVE TOOLKIT / CNTK, MXNET, CHAINER, KERAS, DEEPLEARNING4J, and / or other training frameworks, to implement and perform the operations described herein, thereby enabling the execution of an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API used to indicate two or more first storage locations and two or more second storage locations. In at least one embodiment, as an example, training a neural network model includes using a server (e.g., an NVIDIA DGX server) that also includes at least a GPU (e.g., an AMD MI200, VEGAL10, VEGO20, and ARCTURUS), an optimizer (e.g., an ADAM OPTIMIZER), or a discriminator architecture (e.g., a discriminator architecture from face-vid2vid for training with GAN loss).

[0042] Figure 1 illustrates a bulk memory copy operation of Example 100 according to at least one embodiment. In at least one embodiment, Example 100 includes one or more memory copies 102A-C (e.g., memory copy 102), bulk API calls 104, and / or memory destinations 106A-C (e.g., memory destination 106). In at least one embodiment, Example 100 performs part or all of the bulk memory copy process (e.g., process 600 in Figure 6).

[0043] In at least one embodiment, the processor (e.g., processor 702 in FIG. 7) uses memory copy 102 to indicate information, such as one or more indications to move contiguous memory selections from one set of contiguous memory addresses to another set of contiguous memory addresses. In at least one embodiment, memory copy 102 is a software indication that will be copied to the memory of one or more memory destinations 106, and / or some other indication that will be copied to the memory of one or more memory destinations 106. In at least one embodiment, memory copy 102 is input to bulk API call 104 for collective parsing and execution, thereby copying data from a data source to a data destination in memory.

[0044] In at least one embodiment, the processor (e.g., processor 702 in FIG. 7) uses bulk API call 104 to indicate information, such as information indicating one or more instances of bulk memory copy operations performed via an application programming interface (API). In at least one embodiment, the bulk API call is an instruction to perform a bulk memory copy API (e.g., example 700 in FIG. 7). In at least one embodiment, the bulk API call 104 receives input memory copy 102 as input, thereby performing one or more memory copy operations as a result. In at least one embodiment, the bulk API call 104 is executed upon invocation to copy multiple non-contiguous memory copy operations, wherein a set of start and / or shut-down operations are performed for each group.

[0045] In at least one embodiment, the processor (e.g., processor 702 in FIG. 7) uses memory destination 106 to indicate information, such as information indicating one or more memory addresses, wherein the memory selection indicated by one or more memory copy 102 will be copied. In at least one embodiment, memory destination 106 is one or more memory hardware and / or software addresses whose size is sufficient to copy the indicated data. In at least one embodiment, memory destination 106 is indicated by an associated memory copy 102 to allow copying the appropriate memory.

[0046] In at least one embodiment, Example 100 includes one or more processors for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 100 is the system shown in Figures 1-8, included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 100 performs one or more processes shown in Figures 1-8, such as executing an Application Programming Interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, Example 100 performs one or more processes shown in Figures 9-27C, such as executing an Application Programming Interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein.

[0047] Figure 2 illustrates the input requirements of an example 200 of a bulk memory copy operation according to at least one embodiment. In at least one embodiment, example 200 includes one or more memory copies 202A and / or 202B, a bulk API call 204, and / or a memory destination 206. In at least one embodiment, memory copies 202A and / or 202B include one or more memory sources 208A and / or 208B (e.g., memory source 208), memory sizes 210A and / or 210B (e.g., memory size 210), and / or memory destinations 212A and / or 212B (e.g., memory destination 212). In at least one embodiment, memory copies 202A and / or 202B, bulk API call 204, and / or memory destination 206 have been described in conjunction with the components in the previous figures (e.g., memory copy 102, bulk API call 104, and / or memory destination 106 in Figure 1, respectively) and are fully defined without further description. In at least one embodiment, Example 200 performs part or all of a batched memory copying process (e.g., process 600 in Figure 6).

[0048] In at least one embodiment, the processor (e.g., processor 702 in FIG. 7) uses memory source 208, memory size 210, and / or memory destination 212 to indicate information, such as information indicating one or more memory addresses containing one or more data points to be copied to memory destination 212, including one or more data points indicated to begin from one or more memory addresses indicated by memory source 208 and / or the size of the memory dataset indicated by memory size 210. In at least one embodiment, memory source 208 includes an indication of one or more memory addresses comprising a contiguous set of data points in memory, the size of which is indicated by the associated memory size 210. In at least one embodiment, the memory will be copied to one or more corresponding data addresses indicated by memory destination 212. In at least one embodiment, each instance of memory source 208, memory size 210, and / or memory destination 212 is discontinuous (e.g., any memory address indicated by memory source 208A and / or 208B includes discontinuous data between them, and / or the memory addresses indicated by memory destination 212A and / or 212B are discontinuous between them). In at least one embodiment, all input memory source 208, memory size 210, and / or memory destination 212 are input to bulk API call 204 to perform a data copying process (e.g., process 600 in FIG. 6). In at least one embodiment, memory source 208, memory size 210, and / or memory destination 212 are data stored in memory that indicates one or more parameters of one or more memory copying operations (e.g., memory copy 202).

[0049] In at least one embodiment, Example 200 includes one or more processors for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 200 is the system shown in Figures 1-8, included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 200 performs one or more processes shown in Figures 1-8, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, Example 200 performs one or more processes shown in Figures 9-27C, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein.

[0050] Figure 3 illustrates a differentiated example 300 variant of a batch memory copy operation according to at least one embodiment. In at least one embodiment, example 300 includes one or more memory sources 308, memory destinations 306A and / or 306B, and / or batch memory copy 310. In at least one embodiment, memory source 308 includes one or more memory sources 302A-E. In at least one embodiment, memory destinations 306A and / or 306B include one or more memory destinations 304A-E. In at least one embodiment, batch memory copy 310 includes one or more copy operations 312A-D (e.g., copy operation 312). In at least one embodiment, memory sources 302A-E, memory destinations 306A and / or 306B, and / or memory destination 304 are described in conjunction with one or more components from the previous figures (e.g., memory source 208, memory destination 106, and / or memory destination 212 shown in Figures 1 and / or 2, respectively), and are fully defined without further description. In at least one embodiment, Example 300 performs part or all of a batch memory copying process (e.g., process 600 shown in FIG. 6).

[0051] In at least one embodiment, the processor (e.g., processor 702 shown in FIG. 7) uses memory source 308 to indicate information, such as information indicating one or more selections of memory addresses, said memory addresses including at least all memory addresses covered by memory sources 302A-E. In at least one embodiment, memory source 308 may include one or more data points of memory not indicated by memory sources 302A-E. In at least one embodiment, memory source 308 is one or more physical memory systems (e.g., RAM, DRAM, system memory, and / or any other memory hardware) that include one or more indicated memory addresses indicated by memory sources 302A-E.

[0052] In at least one embodiment, a processor (e.g., processor 702 in FIG. 7) uses bulk memory copy 310 to indicate information, such as information indicating one or more individual memory copy operations to be executed in batches by a bulk memory copy operation. In at least one embodiment, the bulk memory copy performs copy operations 312 in an indicated order. In at least one embodiment, the order of copy operations 312 is determined by one or more processors (e.g., processor 702 in FIG. 7). In at least one embodiment, the bulk memory copy 310 performs reservation of semaphores, memory buffers, and / or any other resources required to perform the copy operation 312 before starting the first copy operation 312, and performs release of all reserved resources after completing the last copy operation 312. In at least one embodiment, the bulk memory copy is performed as a result of an API call (e.g., example 700 in FIG. 7).

[0053] In at least one embodiment, a processor (e.g., processor 702 in FIG. 7) uses copy operation 312 to indicate information, such as information indicating one or more operations to copy one or more data points from an associated memory source 302 to an associated memory destination 304. In at least one embodiment, copy operation 312A is performed because the associated memory source 302 and / or memory destination 304 are not contiguous with other memory locations involved in other copy operations 312. In at least one embodiment, copy operations 312B and / or 312C are performed separately because memory destinations B and / or C are not contiguous. In at least one embodiment, copy operation 312D is performed separately from copy operation AC because its destination is located in separate memory hardware and is therefore not contiguous. In at least one embodiment, copy operation 312D is performed because the associated memory source 302 and / or memory destination 304 involved are contiguous. In at least one embodiment, copy operation 312 is software that performs data copy operations between two sets of memory addresses.

[0054] In at least one embodiment, Example 300 includes one or more processors for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 300 is the system shown in Figures 1-8, included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 300 performs one or more processes shown in Figures 1-8, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, Example 300 performs one or more processes shown in Figures 9-27C, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein.

[0055] Figure 4 illustrates an example 400 of a bulk memory copy operation using a single engine according to at least one embodiment. In at least one embodiment, example 400 includes one or more memory copies 402A-C (e.g., memory copy 402), a bulk API call 404, memory destinations 406A-C (e.g., memory destination 406), and / or engines 408A and / or 408B (e.g., engine 408). In at least one embodiment, memory copy 402, bulk API call 404, and / or memory destination 406 are described in conjunction with one or more components from the previous figures (e.g., memory copy 102, bulk API call 104, and / or memory destination 106 in Figure 1, respectively), and are fully defined without further description. In at least one embodiment, example 400 performs part or all of a bulk memory copy process (e.g., process 600 in Figure 6).

[0056] In at least one embodiment, the processor (e.g., processor 702 in FIG. 7) uses engine 408 to indicate information, such as information indicating one or more software and / or hardware components required to perform one or more memory copy operations. In at least one embodiment, engine 408 can be distinguished by the required memory source and / or destination (which require different memory transfer methods to perform the memory copy operation). In at least one embodiment, bulk API calls 404 separate memory copy 402 using engine 408 to allow memory copy operations using similar engines 408 (e.g., memory copies B and / or C using engine 408B) to be performed together. In at least one embodiment, for example, engine 408 may be a chain of software and / or hardware components required to perform host-to-device copy, while a second engine 408 may be a separate set of software and / or hardware for host-to-host copy. In at least one embodiment, engine 408 may be any combination of hardware and / or software components required to perform a memory copy process (e.g., process 600 in FIG. 6).

[0057] In at least one embodiment, Example 400 includes one or more processors for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 400 is the system shown in Figures 1-8, included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 400 performs one or more processes shown in Figures 1-8, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, Example 400 performs one or more processes shown in Figures 9-27C, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein.

[0058] Figure 5 illustrates an example 500 of a memory copy operation performed over a network according to at least one embodiment. In at least one embodiment, example 500 includes one or more devices 504A and / or 504B (e.g., device 504). In at least one embodiment, device 504A includes one or more memories 508A and / or processors 512A. In at least one embodiment, memory 508A includes one or more memory sources 102A-C (e.g., memory source 502). In at least one embodiment, device 504B includes one or more memories 508B, processor 512B, API 514, and / or driver 516. In at least one embodiment, memory 508B includes one or more memory destinations 506A-C (e.g., memory destination 506). In at least one embodiment, memory 508A, memory source 502, memory 508B, memory destination 506, and / or API 514 are described in conjunction with components in the previous figures (e.g., memory source 308, memory source 302, memory destination 306, memory destination 106, and / or bulk API call 104, respectively, in Figures 1 and / or 3), and are fully defined without further description. In at least one embodiment, example 500 performs part or all of a memory copying process (e.g., process 600 in Figure 6).

[0059] In at least one embodiment, a processor (e.g., processor 702 in FIG. 7) uses device 504 to indicate information, such as information indicating one or more physical devices (e.g., host, device, GPU, CPU, GPGPU, and / or any other designation of a separate device), said devices being indicated by one or more associated memory copy operations to at least partially include memory source 502 and / or memory destination 506. In at least one embodiment, device 504 is a single device, wherein the associated components are also the same components. In at least one embodiment, device 504 may communicate via network 510. In at least one embodiment, device 504 executes one or more APIs 514 to copy data from memory source 502 to memory destination 506.

[0060] In at least one embodiment, the processor (e.g., processor 702 in FIG. 7) uses network 510 to indicate information, such as information indicating one or more nodes including a communication network to allow facilitation between devices 504. In at least one embodiment, network 510 is one or more hardware and / or software systems required to transfer data to perform a memory copy operation (e.g., process 600 in FIG. 6) between memory source 502 and / or memory destination 506. In at least one embodiment, network 510 may not be required.

[0061] In at least one embodiment, a processor (e.g., processor 702 in FIG. 7) uses processor 512 to indicate information, such as one or more processors (e.g., processor 702 in FIG. 7) for facilitating and / or performing a memory copying process (e.g., process 600 in FIG. 6). In at least one embodiment, processor 512 performs one or more other functions related to performing operations within device 504. In at least one embodiment, processor 512 is not the processor performing the memory copying process (e.g., process 600 in FIG. 6), but rather one or more other processors capable of performing part or all of the indicated process.

[0062] In at least one embodiment, a processor (e.g., processor 702 in FIG. 7) uses driver 516 to indicate information, such as information instructing one or more processors (e.g., processor 702 in FIG. 7) and / or software executed by one or more processors to control device 504. In at least one embodiment, driver 516 may include one or more instructions for performing a memory copying process (e.g., process 600 in FIG. 6) and / or facilitating the execution of the memory copying process. In at least one embodiment, driver 516 may be executed by device 504B, device 504A, and / or certain external devices.

[0063] In at least one embodiment, Example 500 includes one or more processors for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 500 is the system shown in Figures 1-8, included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 500 performs one or more processes shown in Figures 1-8, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, Example 500 performs one or more processes shown in Figures 9-27C, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein.

[0064] Figure 6 illustrates an example process 600 for batch memory copying according to at least one embodiment. In at least one embodiment, one or more processors (e.g., processor 702 in Figure 7) initiate process 600 602 (when invoked) to perform a set of one or more batch memory copying operations. In at least one embodiment, the received input uses one or more data formats such that process 600 can iterate to the next feature (e.g., a first feature indicating to begin 602). In at least one embodiment, one or more processors (e.g., processor 702 in Figure 7) use process 600 to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations. In at least one embodiment, process 600 begins at step 602.

[0065] In at least one embodiment, process 600 includes at step 604: verifying the indicated memory to be copied in the indicated memory copy operation. In at least one embodiment, this includes one or more processes for checking whether the memory is valid, whether it has not been reserved by other operations, and whether it can be copied by the memory copy operation.

[0066] In at least one embodiment, process 600 includes at step 606: grouping a group of one or more related memory copy operations by an engine for allowing similar memory copy operations to be executed together.

[0067] In at least one embodiment, process 600 includes at step 608 reserving memory buffers and semaphores required to perform each associated memory copy operation. In at least one embodiment, if multiple engines are selected, this may include multiple sets of buffers and / or semaphores, but memory copy operations executable on the same set of buffers and / or semaphores will be reserved in one set, and the operations will be performed serially.

[0068] In at least one embodiment, process 600 includes, at step 610, a push copy command that performs at least one memory copy operation. In at least one embodiment, if multiple memory copy operations are performed, they will be performed serially in the indicated order unless instructed to be performed in parallel.

[0069] In at least one embodiment, process 600 includes at step 612 determining whether all indicated copy commands have been pushed. In at least one embodiment, if not, process 600 may iterate back to step 610. In at least one embodiment, if yes, process 600 may iterate to step 614.

[0070] In at least one embodiment, process 600 includes at step 614: releasing buffers and semaphores required for and reserved for the previously performed memory copy operation. In at least one embodiment, the resources are released back to the controller for use elsewhere, and, if instructed, the contained information is cleared.

[0071] In at least one embodiment, process 600 may end at step 616 or otherwise terminate. In at least one embodiment, process 600 may terminate if process 600 completes a memory copy operation, indicates completion of a memory copy operation, or otherwise returns an error.

[0072] In at least one embodiment, a processor uses a process 600 comprising one or more steps to execute an application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, as an example, a machine-readable medium storing a set of instructions that, if executed by one or more processors, causes the one or more processors to execute the application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, process 600 includes, is included in, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, process 600 is executed by one or more systems shown in Figures 1-8, for example, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, process 600 is performed by one or more systems shown in Figures 9-27C, for example, to cause the execution of an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein.

[0073] Figure 7 illustrates an example system 700 of a processor according to at least one embodiment. In at least one embodiment, the processor 702 executes one or more processes (e.g., the processes described herein) for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations. In at least one embodiment, the processor 702 executes the process described in conjunction with Figure 1. In at least one embodiment, the processor 702 executes one or more processes (e.g., the processes described in conjunction with Figures 1-7).

[0074] In at least one embodiment, processor 702 includes one or more processors (e.g., processors described in conjunction with Figures 9-27C). In at least one embodiment, processor 802 is any suitable processing unit and / or combination of processing units, such as one or more CPUs, GPUs, GPGPUs, PPUs, and / or variations thereof. In at least one embodiment, processor 802 includes one or more communication modules 704, call parsing modules 706, API modules 708, and / or engine modules 710. In at least one embodiment, one or more communication modules 704, call parsing modules 706, API modules 708, and / or engine modules 710 are part of processor 802 and / or one or more other processors. In at least one embodiment, one or more communication modules 704, call parsing modules 706, API modules 708, and / or engine modules 710 are distributed across multiple processors that communicate via buses, networks, writes to shared memory, and / or any suitable communication process (e.g., the communication processes described herein).

[0075] In at least one embodiment, as used in any implementation described herein, unless the context explicitly states otherwise or otherwise specifies otherwise, a module refers to any combination of software logic, firmware logic, hardware logic, and / or 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 (e.g., individually or in any combination) hardwired circuitry, programmable circuitry, state machine circuitry, fixed-function circuitry, execution unit circuitry, and / or firmware storing instructions executed by the programmable circuitry. In at least one embodiment, a module may be embodied collectively or individually as circuitry forming part of a larger system (e.g., an integrated circuit (IC), a system-on-a-chip (SoC), etc.). In at least one embodiment, a module is combined with any suitable processing unit and / or combination of processing units (e.g., one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof) to perform one or more processes.

[0076] In at least one embodiment, the processor uses communication module 704 to perform and / or facilitate communication operations. In at least one embodiment, communication module 704 provides output to one or more other modules, memories, external hardware, and / or any other data source required to perform a memory copying process (e.g., process 600 in FIG. 6) in the form of data points from memory, data representing information required to perform the memory copying operation, and / or any other data required to perform the memory copying process (e.g., process 600 in FIG. 6). In at least one embodiment, communication module 704 receives input from data points from memory in the form of data representing information required to perform the memory copying operation. In at least one embodiment, communication module 704 is combined with any suitable processing unit and / or combination of processing units (e.g., one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof) to perform and / or facilitate communication operations.

[0077] In at least one embodiment, the processor uses the call parsing module 706 to perform operations to parse one or more batch memory copy operations, thereby allowing additional information required to facilitate resource reservation and / or perform memory copy operations (e.g., process 600 in FIG. 6). In at least one embodiment, the call parsing module 706 provides output to one or more communication modules 704, system memory, and / or other data sources required to perform data parsing, in the form of data indicating parameters and / or resources required for the relevant batch of memory copy operations. In at least one embodiment, the call parsing module 706 receives input in the form of data indicating one or more parameters (e.g., memory size, memory destination, etc.) for the relevant batch of one or more memory copy operations. In at least one embodiment, the call parsing module 706 performs operations to parse one or more batches of memory copy operations, thereby allowing additional information required to facilitate resource reservation and / or perform memory copy operations (e.g., process 600 in FIG. 6) in conjunction with any suitable processing unit and / or processing unit combination (e.g., one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof).

[0078] In at least one embodiment, the processor uses API module 708 to invoke one or more APIs for performing a batch memory copy process (e.g., process 600 in FIG. 6) and / or other APIs required to perform such a process. In at least one embodiment, API module 708 provides output to communication module 704, external memory, and / or any other data source required to perform the memory copy operation (e.g., process 600 in FIG. 6) in the form of data indicating parameters for performing one or more APIs (e.g., example 700 in FIG. 7) and / or one or more invocations of such APIs. In at least one embodiment, API module 708 receives input in the form of parameters indicating one or more APIs and / or data representing one or more parameters required to perform such APIs. In at least one embodiment, API module 708 invokes one or more APIs for performing a batch memory copy process (e.g., process 600 in FIG. 6) and / or other APIs required to perform such a process in combination with any suitable processing unit and / or processing unit combination (e.g., one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof).

[0079] In at least one embodiment, the processor uses engine module 710 to allocate and / or reserve hardware and / or software resources required to perform the relevant memory copying operation using an indicated engine (e.g., engine 408 in FIG. 4). In at least one embodiment, engine module 710 provides output to communication module 704, system memory, and / or any other data source required to perform the memory copying process (e.g., process 600 in FIG. 6) in the form of data indicating one or more resources that need to be reserved to perform the relevant memory copying operation. In at least one embodiment, engine module 710 receives input in the form of parameters for one or more relevant memory copying operations to be performed. In at least one embodiment, engine module 710 uses the indicated engine (e.g., engine 408 in FIG. 4) in conjunction with any suitable processing unit and / or combination of processing units (e.g., one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof) to perform the allocation and / or reservation of hardware and / or software resources required to perform the relevant memory copying operation.

[0080] In at least one embodiment, Example 700 includes one or more processors for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 700 is the system shown in Figures 1-8, included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise perform the operations described herein. In at least one embodiment, Example 700 performs one or more processes shown in Figures 1-8, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein. In at least one embodiment, Example 700 performs one or more processes shown in Figures 9-27C, such as those used to execute an application programming interface (API), to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, and / or otherwise performs the operations described herein.

[0081] In at least one embodiment, performing a batch memory copy operation 802 (“call 802”) is a function call executed by one or more software programs, such as a kernel executed by one or more parallel processing units (PPUs) (e.g., graphics processing units (GPUs)). In at least one embodiment, call 802 is a call to an instruction that causes one or more processors to perform an operation to copy data points in memory within a memory source 808 to generate an indicator 814 of acknowledgment 804 (“response 804”). In at least one embodiment, call 802 is a call to an API that causes one or more processors to perform one or more batch memory copy processes (e.g., process 600 in FIG. 6) to generate one or more indicators 814 of response 804.

[0082] In at least one embodiment, call 802 receives parameters 808, 806, 810, and / or 812 as input, including a memory source 808, a memory destination 806, and / or a memory size 810. In at least one embodiment, the input is described in conjunction with components of one or more of the previous figures (e.g., memory source 208, memory destination 104, and memory size 210 as shown in Figures 1 and / or 2, respectively), and is fully defined without further description.

[0083] In at least one embodiment, call 802 receives parameters 812, including a copy order 812, as input. In at least one embodiment, the copy order 812 includes one or more indications that indicate the order in which a set of associated memory copy operations are performed. In at least one embodiment, the indications may be included in one or more references to the memory size, source, and / or destination provided for each individual memory copy in a batch memory copy operation. For example, in at least one embodiment, a batch memory copy operation includes two host-to-host (H2H) transfers, two device-to-host (D2H) transfers, and three device-to-device (D2D) transfers, which are grouped by transfer type, but still require an order for each type set. In at least one embodiment, the indications may be included in the transfer request such that, for the two indicated H2H transfers, the indicated transfer is performed first. In at least one embodiment, the transfer order may be calculated by one or more processors (e.g., processor 702 in FIG. 7) based on other input parameters of call 802 (e.g., parameters 808, 806, 810). In at least one embodiment, for example, transmissions may be ordered based on the size of the transmitted data, the reliability of the connection between the memory source and the memory destination, and / or any other transmission ordering method.

[0084] In at least one embodiment, a copy order 812 is not required and / or is not provided, wherein memory copy operations are performed serially, in parallel, and / or in combinations thereof, depending on optimization constraints of the relevant engine, dataset source, dataset destination, dataset size, and / or any other parameters indicating the optimized order. In at least one embodiment, the copy order 812 (if present) may include one or more attributes for specifying source ordering, flags for indicating operation ordering requirements, and / or explicitly specifying ordering requirements. In at least one embodiment, other methods for ordering memory copy operations may be used to indicate one or more copy orders.

[0085] In at least one embodiment, call 802 generates parameters 814 as output, including one or more indicators 814. In at least one embodiment, the indicators 814 include one or more indications that one or more batch memory copy operations have been successfully executed, unsuccessfully executed, paused, and / or otherwise terminated. In at least one embodiment, the indicators 814 may include one or more other indications, such as one or more indications of an alternative memory destination, and / or other indications modifying memory transfers.

[0086] In at least one embodiment, the processor uses operation 800, which includes one or more steps for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API used to indicate two or more first storage locations and two or more second storage locations to two or more processors, and / or otherwise perform the operations described herein. In at least one embodiment, as an example, a machine-readable medium stores a set of instructions that, if executed by one or more processors, cause the one or more processors to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API used to indicate two or more first storage locations and two or more second storage locations to two or more processors, and / or otherwise perform the operations described herein. In at least one embodiment, operation 800 includes the system shown in Figures 1-8, is included in the system shown in Figures 1-8, and / or otherwise includes the system shown in Figures 1-8, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API used to indicate two or more first storage locations and two or more second storage locations to two or more processors, and / or performs other operations described herein. In at least one embodiment, operation 800 is performed by one or more systems shown in Figures 9-27C, for example, for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API used to indicate two or more first storage locations and two or more second storage locations to two or more processors, and / or otherwise performs the operations described herein.

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

[0088] Data Center

[0089] Figure 9 illustrates an example data center 900 according to at least one embodiment. The data center 900 may include one or more rooms having racks 902 and auxiliary equipment for housing one or more racks 902 and one or more substrates 904. A rack 902 may include one or more substrates 904. A rack 902 may include a housing for housing and supporting individual substrates 904. Operational aspects of the rack 902 may be adjustable at the rack level (corresponding to a group of substrates 904) or at the substrate level (corresponding to an individual substrate 904), among other options. The rack 902 or substrate 904 may have specific selected maximum operating parameters, such as, but not limited to, power consumption, operating frequency, etc. The data center 900 may be supported by various cooling systems, such as, but not limited to, cooling towers, cooling loops, pumps, and other support systems. The cooling system may include sensors and controllers for monitoring and managing the cooling characteristics of the rack 902. The substrates 904 within the rack 902 may draw operating power from one or more power distribution units (PDUs; not shown). PDUs can be arranged within racks 902, for example, between racks 902 that include substrates 904, or within racks 902 that also house substrates 904.

[0090] Rack 902 and baseboard 904 may include subsystems, modules, add-in cards, and other semiconductor components. Baseboard 904 may include one or more computing units 906, each computing unit 906 may include one or more processors 908, one or more memories 910, and an interface controller 912. Computing units 906 may include any number of processors, such as, but not limited to, central processing units (“CPU”), graphics processing units (“GPU”), or other processors (including accelerators, field-programmable gate arrays (FPGAs), graphics processors, etc.), including any processor described herein, such as, but not limited to, the processors in Figures 10-22B. Computing units 906 may include one or more memory storage devices 910 (e.g., dynamic read-only memory, solid-state storage devices, or disk drives), as well as network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power modules, and cooling modules, etc. One or more computing units 906 may be servers having one or more of the aforementioned computing resources.

[0091] Computing unit 906 may include individual computing unit groups housed in one or more racks (not shown), or in numerous racks within data centers in different geographical locations (also not shown). Individual computing unit groups may include grouped computing, networking, memory, or storage resources that can be configured or allocated to support one or more workloads. Several computing units (e.g., including CPUs and / or other processors) may be grouped within one or more racks to provide computing resources to support one or more workloads. Resource coordinator 914 may configure or otherwise control one or more computing units 906 or groups of computing units. Resource coordinator 914 may include a Software Design Infrastructure (“SDI”) management entity for data center 900. Resource coordinator 914 may include hardware, software, or some combination thereof.

[0092] Data center 900 may include any one or any combination of a framework layer 920, a software layer 930, and an application layer 940. As shown in Figure 9, the framework layer 920 includes a job scheduler 922, a configuration manager 924, a resource manager 926, and a distributed file system 928. The framework layer 920 may include a framework for supporting software 932 of the software layer 930 and / or one or more applications 942 of the application layer 940. Software 932 or application 942 may respectively include web-based service software or applications, such as, but not limited to, software or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. The framework layer 920 may be a type of free and open-source software web application framework, such as, but not limited to, Apache Spark. TM (Hereinafter referred to as "Spark"), which can utilize the distributed file system 928 for large-scale data processing (e.g., "big data"). The job scheduler 922 may include Spark drivers, which facilitate the scheduling of workloads supported by various layers of the data center 900. The configuration manager 924 may be able to configure different layers, such as, but not limited to, the software layer 930 and the framework layer 920 (which includes Spark and the distributed file system 928 for supporting large-scale data processing). The resource manager 926 may be able to manage clustered or grouped compute units 906 mapped to or allocated to support the distributed file system 928 and the job scheduler 922. The resource manager 926 may coordinate with the resource coordinator 914 to manage these mapped or allocated compute resources.

[0093] Software 932 may be included in software layer 930, and may include software used by at least a portion of computing units 906, one or more computing units 906, groups of computing units 906, and / or the distributed file system 928 of framework layer 920. One or more types of software may include, but are not limited to, internet web search software, email virus scanning software, database software, and streaming video content software.

[0094] Application 942 may be included in application layer 940 and may include one or more types of applications used by at least the portions of computing unit 906, one or more computing units 906, groups of computing units 906, and / or the distributed file system 928 of framework layer 920. One or more types of applications may include, but are not limited to, any number of genomics applications, cognitive computing applications, and machine learning applications, including training or inference software, machine learning framework software (e.g., PyTorch, TensorFlow, Caffe, etc.), or other machine learning applications used in conjunction with one or more embodiments.

[0095] Any of the Configuration Manager 924, Resource Manager 926, and Resource Coordinator 914 can implement any number and type of self-modification actions based on any amount and type of data obtained in any technically feasible manner. Self-modification actions can alleviate the burden on data center operators of Data Center 900 to make potentially erroneous configuration decisions and may prevent underutilized and / or poorly performing portions of the data center.

[0096] Data center 900 may include tools, services, software, or other resources for training one or more machine learning models according to one or more embodiments described herein, or for using one or more machine learning models to predict or infer information. For example, a machine learning model can be trained by calculating weight parameters based on a neural network architecture using the software and computing resources described above regarding data center 900. The trained machine learning model corresponding to one or more neural networks can be used with the resources described above regarding data center 900 to infer or predict information using weight parameters calculated through one or more training techniques described herein.

[0097] Data center 900 may use CPUs, application-specific integrated circuits (ASICs), GPUs, FPGAs, or other hardware (e.g., embodiments in Figures 10-22B) to perform some or all of the processes and techniques described elsewhere herein, such as, but not limited to, training and / or inference using the resources described above. Furthermore, one or more of the aforementioned software and / or hardware resources may be configured as a service to allow a user to train or perform information inference, such as, but not limited to, image recognition, speech recognition, or other artificial intelligence services.

[0098] In at least one embodiment, processor 908 may include one or more processors and / or circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. In at least one embodiment, processor 908 is configured by software 932 to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. Data center 900 may use logic, a CPU, an application-specific integrated circuit (ASIC), a GPU, an FPGA, or other hardware (e.g., the embodiments in Figures 10-22B) to perform any of the operations described above or elsewhere herein.

[0099] processor

[0100] The following figures illustrate, without limitation, example processors and processing systems that can be used to execute an application programming interface (API) to copy information from two or more first storage locations to two or more second storage locations, at least in part, based on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations, or otherwise perform some or all of the processes, operations, and / or techniques described elsewhere herein. The example processors and processing systems may be software-configured to execute the application programming interface (API) to copy information from two or more first storage locations to two or more second storage locations, at least in part, based on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. Processors and processing systems may include logic, a central processing unit (CPU), an application-specific integrated circuit (ASIC), a graphics processing unit (GPU), a field-programmable gate array (FPGA), an XPU (i.e., any computing architecture best suited to the needs of the application), or other hardware (e.g., the embodiments in Figures 10-22B) for performing any of the operations described above, below, or elsewhere herein. The processor and / or processing system described herein may include one or more circuits that can be used to execute an application programming interface (API) to copy information from two or more first storage locations to two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. As used herein, one or more circuits may be software-configurable to execute the application programming interface (API) to copy information from two or more first storage locations to two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. Figures 27A and 27B illustrate logic 2715 as described elsewhere herein, which may be used in one or more devices to perform, for example (but not limited to), the operations discussed herein according to at least one embodiment. Logic can refer to any combination of software logic, hardware logic, and / or firmware logic used to provide the functions and / or operations described herein, wherein the logic can be collectively or individually embodied as circuitry forming part of a larger system, such as an integrated circuit (IC), an application-specific integrated circuit (ASIC), a field-programmable gate array (FPGA), a system-on-a-chip (SoC), or one or more processors (e.g., CPU, GPU).

[0101] Figure 10 illustrates a processor according to at least one embodiment, which is a System-on-Chip (SOC) 1000 (also referred to as a system-on-a-chip, superchip, or other names). The SOC 1000 may include processor complex 1010 and processor complex 1040. The SOC 1000 may include any number of processor complexes 1010 and / or processor complexes 1040, which may include any number of processors in any combination as described herein, such as, but not limited to, the processors in Figures 10-22B. For example, processor 1010 may include a central processing unit (CPU), and processor 1040 may include a graphics processor. Alternatively, processor 1010 may include a graphics processor, and processor 1040 may include a graphics processor. The SOC 1000 may include any number of display controllers 1092, any number of multimedia engines 1094, any number of I / O interfaces 1070, any number of memory controllers 1080, and any number of fabrics 1060 in any combination. For ease of explanation, this document uses reference numbers to identify objects and bracket numbers to identify instances (if necessary) to represent multiple instances of similar objects. The SOC 1000 may include processors from Broadcom Inc., Palo Alto, California.

[0102] Processor complex 1010 may include a CPU, processor complex 1040 may include a GPU, and SOC 1000 may include a processing unit integrating processor complex 1010 and processor complex 1040 onto a single chip. Certain tasks may be assigned to processor complex 1010, while other tasks may be assigned to processor complex 1040. Processor complex 1010 may be configured to execute main control software associated with SOC 1000, such as, but not limited to, an operating system. Processor complex 1010 may be the main processor of SOC 1000, controlling and coordinating the operation of other processors. Processor complex 1010 may issue commands that control the operation of processor complex 1040 to perform some or all of the operations described herein. Processor complex 1010 may be configured to execute host-executable code derived from CUDA or other source code (e.g., HIP source code), while processor complex 1040 may be configured to execute device-executable code derived from CUDA or other source code to perform any of the operations described herein.

[0103] The processor complex 1010 may include cores 1020(1)-1020(4) and a cache (e.g., L3 cache) 1030 for storing information for performing the operations described herein. The processor complex 1010 may include any number of cores 1020 in any combination and any number and type of cache. The cores 1020 may be configured to execute instructions of a specific instruction set architecture (“ISA”) to perform some or all of the operations described herein. Each core 1020 may include a CPU core. Cores 1020(1)-1020(4) may be referred to as compute units or arithmetic units. The SOC 1000 may include any number of processor complexes 1010, architecture 1060, I / O interface 1070, and memory controller 1080.

[0104] Each core 1020 may include a fetch / decode unit 1022, an integer execution engine 1024, a floating-point execution engine 1026, and an L2 cache 1028. The fetch / decode unit 1022 may fetch instructions to perform some or all of the operations described herein (e.g., but not limited to APIs compiled into instructions) and decode those instructions, generate micro-operations, and dispatch individual micro-instructions to the integer execution engine 1024 and / or the floating-point execution engine 1026. The fetch / decode unit 1022 may concurrently dispatch one micro-instruction to the integer execution engine 1024 and another micro-instruction to the floating-point execution engine 1026. The integer execution engine 1024 may perform integer and memory operations. The floating-point engine 1026 may perform floating-point and vector operations. The fetch / decode unit 1022 may dispatch micro-instructions to one or more execution engines, which may replace both the integer execution engine 1024 and the floating-point execution engine 1026.

[0105] Each core 1020(i) (where i is an integer representing a specific instance of core 1020) can access the L2 cache 1028(i) included in core 1020(i). Each core 1020 included in core complex 1010(j) (where j is an integer representing a specific instance of core complex 1010) can be connected to other cores 1020 included in core complex 1010(j) via the L3 cache 1030(j) included in core complex 1010(j). The cores 1020 included in core complex 1010(j) (where j is an integer representing a specific instance of core complex 1010) can access all L3 caches 1030(j) included in core complex 1010(j). The L3 cache 1030 can include any number of slices.

[0106] Processor complex 1040 may be a graphics complex that can be configured to perform computational operations (e.g., the computational operations described herein) in a highly parallel manner. Processor complex 1040 may be configured to perform graphics pipeline operations, such as, but not limited to, drawing commands, pixel operations, geometric calculations, and other operations associated with rendering an image to a display. Processor complex 1040 may be configured to perform graphics-independent operations, such as, but not limited to, neural network training and / or simulation. Processor complex 1040 may be configured to perform both graphics-related and graphics-independent operations.

[0107] The processor complex 1040 may include any number of compute units 1050(1)-1050(N) (where N is any integer greater than 1) and an L2 cache 1042. The compute units 1050 may share the L2 cache 1042, which may store information that will be used to perform some or all of the operations described herein. The L2 cache 1042 may be partitioned. The processor complex 1040 may include any number of compute units 1050 and any number (including zero) and type of cache. The processor complex 1040 may include any number of dedicated graphics hardware.

[0108] Each compute unit 1050 may include any number of SIMD units 1052(1)-1052(N) (where N is any integer greater than 1) and shared memory 1054. Each SIMD unit 1052 may implement a SIMD architecture and may be configured to perform some or all of the operations described herein in parallel. Each compute unit 1050 may execute any number of thread blocks, but each thread block may execute on a single compute unit 1050, although in some embodiments, the thread block may execute on multiple compute units. A thread block may include any number of execution threads. A workgroup may be a thread block. Each SIMD unit 1052 may execute a set of threads. A set of threads (e.g., 16 threads), also referred to as a warp, subgroup, or wavefront (e.g., used by AMD and Intel), may belong to a single thread block and be configured to process different datasets based on a single instruction set. Prediction may be used to disable one or more threads in a warp, subgroup, or wavefront. A lane may be a thread. A work item can be a thread, such as (but not limited to) an OpenCL thread. Different thread bundles, subgroups, or wavefronts within a thread block can be synchronized together and communicate via shared memory 1054. Each compute unit 1050 can include one or more thread block clusters, where thread block clusters can implement programmable control over locality at a larger granularity than a single thread block in a single streaming multiprocessor (SM). Thread block clusters (also referred to as “clusters”) can support multiple thread blocks running concurrently across streaming multiprocessors, thereby synchronously and cooperatively acquiring, exchanging, or otherwise using data. In at least one embodiment, a streaming multiprocessor (“SM”) can refer to a streaming microprocessor, a streaming processor (“SP”), a streaming processing unit (“SPU”), a compute unit (“CU”), an execution unit (“EU”), and / or a slice, where a slice in this context can refer to a portion of the processing resources within a processing unit (e.g., 16 cores, a ray tracing unit, a thread bootstrap, or a scheduler).

[0109] Structure 1060 may be a system interconnect that facilitates data and control transfers across processor complex 1010, processor complex 1040, I / O interface 1070, memory controller 1080, display controller 1092, and multimedia engine 1094, for example, to perform some or all of the operations described herein. SOC 1000 may include any number and type of system interconnects other than or replacing structure 1060, facilitating data and control transfers across any number and type of directly or indirectly linked components within or outside SOC 1000. I / O interface 1070 may represent any number and type of I / O interfaces (e.g., PCI, PCI extensions (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). Various types of peripheral devices may be coupled to I / O interface 1070. Peripheral devices that may be coupled to I / O interface 1070 may include keyboards, mice, printers, scanners, joysticks or other types of game controllers, media recording devices, external storage devices, network interface cards, etc.

[0110] Display controller 1092 can display images on one or more display devices, such as, but not limited to, liquid crystal displays (“LCD”) devices. Multimedia engine 1094 can include any number and type of circuitry related to multimedia, such as, but not limited to, video decoders, video encoders, image signal processors, etc. Memory controller 1080 can facilitate data transfer between SOC 1000 and unified system memory 1090. Processor complex 1010 and processor complex 1040 can share unified system memory 1090. Unified system memory 1090 can include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as, but not limited to, synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory. Unified system memory 1090 can include 3D stacked memory, including but not limited to high bandwidth memory (HBM), HBM2e, or HDM3.

[0111] The SOC 1000 can implement a memory subsystem including any number and type of memory controllers 1080 and memory devices (e.g., shared memory 1054), which may be dedicated to a single component or shared among multiple components to perform any of the operations described herein. The SOC 1000 can implement a cache subsystem including one or more cache memories (e.g., L2 cache 1028, L3 cache 1030, and L2 cache 1042), each cache memory may be dedicated to any number of components (e.g., core 1020, core complex 1010, SIMD unit 1052, compute unit 1050, and processor complex 1040), or may be shared among any number of components (e.g., core 1020, core complex 1010, SIMD unit 1052, compute unit 1050, and processor complex 1040).

[0112] In at least one embodiment, the SOC 1000 may include one or more circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API that indicate the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. One or more circuitry may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API that indicate the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0113] Figure 11A illustrates a parallel processor 1100 according to at least one embodiment. The parallel processor 1100 may be implemented using one or more circuits and may be referred to as a programmable processor (e.g., CPU and / or GPU), logic, application-specific integrated circuit (ASIC), field-programmable gate array (FPGA), or other hardware (e.g., embodiments in Figures 10-22B) for performing any of the operations described above or elsewhere herein.

[0114] Parallel processor 1100 may include parallel processing unit 1102 for performing any of the operations described above or elsewhere herein. Parallel processing unit 1102 may include I / O unit 1104 that enables communication with other devices, including other instances of parallel processing unit 1102. I / O unit 1104 may be directly connected to other devices. I / O unit 1104 may be connected to other devices via the use of a hub or switch interface, such as, but not limited to, memory hub 1105. The connection between memory hub 1105 and I / O unit 1104 may form a communication link 1113. I / O unit 1104 may be connected to host interface 1106 and memory crossbar switch 1116, wherein host interface 1106 receives commands directed to perform processing operations, and memory crossbar switch 1116 receives commands directed to perform memory operations.

[0115] When host interface 1106 receives a command buffer via I / O unit 1104, host interface 1106 can route the work operations that execute these commands to front-end 1108. Front-end 1108 can be coupled to scheduler 1110 (which may be referred to as sequencer), which is configured to distribute commands or other work items to processing cluster array 1112. Scheduler 1110 can ensure that processing cluster array 1112 is correctly configured and in an active state before tasks are distributed to the cluster of processing cluster array 1112. Scheduler 1110 can be implemented via firmware logic executed on a microcontroller. The microcontroller-implemented scheduler 1110 can be configured to perform complex scheduling and work distribution operations at both coarse and fine granular levels, thereby enabling fast preemption and context switching of threads executing on processing array 1112. Host software can validate workloads scheduled on processing cluster array 1112 via one of multiple graphics processing paths. The workload can then be automatically distributed to the processing array cluster 1112 by the scheduler 1110 logic within the microcontroller, which includes the scheduler 1110.

[0116] Processing cluster array 1112 can perform any of the operations described above or elsewhere herein, and may include up to “N” processing clusters (e.g., clusters 1114A, 1114B through 1114N), where “N” represents a positive integer (which may be a different integer “N” than used in other diagrams). Each cluster 1114A-1114N in processing cluster array 1112 can execute a large number of concurrent threads. Scheduler 1110 may use various scheduling and / or work distribution algorithms to distribute work to clusters 1114A-1114N in processing cluster array 1112, which may vary depending on the workload generated by each type of program or computation. Scheduling may be dynamically handled by scheduler 1110 or may be assisted by compiler logic during the compilation of program logic configured to be executed by processing cluster array 1112. Different clusters 1114A-1114N of processing cluster array 1112 may be assigned to process different types of programs or perform different types of computations.

[0117] The processing cluster array 1112 can be configured to perform various types of parallel processing operations, such as, but not limited to, any of the operations described above or elsewhere herein. The processing cluster array 1112 can be configured to perform general-purpose parallel computing operations. For example, the processing cluster array 1112 may include logic for performing processing tasks, including filtering video and / or audio data, performing modeling operations (including physical operations), and performing data transformations.

[0118] Processing cluster array 1112 can be configured to perform parallel graphics processing operations. Processing cluster array 1112 may include additional logic for supporting the execution of such graphics processing operations, including but not limited to texture sampling logic for performing texture operations, as well as tessellation logic and other vertex processing logic. Processing cluster array 1112 can be configured to execute shader programs related to graphics processing, such as, but not limited to, vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. Parallel processing unit 1102 can transfer data from system memory via I / O unit 1104 for processing. During processing, the transferred data may be stored in on-chip memory (e.g., parallel processor memory 1122) during processing and then written back to system memory.

[0119] When the parallel processing unit 1102 is used to perform graphics processing, the scheduler 1110 can be configured to divide the processing workload into tasks of approximately equal size to better distribute graphics processing operations to multiple clusters 1114A-1114N of the processing cluster array 1112. Each part of the processing cluster array 1112 can be configured to perform different types of processing. For example, a first part can be configured to perform vertex shading and topology generation, a second part can be configured to perform tessellation and geometry shading, and a third part can be configured to perform pixel shading or other screen-space operations to produce a rendered image for display. Intermediate data generated by one or more clusters 1114A-1114N can be stored in a buffer to allow intermediate data to be transferred between clusters 1114A-1114N for further processing.

[0120] Processing cluster array 1112 can receive processing tasks to be executed via scheduler 1110, which receives commands defining the processing tasks from front end 1108. Processing tasks may include indexes of data to be processed, such as surface (patch) data, primitive data, vertex data, and / or pixel data, as well as state parameters and commands defining how to process the data (e.g., which program to execute). Scheduler 1110 can be configured to retrieve the index corresponding to the task, or can receive the index from front end 1108. Front end 1108 can be configured to ensure that processing cluster array 1112 is configured to be active before the workload specified by the incoming command buffer (e.g., batch buffer, push buffer, etc.) is initiated.

[0121] Each instance of one or more instances of parallel processing unit 1102 may be coupled to parallel processor memory 1122 to perform any of the operations described above or elsewhere herein. Parallel processor memory 1122 may be accessed via memory crossbar switch 1116, which may receive memory requests from processing cluster array 1112 and I / O unit 1104. Memory crossbar switch 1116 may access parallel processor memory 1122 via memory interface 1118. Memory interface 1118 may include multiple partition units (e.g., partition units 1120A, 1120B through 1120N), each partition unit may be coupled to a portion (e.g., a memory cell) of parallel processor memory 1122. The number of partition units 1120A-1120N can be configured to be equal to the number of memory units, such that the first partition unit 1120A has a corresponding first memory unit 1124A, the second partition unit 1120B has a corresponding memory unit 1124B, and the Nth partition unit 1120N has a corresponding Nth memory unit 1124N. The number of partition units 1120A-1120N may not be equal to the number of memory units.

[0122] Memory cells 1124A-1124N may include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as, but not limited to, synchronous graphics random access memory (SGRAM), which includes graphics double data rate (GDDR) memory. Memory cells 1124A-1124N may also include 3D stacked memory, including but not limited to high-bandwidth memory (HBM), HBM2e, or HDM3. Render targets (e.g., but not limited to framebuffers or texture maps) may be stored in memory cells 1124A-1124N, allowing partitioning cells 1120A-1120N to write portions of each render target in parallel to efficiently utilize the available bandwidth of the parallel processor memory 1122. A local instance of the parallel processor memory 1122 may not be included to support a unified memory design that combines system memory with local cache memory.

[0123] Any cluster 1114A-1114N in the processing cluster array 1112 can process data to be written to any memory cell 1124A-1124N within the parallel processor memory 1122. The memory crossbar switch 1116 can be configured to transfer the output of each cluster 1114A-1114N to any partition cell 1120A-1120N, or to another cluster 1114A-1114N where additional processing operations can be performed on the output. Each cluster 1114A-1114N can communicate with the memory interface 1118 via the memory crossbar switch 1116 to read from or write to various external memory devices. The memory crossbar switch 1116 can be connected to the memory interface 1118 to communicate with the I / O unit 1104, or to a local instance of the parallel processor memory 1122, enabling processing units within different processing clusters 1114A-1114N to communicate with system memory or other memory local to the non-parallel processing unit 1102. The memory crossbar switch 1116 can use virtual channels to separate traffic flows between clusters 1114A-1114N and partition units 1120A-1120N.

[0124] Multiple instances of the parallel processing unit 1102 can be mounted on a single add-in card, or multiple add-in cards can be interconnected. Even if different instances of the parallel processing unit 1102 have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences, these different instances can be configured to interoperate. For example, some instances of the parallel processing unit 1102 may include higher-precision floating-point units relative to other instances. Systems including one or more instances of the parallel processing unit 1102 or the parallel processor 1100 can be implemented in a variety of configurations and form factors, including but not limited to desktop, laptop, or handheld personal computers, servers, workstations, game consoles, and / or embedded systems.

[0125] Figure 11A also includes a block diagram of a partitioning unit 1120 according to at least one embodiment. Partitioning unit 1120 is an example of one of partitioning units 1120A-1120N in Figure 11A. Partitioning unit 1120 may include an L2 cache 1121, a frame buffer interface 1125, and a ROP 1126 (raster operation unit). The L2 cache 1121 may be a read / write cache configured to perform load and store operations received from memory crossbar switch 1116 and ROP 1126. Read misses and urgent write-back requests may be output from the L2 cache 1121 to the frame buffer interface 1125 for processing. Updates may also be sent to the frame buffer via the frame buffer interface 1125 for processing. The frame buffer interface 1125 may engage with one of the memory cells in the parallel processor memory, such as, but not limited to, memory cells 1124A-1124N in Figure 11A (e.g., located within parallel processor memory 1122).

[0126] ROP 1126 can be a processing unit that performs raster operations, such as, but not limited to, stenciling, z-testing, blending, etc. ROP 1126 can then output processed graphics data stored in graphics memory. ROP 1126 may include compression logic for compressing depth or color data written to memory and decompressing depth or color data read from memory. The compression logic can be lossless compression logic that utilizes one or more compression algorithms. The type of compression performed by ROP 1126 can vary based on the statistical characteristics of the data to be compressed. For example, incremental color compression is performed on depth and color data on a per-tile basis.

[0127] ROP 1126 can be included within each processing cluster (e.g., clusters 1114A-1114N in Figure 11A) instead of within partition unit 1120. Read and write requests for pixel data (instead of pixel fragment data) can be transferred via memory crossbar switch 1116. Processed graphics data can be displayed on a monitor and routed for further processing by a processor, or routed for further processing by one of the processing entities within parallel processor 1100 in Figure 11A.

[0128] In at least one embodiment, the parallel processor 1100 may include one or more circuitry configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuitry may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0129] Figure 11B includes a block diagram of a processing cluster 1114 within a parallel processing unit according to at least one embodiment. The processing cluster may be an instance of one of the processing clusters 1114A-1114N in Figure 11A, which can be used to perform any of the operations described above or elsewhere herein. The processing cluster 1114 can be configured to execute a number of threads in parallel, where a “thread” refers to an instance of a specific program executed on a specific input dataset. Single Instruction Multiple Data (SIMD) instruction issuing techniques can be used to support the parallel execution of a large number of threads without providing multiple independent instruction units. Single Instruction Multiple Threading (SIMT) techniques can be used to support the parallel execution of a large number of typically synchronous threads using a common instruction unit configured to issue instructions to a set of processing engines within each processing cluster.

[0130] The operation of processing cluster 1114 can be controlled via pipeline manager 1132, which distributes processing tasks to SIMT parallel processors. Pipeline manager 1132 can receive instructions from scheduler 1110 in FIG11A and manage the execution of these instructions via graphics multiprocessor 1134 and / or texture unit 1136. Graphics multiprocessor 1134 can be an example instance of a SIMT parallel processor. However, processing cluster 1114 can include various types of SIMT parallel processors with different architectures. Processing cluster 1114 can include one or more instances of graphics multiprocessor 1134. Graphics multiprocessor 1134 can process data and can distribute the processed data to one of several possible destinations, including other shader units, using data cross switch 1140. Pipeline manager 1132 can facilitate the distribution of processed data by specifying the destination of the processed data to be distributed via data cross switch 1140.

[0131] Each graphics multiprocessor 1134 within the processing cluster 1114 may include a set of identical functional execution logic (e.g., arithmetic logic units, load-memory units, etc.) for performing computations for any of the operations described above or elsewhere herein. The functional execution logic may be configured in a pipelined manner, where new instructions can be issued before previous instructions complete. The functional execution logic may support a wide range of operations, including integer and floating-point arithmetic, comparison operations, Boolean operations, bit shifting, and computation of various algebraic functions. Different operations can be performed using the same functional unit hardware, and arbitrary combinations of functional units are possible.

[0132] Instructions transmitted to the processing cluster 1114 can form threads, which may also be called thread bundles, subgroups, waves, or wavefronts. A group of threads executing across a set of parallel processing engines can be called a thread group. Thread groups can execute a common program on different input data. Each thread within a thread group can be assigned to a different processing engine within the graphics multiprocessor 1134. The number of threads in a thread group can be less than the number of processing engines within the graphics multiprocessor 1134. When the number of threads in a thread group is less than the number of processing engines, one or more processing engines may be idle during the processing cycle of that thread group. The number of threads in a thread group can also be more than the number of processing engines within the graphics multiprocessor 1134. When the number of threads in a thread group is more than the number of processing engines within the graphics multiprocessor 1134, processing can be performed in consecutive clock cycles. Multiple thread groups can execute concurrently on the graphics multiprocessor 1134.

[0133] The graphics multiprocessor 1134 includes an internal cache memory for performing load and store operations, such as, but not limited to, any of the operations described above or elsewhere herein. The graphics multiprocessor 1134 may forgo the internal cache and instead use a cache memory within the processing cluster 1114 (e.g., L1 cache 1148). Each graphics multiprocessor 1134 may also access an L2 cache within a partition unit (e.g., partition units 1120A-1120N in FIG. 11A) that can be shared across all processing clusters 1114 and can be used to transfer data between threads. The graphics multiprocessor 1134 may also access off-chip global memory, which may include one or more of local parallel processor memory and / or system memory. Any memory outside of the parallel processing unit 1102 may be used as global memory. The processing cluster 1114 may include multiple instances of the graphics multiprocessor 1134 and may share common instructions and data, which may be stored in the L1 cache 1148.

[0134] Each processing cluster 1114 may include an MMU 1145 (Memory Management Unit) that can be configured to map virtual addresses to physical addresses. One or more instances of the MMU 1145 may reside within the memory interface 1118 of Figure 11A. The MMU 1145 may include: a set of page table entries (PTEs) for mapping virtual addresses to physical addresses of tiles; and optional cache line indexes. The MMU 1145 may include address translation back buffers (TLBs) or caches that may reside within the graphics multiprocessor 1134 or the L1 1148 cache or within the processing cluster 1114. Physical addresses can be processed to distribute surface data access locally, allowing for efficient request interleaving between partition units. The cache line indexes can be used to determine whether a request for a cache line is a hit or a miss.

[0135] Processing cluster 1114 can be configured such that each graphics multiprocessor 1134 is coupled to a texture unit 1136 to perform texture mapping operations, such as determining texture sample locations, reading texture data, and filtering texture data. Texture data can be read from an internal texture L1 cache (not shown) or an L1 cache within the graphics multiprocessor 1134, and can be retrieved as needed from an L2 cache, local parallel processor memory, or system memory. Each graphics multiprocessor 1134 can output a processed task to a data crossbar switch 1140 to provide the processed task to another processing cluster 1114 for further processing, or store the processed task in an L2 cache, local parallel processor memory, or system memory via a memory crossbar switch 1116. Pre-Raster Operation Unit 1142 can be configured to receive data from the graphics multiprocessor 1134 and direct the data to ROP units, which can be located together with partitioning units described herein (e.g., partitioning units 1120A-1120N in FIG. 11A). The PreROP 1142 unit can perform color blending optimization, organize pixel color data, and perform address translation.

[0136] In at least one embodiment, the processing cluster 1114 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0137] Figure 11C illustrates a graphics multiprocessor 1134 according to at least one embodiment, for example, to perform any of the operations described above or elsewhere herein. The graphics multiprocessor 1134 may be coupled to a pipeline manager 1132 of a processing cluster 1114. The graphics multiprocessor 1134 may include an execution pipeline including, but not limited to, an instruction cache 1152 (e.g., which may store instructions, such as, but not limited to, compiled API instructions), instruction units 1154, an address mapping unit 1156, a register file 1158, one or more general-purpose graphics processing unit (GPGPU) cores 1162, and one or more load / store units 1166, one or more of which may perform load / store operations to load / store instructions corresponding to the execution operation. The GPGPU cores 1162 and load / store units 1166 may be coupled to cache memory 1172 and shared memory 1170 via a memory and cache interconnect 1168. The GPGPU cores 1162 may be part of a SoC, such as, but not limited to, the integrated circuit 1000 in Figure 10.

[0138] Instruction cache 1152 can receive a stream of instructions to be executed from pipeline manager 1132 (e.g., to perform any of the operations described above or elsewhere herein). Instructions can be cached in instruction cache 1152 and dispatched for execution by instruction unit 1154. Instruction unit 1154 can dispatch instructions into thread groups (e.g., thread bundles, subgroups, wavefronts, or waves), with each thread in the thread group assigned to a different execution unit within GPGPU core 1162. Instructions can access any of the local, shared, or global address spaces by specifying an address within a unified address space. Address mapping unit 1156 can be used to translate addresses in the unified address space into different memory addresses accessible by load / store unit 1166.

[0139] Register file 1158 provides a set of registers for the functional units of graphics multiprocessor 1134. Register file 1158 provides temporary storage for operands on data paths connected to functional units of graphics multiprocessor 1134 (e.g., GPGPU core 1162, load / store unit 1166). Register file 1158 can be partitioned among the functional units such that each functional unit is allocated a dedicated portion of register file 1158. Register file 1158 can be partitioned among different thread bundles (which may be referred to as wavefronts, subgroups, and / or waves or threads) executed by graphics multiprocessor 1134.

[0140] Each GPGPU core 1162 may include a floating-point unit (FPU) and / or an integer arithmetic logic unit (ALU) for executing instructions of the graphics multiprocessor 1134. The architectures of the GPGPU cores 1162 may be similar or different. A first part of the GPGPU core 1162 may include a single-precision FPU and an integer ALU, while a second part of the GPGPU core may include a double-precision FPU. The FPU may implement IEEE 754-2008 standard floating-point arithmetic or enable variable-precision floating-point arithmetic. The graphics multiprocessor 1134 may also include one or more fixed-function or special-function units for performing specific functions, such as, but not limited to, copying rectangles or pixel blending operations. One or more of the GPGPU cores 1162 may also include fixed-function or special-function logic.

[0141] The GPGPU core 1162 may include SIMD logic capable of executing a single instruction on multiple sets of data. The GPGPU core 1162 can physically execute SIMD4, SIMD8, and SIMD16 instructions, and logically execute SIMD1, SIMD2, and SIMD32 instructions. The SIMD instructions for the GPGPU core may be generated by the shader compiler at compile time, or may be automatically generated when executing programs written and compiled for Single Program Multiple Data (SPMD) or SIMT architectures. Multiple threads of the program can be configured for a SIMT execution model that can be executed via a single SIMD instruction. For example, eight SIMT threads performing the same or similar operations can be executed in parallel via a single SIMD8 logic unit.

[0142] The memory and cache interconnect 1168 may include an interconnect network that connects each functional unit of the graphics multiprocessor 1134 to the register file 1158 and shared memory 1170. The memory and cache interconnect 1168 may be a cross-switch interconnect that allows the load / store unit 1166 to perform load and store operations between the shared memory 1170 and the register file 1158. The register file 1158 may operate at the same frequency as the GPGPU core 1162, thus data transfer between the GPGPU core 1162 and the register file 1158 can have very low latency. The shared memory 1170 can be used to implement communication between threads executing on functional units within the graphics multiprocessor 1134. The cache memory 1172 can be used as a data cache, for example, for caching texture data transferred between functional units and texture units 1136. The shared memory 1170 can also be used as a program-managed cache. In addition to automatically caching the data stored in the cache memory 1172, threads executing on the GPGPU core 1162 can also programmatically store data in shared memory.

[0143] The parallel processor or GPGPU described herein can be communicatively coupled to a host / processor core to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general-purpose GPU (GPGPU) functions. The GPU can be communicatively coupled to the host processor / core via a bus or other interconnect (e.g., high-speed interconnects, such as, but not limited to, PCIe or NVLink). A System-on-a-Chip (SoC) may include the parallel processor or GPGPU described herein, which executes on the SoC. The GPU may be integrated as a core on a package or chip and communicatively coupled to the core via an internal processor bus / interconnect within the package or chip. Regardless of the GPU's connection method, the processor core can assign work to the GPU in the form of a sequence of commands / instructions contained in a job descriptor. The GPU can then use dedicated circuitry / logic to efficiently process these commands / instructions to perform any of the operations described above or elsewhere herein.

[0144] In at least one embodiment, the graphics multiprocessor 1134 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0145] Figure 12 illustrates a processor 1200 according to at least one embodiment. The processor 1200 may include a hybrid architecture processor (e.g., Lunar Lake or Meteor Lake) from Intel Corporation, Santa Clara, California, or other processors sharing at least some of the components described herein. The processor 1200 may include one or more central processing units (CPU 1202), one or more graphics processing units (GPU 1206), and / or one or more neural processing units (NPU 1208), which may be, for example, dedicated AI accelerators for offloading artificial intelligence (AI) workloads from the CPU 1202 and GPU 1206. The processor 1200 may use instructions that, if executed, cause the processor 1200 and / or any of its components to perform some or all of the processes and techniques described elsewhere herein. The processor 1200 may include any number of memory and cache units 1210 for facilitating processing between different components of the processor 1200. The memory and cache 1210 on processor 1200 may include one or more levels of cache (e.g., L1, L2, L3, and / or last-level cache) and high-bandwidth memory (e.g., HBM2e or HBM3) in any combination. Regarding processor 1200 and any components described above or elsewhere herein, one or more APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) internal or external to processor 1200. The results of APIs may be stored in storage devices internal or external to processor 1200, including registers, DRAM, flash memory, SRAM, cache, or other memory. One or more APIs described herein may include calls.

[0146] Processor 1200 may include a computing engine as CPU 1202, and may include any number of cores, such as, but not limited to, up to 16 cores / 22 threads. The cores in CPU 1202 may include P-cores (performance), E-cores (high efficiency), and LP-E cores (low-power, high-efficiency). Performance cores can be used for low-latency, single-threaded, computationally intensive workloads, while high-efficiency cores can be used for multi-threaded, less computationally intensive workloads. Low-power, high-efficiency cores can be used for scalable multi-threaded execution and offloading background tasks. P-cores can be used for single-threaded and limited-threaded execution, while E-cores and LP-E cores are used for multi-threaded throughput and power efficiency.

[0147] The GPU 1206 can include any number of graphics engines, such as, but not limited to, those with 8 Xe cores (up to 128 execution units or EUs). Arc TM Graphics Engine (Xe LPG). As shown in Figure 12, GPU 1206 may include a vector engine 1210 and a matrix engine 1212, which can, for example, run FP, INT, and matrix operation tasks simultaneously, individually, or in batches. GPU 1206 may include a load / store unit 1214, as well as other memories, such as, but not limited to, an instruction cache (I$) 1216 and an L1 cache / subsystem local memory (SLM) 1218, which may, for example, store instructions for performing any of the operations described above or elsewhere herein.

[0148] The NPU 1204 may include one or more AI Boost integrates a Neural Processing Unit (NPU). The NPU 1204 can be enumerated as an integrated PCIe device to the host processor. The NPU 1204 may include one or more (e.g., two) Neural Computation Engine (NCE) tiles 1230. Each tile may be configured with any combination of, but not limited to, the following: (e.g., 2000) Multiply-Accumulate (MAC) engines 1234, a post-processing engine (not shown), an AEP processor (not shown), and memory per tile (2MB dedicated SRAM), as shown in Figure 12. For general computing needs, the Neural Computation Engine 1230 may include a streaming hybrid architecture vector engine (SHAVE) 1228 for high-performance parallel computing, which may include a DMA (Direct Memory Access) engine 1224 for transporting data between system memory DRAM (Dynamic Random Access Memory) 1226 and a software-managed cache. The built-in device MMU (Memory Management Unit) 1222, along with the IOMMU (Input-Output Memory Management Unit) (not shown), can support multiple concurrent hardware contexts and provide secure isolation between execution contexts according to the MCDM (Microsoft Computation Driver Model) architecture. The processor 1200 may also include a media unit (not shown), which may be included on or separate from the XCD or other components of the processor 1200 to enable video playback and processing of compressed or uncompressed data, such as hardware-accelerated decoding support for HEVC, AV1, VP9, ​​and AVC, as well as hardware-accelerated encoding support for HEVC, VP9, ​​and AVC.

[0149] Thread bootstrap ( The Thread Director (which includes firmware built into the processor 1200) can prioritize and manage the distribution of workloads, thereby sending tasks to optimized cores. For example, the Thread Director can tie P cores, E cores, and / or LP-E cores (as described above) together with task scheduling capabilities and the ability to send less demanding tasks to E cores or LP-E cores. Deep learning acceleration ( DLBoost (not shown) can provide built-in AI acceleration for training and inference workloads and may include support for VNNI (for CPU) and DP4a (for GPU) instruction sets. This instruction set can be used with OpenVINO. TM The toolkit and oneAPI are optimized to accelerate INT8 inference. For example, the software stack described elsewhere in this document can be used to leverage OpenVINO. TM The toolkit enables AI inference. The processor 1200 can be configured to execute applications, such as, but not limited to, CUDA programs.

[0150] In at least one embodiment, the processor 1200 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0151] Processor 1200 may alternatively include a processor based on Qualcomm's AIEngine Direct architecture from Santa Clara, California, or other processors sharing at least some of the components described herein. It may include any number of NPUs, GPUs, CPUs, and other associated components, such as, but not limited to, an NPU 1204 as a Hexagon NPU, a GPU 1206 as an Adreno GPU, a CPU 1202 as a Kryo or Qualcomm Oryon CPU, and a Qualcomm Sensing Hub (not shown) and a memory subsystem 1210, in any combination. Hexagon NPU 1204 may include power rails, micro-tile inference units, hardware acceleration units, tensor units, scalar units, and vector units (all not shown), which may have dedicated or shared memory (e.g., cache or memory, such as HBM3) for storing, for example, instructions for performing any of the operations described above or elsewhere herein. The Adreno GPU 1206 can provide graphics and parallel processing for AI, in formats including but not limited to 32-bit floating-point (FP32), 16-bit floating-point (FP16), and 8-bit integer (INT8). The Kryo or Qualcomm Oryon CPU 1202 can execute AI workloads and handle the contextualization of ubiquitous generative AI applications. The CPU 1202 may also include an instruction fetch unit, a renaming and deprecation unit, a memory management unit, a vector execution unit, an integer execution unit, and a load and store unit for processing and instruction management. Regarding the processor 1200 and any of its components described above or elsewhere herein, one or more APIs described herein can, for example, be compiled into instructions that can be fetched by the instruction fetch unit, decoded by the processor decoder or equivalent, scheduled (e.g., sequentially or out of order) for execution by the scheduler or equivalent, executed by execution logic or equivalent, reordered, and then deprecated by the renaming and deprecation unit. The API (and / or compiled instructions including the API) can be stored in any storage device (e.g., cache and / or memory) inside or outside the processor 1200. An arbitrary number of CPU cores 1202 can be included in an arbitrary number of CPU clusters, which can be coupled to memory and / or cache, such as, but not limited to, a shared L2 cache. Memory can be separate or shared; for example, the CPU clusters of CPU cores 1202 can be coupled to a memory subsystem 1210, which can include structures capable of reading and writing to memory (e.g., DRAM), system-level caches, and an arbitrary number of memory management units.The Qualcomm sensing hub (not shown) includes a miniature NPU, power rails, and conventional sensors (such as gyroscopes, accelerometers, or even barometers) that support voice and data streaming. The memory subsystem 1210 may include memory and cache on the processor 1200, which may include L1 or more levels of cache (e.g., L1, L2, L3, and / or last-level cache) and high-bandwidth memory (e.g., HBM2e or HBM3) in any combination, for example, for storing information and / or instructions for performing any of the operations described above or elsewhere herein. All or part of the memory and / or cache in the memory subsystem 1210 may be shared or used individually by any component or combination of components on the processor 1200 (e.g., GPU 1206, NPU 1204, and CPU 1202).

[0152] The Qualcomm AI Engine 1200 can be programmed and controlled using a software stack to perform some or all of the operations described herein, including, for example... A neural processing SDK is provided for inference on Android, Linux, and Windows. Developer libraries and services support programming languages, virtual platforms, and compilers. At lower levels of the software stack, system software includes a basic real-time operating system (RTOS), system interfaces, and drivers. The software stack supports various operating systems, including Android, Windows, Linux, and QNX, as well as deployment and monitoring infrastructures such as Prometheus, Kubernetes, and Docker. OpenCL and DirectML are supported for direct cross-platform access to the GPU 1206. For the CPU 1202, LLVM compiler infrastructure optimizations enable accelerated and efficient AI inference. Regarding the Qualcomm AI Engine 1200 and any of its components described above or elsewhere herein, one or more APIs described herein can, for example, be compiled into instructions that can be fetched by instruction fetching logic or equivalents, decoded by processor decoders or equivalents, scheduled (e.g., sequentially or out of order) for execution by a scheduler or equivalent, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. The API (and / or compiled instructions including the API) can be stored in any storage device (e.g., cache and / or memory) inside or outside the Qualcomm AI Engine 1200. The results of the API can be stored in storage devices inside or outside the Qualcomm AI Engine 1200, including registers, DRAM, flash memory, SRAM, cache, or other memory.

[0153] In at least one embodiment, the processor 1200 or the Qualcomm AI engine 1200 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0154] Figure 13A illustrates a processor 1300 according to at least one embodiment. Processor 1300 may include a Scalable family of processors from Intel Corporation, Santa Clara, California, or other processors that share at least some of the components described herein. Processor 1300 may include one or more cores 1312(1)-1312(N) capable of performing the operations described elsewhere herein, where N is any integer greater than 1. Cores 1312(1)-1312(N) may be interconnected using ring and / or mesh interconnects. Utilizing a mesh interconnect architecture, arrays of vertical and horizontal communication paths may allow traversal from one core to another 1312(1)-1312(N) via the shortest path (jumping to the correct row along the vertical path and to the correct column along the horizontal path). For mesh interconnects, dies may accommodate cores 1312(1)-1312(N) and may include aggregated mesh stop point (CMS) grids that may be associated with cores 1312(1)-1312(N) (e.g., 1:1). Each core can be associated with a low-level cache (LLC) slice 1314(1)-1314(N), or cores 1312(1)-1312(N) can share a cache, such as a low-level cache. LLC 1314(1)-1314(N) can be inclusive or non-inclusive (having blocks not present in the high-level cache) by merging blocks from a higher-level cache (e.g., L2 cache). Each core and LLC slice can include a caching and home agent (CHA) (not shown), which can maintain cache consistency by providing resource scalability via the mesh interconnect. Super Path Interconnect ( UPI 1316 provides cache coherency capabilities. UPI 1316 enables coherent interconnects for scalable systems and allows multiple processors to share a single shared address space via links, such as, but not limited to, two or three UPI links per processor.

[0155] The processor 1300 may also include a system agent 1310, which may house and / or perform various functions, such as, but not limited to, memory management, display functions, and / or input / output (I / O) functions. For example, the processor 1300 may include one or more integrated memory controllers (IMCs) 1308. IMCs 1308 may control and manage memory, such as, but not limited to, different memory types, such as DDR RAM, such as DDR4, or other memory described elsewhere herein. The system agent 1310 may include a display controller (not shown) for supporting one or more displays. The system agent 1310 may also integrate a PCIe 1304 (e.g., up to 20 PCIe lanes), which may, for example, be connected to an external dedicated graphics connector via a DMI bus (e.g., Intel's DMI 3.0 bus) 1306. The system agent 1310 may include an image processing unit (IPU) (not shown) that integrates an on-die image signal processor (ISP). Structure 1302 provides scalability for connecting to other nodes (e.g., processors, such as processor 1300) and can, for example, be connected to Cornelis Networks (…). It can be used together with elements of a scalable system framework that provides performance for high-performance computing (HPC) workloads and the ability to scale to tens of thousands of nodes.

[0156] Figure 13B illustrates the components within a core 1312 according to at least one embodiment. The core 1312 may include a front-end 1318, a back-end or execution engine 1332, and a memory subsystem 1342. The front-end 1318 can provide operations (e.g., operations described elsewhere herein) to the execution engine 1332 by decoding instructions stored in memory. For example, the front-end 1318 may include micro-operation (μOps) cache paths and / or conventional paths, and a branch prediction unit 1320 that can determine path instructions. A conventional path of instructions might involve fetching variable-length (e.g., x86) instructions from an L1 instruction cache, queuing these instructions into an instruction queue 1324, and decoding the instructions into μOps that can be provided to the allocation queue 1328 using a decoder 1326. Alternatively, the μOps cache path may include a cache that includes decoded μOps (μOps 1330) that can be sent to the allocation queue 1328. The allocation queue 1328 can act as an interface between the front end 1318 and the execution engine 1332, and can provide instructions to the execution engine 1332. For example, one or more APIs described herein can be compiled into instructions that can be stored, processed, and executed by the front end 1318 and the execution engine 1332, and stored in the memory subsystem 1342.

[0157] Execution engine 1332 can receive micro-operations into reordering buffer 1334, which can register, rename, and deregister μOPs. μOPs can be sent from the reordering buffer to scheduler 1336, which can be connected to one or more different execution units 1338. Execution units 1338 can perform operations such as basic arithmetic logic unit (ALU) operations, multiplication, division, and / or more complex operations, such as, but not limited to, various vector operations. Scheduler 1336 can manage the queuing of μOPs for one or more execution units 1338 based on, for example, the operations that need to be performed.

[0158] The memory subsystem 1342 can handle load and store requests as well as sorting operations. For example, μOPs may be associated with memory accesses (e.g., load and store), and these μOPs can be sent through dedicated scheduler ports that can perform these memory operations. For example, store and load operations can be sent to load and store buffers 1344. The memory subsystem 1342 may also include shared or separate L1 data and instruction caches 1346, and an L2 cache 1348 that can be used and shared by the L1 data and instruction caches 1346. As described above with respect to Figure 13A, each core 1312 may be connected to a slice of a third-level cache (e.g., LLC 1314) that can be shared by all cores 1312.

[0159] In at least one embodiment, the processor 1300 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0160] Figure 14 illustrates an AI accelerator 1400 according to at least one embodiment. The processor 1400 may include a processor with an AI accelerator architecture manufactured by Intel Corporation, Santa Clara, California, or other processors sharing at least some of the components described herein. The AI ​​accelerator 1400 may use instructions that, if executed by the AI ​​accelerator 1400, cause the AI ​​accelerator 1400 to perform some or all of the processes and techniques described elsewhere herein. For example, with respect to the AI ​​accelerator 1400 and any components described above or elsewhere herein, one or more APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetching logic or equivalents, decoded by processor decoders or equivalents, scheduled (e.g., sequentially or out of order) for execution by a scheduler or equivalent, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. The APIs (and / or compiled instructions including the APIs) may be stored in any storage device internal or external to the AI ​​accelerator 1400 (e.g., in cache and / or memory). The results of the API can be stored in internal or external storage devices of the AI ​​accelerator 1400, including registers, DRAM, flash memory, SRAM, cache, or other memory. The AI ​​accelerator 1400 may include one or more compute dies, which may include homogeneous or heterogeneous processors. The compute dies may include one or more central processing units (CPUs), one or more graphics processing units (GPUs), or a combination of both.

[0161] In at least one embodiment, the computational die may include a computational engine for performing AI computations. In at least one embodiment, the computational die of the AI ​​accelerator 1400 may be split into any number (e.g., four) clusters, which may be referred to as DCORE (Deep Learning Core) 1406, and include any number of matrix multiplication engines (MME) 1408, tensor processor cores (TPC) 1410, and L2 cache 1414 in any combination. The MME 1408 may perform operations using matrix multiplication, such as fully connected layers, convolutions, and batch general matrix multiplication (GEMM). The MME 1408 may be equipped with a multiplication-accumulation unit (MAC) (not shown), which may perform general matrix multiplication (GEMM) operations, such as, but not limited to, AxB multiplication, which involves generating a tensor C [NxM] from two input tensors A [NxK] and B [KxN]. The MME 1408 may be programmed with array dimensions, positions, data types, and various operands. The MME 1408 can retrieve tensors A and B from memory and pull them into its streaming buffer for parallel matrix multiplication by the MAC. After completion, the MME 1408 can push tensor C back to memory. The TPC1410 may include any number of scalar units for performing scalar operations, any number of vector units for performing vector operations, any number of register files or local memory units (e.g., vector local memory), and load and store components for instructions, which may be coupled to memory or caches (e.g., HBM, L3 cache, and / or L2 cache) (all not shown). The TPC can support different types of parallel processing, such as Very Long Instruction Word (VLIW) Single Instruction Multiple Data (SIMD) data types such as, but not limited to, FP32, BF16, FP16, and FP8 (both E4M3 and E5M2), UINT32, INT32, UINT16, INT16, UINT8, and INT8 data types. Any number of computational dies can be interconnected. The interconnection of connectable computing dies can be via an intermediary bridge, which is transparent to the software, for example.

[0162] The memory on the AI ​​accelerator 1400 may include one or more levels of cache (e.g., L1, L2, L3, and / or last-level cache) and high-bandwidth memory (e.g., HBM2e or HBM3) in any combination. The memory and / or cache system may be unified or separate. The compute die of the AI ​​accelerator 1400 may include on-chip memory comprising one or more levels (e.g., two levels) of cache. On-chip SRAM or other memory described elsewhere herein may be used as a unified last-level cache (L3) or split into multiple slices of L2 cache accessible to the MME 1408 and TPC 1410 groups. Using on-chip memory as an L2 or L3 cache is entirely software-configurable, and the software can dynamically determine its optimal cache allocation based on I / O tensors. AI accelerator 1400 may include one or more memory management units (MMUs) 1422 for managing memory, such as allowing the AI ​​accelerator 1400 memory subsystem to run in virtual space when accessing VRAM.

[0163] AI accelerator 1400 may include a communication port (e.g., a PCIe Gen5 x16 port) 1402 for communicating with a host and scheduling and synchronization unit 1404. AI accelerator 1400 may include a media unit 1416, which may include any number or combination of media decoder engines (DECs) 1420 and rotation engines (ROTs) 1418. AI accelerator 1400 may include a network unit 1424, which may include any number or combination of network ports 1426 and an accompanying RDMA engine 1428, L2 cache, and memory (e.g., HBM2e or HBM3) stack. AI accelerator 1400 may include a programmable control path entity (not shown) for managing the parallel and efficient execution of the various engines. The control path may include a submission queue (SQ) that can be issued by the runtime system, a completion queue (CQ) that can be used for job completion reporting, a programmable scheduling mechanism that can be used for task scheduling, a programmable hardware synchronization mechanism or "synchronization manager (SM)" that can be used for hardware synchronization, and a programmable interrupt service mechanism or "interrupt manager (INTR)" that can pass asynchronous events to drivers.

[0164] AI Accelerator 1400 may include media decoding units supporting video formats such as, but not limited to, HEVC, Progressive H.264, SVC base layer, MVC, VP9, ​​JPEG, and Progressive JPEG. AI Accelerator 1400 may support post-processing of the decoded media stream, such as, but not limited to, image downsizing (image resizing), vertical and horizontal scaling at different scaling ratios, image enlargement, image cropping, bilinear scaling, and Lancos scaling. AI Accelerator 1400 may implement two post-processing channels per decoder unit, one for scalar (up and down) and the other solely for outputting the original image. AI Accelerator 1400 may include a hardware rotation engine that performs the following transformations on the input image: 2D rotation, 3D rotation, projection, image warping and de-warping, resampling of the input data at user-defined coordinates, and rescaling.

[0165] The RDMA 1428 based on converged Ethernet on the AI ​​accelerator 1400 enables scaling from a single node (i.e., from a single AI accelerator 1400 to hundreds or thousands of nodes or AI accelerators 1400). The network subsystem 1424 may include... The accelerator 1400 includes an In-Gigabit Ethernet Communication Library (IGCL), a master controller coordinating data movement, and a programmable scheduling mechanism that enables smooth engine activation while maintaining task dependencies. The accelerator network subsystem may include a Gigabit Ethernet NIC port 1426, a Layer 2 MAC (not shown), and an RDMA engine 1428. The AI ​​accelerator 1400 may include an aggregation engine for performing summation activities. All engines in the processor 1400 can run in parallel; for example, the MME 1408, TPC 1410, and NIC 1426 can all operate simultaneously. Dependencies may exist between operations running on different engines; for example, the output of one engine may be used as the input of another engine, and / or the MME, TPC, and NIC may be scheduled to run in parallel. When one engine completes its execution, another engine can be scheduled to begin working on the next operation (executed immediately after its input is ready).

[0166] The AI ​​accelerator 1400 can be operated and controlled using a software layer 1428, which may include low-level components such as, but not limited to, a graph compiler, an automatic kernel fusionist and pre-compiled kernel libraries, and integrations with the AI ​​ecosystem such as, but not limited to, PyTorch, DeepSpeed, Hugging Face, vLLM, Ray, etc., or as described elsewhere in this document regarding software and programming platforms. The software layer 1428 may include implementations of algorithms such as, but not limited to, paged attention, flash attention, etc. The software layer 1428 can generate optimized binary code that implements a given model topology, such as, but not limited to, performing operator fusion, data layout management, parallelization, pipeline and memory management, and graph-level optimization.

[0167] In at least one embodiment, the AI ​​accelerator 1400 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0168] This document describes a neuromorphic computing system employing a multi-core architecture, wherein each core houses computing elements including neurons, synapses with on-chip learning capabilities, and local memory for storing synaptic weights and routing tables. Figure 15 is a simplified block diagram 1500 illustrating at least a portion of an example of such a neuromorphic computing device 1505 according to at least one embodiment. The neuromorphic computing device 1505 may include a neuromorphic processor from Intel Corporation, Santa Clara, California, or other processors that include at least a portion of the components described herein. As shown in this example, the device 1505 may be equipped with a network 1510 consisting of multiple neural network cores interconnected by an on-device network, thereby potentially defining multiple distinct connections between the cores. For example, the device 1505 may provide a network 1510 of spiking neural network cores, each core capable of communicating via short packet pulse messages sent from one core to another through network channels. Each core (e.g., 1515) may have processing and memory resources, as well as logic, for implementing a number of primitive nonlinear time computation elements, such as, but not limited to, multiple (e.g., more than 1000) different artificial neurons (referred to herein as “neurons”). For example, each core may be able to implement multiple neurons concurrently, such that a neuromorphic core can implement many, many neurons using device 1505. Regarding neuromorphic computing device 1505 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) internal or external to neuromorphic computing device 1505. The results of the API can be stored in internal or external storage devices of the neuromorphic computing device 1505, including registers, DRAM, flash memory, SRAM, cache or other equivalent memory.

[0169] Continuing with the example in Figure 15, the neuromorphic computing device 1505 may further include a processor 1520 and a system memory 1525 for implementing one or more components to manage and provide the functionality of the neuromorphic computing device 1505. For example, a system manager 1530 may be provided to manage the global attributes and operations of the neuromorphic computing device 1505 (e.g., attributes affecting the core network 1510, multiple cores in network 1510, interconnection of the neuromorphic computing device 1505 with other devices, managing access to the global system memory 1525, and other potential examples). In one example, the system manager 1530 may manage the definition and configuration of specific routing tables for individual routers in network 1510, orchestration of network definitions and attributes to be applied to network 1510 (e.g., weights, attenuation rates, etc.), core synchronization and time multiplexing management, routing input to the appropriate core, and other potential functions.

[0170] As another example, the neuromorphic computing device 1505 may also include a programming interface 1535 through which a user or system can specify the neural network definition to be applied (e.g., via routing tables and individual neuron attributes), implemented by the neuromorphic core grid 1510. A software-based programming tool may be provided with or separately from the neuromorphic computing device 1505, through which a user can provide a definition for a specific neural network to be implemented using the neuromorphic core network 1510. The programming interface 1535 can receive input from a programmer, then generate the corresponding routing table and populate the specified parameters into the local memory of each neuromorphic core (e.g., 1515) to implement the corresponding custom artificial neural network implemented by the neuromorphic core 1515.

[0171] In certain circumstances, the neuromorphic computing device 1505 can advantageously engage and interoperate with other devices, including general-purpose computing devices, to enable specific applications and use cases. Therefore, in some cases, external interface logic 1540 may be provided to communicate with one or more other devices (e.g., via one or more defined communication protocols). External interface 1540 may be used to accept input data from another device or an external memory controller used as an input data source. External interface 1540 may additionally or alternatively be used to allow the results or outputs of computations performed using the neural network implemented using the neuromorphic computing device 1505 to be provided to another device (e.g., another general-purpose processor implementing machine learning algorithms) to enable additional applications and enhancements, among other examples.

[0172] Figure 15 illustrates a network 1510 of multiple neural network cores interconnected by a network on the device, showing a portion of a network structure interconnecting multiple neuromorphic cores (e.g., 1515a-d). For example, several neuromorphic cores (e.g., 1515a-d) may be provided in a mesh, each core interconnected via a network including multiple routers (e.g., 1550). In one implementation, each neuromorphic core (e.g., 1515a-d) may be connected to a single router (e.g., 1550) in the routers, and the router may be connected to at least one other router (as shown at 1510 in Figure 15). As an example, in a particular implementation, four neuromorphic cores (e.g., 1515a-d) may be connected to a single router (e.g., 1550), and each router 1550 may be connected to two or more other routers to form a multi-core mesh, thereby allowing each neuromorphic core to interconnect with each other neuromorphic core in the neuromorphic computing device 1505. Furthermore, since each neuromorphic core can be configured to implement multiple different neurons, the router network of the neuromorphic computing device 1505 can similarly implement connections or artificial synapses (or simply "synapses") defined between any two of the potential many (e.g., 30,000+) neurons defined using the neuromorphic core network 1510 provided in the neuromorphic computing device 1505.

[0173] Figure 15 shows a block diagram of the internal components of an example implementation of the illustrated neuromorphic core 1515. In one example, a single neuromorphic core may implement a number of neurons (e.g., 1024), which share the architectural resources of the neuromorphic core 1515 in a time-division multiplexing manner. In one example, each neuromorphic core 1515 may include a processor block 1555 capable of executing arithmetic functions and routing related to the implementation of the digitally implemented artificial neurons, such as, but not limited to, those explained herein. Each neuromorphic core 1515 may also provide local memory in which routing tables of the neural network can be stored and accessed, accumulated potentials of each cell body of each neuron implemented using core 1515 can be tracked, parameters of each neuron implemented by core 1515 can be recorded, and other data and usage can be recorded. Components or architectural resources of the neuromorphic core 1515 may also include: an input interface 1565 for receiving input spike messages generated by other neurons on other neuromorphic cores; and an output interface 1570 for sending spike messages to other neuromorphic cores via a mesh network 1510. In some instances, the routing logic of the neuromorphic core 1515 can be implemented at least partially using the output interface 1570. Furthermore, in some cases, the core (e.g., 1515) can implement multiple neurons within an example SNN, and some of these neurons can be interconnected. In this case, spiking messages sent between neurons hosted on the core 1515 can forgo communication via the routing structure of the neuromorphic computing device 1505 and can be managed locally within the specific neuromorphic core 1515.

[0174] Each neuromorphic core may also include logic for implementing artificial dendrites 1580 and artificial cell bodies 1585 (hereinafter referred to as “dendrites” and “cell bodies”, respectively) for each neuron 1575. Dendrite 1580 may be a hardware-implemented process for receiving impulses from network 1510. Cell body 1585 may be a hardware-implemented process for receiving the current time-accumulated neurotransmitter mass of each dendrite and evolving the potential states of each dendrite and cell body to generate outgoing impulse messages at appropriate times. Dendrite 1580 may be defined for each connection receiving input from another source (e.g., another neuron). In one implementation, the dendritic process 1580 may receive and process the impulse message as it arrives serially from network 1510 in a time-division multiplexed manner. With the reception of impulses, neuronal activation (tracked using cell body 1585 (and local memory 1560)) may increase. When the activation of a neuron 1575 exceeds a threshold set for neuron 1575, neuron 1575 generates a spike message, which is propagated via output interface 1570 to a fixed set of fan-out neurons. The network distributes the spike messages to all destination neurons, which in turn can update their activation in a transient, time-dependent manner in response. This can lead to some of the destination neurons also exceeding their corresponding thresholds and triggering further spike messages, just as in real biological neural networks.

[0175] As described above, the neuromorphic computing device 1505 can reliably implement spiking-based neural computing models. Such models are also referred to as spiking neural networks (SNNs). In addition to neuronal and synaptic states, SNNs incorporate temporal concepts. For example, in SNNs, communication occurs via event-driven action potentials or impulses that convey no explicit information other than the impulse duration and the implicit source and destination neuron pairs corresponding to the impulse transmission. The computation of the result of a dynamic nonlinear integral as a weighted impulse input occurs in each neuron. In some implementations, loops and dynamic feedback can be incorporated into the SNN computation model. Furthermore, various network connectivity models can be employed to model a wide range of real-world networks or relationships, including fully connected (all-to-all) networks, feedforward trees, fully random projections, “small-world” networks, and other examples. Isomorphic two-dimensional networks of the neuromorphic core (e.g., but not limited to the network shown in the example of Figure 15) can advantageously support all these network models. Since some or all of the core of the neuromorphic computing device 1505 can be connected, some or all of the neurons defined in the core can also be fully connected via a certain number of router hops. The neuromorphic computing device 1505 may also include a fully configurable routing table for defining a variety of different neural networks by allowing neurons in each core to distribute their spiking to any number of cores in the grid 1510 to achieve a completely arbitrary connection graph.

[0176] Improved implementations of systems capable of supporting SNNs, such as, but not limited to, VLSI hardware devices as shown in the example of Figure 15, can provide high-speed, reliable circuitry to model the information processing algorithms employed by the brain, but in a more programmable manner. For example, while a biological brain can only perform a specific set of defined behaviors (a consequence of years of development), neuromorphic processor devices can provide the ability to rapidly reprogram all neural parameters. Therefore, a single neuromorphic processor can be used to implement a wider range of behaviors than a single slice of biological brain tissue. This distinction can be achieved by employing neuromorphic processors with neuromorphic designs that are distinctly different from those found in natural neural circuitry.

[0177] As an example, a neuromorphic processor can implement a spontaneous neural network (SNN) using time-multiplexed computation in both a spiking communication network and the neuronal mechanism of the neuromorphic computing device 1505. Therefore, the physical circuitry of the neuromorphic computing device 1505 can be shared by many neurons to achieve a higher neuron density. Through time multiplexing, the network can connect N cores with a total wiring length of O(N), while the length of discrete point-to-point wiring will be extended to O(N). 2 This significantly reduces wiring resources to accommodate planar and non-plastic VLSI routing techniques, among other examples. In the neuromorphic core, time multiplexing can be implemented through dense memory allocation, for example, using static random access memory (SRAM) with a shared bus, address decoding logic, and other multiplexed logic elements. The state of each neuron can be stored in the processor's memory, where data describing the state of each neuron includes the state of the collective synapse of each neuron, all currents and voltages on its membrane, and other example information (e.g., but not limited to configuration and other information).

[0178] Neuromorphic processors can be implemented in a “digital” manner, unlike other processors that employ more “analog” or “isomorphic” neuromorphic approaches. For example, a digital implementation can use digital adder and multiplier circuitry to integrate synaptic currents, in contrast to an analog isomorphic neuromorphic approach that accumulates charge on capacitors in a manner similar to how neurons accumulate synaptic charge on their lipid membranes. For instance, the accumulated synaptic charge for each neuron can be stored in the local memory of the corresponding core. Furthermore, at the architectural level of an example digital neuromorphic processor, reliable and deterministic operation can be achieved through time synchronization across the core network, ensuring that any two executions of the design, given the same initial conditions and configuration, will produce the same results. Asynchronicity can be reserved at the circuit level to allow individual cores to operate as quickly and freely as possible while maintaining determinism at the system level. Therefore, in neural computing, the concept of time as a time variable can be abstracted away from the “wall clock” time used by the hardware to perform computations. Thus, in some implementations, a time synchronization mechanism can be provided that globally synchronizes the neuromorphic cores at discrete time intervals. Synchronization mechanisms allow neural computation to be completed at the fastest speed allowed by the circuitry, and there is a difference between the runtime and the biological time for modeling neuromorphic systems.

[0179] In operation, the neuromorphic computing device 1505 can start in an idle state when all neuromorphic cores are inactive. As each core asynchronously loops through its neurons, it generates impulse messages, which are routed by the mesh interconnect to the appropriate destination core containing all destination neurons. The implementation of multiple neurons on a single neuromorphic core can be time-multiplexed, and time steps can be defined, where all impulses involving multiple neurons can be processed and considered using the shared resources of the respective cores. When each core completes its service to its neurons within the corresponding time step, in some implementations, the core can communicate with neighboring cores using synchronization messages (e.g., using a handshake) to refresh the mesh of all transmitted impulse messages, allowing the core to safely determine that all impulses have been serviced within a certain time step. At this point, all cores can be considered synchronized, allowing them to advance their time steps and return to the initial state to begin the next time step.

[0180] Given this context, as described above, a device (e.g., 1505) can be provided to realize an interconnected neuromorphic core network 1510, wherein the core 1515 can realize multiple artificial neurons capable of interconnecting to realize an SNN. Each neuromorphic core (e.g., 1515) can provide two loosely coupled asynchronous processes: an input dendrite process (e.g., 1580) that receives impulses from the network 1510 and applies them to the appropriate destination dendritic chamber at an appropriate future time; and an output cell body process (e.g., 1585) that receives the current-time accumulated neurotransmitter mass of each dendritic chamber and evolves the membrane potential state of each dendrite and cell body to generate an outgoing impulse message at an appropriate time (e.g., when the threshold potential of the cell body is reached). It should be noted that, from a biological perspective, the names of dendrites and cell bodies used herein are only approximate to the function of these features and should not be interpreted too literally.

[0181] In at least one embodiment, the neuromorphic computing device 1505 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0182] Figure 16 is a block diagram of an embodiment of a multi-node network capable of enabling remote memory computing according to any embodiment. System 1600 may represent the node network described herein, for example, the node network may be used to perform some or all of the operations described herein. System 1600 may represent a data center. System 1600 may represent a server farm. System 1600 may represent a data cloud or processing cloud. System 1600 may represent a supercomputer. System 1600 may include tens, hundreds, or thousands of nodes. The nodes of System 1600 may include processors, such as, but not limited to, a central processing unit (CPU), a graphics processing unit (GPU), or any combination of processors described herein, such as, but not limited to, other processors shown in Figures 10-22B. For any processor in System 1600 and any component described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) inside or outside the processor or node. The results of the APIs may be stored in storage devices inside or outside the processor or node, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents. System 1600 may include more than nine thousand nodes, each node comprising two Intel Xeon Max processors, six Intel Max series GPUs, and a unified memory architecture, such as, but not limited to, the architecture used in Intel's Aurora supercomputer in Santa Clara, California, or other supercomputers that share at least some of the components described herein.

[0183] One or more clients 1602 send requests to system 1600 via network 1604. Network 1604 represents one or more local area networks, wide area networks, or a combination of both. Clients 1602 can be human or machine clients that generate requests for operations to be performed by system 1600. System 1600 executes the application or data computation task requested by client 1602.

[0184] System 1600 may include one or more racks, which represent structural and interconnect resources for housing and interconnecting multiple computing nodes. Rack 1610 may include multiple nodes 1630. Rack 1610 may carry multiple blade assemblies 1620. Carrying can refer to providing power, structural or mechanical support, and interconnection. Blade 1620 may refer to computing resources on a printed circuit board (PCB), where the PCB houses the hardware components of one or more nodes 1630. Blade 1620 may or may not include a chassis, housing, or other “box” besides those provided by rack 1610. Blade 1620 may include a housing with exposed connectors for connection to rack 1610. System 1600 may or may not include rack 1610, and each blade 1620 may include a chassis or housing that may be stacked or otherwise closely proximate with other blades, allowing nodes 1630 to interconnect. The System 1600 can include 10,624 compute blades, comprising 63,744 Intel Max series GPUs and 21,248 Intel Xeon Max CPUs across 166 racks.

[0185] System 1600 may include architecture 1670, which represents one or more interconnectors of nodes 1630. Architecture 1670 may include multiple switches 1672 or routers or other hardware for routing signaling between nodes 1630. Furthermore, architecture 1670 may couple system 1600 to network 1604 for access by client 1602. In addition to routing devices, architecture 1670 may also be considered to include cables or ports or other hardware devices for coupling nodes 1630 together. Architecture 1670 may have one or more associated protocols for managing signaling routing through system 1600. One or more protocols are at least partially dependent on the hardware devices used in system 1600.

[0186] As shown in the figure, rack 1610 may include N blades 1620. In addition to rack 1610, system 1600 may also include rack 1650. As shown in the figure, rack 1650 may include M blades 1660. M is not necessarily the same as N; therefore, it is understood that various different hardware device components can be used and coupled together into system 1600 via structure 1670. Blade 1660 may be the same as or similar to blade 1620. Node 1630 may be any type of node described herein and is not necessarily of the same type. System 1600 is not limited to homogeneous or non-homogeneous structures.

[0187] The nodes in blade 1620(0) are shown in detail. However, other nodes in system 1600 may be the same or similar. At least some nodes 1630 may be compute nodes, having processor 1632 and memory 1640. A compute node is a node having processing resources (e.g., one or more processors) that executes an operating system and can receive and process one or more tasks. At least some nodes 1630 may include storage server nodes, which have servers as processing resources 1632 and memory 1640. A storage server is a node having more storage resources than a compute node, and instead of having processors for performing tasks, a storage server includes processing resources for managing access to storage nodes within the storage server.

[0188] Node 1630 may include interface controller 1634, which may represent logic for controlling node 1630's access to structure 1670. The logic may include hardware resources for interconnecting to physical interconnect hardware. The logic may include software or firmware logic for managing the interconnect. Interface controller 1634 may include a host structure interface, which may include a structure interface according to any embodiment described herein.

[0189] Node 1630 may include a memory subsystem 1640. Memory 1640 may include a memory computation resource (comp) 1642, which represents the ability of memory 1640 to perform one or more memory computations. System 1600 supports remote memory operations, such as, but not limited to, those described elsewhere herein. Therefore, node 1630 may request a remote node to perform a memory computation, wherein the data used for the computation remains local to the executing node, and is not sent via structure 1670 or from memory to the structure interface. In response to the execution of the memory computation, the executing node may provide the result to the requesting node.

[0190] Processor 1632 may include one or more individual processors. Each individual processor may include a single processing unit, a multi-core processing unit, or a combination thereof. A processing unit may include a main processor, such as, but not limited to, a CPU (Central Processing Unit), a peripheral processor (such as, but not limited to, a GPU (Graphics Processing Unit)), or a combination thereof. Memory 1640 may be or include memory devices and a memory controller.

[0191] The term "memory device" can refer to different types of memory. Memory devices generally refer to volatile memory technology. Volatile memory is memory whose state (and the data stored within it) is uncertain if power is interrupted. Non-volatile memory is memory whose state is deterministic even if power is interrupted. Dynamically volatile memory can refresh the data stored in the device to maintain its state. An example of dynamically volatile memory includes DRAM (Dynamic Random Access Memory) or variations thereof, such as, but not limited to, Synchronous DRAM (SDRAM). The memory subsystem described in this article is compatible with a variety of memory technologies, such as, but not limited to, DDR3 (Double Data Rate version 3, originally released by JEDEC (Joint Electron Device Engineering Committee) on June 27, 2007, currently version 21), DDR4 (DDR version 4, initial specification released by JEDEC in September 2012), DDR4E (DDR version 4, extended version, currently under discussion by JEDEC), LPDDR3 (Low Power DDR version 3, JESD209-3B, released by JEDEC in August 2013), and LPDDR4 (Low Power Double Data Rate (LPDDR) version 4). JESD209-4 (originally released by JEDEC in August 2014), WIO2 (Wide I / O2), JESD229-2 (originally released by JEDEC in August 2014), HBM (High Bandwidth DRAM), JESD235 (originally released by JEDEC in October 2013), DDR5 (DDR version 5, currently under discussion by JEDEC), LPDDR5 (currently under discussion by JEDEC), HBM2 (HBM version 2, currently under discussion by JEDEC) or combinations of other memory technologies, as well as technologies derived from or extended based on such specifications.

[0192] In addition to or as an alternative to volatile memory, in one embodiment, a reference to a memory device may refer to a non-volatile memory device whose state is deterministic even when power is interrupted. In one embodiment, a non-volatile memory device is a block-addressable memory device, such as, but not limited to, NAND or NOR technology. Therefore, the memory device may also include future-generation non-volatile devices, such as, but not limited to, three-dimensional cross-point (3DXP) memory devices, other byte-addressable non-volatile memory devices, or memory devices using chalcogenide phase change materials (e.g., chalcogenide glasses). In one embodiment, the memory device may be or include multi-threshold NAND flash memory, NOR flash memory, single-level or multi-level phase change memory (PCM) or switched phase change memory (PCMS), resistive memory, nanowire memory, ferroelectric transistor random access memory (FeTRAM), magnetoresistive random access memory (MRAM) incorporating memristor technology, or spin-transfer torque (STT)-MRAM, or any combination of the foregoing, or other memories.

[0193] In at least one embodiment, system 1600 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0194] Figure 17 illustrates an accelerated processing unit 1700 according to at least one embodiment. The accelerated processing unit 1700 may include a processor based on the CDNA architecture of AMD, Inc., Santa Clara, California, or other processors sharing at least some of the components described herein. The accelerated processing unit 1700 may include one or more accelerator complex dies (XCDs) 1704 for performing the operations described elsewhere in this document, such as, but not limited to, graphics processing and / or parallel processing and instruction-level parallel computing, including support for multiple precisions (INT8, FP8, BF16, FP16, TF32, FP32, and FP64) and sparse matrix data (i.e., sparsity). In some cases, the XCD may be referred to as a graphics computing die (GCD). The accelerated processing unit 1700 may include one or more complex computing dies (CCDs) 1706 for performing the operations described elsewhere in this document, such as, but not limited to, operations performed by a host processor. In some cases, the CCD may be referred to as a core complex or CCX, such as, but not limited to, the CCX used in AMD Ryzen processors. XCDs and CCDs can share any type of cache or memory (e.g., one or more memory cells 1702), or a cache or memory can be allocated to each XCD or CCD or group of XCDs or CCDs. For example, AMD Infinity Fabric within the package connects XCDs and CCDs to a shared AMD Infinity Cache 1708, and in some embodiments, to high-bandwidth memory (e.g., HMB3). Accelerated processing unit 1700 may include an AMD MI300a processor comprising three CPU die (or CCD) and six accelerator die (XCD) on top of four input-output dies (IODs) that may be layered on a single silicon die (e.g., via AMD Infinity Fabric) and linked together to eight high-bandwidth DRAM stacks in a ring to form a superchip. For systems using only accelerators, the AMD MI300x processor replaces the CCD with two or more XCDs.

[0195] Accelerated processing unit 1700 may include one or more input / output (I / O) interfaces. For example, XCD 1704 and CCD 1706 may coexist on one or more input-output dies (IODs) 1710, which may include one or more I / O interfaces. IOD 1710 may include any number and type of I / O interfaces (e.g., PCI, PCI expansion (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). Various types of peripheral devices may be coupled to I / O interfaces 1770. The I / O interfaces of IOD 1710 may also be used to connect one or more accelerated processing units 1700, for example, in a server architecture.

[0196] Accelerated processing unit 1700 may include one or more memory units 1702 for storing instructions and other information for performing the operations described in other parts of this document. Memory units 1702 may include any volatile memory, such as, but not limited to, the memory types described in other parts of this document, and may include, for example, high-bandwidth memory (e.g., HMB3) or high-bandwidth DRAM. The memory associated with accelerated processing unit 1700 (e.g., memory unit 1702) may include system memory, which can be used for, for example, commands, instructions, and constants, as well as input and output. Memory unit 1702 may also include device memory, which can be used for storage and, for example, for commands, instructions, and constants, as well as input and output, as a return buffer, and for private data. Memory unit 1702 may be linked to one or more IODs 1710. In at least one embodiment, L1 cache 1720 initiates a memory hierarchy including a shared L2 cache 1728 (e.g., within an XCD). AMD Infinity Cache TM It is the last-level cache (LLC) located on the active I / O die (IOD). The CCD 1706 and XCD1704 can have independent or shared memory. AMD Infinity architecture and AMD Infinity Fabric TM The technology enables consistent, high-throughput unification of GPU and CPU chip technologies (such as XCD, CCD, and / or CCX) with memory (such as stacked HBM3 memory) in a single device and across multiple device platforms.

[0197] As shown in Figure 17, the XCD 1704 may include a set of shared global resources 1730, which may include a hardware scheduler 1732 and an asynchronous compute engine (ACE) 1724. The ACE 1724 sends tasks (e.g., compute shader workgroups) to compute units (CUs or cores) 1734. Each of the ACEs 1724 (e.g., four) may be associated with a CU 1734 (e.g., 40 CUs), and some CUs 1734 may be disabled for yield management. CUs 1734 may have dedicated caches or shared caches (e.g., L2 caches) 1728 for consolidating all memory traffic on a single die. CU1734 may include threaded and parallel processor cores, including instruction fetching and scheduling using a scheduler (S) 1712, a matrix core unit (MCU) 1716, and a shader core (SC) 1718 (e.g., execution units for scalar, vector, and matrix data types), and a load / store pipeline with an L1 cache 1720 and a local data share (LDS) 1714. The local data share may include, for example, sticky-note RAM with built-in arithmetic capabilities, allowing data sharing between threads in a workgroup. Instruction cache 1740 (e.g., for storing and providing instructions for performing the operations described elsewhere in this document) and constant cache 1738 may be connected to one or more CUs and may be shared between two CUs. Matrix core 1716 can handle various data types, such as, but not limited to, INT8, FP8, FP16, BF16, and TF32 data types. Accelerated processing unit 1700 may include computation units 1734, which may be arranged in an array format, such as as a data parallel processor (DPP) array. The hyper-threaded dispatch processor 1742 can communicate with the compute unit 1734, and the command processor 1744 can read commands written by the host to memory-mapped registers in the system memory address space (not shown). When a command is completed, the command processor 1744 can send a hardware-generated interrupt to the host processor (e.g., a CCD). The memory controller 1736 can also directly access all device memories and system memory regions specified by the host. To satisfy read and write requests, the memory controller 1736 can perform the functions of a direct memory access (DMA) controller, including calculating the memory address offset based on the format of the requested data in memory. For example, one or more APIs described herein can be compiled into instructions that can be stored in the instruction cache 1740, then fetched by the instruction fetch logic in the processor 1740, decoded by the processor decoder or equivalent, scheduled (e.g., sequentially or out of order) by the scheduler or equivalent for execution, executed by the execution logic or equivalent, reordered, and then retired by the retirement logic or equivalent.The API (and / or compiled instructions including the API) can be stored in any storage device, either inside or outside the processor 1700 (e.g., in cache and / or memory). The results of the API can be stored in storage devices, either inside or outside the processor 1700, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0198] Applications may include programs running on the main processor (e.g., a CCD) and programs running on one or more XCDs (referred to as kernels). Programs can be controlled by host commands that set internal base addresses and other configuration registers, specify data fields on which the accelerator processing unit 1700 can run, invalidate and flush caches on the accelerator processing unit 1700, and cause the accelerator processing unit 1700 to begin executing a program. A kernel can be referred to as a program executed by the accelerator processing unit 1700. Kernels can execute independently on each work item or as a group of work items, referred to as a wavefront, which can execute kernels on all (e.g., 64) work items in a single pass. The computation unit 1734 may include: a scalar arithmetic logic unit (ALU) that can operate on a single value for each wavefront (shared by all work items); a vector ALU that can operate on a unique value for each work item; a local data share 1714 that allows work items within a workgroup to communicate and share data; a scalar memory (not shown) that can transfer data between the scalar general-purpose registers (SGPRs) and memory via cache; and a vector memory that can transfer data between the vector general-purpose registers (VGPRs) and memory, including sampling texture maps. Kernel control flow can be manipulated using scalar ALU instructions, which may include if / else statements, branches, and loops. Scalar ALU (SALU) and memory instructions can operate on the entire wavefront and operate on one or more SGPRs. Vector memory and ALU instructions can operate on all work items in the wavefront simultaneously.

[0199] In at least one embodiment, the acceleration processing unit 1700 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0200] Figure 18 illustrates a processor 1800, such as, but not limited to, a Zen architecture-based processor (e.g., Zen1, 2, 3, 4, 5, or other architectures) from AMD Inc. of Santa Clara, California, or other processors that share at least some of the components described herein. The processor 1800 includes one or more CPU dies 1802(1)–1802(N), where N is any integer greater than 1. The CPU die 1802 may include any number of processor cores 1816 (e.g., for performing any operations described elsewhere herein) and any number of cache memories (e.g., for storing instructions and other information to perform any operations described elsewhere herein). For example, an L2 cache unit 1818 may be coupled to a processor core 1816, and the processor cores 1816 may share and / or be individually coupled to the L2 cache unit 1818. Processor core 1816 can be coupled to L3 cache 1822 and / or a shared L3 cache, which can be the lowest level cache (LLC) 1822 used to access data and other information used by processor core 1816. One or more processor cores 1816 and one or more L2 cache units 1818 can be included in a core complex (CCX) 1820, which can include (e.g., 32MB) a shared cache (e.g., L3 cache 1822). Core complex 1820 can be manufactured onto a die (CCD or CPU die) 1802. For example, up to 12 core complexes 1820 can be configured in a processor along with 8 CPU dies 1802, providing up to 96 processor cores 1816 for processor 1800. For example, a "Zen 4c" core complex 1820 can include up to 8 cores 1816 and a shared 16MB L3 cache 1822. Two core complexes from these core complexes 1820 can be combined onto a single CPU die 1802, resulting in 16 cores per die and a total of 32MB of L3 cache per die 1822. Up to eight CPU dies 1802 can be combined with I / O units 1804 to provide up to 128 processor cores 1816 for the CPU. Up to four "Zen 4c" dies mentioned above can be combined to provide up to 64 processor cores 1816 for the CPU.

[0201] Processor 1800 may include various configurations for input / output operations, which will be further described herein. I / O unit 1804 may include one or more memory controllers 1806 capable of managing the memory usage of processor 1800 (e.g., DDR5 memory). I / O unit 1804 may include one or more SATA disk controllers for managing storage device 1812, and one or more Compute Express Links (CXL) for providing CPU-to-device and CPU-to-memory connectivity, which can be flexibly assigned to specific functions during server design. TM 1.1+ Memory controller 1814. I / O unit 1804 may include PCIe controller 1808 for connecting peripherals and other components connected to processor 1800. I / O unit 1804 may also include USB port 1810 for connecting to other components separate from processor 1800. CPU die 1802 may support any number of connections to I / O unit 1804, for example, one or two connections. As shown, I / O unit 1804 may include components further described herein, and I / O unit 1804 may be an I / O die accommodating several different components. Memory controller 1806, PCIe controller 1808, USB port 1810, SATA controller 1812 and / or CXL controller 1814 may be individually integrated into any location within processor 1800, or integrated in any group or combination.

[0202] Processor 1800 may include an Infinity Fabric 1824 interconnect (which may be similar to or based on a PCIe architecture) that provides connectivity between the CPU (e.g., CPU die 1802(1)-1802(N)), graphics processor 1826, inference engine 1832, and other components in a multi-chip architecture (e.g., security processor 1828 and I / O unit 1804). One or more AMD Infinity Fabrics TM Interconnect 1810 can be connected to CPU dies 1802(1)-1802(N) and used as a connection between CPUs. One or more Infinity Fabric connections 1810 can connect each CPU die 1802 to the I / O unit 1810.

[0203] In at least one embodiment, processor 1800 may include a central processing unit (CPU) and other related hardware and software described above and further herein. Processor 1800 may also include a graphics processor 1826. Graphics processor 1826 may be used for image generation and processing, as well as other computations and operations described further herein. Graphics processor 1826 may be based on AMD's RDNA 3 or 3.5 architecture, located in Santa Clara, California. Graphics processor 1826 may include a graphics computing die (GCD) and a memory cache die (MCD). The GCD may include any number of computing units (CUs) for graphics or other processing, such as operations performed by an arithmetic logic unit (ALU) described further herein. Graphics processor 1826 may include an L2 cache available for use by the computing units. The MCD (not shown) may include any number of memory cells and may include a cache (e.g., an L3 cache) and a memory interface for coupling to memory (e.g., memory 1842(1)-(N), where N is an integer). Components within the graphics processor 1826 can be connected using various methods, such as using Infinity Fabric 1824 interconnects, either internally or externally to the graphics processor 1826.

[0204] Inference engine 1832 can provide neural processing capabilities to processor 1800 for computational processes used in neural networks, deep learning, and other AI-related operations, which will be further described herein. Processor 1800 may include: a security processor 1828 for managing the security of processor 1800; a display controller 1830 for controlling the display; a system management unit 1834 for managing and operating some or all components on processor 1800; a multimedia engine 1836 for audio and video operations; a fusion controller hub 1838 for managing USB, SATA, and PCIe connections to processor 1800; and a sensor fusion hub 1840 for managing sensors (e.g., accelerometers). Processor 1800 may also include memory 1842(1)-(N), where N is any integer. Memory may include different memory types, such as LPDDR5 and / or DDR5, or other memory described elsewhere herein.

[0205] To perform the operations further described herein, processor 1800 may include an execution pipeline including a front end that may include a cache for storing instructions (e.g., an L1 cache) (not shown). A branch predictor may modify the instruction stream. Instructions may be decoded by a decoder, dispatched to a back end for execution, and renamed. For example, the instruction fetch and decode pipeline may be dispatched to integer or floating-point execution operations, which may be scheduled by a scheduler and passed to vectors and / or general-purpose registers. Floating-point multipliers and / or addition operations may be processed, and an arithmetic logic unit (ALU) may also be used to perform computations, such as arithmetic and logical operations. The output of the computation unit may be coupled to a load / store queue that may be connected to a cache, such as an L1 cache and / or an L2 cache.

[0206] With respect to processor 1800 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may be compiled, for example, into instructions or equivalents (e.g., AVX-512 instructions based on a SIMD model). These instructions or equivalents may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) internal or external to processor 1800. The results of the APIs may then be stored in storage devices internal or external to processor 1800, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0207] In at least one embodiment, the processor 1800 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0208] Figure 19 illustrates an example of a processing core 1900, which may implement an Arm architecture (e.g., v9.0-A) or another processor that shares at least some of the components described herein. TM The V2 core 1900 can be implemented within a DynamIQ shared unit (DSU) cluster via a DSU-110 interconnect 1954 for connecting one or more cores, for example, for parallel processing. Neoverse TM The V2 core can be implemented as a single core in a DSU cluster configured for direct interconnection, with or without L3 cache, listener filter, or listener control unit (SCU) logic (not shown). Neoverse TMThe V2 core may include a CPU bridge 1952 for connecting core 1900 to the DSU-110 interconnect. This bridge may also connect core 1900 to an external memory system and the remainder of the on-chip system. The L1 instruction memory system 1902 may fetch instructions from instruction cache 1904 and deliver instructions (e.g., one or more compileable APIs described herein) to instruction decoding unit 1910, for example, to perform some or all of the operations described above or elsewhere herein. The L1 instruction memory system 1902 may include L1 instruction cache 1904 (e.g., with 64-byte cache lines), L1 instruction translation back buffer (TLB) 1906 (e.g., natively supporting 4KB, 16KB, 64KB, and 2MB page sizes), and macro operation cache (MOP) 1908 (e.g., a 1536-entry, 4-way skewed associated L0 MOP cache), which may include decoded and optimized instructions for higher performance. Instruction decoding unit 1910 may decode AArch64 instructions into its internal format. The register renaming unit 1912 can perform register renaming to facilitate out-of-order execution and dispatch decoded instructions to various issue queues. The instruction issuing unit 1914 can control when decoded instructions are dispatched to the execution pipeline, and it can include an issue queue for storing instructions to be dispatched to the execution pipeline. The integer execution pipeline 1916 can be included in the execution pipeline and includes an integer execution unit 1918 that can perform arithmetic and logical data processing operations. The vector execution unit 1920 can be included in the execution pipeline and can execute advanced SIMD and floating-point arithmetic (FPU) 1922, execute Scalable Vector Extension (SVE) and Scalable Vector Extension 2 (SVE2) instructions 1924, and can also selectively execute cryptographic instructions 1926. The advanced SIMD can include a media and signal processing architecture that primarily adds instructions for audio, video, 3D graphics, image, and speech processing. The floating-point architecture provides support for single-precision and double-precision floating-point operations. The L1 data memory system 1930 executes load and store instructions, as well as service memory coherence requests. The L1 data memory system 1930 may include an L1 data cache 1932 and a fully associative L1 data TLB 1934, natively supporting 4KB, 16KB, and 64KB page sizes and 2MB and 512MB block sizes. The memory management unit (MMU) 1928 provides fine-grained memory system control through a set of virtual-to-physical address mappings and memory attributes, which are stored in a translation table and saved to the TLB 1934 after address translation. The L2 memory system 1936 may include an L2 cache 1938 and can be connected to the DSU-110 1954 via an asynchronous CPU bridge 1952. TMThe V2 core 1900 supports a range of debugging, testing, and tracing options, including the tracing unit 1942, the tracing buffer 1940, and the embedded logic analyzer (ELA) 1948. Neoverse TM The V2 core 1900 implements the Statistical Analysis Extension (SPE) 1944, which provides a statistical view of the performance characteristics of executed instructions. Software writers can leverage these views to optimize code for better performance. The Performance Monitoring Unit (PMU) 1946 provides a performance monitor that can be configured to collect statistics on the operation of each core and memory system. This information can be used for debugging and code analysis. The General Purpose Interrupt Controller (GIC) CPU interface 1950, when integrated with external allocator components, serves as a resource for supporting and managing interrupts in a cluster system. In a cluster, each Neoverse... TM There can be a CPU bridge 1952 between the V2 core 1900 and the DSU-110 1954. The CPU bridge 1952 can control the buffering and synchronization between the core 1900 and the DSU-110 1954. The CPU bridge 1952 can be asynchronous to allow each core 1900 to use a different frequency, power, and area implementation point. The CPU bridge 1952 can operate synchronously without affecting other interfaces, such as, but not limited to, asynchronous debug and tracing interfaces.

[0209] In at least one embodiment, core 1900 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0210] Figure 20 illustrates one or more chips including one or more Tensor Processing Units (TPUs) 2000 according to at least one embodiment. The TPU 2000 in Figure 20 may include an application-specific integrated circuit (ASIC), for example, for performing some or all of the operations described above or elsewhere herein, such as, but not limited to, machine learning workloads that accelerate matrix operations. The TPU 2000 may be an ASIC from Alphabet Corporation, Mountain View, California. Cloud TPUs include a cloud service that enables the TPU to be used as a scalable resource for processing tasks, such as, but not limited to, machine learning workloads that can run on frameworks such as, but not limited to, TensorFlow, PyTorch, and JAX.

[0211] Chip 2000 may include any number of TPUs, which may include Tensor Core 2006. Tensor Core 2006 may include one or more core sequencers 2008, vector processing units (VPUs) 2010, matrix multiplication units (MXUs) 2012(A)-2014(N) (where N is any integer greater than 1), and transpose permutation units 2016. Core sequencer 2008 may fetch instructions (e.g., VLIW (Very Long Instruction Word)) from the instruction memory (Imem) of core 2006, perform scalar operations using scalar data memory (Smem) and scalar registers (Sregs) (not shown), and forward vector instructions to the vector processing units (VPUs) 2010. For example, an instruction may initiate eight operations: two scalars, two vector ALUs, vector loading and storing, and queuing data into the matrix multiplication and transpose units and a pair of slots for queuing data from them. The VPU2010 can perform vector operations using a large on-chip vector memory (Vmem) and vector registers (Vregs). The VPU 2010 can stream data to or from the MXU via decoupled FIFOs. The VPU 2010 can collect and distribute data to the Vmem using both data-level parallelism (2D matrix and vector functional units) and instruction-level parallelism (8 operations per instruction). Large two-dimensional matrix multiplication units (MXUs) 2012(A)–2012(N) can, for example, use systolic arrays to reduce area and power consumption, and use large software-controlled on-chip memory instead of caches. The transpose-reduction-permute unit 2016 can perform matrix transpose, reduction, and permute operations (e.g., 128x128) on VPU 2010 channels. A high-bandwidth memory 2004 can be used for on-chip applications. One or more chips 2000 can be interconnected for computation. For example, one or more chips 2000 can be connected to form a torus, such as a 2D torus. Chip 2000 may also include any number (e.g., 4) of inter-core interconnect (ICI) links 2018, which can enable direct connections between chips to form a supercomputer.

[0212] For any processor in the chip 2000 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) external to or internal to any processor in the chip 2000. The results of the APIs may then be stored in any storage device internal to or external to any processor in the chip 2000, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0213] In at least one embodiment, chip 2000 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0214] Figure 21 illustrates a vector processor according to at least one embodiment. The vector processor 2100 may support the RISC-V standard. The vector processor 2100 may include one or more cores 2110 (e.g., scalar units) and one or more vector processing units (VPUs) 2142 (e.g., vector units), which may, for example, perform some or all of the operations described above or elsewhere herein. The core 2110 may include an Andes Custom Extension (ACE) 2116, which can be used to deliver custom instructions to the processor 2100. The core 2110 may include a 1-cycle multiplier and a 1-cycle instruction / data local memory (ILM / DLM) for improving parallelism by allowing simultaneous instruction fetching and data access. A memory management unit (MMU) 2124 manages system memory and cache, and provides branch execution, instruction pair issuance, L1 instruction / data cache, and local memory storage. The core 2110 may include a physical memory protection and programmable physical memory attribute unit (PMP / PPMA) 2122. Core 2110 may include a digital signal processor (DSP) 2128 and a floating-point unit (FPU) 2126, as well as a load-memory unit (LSU) 2132 for interaction with memory hierarchies (D$2134 and I$2130). Core 2110 may include a branch prediction unit 2118 and a multiplier unit 2120.

[0215] The vector processing unit (VPU) 2142 may include one or more vector function units (FU) 2146(A)-2146(N) (these vector function units may be linked together for parallel processing), a separate memory path for loading / storing RISC-V vectors (RVVs) via ACE-RVV 2148 and AndesStreaming port (ASP) 2144, and a vector load / store unit (VLSU) 2150.

[0216] Vector processor 2100 may include bus interfaces, such as, but not limited to, a cache-accessible L2 cache port 2156, a non-cache-accessible MMIO port 2154, an input-output coherence port (IOCP) 2158 for a cacheless bus master, a local memory access port for accessing ILM / DLM 2112 and High Bandwidth Vector Memory (HVM) 2136, and a shared peripheral port (SPP) 2152 for external peripherals. Other memory ports include an LM slave port AXI 2102 and an HVM subordinate port AXI 2104.

[0217] For any processor in processor 2100 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device external to or internal to processor 2100 (e.g., in cache and / or memory). The results of the APIs may then be stored in storage devices internal to or external to processor 2100, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0218] In at least one embodiment, the vector processor 2100 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0219] Figure 22A illustrates a schematic diagram of an example multi-core tiled processor microarchitecture. The multi-core tiled processor in Figure 22A may include a language processing processor. As shown in Figure 22A, each “tile” in the processor architecture is a processing element bundled together using an on-chip network (NoC) that can be used to perform some or all of the operations described above or elsewhere in this document. For example, each tile may have an instruction dispatch 2204 and integer (INT) units 2206 and floating-point (FP) units 2208, a load-memory unit (LSU) 2212 for engaging with a memory hierarchy (data cache (D$) 2210 and instruction cache (I$) 2214), and a network (NET) interface 2216 for communicating with other tiles. Some tiles in processor 2200 may include a memory controller 2202 for managing and controlling memory, as further described herein. Processor 2200 may have a functionally sliced ​​architecture. Processor 2200 may reside on an application-specific integrated circuit (ASIC), and Figure 22A may represent the layout of an ASIC. Processor 2200 may include a coprocessor designed to execute instructions for a predictive model. A predictive model is any model configured to make predictions based on input data. The predictive model may use a classifier for classification predictions. The predictive model may be a machine learning model, such as, but not limited to, a tensor flow model, and processor 2200 is a tensor flow processor.

[0220] Processor 2200 may employ different microarchitectures that decompose the functional units shown in each tile of Figure 22B. Instead, the functional tiles of processor 2200 may be aggregated into multiple functional processing units (hereinafter referred to as "slices") 2204, each corresponding to a specific function type (e.g., FP / INT 2218, NET 2220, MEM 2222). For example, as shown in Figure 22B, each slice may correspond to a row of functional tiles extending in a north-south direction. Furthermore, processor 2200 may also include communication channels for carrying data between tiles of different slices, each communication channel extending horizontally in an east-west direction. Each communication channel may be connected to each slice 2204 of processor 2200.

[0221] The slices 2204 of processor 2200 may each correspond to different functions and may include arithmetic logic slices (e.g., FP / INT 2218), channel switching slices (e.g., NET 2220), and memory slices (e.g., MEM 2222). Arithmetic logic units can perform one or more arithmetic and / or logical operations on data received via communication channels to generate output data. Examples of arithmetic logic units may be matrix multiplication units and vector multiplication units. Memory slices include memory cells that store data. Memory slices can provide data to other slices via communication channels. Memory slices can also receive data from other slices via communication channels. Channel switching slices can configurably route data from one communication channel to any other communication channel. For example, data from a first channel can be provided to a second channel via a channel switching slice. In some embodiments, a channel switching slice can be implemented as a crossbar switch. Each slice 2204 also includes its own instruction queue (not shown) for storing instructions and an instruction control unit (ICU) for controlling instruction execution. Instructions in a given instruction queue can only be executed by a tile in its associated functional slice and not by other slices of processor 2200.

[0222] By arranging the tiles of processor 2200 into different functional slices 2204, the on-chip instruction and control flow of processor 2200 can be separated from the data flow. For example, according to some embodiments, one arrow in FIG22B illustrates the instruction flow within the processor architecture. According to at least one embodiment, another arrow in FIG22B illustrates the data flow within the processor architecture. As shown, instruction and control flow can flow across the tiles of processor 2200 in a first direction (e.g., along the length of the functional slice in a north-south direction, as indicated by the first arrow), while data flow can flow across the tiles of processor 2200 in a second direction (e.g., across the functional slice in an east-west direction, as indicated by the second arrow), which is perpendicular to the first direction.

[0223] Different functional slices of processor 2200 can correspond to MEM 2222 (memory), VXM (vector execution module), MXM (matrix execution module), NIM (numerical interpretation module), and SXM (swapping and permutation module). Each slice can include N tiles, all of which can be controlled by the same instruction control unit (ICU) (not shown). Each slice can operate completely independently and can only be coordinated using barrier-like synchronization primitives or by the compiler using tractable determinism. Each tile of processor 2200 can correspond to an execution unit organized as ×MSIMD tiles. For example, each tile of the on-chip memory of processor 2200 can be organized to atomically store L-element vectors. Therefore, MEM slices with N tiles can work together to store or process large vectors (e.g., with a total of N×M elements).

[0224] Tiles within a slice can execute instructions in an "interleaved" manner, where instructions can be issued tile-by-tile within the slice over N cycle periods. Functional slices can be physically arranged on the chip to allow for efficient data flow for pipelined execution over hundreds of cycles for common patterns. The data flow can perform a single "u-turn" (direction change) corresponding to a single matrix operation before being written back to memory; in some embodiments, a particular data flow can change direction multiple times before writing the resulting data back to memory (due to multiple matrix and vector operations).

[0225] When using a processor 2200 with a function slice architecture (e.g., a TSP), the TSP compiler (not shown) generates an explicit plan of how the processor 2200 executes programs (e.g., microprograms). The compiler can specify when each operation will be executed, which function slices will perform the work, and which STREAM registers will hold operands. The compiler can maintain a high-fidelity (cycle-accurate) model of the processor 2200 (e.g., TSP) hardware state so that the microprogram can coordinate data flow.

[0226] Processor 2200 (e.g., TSP) can use a web-hosted compiler that takes a model (e.g., an ML model, such as but not limited to a TensorFlow model) as input and issues a proprietary stream of instructions for processor 2200 (e.g., TSP). The compiler is responsible for coordinating the control and data flow of the program and specifying any instruction-level parallelism by explicitly bundling instructions that can and should be executed concurrently so that they can be dispatched together. The main hardware architecture includes the architecture-visible streaming register file (STREAM), which will be described in more detail below, and acts as a conduit for operands to flow from MEM slices (e.g., SRAM) to functional slices (and vice versa).

[0227] The MEM 2222 of processor 2200 can be used as: (1) a storage for model parameters, microprograms, and data on which they operate; and (2) an on-chip network (NoC) for transferring data operands from the MEM to functional slices and returning computation results to the MEM. In some embodiments, the on-chip memory may consume approximately 75% of the chip area of ​​processor 2200. In some embodiments, the on-chip memory of the MEM tile may include SRAM instead of DRAM due to the bandwidth requirements of processor 2200. The on-chip memory capacity of processor 2200 may be determined by: (i) the number of ML models that can reside on the chip simultaneously, (ii) the size of any given model, and (iii) partitioning of large models for adaptation to a multi-chip system. In some embodiments, the MEM system of processor 2200 may provide multiple memory slices organized into two distinct hemispheres (referred to as “MEM WEST” and “MEM EAST”, respectively).

[0228] The memory slices in each hemisphere can be mirrored such that the slices are physically numbered {0,...L} in the eastern hemisphere and {L,...0} in the western hemisphere, such that memory slice 0 in each hemisphere corresponds to the slice of the VXM slice closest to the hemisphere, where each hemisphere comprises L slices. Data transfer towards the chip center can be referred to as inward, while data transfer towards the outer edge of the chip (easternmost or westernmost) can be referred to as outward. Although the memory hemispheres of processor 2200 can be referred to as east and west, it is understood that other names may be used to refer to different memory hemispheres in other embodiments.

[0229] In some embodiments, streaming register files (referred to as STREAMS) transfer operands and results between the SRAM of the MEM slices of the processor 2200 and the functional slices. In some embodiments, multiple MEM slices (e.g., 2 to 10 adjacent MEM slices) can be physically organized into sets. Each slice set can be located between a pair of STREAMS register files, allowing each slice to read from or write to the STREAMS registers in either direction. By placing the STREAMS register files between sets of MEM slices, the number of cycles required to transfer data operands across hemispheres can be reduced (e.g., reduced by a factor corresponding to the number of slices per set). The number of slices per set can be configured based on the distance of data transfer within a single clock cycle.

[0230] For any processor in Figure 22A and any component described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. The API (and / or compiled instructions including the API) may be stored in any storage device internal or external to processor 2200 (e.g., in cache and / or memory). The results of the API may then be stored in storage devices internal or external to processor 2200, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0231] In at least one embodiment, processor 2200 may include one or more circuits configured to execute an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0232] Software Structure

[0233] The following figures illustrate, in a non-limiting manner, examples of software structures for implementing at least one embodiment.

[0234] Figure 23 illustrates the software stack of a programming platform according to at least one embodiment. The programming platform may include a platform for leveraging hardware on a computing system to accelerate computational tasks. In at least one embodiment, software developers can access the programming platform through libraries, compiler instructions, and / or extensions to the programming language. The programming platform may be CUDA, Radeon Open Computing Platform (“ROCm”), OpenCL (OpenCL... TM Developed by the Khronos Group), SYCL, or Intel OneAPI.

[0235] The software stack 2300 of the programming platform can provide an execution environment for the application 2301. The application 2301 may include any computer software that can be launched on the software stack 2300. The application 2301 may include artificial intelligence (“AI”) / machine learning (“ML”) applications, high-performance computing (“HPC”) applications, virtual desktop infrastructure (“VDI”), or data center workloads.

[0236] Application 2301 and software stack 2300 run on hardware 2308. Hardware 2308 may include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices supporting programming platforms. Software stack 2300 may be vendor-specific and compatible only with vendor-specific devices, such as CUDA, ROCm, OneAPI, OpenCL, or other implementations. Hardware 2308 may include a host connected to one or more devices that can be accessed via application programming interface (“API”) calls to perform computational tasks. In at least one embodiment, the devices within hardware 2308 may include GPUs, FPGAs, AI engines, or other computing devices (but may also include CPUs) and their memory, while the host within hardware 2308 may include CPUs (but may also include computing devices) and their memory. For any hardware 2308 described above or elsewhere herein, the one or more APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetching logic, decoded by a processor decoder, scheduled (e.g., sequentially or out of order) for execution by a scheduler, executed by execution logic, reordered, and then retired by retirement logic. The API (and / or compiled instructions including the API) can be stored in any storage device (e.g., cache and / or memory) internal or external to the hardware 2308. The results of the API can be stored in storage devices internal or external to the hardware 2308, including registers, DRAM, flash memory, SRAM, cache, or other memory. One or more APIs described herein can receive calls. One or more APIs described herein can communicate with a library or a portion of a library to perform the function described by the call. One or more APIs described herein can receive calls and communicate with a library or a portion of a library to perform the function described by the call.

[0237] The software stack 2300 of the programming platform may include multiple libraries 2303, a runtime 2305, optional drivers / interfaces 2307, and device kernel drivers 2308. Each library 2303 may include data and programming code that can be used by computer programs and utilized during software development. Library 2303 may include pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. Library 2303 may include functionality optimized for execution on one or more types of devices. Library 2303 may include functionality for performing mathematical, deep learning, and / or other types of operations on the device. Library 2303 may be associated with corresponding APIs 2302, which may include one or more APIs for exposing the functionality implemented in library 2303. A processor (e.g., CPU, GPU) may execute, call, or otherwise use one or more APIs to determine kernel priority. For example, a first kernel (e.g., a parent kernel) may launch a second kernel (e.g., a child kernel), and the processor may use the second kernel to launch an additional kernel (e.g., a grandchild kernel) independent of the first kernel. The processor can execute APIs or call APIs from memory to support dynamic stream priorities (e.g., updating priorities when performing operations using streams). For example, when the processor executes the API, it allows the programmer to copy stream priorities from one stream to one or more other streams.

[0238] Software stack 2300 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to set the stream's priority at any time after the stream is created. Software stack 2300 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to obtain the stream's current priority, where the priority is one of several attributes of the stream. Software stack 2300 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to obtain the stream's current priority as a single attribute. Software stack 2300 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to start the kernel to perform operations on the stream at a set priority, which may be different from the stream priority. Software stack 2300 may include an API for indicating whether an object (e.g., a thread synchronization object, such as, but not limited to, a barrier) tracks whether all data movement operations of a set of threads running on the GPU have a specified state after a specified time period, wherein the specified state may be a state indicating that data has been moved and is ready for use, and is specified using expected parity values ​​as input to the API.

[0239] Software stack 2300 may include one or more APIs for updating the kernel. The processor may execute APIs or call APIs from memory to update existing APIs, thereby supporting a context-independent kernel. This allows programmers to add kernel nodes to a graph without a graph context, so that the graph context can be dynamically associated with the kernel at runtime. Software stack 2300 may include one or more APIs that allow programmers to obtain kernel identifiers and graph contexts as separate parameters from kernel nodes, thereby enabling parameter retrieval from both the kernel and the context-independent kernel. Software stack 2300 may include one or more APIs that use parallel processors (e.g., but not limited to one or more graphics processing units) to initiate task graphs (e.g., task graphs) and execute one or more task graphs (e.g., including one or more programs).

[0240] Software stack 2300 may include one or more APIs for associating one or more instructions with one or more memory sorting operations (e.g., but not limited to fence or memory barrier operations). Instructions may be associated with one or more domains, causing memory sorting operations to execute in association with one or more specific domains without interfering with instructions in other domains. APIs may indicate that a thread has reached (e.g., reached a thread synchronization barrier) or completed a certain phase of work associated with an asynchronous data movement operation on the GPU. Software stack 2300 may include one or more APIs that allow programmers to manually indicate an expected transaction count when a thread completes a certain phase of work; this transaction count can be used to update an object used to track whether all data movement operations for a set of threads have been completed.

[0241] Application 2301 can be written as source code and then compiled into executable code, as discussed in more detail below with reference to Figures 24 and 25. The executable code of application 2301 can run, at least partially, within the execution environment provided by software stack 2300. During the execution of application 2301, code that needs to run on a device (rather than a host) may be encountered. In this case, runtime 2305 can be invoked to load and launch the required code on the device. Runtime 2305 can include any technically feasible runtime system capable of supporting the execution of application 2301.

[0242] Runtime 2305 can be implemented as one or more runtime libraries associated with the corresponding API (which is shown as API 2304). One or more such runtime libraries may include functions for memory management, execution control, device management, error handling, and / or synchronization, etc. Memory management functions may include functions for allocating, dealing with, and copying device memory, as well as functions for transferring data between host memory and device memory. Execution control functions may include functions for starting functions on the device (sometimes referred to as the "kernel" when the function is a global function that can be called from the host) and setting attribute values ​​in buffers maintained by the runtime libraries so that a given function can be executed on the device.

[0243] Runtime libraries and corresponding API 2304 can be implemented in any technically feasible manner. One (or any number) APIs can expose a set of low-level functions for fine-grained control of the device, while another (or any number) APIs can expose a set of high-level functions for such functions. High-level runtime APIs can be built on top of low-level APIs. One or more runtime APIs can be language-specific APIs, which can be layered on top of language-independent runtime APIs.

[0244] Optional drivers or interfaces 2307 can be implemented, for example, for CUDA and ROCm implementations, which will be described further below. Optional drivers / interfaces 2307 can be associated with optional driver or interface APIs, such as, but not limited to, the CUDA and / or ROCm APIs.

[0245] One or more processors disclosed in the “processing system” may execute, access, or otherwise use the software stack 2300. For example, the system-on-a-chip 1000, parallel processor 1100, graphics multiprocessor 1134, processor 1200, processor 1300, accelerator 1400, neuromorphic processor 1505, supercomputer 1600, acceleration processing unit 1700, processor 1800, processor 1900, tensor processing unit 2000, processor 2100, and language processing unit 2200 may execute, use, call, or otherwise implement (e.g., by accessing memory) one or more APIs included in the software stack 2300.

[0246] Device kernel driver 2308 can be configured to facilitate communication with underlying devices. Device kernel driver 2308 can provide low-level functionality for APIs (such as, but not limited to, API 2304) and / or other software. Device kernel driver 2308 can be configured to compile intermediate representation (“IR”) code into binary code at runtime. For CUDA or other implementations (such as, but not limited to, ROCm, OneAPI, or OpenCL), device kernel driver 2308 can compile non-hardware-specific parallel thread execution (“PTX”) IR code into binary code for a specific target device at runtime (and cache the compiled binary code), which is sometimes referred to as “finalized” code. Doing so allows the finalized code to run on a target device that may not have existed when the source code was initially compiled into PTX code. Alternatively, device source code can be compiled into binary code offline without device kernel driver 2308 compiling the IR code at runtime.

[0247] The processor described elsewhere in this document (e.g., but not limited to the processor in Figures 10-22B) may include one or more circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere in this document. The one or more circuitry may be configured by software (e.g., software stack 2300) to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere in this document.

[0248] According to at least one embodiment, the software stack 2300 in FIG23 can be executed in a CUDA implementation. The CUDA software stack 2300, on which application 2301 can be launched, may include a CUDA library 2303, a CUDA runtime 2305, a CUDA driver 2307, and a device kernel driver 2308. The CUDA software stack 2300 can be executed on hardware (e.g., a graphics multiprocessor 1134, which may include a CUDA-enabled GPU developed by NVIDIA Corporation, Santa Clara, California).

[0249] Application 2301, CUDA runtime 2305, and device kernel driver 2308 can perform the functions described above and elsewhere in this document. CUDA driver 2307 may include a library (libcuda.so) that implements CUDA driver API 2306. Similar to CUDA runtime API 2304 implemented by the CUDA runtime library (cudart), CUDA driver API 2306 exposes functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability. CUDA driver API 2306 differs from CUDA runtime API 2304 in that CUDA runtime API 2304 simplifies device code management by providing implicit initialization, context (similar to processes) management, and module (similar to dynamically loaded libraries) management. Compared to the high-level CUDA runtime API 2304, CUDA driver API 2306 can serve as a low-level API, providing finer-grained device control, especially in terms of context and module loading. CUDA driver API 2306 can expose context management functions not exposed in CUDA runtime API 2304. CUDA driver API 2306 can also be language-agnostic, supporting technologies such as OpenCL in addition to CUDA runtime API 2304. Furthermore, development libraries, including CUDA runtime 2305, can be considered separate from driver components, including user-mode CUDA driver 2307 and kernel-mode device driver 2308 (sometimes referred to as the "display" driver).

[0250] CUDA library 2303 may include mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries, which parallel computing applications (such as, but not limited to, application 2301) may utilize. CUDA library 2303 may include mathematical libraries, such as, but not limited to, the cuBLAS library (an implementation of the basic linear algebra subroutine (“BLAS”) for performing linear algebra operations), the cuFFT library (for computing the Fast Fourier Transform (“FFT”)), and the cuRAND library (for generating random numbers), etc. CUDA library 2303 may include deep learning libraries, such as, but not limited to, the cuDNN primitive library for deep neural networks and the TensorRT platform for high-performance deep learning inference, etc.

[0251] In at least one embodiment, the processor described elsewhere herein (e.g., but not limited to the processor in Figures 10-22B) may include one or more circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuitry may be configured by software (e.g., software stack 2300) to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0252] According to at least one embodiment, the software stack 2300 in Figure 23 can be executed in a ROCm implementation. An application 2301 can be launched on the ROCm software stack 2300, which includes a language runtime 2303, a system runtime 2305, a thunk 2307, and a ROCm kernel driver 2308. The ROCm software stack 2300 executes on hardware 2309, which may include a ROCm-enabled GPU developed by AMD Inc. of Santa Clara, California.

[0253] Application 2301 can perform functions similar to those discussed above in conjunction with Figure 23. Furthermore, language runtime 2303 and system runtime 2305 can perform functions similar to those discussed above in conjunction with Figure 23. The difference between language runtime 2303 and system runtime 2305 is that system runtime 2305 is a language-independent runtime that implements the ROCr System Runtime API 2304 and uses the Heterogeneous System Architecture (“HSA”) runtime API. The HSA runtime API may include a streamlined user-mode API that exposes interfaces for accessing and interacting with the AMD GPU, including functions for memory management, execution control via kernel architecture dispatch, error handling, system and agent information, and runtime initialization and shutdown. Unlike system runtime 2305, language runtime 2303 may be an implementation of language-specific runtime API 2302, which resides in a layer above the ROCr System Runtime API 2304. The language runtime API may include the Heterogeneous Computing Portable Interface (“HIP”) language runtime API, the Heterogeneous Computing Compiler (“HCC”) language runtime API, or the OpenCL API, etc. The HIP language is an extension of the C++ programming language, having a functionally similar version to the CUDA mechanism, and the HIP language runtime API may include functions similar to the CUDA runtime API discussed above in conjunction with Figure 23, such as, but not limited to, memory management, execution control, device management, error handling, and synchronization functions.

[0254] The Thunk(ROCt)2307 can be interface 2306, which is used to interact with the underlying ROCm driver 2308. The ROCm driver 2308 can be the ROCk driver, a combination of the AMD GPU driver and the HSA core driver (amdkfd). The AMD GPU driver can be a device core driver developed by AMD for GPUs, performing functions similar to the device core driver 2309 discussed above in conjunction with Figure 23. The HSA core driver can be a driver that allows different types of processors to share system resources more efficiently through hardware features.

[0255] Various libraries (not shown) may be included in the ROCm software stack 2300 on top of the language runtime 2303, providing functionality similar to the CUDA library 2303 discussed above in conjunction with Figure 23. These libraries may include mathematical libraries, deep learning libraries, and / or other libraries, such as, but not limited to, the hipBLAS library which implements functionality similar to CUDA cuBLAS, the rocFFT library for computing FFTs similar to CUDA cuFFT, etc.

[0256] The processor described elsewhere in this document (e.g., but not limited to the processor in Figures 10-22B) may include one or more circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere in this document. The one or more circuitry may be configured by software (e.g., software stack 2300) to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere in this document.

[0257] According to at least one embodiment, the software stack 2300 in FIG23 can be executed in an OpenCL implementation. The OpenCL software stack 2300, on which the application 2301 can be launched, may include an OpenCL framework 2303, an OpenCL runtime 2305, and a driver 2308. The OpenCL software stack 2300 can be executed on non-vendor-specific hardware 2309. Because devices developed by different vendors support OpenCL, specific OpenCL drivers may be required for interoperability with the hardware of these vendors.

[0258] Application 2301, OpenCL runtime 2305, device kernel driver 2308, and hardware 2309 can perform functions similar to other implementations of application 2301, runtime 2305, device kernel driver 2308, and hardware 2309 discussed above in conjunction with Figure 23. Application 2301 may also include an OpenCL kernel (not shown), whose code will be executed on the device.

[0259] OpenCL can define a "platform" that allows a host to control devices connected to it. The OpenCL framework provides platform-level APIs and runtime APIs, shown as Platform API 2302 and Runtime API 2304, respectively. Runtime API 2304 uses contexts to manage kernel execution on devices. Each identified device can be associated with a corresponding context, which Runtime API 2304 uses to manage the device's command queue, program objects, kernel objects, shared memory objects, and so on. Platform API 2302 exposes functions that allow the use of device contexts to select and initialize devices, submit work to devices via command queues, and enable data transfer with devices. In addition, the OpenCL framework provides various built-in functions (not shown), including mathematical functions, relational functions, and image processing functions.

[0260] The OpenCL framework 2303 may also include a compiler (not shown). Source code can be compiled offline before application execution or online during application execution. Unlike CUDA and ROCm, OpenCL applications can be compiled online by a compiler representing any number of compilers that can be used to compile source code and / or IR code (e.g., but not limited to Standard Portable Intermediate Representation (“SPIR-V”) code) into binary code. Alternatively, OpenCL applications can be compiled offline before execution.

[0261] In at least one embodiment, the processor described elsewhere herein (e.g., but not limited to the processor in Figures 10-22B) may include one or more circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuitry may be configured by software (e.g., software stack 2300) to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0262] According to at least one embodiment, the software may be supported by a programming platform configured to support various programming models, middleware and / or libraries, and frameworks that the application may rely on. The application may be an AI / ML application implemented using, for example, a deep learning framework (e.g., but not limited to, MXNet, PyTorch, or TensorFlow), which may rely on libraries such as, but not limited to, cuDNN, the NVIDIA Collective Communication Library (“NCCL”), and / or the NVIDIA Developer Data Loading Library (“DALI”) CUDA library to provide accelerated computation on the underlying hardware.

[0263] The programming platform can be one of the CUDA, ROCm, or OpenCL platforms described above in conjunction with Figure 23. The programming platform can support various programming models, which can be abstractions of the underlying computing system that allow the expression of algorithms and data structures. Programming models can expose features of the underlying hardware to improve performance. Programming models may include CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++AMP”), Open Multiprocessing (“OpenMP”), Open Accelerators (“OpenACC”), and / or Vulcan Compute.

[0264] Libraries and / or middleware can provide abstract implementations of programming models. Such libraries may include data and programming code that computer programs can use and leverage during software development. Such middleware may include software that provides services to applications beyond those offered by the programming platform. Libraries and / or middleware may include cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Furthermore, libraries and / or middleware may include the NCCL and ROCm communication collection library (“RCCL”) libraries that provide communication routines for GPUs, the MIOpen library for accelerating deep learning, and / or the Eigen library for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0265] Application frameworks may depend on libraries and / or middleware. Each application framework can be a software framework that provides a standard structure for implementing application software. Returning to the AI / ML example discussed above, AI / ML applications can be implemented using frameworks such as, but not limited to, deep learning frameworks like Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet.

[0266] In at least one embodiment, the processor described elsewhere herein (e.g., but not limited to the processor in Figures 10-22B) may include one or more circuitry for executing an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuitry may be configured by software (e.g., the programming platform described herein) to execute the application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations, or otherwise perform any of the operations described above or elsewhere herein.

[0267] Figure 24 illustrates compiled code for execution on one of the programming platforms shown in Figure 23 above, according to at least one embodiment. Compiler 2401 is configured to receive source code 2400, compile source code 2400, and output an executable file 2410. Compiler 2401 can be configured to convert source code 2400 into host executable code 2407 for execution on a host and device executable code 2408 for execution on a device. Source code 2400 can be compiled offline before executing the application or online during application execution. Source code 2400 can include code in any programming language supported by compiler 2401, such as, but not limited to, C++, C, Fortran, etc. Source code 2400 can be included in a single source file containing both host code and device code, with the location of the device code indicated. The single source file can be a .cu file including CUDA code, or a .hip.cpp file including HIP code, or a file in other formats including both host code and device code. Alternatively, source code 3300 may include multiple source code files instead of a single source file, with host code and device code separated into these files. Compiler 2401 includes or can access one or more libraries to identify API call sequences to execute a single fused API, where the single fused API is a combination of two or more APIs. In at least one embodiment, compiler 2401 may be an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code in .cu files, or an HCC compiler for compiling HIP code in .hip.cpp files, or other compilers.

[0268] Compiler 2401 can be configured to compile source code 2400 into host executable code 2407 for execution on a host and device executable code 2408 for execution on a device. The operations performed by compiler 2401 include parsing source code 2400 into an abstract system tree (AST), performing optimizations, and generating executable code. When source code 2400 comprises a single source file, compiler 2401 can separate the device code from the host code in that single source file, compile the device code and host code into device executable code 2408 and host executable code 2407 respectively, and link the device executable code 2408 and host executable code 2407 together to form a single file.

[0269] Compiler 2401 may include compiler front-end 2402, host compiler 2405, device compiler 2406, and linker 2409. Compiler front-end 2402 may be configured to separate device code 2404 from host code 2403 in source code 2400. In at least one embodiment, device code 2404 may be compiled by device compiler 2406 into device executable code 2408, which may include binary code or IR code as described above. Separately, host code 2403 may be compiled by host compiler 2405 into host executable code 2407. For other compilers such as NVCC (e.g., but not limited to oneAPI, ROCm, and OpenCL compilers), host compiler 2405 may be a general-purpose C / C++ compiler that outputs native object code, while device compiler 2406 may be a low-level virtual machine (“LLVM”) based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code. For HCC, both host compiler 2405 and device compiler 2406 may be LLVM based compilers that output object binary code.

[0270] After compiling source code 2400 into host executable code 2407 and device executable code 2408, linker 2409 can link the host executable code 2407 and device executable code 2408 together to form executable file 2410. The host's native object code and the device's PTX or binary code can be linked together in an executable and linkable format (“ELF”) file, a container format for storing object code. Host executable code 2407 and device executable code 2408 can take any suitable format, such as, but not limited to, binary code and / or IR code. In at least one embodiment, for CUDA, host executable code 2407 may include native object code, while device executable code 2408 may include code in a PTX intermediate representation. In at least one embodiment, for ROCm, both host executable code 2407 and device executable code 2408 can include object binary code. Other implementations (e.g., but not limited to oneAPI, OpenCL) are considered and can be performed similarly to the CUDA and ROCm implementations described above.

[0271] Source code 2400 can be transformed before compilation. The source code is passed through a transformation tool (not shown) that transforms source code 2400 into transformed source code. Compiler 2401 can be used to compile the transformed source code into host executable code 2407 and device executable code 2408, in a process similar to compiler 2401 compiling source code 2400 into host executable code 2407 and device executable code 2408, as described above in conjunction with Figure 24.

[0272] The conversions performed by the conversion tool can be used to port source code 2400 to an environment different from the one originally intended to run in. The conversion tool may include a HIP converter, which "hipify" CUDA code intended for the CUDA platform into HIP code that can be compiled and executed on the ROCm platform. The conversion of source code 2400 may include parsing source code 2400 and converting calls to APIs provided by one programming model (e.g., CUDA) into corresponding calls to APIs provided by another programming model (e.g., HIP), as discussed in more detail below with reference to Figure 25. Returning to the example of HIPing CUDA code, calls to CUDA runtime APIs, CUDA driver APIs, and / or CUDA libraries can be converted into corresponding HIP API calls. The automatic conversions performed by conversion tool 2401 may sometimes be incomplete, requiring additional manual intervention to fully port source code 2400.

[0273] One or more techniques described herein can leverage other methods of converting one type of code into another to achieve interoperability between different device architectures. In at least one embodiment, an application for one platform (e.g., a CUDA application) can be compiled into code for implementation on another platform (e.g., an AMD processor, an Intel processor, or another processor). For example, source code 2400 may include source code for one platform (e.g., CUDA). Compiler 2401 can compile source code 2400 into an executable file 2410 that can be used by another platform (e.g., AMD or Intel). Programming toolkits can allow applications for one platform (e.g., CUDA) to be compiled (e.g., natively compiled) against another platform (e.g., AMD or Intel). For example, a GPGPU programming toolkit can allow CUDA applications to be natively compiled against AMD GPUs. Programs (e.g., CUDA programs) or their build systems do not require modification or conversion into other languages ​​before being compiled into code for another platform. The compiler can accept the same command-line options and programming cognates (e.g., CUDA cognates) as another compiler (e.g., nvcc for CUDA), serving as a direct replacement for the installation of emulation toolkits (e.g., the NVIDIA CUDA Toolkit), so existing build tools and scripts (e.g., cmake) will work without further modification. In at least one embodiment, an nvcc-cognate CUDA can be compiled for AMD GPUs (including PTX asm) using an nvcc-compatible compiler. Implementations of the CUDA runtime and driver APIs for AMD GPUs can be used. Libraries (e.g., open-source wrapper libraries) can provide APIs by delegating to the corresponding ROCm libraries, such as the "CUDA-X" API. Example implementations include SCALE from Spectral Compute in London, UK. SCALE does not provide a new method for writing GPGPU software, but rather allows programs written in the popular CUDA language to be directly compiled for AMD GPUs. Additional implementations could include the Clang compiler, which provides language front-ends and tooling infrastructure for languages ​​in the C family of languages ​​(C, C++, Objective C / C++, OpenCL, CUDA, and RenderScript).In at least one embodiment, the compiler described herein (e.g., but not limited to compiler 2401, compiler 2405, and / or compiler 2406) may include one or more circuitry for compiling code (e.g., CUDA, HIP, OpenCL, OneAPI, or others) to execute an application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations based at least in part on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations, and / or any of the operations described above or elsewhere herein.

[0274] Figure 25 illustrates a system 2500 configured to compile and execute CUDA source code 2510 using different types of processing units according to at least one embodiment. The system 2500 includes CUDA source code 2510, CUDA compiler 2550, host executable code 2570(1), host executable code 2570(2), CUDA device executable code 2584, CPU 2590, CUDA-enabled GPU 2594, GPU 2592, CUDA to HIP conversion tool 2520, HIP source code 2530, HIP compiler driver 2540, HCC 2560, and HCC device executable code 2582.

[0275] CUDA source code 2510 can be a collection of human-readable code in the CUDA programming language. The CUDA programming language can be an extension of the C++ programming language, including mechanisms for defining device code and distinguishing between device code and host code. Device code can include source code that, after compilation, can be executed in parallel on a device. A device can be a processor optimized for parallel instruction processing, such as, but not limited to, a CUDA-enabled GPU 2590, GPU 2592, or other GPGPUs. Host code is source code that, after compilation, can be executed on a host machine. A host machine is a processor optimized for sequential instruction processing, such as, but not limited to, a CPU 2590.

[0276] CUDA source code 2510 may include any number (including zero) of global functions 2512, any number (including zero) of device functions 2514, any number (including zero) of host functions 2516, and any number (including zero) of host / device functions 2518. Global functions 2512, device functions 2514, host functions 2516, and host / device functions 2518 can be mixed within CUDA source code 2510. Each global function 2512 can be executed on the device and called from the host. Therefore, one or more global functions 2512 can act as entry points for the device. Each global function 2512 can be a kernel. In a technique called dynamic parallelism, one or more global functions 2512 can define a kernel that can be executed on the device and called from the device. During execution, the kernel can be executed in parallel N times by N different threads on the device (where N is any positive integer).

[0277] Each device function 2514 can be executed on a device and can only be called from that device. Each host function 2516 can be executed on a host and can only be called from that host. Each host / device function 2516 can define a host version of a function that can be executed on a host and called only from that host, and a device version of a function that can be executed on a device and called only from that device.

[0278] CUDA source code 2510 can also include any number of calls to any number of functions, which can be defined through CUDA runtime API 2502. CUDA runtime API 2502 can include any number of functions that execute on the host to allocate and deallocate device memory, transfer data between host memory and device memory, manage systems with multiple devices, etc. CUDA source code 2510 can also include any number of calls to any number of functions, which can be specified in any number of other CUDA APIs. CUDA APIs can be any API designed for use by CUDA code. CUDA APIs can include CUDA runtime API 2502, CUDA driver APIs, APIs for any number of CUDA libraries, etc., including any APIs described elsewhere in this document. Compared to CUDA runtime API 2502, CUDA driver APIs can be lower-level APIs but can provide more fine-grained control over devices. Examples of CUDA libraries include cuBLAS, cuFFT, cURAND, cuDNN, etc.

[0279] The CUDA compiler 2550 can compile input CUDA code (e.g., CUDA source code 2510) to generate host executable code 2570(1) and CUDA device executable code 2584. The CUDA compiler 2550 can be, but is not limited to, an NVCC. The host executable code 2570(1) can be a compiled version of the host code included in the input source code, which can be executed on the CPU 2590. The CPU 2590 can be any processor optimized for sequential instruction processing.

[0280] CUDA device executable code 2584 may be a compiled version of the device code included in the input source code, which can be executed on a CUDA-enabled GPU 2594. CUDA device executable code 2584 may include binary code. CUDA device executable code 2584 may include IR code (e.g., but not limited to PTX code), which is further compiled at runtime by the device driver into binary code for a specific target device (e.g., a CUDA-enabled GPU 2594). CUDA-enabled GPU 2594 may include any processor optimized for parallel instruction processing and supporting CUDA. CUDA-enabled GPU 2594 may be developed by NVIDIA Corporation, located in Santa Clara, California.

[0281] The CUDA to HIP conversion tool 2520 can be configured to convert CUDA source code 2510 into functionally similar HIP source code 2530. The HIP source code 2530 may include a collection of human-readable code written in the HIP programming language. The HIP code may include human-readable code written in the HIP programming language. The HIP programming language may include extensions to the C++ programming language that include functionally similar versions of the CUDA mechanisms for defining device code and distinguishing it from host code. The HIP programming language may include a subset of the functionality of the CUDA programming language. For example, the HIP programming language may include a mechanism for defining global functions 2512, but such HIP programming languages ​​may lack support for dynamic parallelism, so the global function 2512 defined in the HIP code may only be called from the host.

[0282] HIP source code 2530 may include any number (including zero) of global functions 2512, any number (including zero) of device functions 2514, any number (including zero) of host functions 2516, and any number (including zero) of host / device functions 2518. HIP source code 2530 may also include any number of calls to any number of functions specified in the HIP runtime API 2532. The HIP runtime API 2532 may include functionally similar versions of a subset of functions contained in the CUDA runtime API 2502. HIP source code 2530 may also include any number of calls to any number of functions specified in any number of other HIP APIs. The HIP API can be any API designed for use by HIP code and / or ROCm. The HIP API may include the HIP runtime API 2532, the HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, etc.

[0283] The CUDA to HIP conversion tool 2520 can convert each kernel call in CUDA code from CUDA syntax to HIP syntax, and can convert any number of other CUDA calls in CUDA code into any number of other functionally similar HIP calls. CUDA calls can include calls to functions specified in the CUDA API, and HIP calls can include calls to functions specified in the HIP API. The CUDA to HIP conversion tool 2520 can convert any number of calls to functions specified in the CUDA runtime API 2502 into any number of calls to functions specified in the HIP runtime API 2532.

[0284] The CUDA to HIP conversion tool 2520 may include a tool called hipify-perl, which performs a text-based conversion process. The CUDA to HIP conversion tool 2520 may also include a tool called hipify-clang, which, compared to hipify-perl, performs a more complex and robust conversion process, including parsing the CUDA code using clang (a compiler front-end) and then converting the resulting symbols. Converting CUDA code to HIP code may include modifications beyond those performed by the CUDA to HIP conversion tool 2520 (e.g., manual editing).

[0285] HIP compiler driver 2540 may include a front-end that determines target device 2546 and then configures a compiler compatible with target device 2546 to compile HIP source code 2530. Target device 2546 may include a processor optimized for parallel instruction processing. HIP compiler driver 2540 may determine target device 2546 in any technically feasible manner.

[0286] If the target device 2546 is CUDA compatible (e.g., a CUDA-enabled GPU 2594), the HIP compiler driver 2540 can generate HIP / NVCC compilation commands 2542. The HIP / NVCC compilation commands 2542 can configure the CUDA compiler 2550 to compile HIP source code 2530 using a HIP-to-CUDA translation header and the CUDA runtime library. In response to the HIP / NVCC compilation commands 2542, the CUDA compiler 2550 can generate host executable code 2570(1) and CUDA device executable code 2584.

[0287] If the target device 2546 is incompatible with CUDA, the HIP compiler driver 2540 can generate HIP / HCC compilation command 2544. HIP / HCC compilation command 2544 can configure HCC 2560 to compile HIP source code 2530 using the HCC header and HIP / HCC runtime library. In response to HIP / HCC compilation command 2544, HCC 2560 can generate host executable code 2570(2) and HCC device executable code 2582. HCC device executable code 2582 can be a compiled version of the device code included in HIP source code 2530, which can be executed on GPU 2592. GPU 2592 can be any processor optimized for parallel instruction processing, incompatible with CUDA, and compatible with HCC. GPU 2592 can be developed by AMD Inc., located in Santa Clara, California. GPU 2592 can include GPU 2592 that does not support CUDA.

[0288] For illustrative purposes only, Figure 25 illustrates three different flows that can be implemented in at least one embodiment for compiling CUDA source code 2510 for execution on CPU 2590 and various devices. A direct CUDA stream can compile CUDA source code 2510 for execution on CPU 2590 and CUDA-enabled GPU 2594 without converting CUDA source code 2510 to HIP source code 2530. An indirect CUDA stream can convert CUDA source code 2510 to HIP source code 2530 and then compile HIP source code 2530 for execution on CPU 2590 and CUDA-enabled GPU 2594. A CUDA / HCC stream can convert CUDA source code 2510 to HIP source code 2530 and then compile HIP source code 2530 for execution on CPU 2590 and GPU 2592.

[0289] The achievable direct CUDA stream is represented by dashed lines and a series of bubbles labeled A1-A3. As shown in bubble A1, CUDA compiler 2550 can receive CUDA source code 2510 and CUDA compilation command 2548, which configures CUDA compiler 2550 to compile CUDA source code 2510. The CUDA source code 2510 available for direct CUDA stream can be written in a CUDA programming language based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In response to CUDA compilation command 2548, CUDA compiler 2550 can generate host executable code 2570(1) and CUDA device executable code 2584 (shown in bubble A2). As shown in bubble A3, host executable code 2570(1) and CUDA device executable code 2584 can be executed on CPU 2590 and CUDA-enabled GPU 2594, respectively. CUDA device executable code 2584 may include binary code. CUDA device executable code 2584 can include PTX code and can be further compiled at runtime into binary code for a specific target device.

[0290] The achievable indirect CUDA streams are represented by dashed lines and a series of bubbles labeled B1-B6. As shown in bubble B1, the CUDA-to-HIP conversion tool 2520 can receive CUDA source code 2510. As shown in bubble B2, the CUDA-to-HIP conversion tool 2520 can convert CUDA source code 2510 into HIP source code 2530. As shown in bubble B3, the HIP compiler driver 2540 can receive HIP source code 2530 and determine that the target device 2546 supports CUDA.

[0291] As indicated by bubble B4, the HIP compiler driver 2540 can generate HIP / NVCC compilation command 2542 and transfer HIP / NVCC compilation command 2542 and HIP source code 2530 to CUDA compiler 2550. HIP / NVCC compilation command 2542 can configure CUDA compiler 2550 to compile HIP source code 2530 using a HIP-to-CUDA translation header and a CUDA runtime library. The HIP-to-CUDA translation header can convert 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. CUDA compiler 2550 can combine the HIP-to-CUDA translation header with the CUDA runtime library corresponding to CUDA runtime API 2502 to generate host executable code 2570(1) and CUDA device executable code 2584. In response to HIP / NVCC compilation command 2542, CUDA compiler 2550 can generate host executable code 2570(1) and CUDA device executable code 2584 (represented by bubble B5). As shown in bubble B6, host executable code 2570(1) and CUDA device executable code 2584 can be executed on CPU 2590 and CUDA-enabled GPU 2594, respectively. CUDA device executable code 2584 may include binary code. CUDA device executable code 2584 may include PTX code and can be further compiled at runtime into binary code for a specific target device.

[0292] The implementable CUDA / HCC streams are represented by solid lines and a series of bubbles labeled C1-C6. As shown in bubble C1, the CUDA-to-HIP conversion tool 2520 can receive CUDA source code 2510. As shown in bubble C2, the CUDA-to-HIP conversion tool 2520 can convert CUDA source code 2510 into HIP source code 2530. As shown in bubble C3, the HIP compiler driver 2540 can receive HIP source code 2530 and can determine that the target device 2546 does not support CUDA.

[0293] HIP compiler driver 2540 can generate HIP / HCC compilation command 2544 and transfer HIP / HCC compilation command 2544 and HIP source code 2530 to HCC 2560 (shown as bubble C4). HIP / HCC compilation command 2544 can configure HCC 2560 to compile HIP source code 2530 using HCC header files and HIP / HCC runtime library. HIP / HCC runtime library can correspond to HIP runtime API 2532. HCC header can include any number and type of HIP and HCC interoperability mechanisms. In response to HIP / HCC compilation command 2544, HCC 2560 can generate host executable code 2570(2) and HCC device executable code 2582 (shown as bubble C5). As shown in the bubble labeled C6, host executable code 2570(2) and HCC device executable code 2582 can be executed on CPU 2590 and GPU 2592, respectively.

[0294] After converting CUDA source code 2510 to HIP source code 2530, executable code for a CUDA-enabled GPU 2594 or GPU 2592 can then be generated using HIP compiler driver 2540 without re-executing the CUDA to HIP conversion tool 2520. CUDA to HIP conversion tool 2520 can convert CUDA source code 2510 to HIP source code 2530 and then store the HIP source code 2530 in memory. HIP compiler driver 2540 can then configure HCC 2560 to generate host executable code 2570(2) and HCC device executable code 2582 based on the HIP source code 2530. In at least one embodiment, HIP compiler driver 2540 subsequently configures CUDA compiler 2550 to generate host executable code 2570(1) and CUDA device executable code 2584 based on the stored HIP source code 2530.

[0295] According to at least one embodiment, the example kernel can be converted by the CUDA to HIP conversion tool 2520 in Figure 25. The CUDA source code 2510 divides the overall problem that a given kernel is designed to solve into relatively coarse subproblems, which can be solved independently using thread blocks. Each thread block includes any number of threads. Each subproblem can be divided into relatively fine pieces, which can be solved collaboratively in parallel by threads within the thread block. Threads within a thread block can cooperate by sharing data via shared memory and by synchronizing execution to coordinate memory access.

[0296] CUDA source code 2510 can organize thread blocks associated with a given kernel into a one-dimensional, two-dimensional, or three-dimensional thread block grid. Each thread block contains any number of threads, and the grid contains any number of thread blocks.

[0297] A kernel can be a function defined in device code using the "__global__" declaration specifier. The dimensions of the raster of kernels and their associated streams that execute a given kernel call can be specified using the CUDA kernel startup syntax. The CUDA kernel startup syntax is specified as "KernelName <<<GridSize,BlockSize,SharedMemorySize,Stream> >>

[0298] (KernelArguments);". The execution configuration syntax can include a "<<<...>>>" structure between the kernel name ("KernelName") and the parenthesized list of kernel parameters ("KernelArguments"). The CUDA kernel boot syntax can include the CUDA boot function syntax, instead of the execution configuration syntax.

[0299] "GridSize" can be of type dim3 and specifies the dimensions and size of the grid. The dim3 type can be a CUDA-defined structure containing unsigned integers x, y, and z. If z is not specified, it defaults to 1. If y is not specified, it defaults to 1. The number of thread blocks in the grid can be equal to the product of GridSize.x, GridSize.y, and GridSize.z. "BlockSize" can be of type dim3 and specifies the dimensions and size of each thread block. The number of threads per thread block can be equal to the product of BlockSize.x, BlockSize.y, and BlockSize.z. Each thread executing the kernel can be assigned a unique thread ID, which can be accessed in the kernel via built-in variables such as "threadIdx".

[0300] Regarding the CUDA kernel startup syntax, "SharedMemorySize" is an optional argument that, in addition to statically allocated memory, can specify the number of bytes in shared memory dynamically allocated per thread block for a given kernel call. The default value for SharedMemorySize is zero. Regarding the CUDA kernel startup syntax, "Stream" is also an optional argument that specifies the associated stream; the default value is zero to specify the default stream. A stream can be a sequence of commands executed sequentially (which may be emitted by different host threads). Different streams can execute commands out of order or concurrently.

[0301] CUDA source code 2510 may include the kernel definition and main function of the example kernel "MatAdd". The main function may be host code that executes on the host machine and includes a kernel call that causes the MatAdd kernel to execute on the device. The MatAdd kernel can add two NxN matrices A and B, where N is a positive integer, and store the result in matrix C. The main function can define the threadsPerBlock variable as 16x16 and the numBlocks variable as N / 16 x N / 16. Then, the main function can specify the kernel call "MatAdd <<<numBlocks,threadsPerBlock> >>(A,B,C);”. According to the CUDA kernel startup syntax, the kernel MatAdd can be executed using a thread block grid of dimensions N / 16 x N / 16, where each thread block is 16x16. Each thread block can contain 256 threads, and a grid with enough thread blocks can be created so that there is one thread per matrix element, and each thread in such a grid can execute the kernel MatAdd to perform a pairwise addition.

[0302] When converting CUDA source code 2510 to HIP source code 2530, the CUDA-to-HIP conversion tool 2520 can convert each kernel call in CUDA source code 2510 from CUDA kernel startup syntax to HIP kernel startup syntax, and can convert any number of other CUDA calls in source code 2510 into any number of other functionally similar HIP calls. The HIP kernel startup syntax can be specified as "hipLaunchKernelGGL(KernelName,GridSize,BlockSize,SharedMemorySize,Stream,KernelArguments);". Each of KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments has the same meaning in the HIP kernel startup syntax as it does in the CUDA kernel startup syntax (described earlier in this document). The arguments SharedMemorySize and Stream can be required in the HIP kernel startup syntax, but optional in the CUDA kernel startup syntax.

[0303] A portion of HIP source code 2530 can be identical to a portion of the CUDA source code 2510 shown, except for the kernel call that causes the kernel MatAdd to execute on the device. The kernel MatAdd can be defined in HIP source code 2530 using the same "__global__" declaration specifier as used in defining the kernel MatAdd in CUDA source code 2510. The kernel call in HIP source code 2530 could be "hipLaunchKernelGGL(MatAdd,numBlocks,threadsPerBlock,0,0,A,B,C);", while the corresponding kernel call in CUDA source code 2510 is "MatAdd<<<numBlocks,threadsPerBlock> >>(A,B,C);

[0304] Other implementations are conceivable, and these can be implemented similarly to the CUDA and HIP implementations described above, such as oneAPI, OpenCL, and other programming platforms. Code can be converted in any direction. For example, CUDA can be converted to HIP, and CUDA can be converted to OpenCL. SnuCL-Tr and CUCL can be used to convert OpenCL to CUDA or CUDA to OpenCL, respectively. Compiled code or intermediate representations (such as CUDA PTX code) can also be converted to run on other processor platforms (such as AMD or Intel). For example, conversion tools (such as ZLUDA) can be used to convert PTX code to run on Intel or AMD processors.

[0305] The techniques described herein can utilize the oneAPI programming model. The oneAPI programming model can refer to a programming model used to interact with various computing accelerator architectures. OneAPI can refer to an application programming interface (API) designed to interact with various computing accelerator architectures. The oneAPI programming model can use the DPC++ programming language. The DPC++ programming language can refer to a high-level language used for data-parallel programming productivity. The DPC++ programming language can be at least partially based on the C and / or C++ programming languages. The oneAPI programming model can be, but is not limited to, a programming model developed by Intel Corporation of Santa Clara, California.

[0306] OneAPI and / or the oneAPI programming model can be used to interact with a variety of accelerators, GPUs, processors, and / or their variant architectures. OneAPI may include a set of libraries that implement various functions. OneAPI may include at least the oneAPIDPC++ library, the oneAPI math kernel library, the oneAPI data analysis library, the oneAPI deep neural network library, the oneAPI collection communication library, the oneAPI thread building block library, the oneAPI video processing library, and / or their variants.

[0307] The oneAPIDPC++ library (also known as oneDPL) can be a library that implements algorithms and functions to accelerate DPC++ kernel programming. oneDPL can implement one or more Standard Template Library (STL) functions. OneDPL can implement one or more parallel STL functions. OneDPL can provide a set of library classes and functions, such as, but not limited to, parallel algorithms, iterators, function object classes, range-based APIs, and / or variations thereof. OneDPL can implement one or more classes and / or functions from the C++ Standard Library. OneDPL can implement one or more random number generator functions.

[0308] The oneAPI math kernel library (also known as oneMKL) can be a library that implements various optimized and parallelized routines for various mathematical functions and / or operations. OneMKL can implement one or more Basic Linear Algebra Subroutines (BLAS) and / or Linear Algebra Packages (LAPACK) dense linear algebra routines. OneMKL can implement one or more sparse BLAS linear algebra routines. OneMKL can implement one or more random number generators (RNGs). OneMKL can implement one or more vector math (VM) routines for performing mathematical operations on vectors. OneMKL can implement one or more Fast Fourier Transform (FFT) functions.

[0309] The OneAPI data analysis library (also known as oneDAL) can include libraries for implementing various data analysis applications and distributed computing. OneDAL can implement various algorithms for data analysis preprocessing, transformation, analysis, modeling, validation, and decision-making, including batch, online, and distributed computing processing modes. OneDAL can implement various C++ and / or Java APIs as well as various connectors for connecting to one or more data sources. OneDAL can implement the DPC++ API extension to the traditional C++ interface and support GPU usage for various algorithms.

[0310] The OneAPI Deep Neural Network Library (also known as oneDNN) can include libraries that implement various deep learning functions. OneDNN can implement various neural networks, machine learning and deep learning functions, algorithms and / or variations thereof.

[0311] The OneAPI collection communication library (also known as oneCCL) can include libraries that implement a variety of applications for deep learning and machine learning workloads. OneCCL can be built on lower-level communication middleware, such as, but not limited to, message passing interfaces (MPI) and libfabrics. OneCCL can support a set of deep learning-specific optimizations, such as, but not limited to, priority ordering, persistence operations, out-of-order execution, and / or variations thereof. OneCCL can implement a variety of CPU and GPU functionalities.

[0312] The OneAPI Thread Building Blocks library (also known as oneTBB) can include libraries that implement various parallel processes for a wide range of applications. OneTBB can be used for task-based shared parallel programming on a host machine. OneTBB can implement general-purpose parallel algorithms. OneTBB can implement concurrent containers. OneTBB can implement scalable memory allocators. OneTBB can implement work-stealing task schedulers. OneTBB can implement low-level synchronization primitives. OneTBB can be compiler-independent and can be used on a variety of processors, such as, but not limited to, GPUs, PPUs, CPUs, and / or variants thereof.

[0313] The OneAPI video processing library (also known as oneVPL) can include libraries for accelerating video processing in one or more applications. OneVPL can implement a variety of video decoding, encoding, and processing functions. OneVPL can implement various functions for media pipelines on CPUs, GPUs, and other accelerators. OneVPL can implement device discovery and selection in media-centric and video analytics workloads. OneVPL can implement API primitives for zero-copy buffer sharing.

[0314] The oneAPI programming model can use the DPC++ programming language. The DPC++ programming language can include a version of a programming language that is functionally similar to the CUDA mechanism to define device code and distinguish it from host code. The DPC++ programming language can include a subset of the functionality of the CUDA programming language. One or more CUDA programming model operations can be performed using the oneAPI programming model utilizing the DPC++ programming language.

[0315] Any application programming interface (API) described herein may be compiled by a compiler, interpreter, or other software tool into one or more instructions, operations, or other signals. Compilation may include generating one or more machine-executable instructions, operations, or other signals from source code. When executed, an API compiled into one or more instructions, operations, or other signals may cause one or more processors (such as, but not limited to, the processor described in Figure 10-22B or any other logic circuitry further described herein) to perform one or more computational operations.

[0316] In at least one embodiment, the conversion tools described elsewhere herein (e.g., but not limited to) may include one or more circuitry for converting CUDA code into HIP, oneAPI, OpenCL, or any other language for performing any of the operations described above or elsewhere herein, wherein the CUDA code executes an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations. One or more circuitry may be software-configurable for converting CUDA code into HIP, oneAPI, OpenCL, or any other language for performing any of the operations described above or elsewhere herein, wherein the CUDA code executes an application programming interface (API) to copy information from the two or more first storage locations to the two or more second storage locations, at least in part based on one or more parameters of the API indicating the two or more first storage locations and the two or more second storage locations.

[0317] Autonomous vehicles

[0318] Figure 26 illustrates an example of an autonomous vehicle 2600 according to at least one embodiment. The autonomous vehicle 2600 (alternately referred to herein as "vehicle 2600") can be a passenger vehicle, such as, but not limited to, a sedan, truck, bus, and / or other type of vehicle capable of accommodating one or more passengers. In at least one embodiment, vehicle 2600 can be a semi-trailer tractor for transporting goods. Vehicle 2600 can be an aircraft, a robotic vehicle, or other type of vehicle.

[0319] Autonomous vehicles can be described according to their automation levels, which are defined by the National Highway Traffic Safety Administration (NHTSA) (a division of the U.S. Department of Transportation) and the Society of Automotive Engineers (SAE) in their "Taxonomy and Definitions for Terms Related to Driving Automation Systems for On-Road Motor Vehicles" (e.g., standard number J3016-201806 published June 15, 2018; standard number J3016-201609 published September 30, 2016; and previous and future versions of this standard). In at least one embodiment, vehicle 2600 can achieve functionality at one or more levels of autonomy, from Level 1 to Level 5. For example, in at least one embodiment, depending on the embodiment, vehicle 2600 can achieve conditional automation (Level 3), high automation (Level 4), and / or full automation (Level 5).

[0320] Vehicle 2600 may include components such as, but not limited to, chassis, body, wheels (e.g., 2, 4, 6, 8, 18, etc.), tires, axles, and other vehicle parts. Vehicle 2600 may include a propulsion system 2650, such as, but not limited to, an internal combustion engine, a hybrid power plant, an all-electric motor, and / or other propulsion system types. Propulsion system 2650 may be connected to the drivetrain of vehicle 2600 (which may include a transmission) to achieve propulsion of vehicle 2600. Propulsion system 2650 may be controlled based on signals received from throttle / accelerator 2652.

[0321] A steering system 2654 (which may include a steering wheel) is used to maneuver the vehicle 2600 (e.g., to travel along a desired path or route) while the propulsion system 2650 is in operation (e.g., when the vehicle 2600 is in motion). The steering system 2654 may receive signals from a steering actuator 2656. The steering wheel is optional for fully automated (Level 5) functionality. A brake sensor system 2646 may be used to operate the vehicle brakes in response to signals received from a brake actuator 2648 and / or brake sensors.

[0322] Controller 2636 may include one or more system-on-a-chip (“SoC”) and / or graphics processing units (“GPUs”) that can provide signals (e.g., signals representing commands) to one or more components and / or systems of vehicle 2600. For example, controller 2636 may send signals to operate vehicle brakes via brake actuator 2648, steering system 2654 via steering actuator 2656, and propulsion system 2650 via throttle / accelerator 2652. Controller 2636 may include one or more onboard (e.g., integrated) computing devices for processing sensor signals and outputting operational commands (e.g., signals representing commands) to enable autonomous driving and / or assist a human driver in driving vehicle 2600. Controller 2636 may include a first controller for autonomous driving functions, a second controller for functional safety functions, a third controller for artificial intelligence functions (e.g., computer vision), a fourth controller for infotainment functions, a fifth controller for redundancy in emergency situations, and / or other controllers. A single controller may handle two or more of the above functions, two or more controllers may handle a single function, and / or any combination thereof.

[0323] The controller 2636 may provide signals for controlling one or more components and / or systems of the vehicle 2600 in response to sensor data (e.g., sensor inputs) received from one or more sensors. Sensor data can be received from, for example, a Global Navigation Satellite System (“GNSS”) sensor 2658 (e.g., a Global Positioning System sensor), a RADAR (radar) sensor 2660, an ultrasonic sensor 2662, a LIDAR (light radar) sensor 2664, an Inertial Measurement Unit (“IMU”) sensor 2666 (e.g., an accelerometer, a gyroscope, one or more magnetic compasses, a magnetometer, etc.), a microphone 2696, a stereo camera 2668, a wide-angle camera 2670 (e.g., a fisheye camera), an infrared camera 2672, a surround-view camera 2674 (e.g., a 360-degree camera), a long-range camera 2698, a medium-range camera 2676, a speed sensor 2644 (e.g., for measuring the speed of vehicle 2600), a vibration sensor 2642, a steering sensor 2640, a braking sensor (e.g., as part of a braking sensor system 2646), and / or other types of sensors.

[0324] One or more controllers 2636 may receive input (e.g., represented by input data) from the instrument panel 2632 of the vehicle 2600 and provide output (e.g., represented by output data, display data, etc.) via a human-machine interface (“HMI”) display 2634, an audible annunciator, a speaker, and / or via other components of the vehicle 2600. Output may include, but is not limited to, vehicle speed, rate, time, map data (e.g., a high-resolution map (not shown), location data (e.g., the location of the vehicle 2600, such as, but not limited to, its location on a map), direction, the location of other vehicles (e.g., occupying a grid), information about objects, and the status of objects perceived by the controllers 2636. For example, the HMI display 2634 may display information about the presence of one or more objects (e.g., street signs, warning signs, traffic light changes, etc.) and / or information about driving actions that the vehicle has performed, is performing, or will perform (e.g., changing lanes now, exiting from exit 34B in two miles, etc.).

[0325] Each component, feature, and system of vehicle 2600 in Figure 26 can be connected via bus 2602. Bus 2602 may include a CAN data interface (also referred to herein as the "CAN bus"). CAN can be a network within vehicle 2600 used to assist in the control of various features and functions of vehicle 2600, such as, but not limited to, the activation of brakes, acceleration, braking, steering, windshield wipers, etc. Bus 2602 can be configured to have dozens or even hundreds of nodes, each with its own unique identifier (e.g., a CAN ID). Bus 2602 can be read to locate steering wheel angle, ground speed, engine revolutions per minute ("RPM"), button positions, and / or other vehicle status indicators. Bus 2602 can be an ASIL B compliant CAN bus.

[0326] In addition to CAN, or as an alternative to CAN, FlexRay and / or Ethernet protocols may be used. Bus 2602 may consist of any number of buses, which may include zero or more CAN buses, zero or more FlexRay buses, zero or more Ethernet buses, and / or zero or more other types of buses using different protocols. Two or more buses may be used to perform different functions and / or for redundancy. For example, a first bus may be used for collision avoidance functions, and a second bus may be used for actuation control. Each bus in bus 2602 may communicate with any component of vehicle 2600, and two or more buses of bus 2602 may communicate with corresponding components. Any number of System-on-Chip (“SoC”) 2604 (e.g., but not limited to SoC 2604(A) and SoC 2604(B)), each controller 2636, and / or each computer within the vehicle may access the same input data (e.g., input from sensors of vehicle 2600) and may be connected to a common bus, such as a CAN bus.

[0327] According to at least one embodiment, for the autonomous vehicle 2600 in FIG26, any number of cameras can be placed at any selected camera locations and within any chosen field of view. The cameras and corresponding fields of view are exemplary embodiments and are not intended to be limiting. For example, additional and / or alternative cameras may be included, and / or the cameras may be located at different locations within the vehicle 2600.

[0328] The camera type may include a digital camera suitable for components and / or systems in vehicle 2600. The camera may operate under Automotive Safety Integrity Level (“ASIL”) B and / or other ASILs. According to embodiments, the camera type may support any image capture rate, such as, but not limited to, 60 frames per second (fps), 1220 fps, 240 fps, etc. The camera may be able to use a rolling shutter, a global shutter, other types of shutters, or combinations thereof. In at least one embodiment, the color filter array may include a red-to-clear-to-clear (“RCCC”) color filter array, a red-to-clear-to-blue (“RCCB”) color filter array, a red-blue-green (“RBGC”) color filter array, a Foveon X3 color filter array, a Bayer sensor (“RGGB”) color filter array, a monochrome sensor color filter array, and / or other types of color filter arrays. Transparent pixel cameras (e.g., but not limited to cameras with RCCC, RCCB, and / or RBGC color filter arrays) may be used to enhance sensitivity.

[0329] One or more cameras may be used to perform advanced driver assistance system (“ADAS”) functions (e.g., as part of a redundancy or fail-safe design). For example, a multi-function monochrome camera may be installed to provide functions including lane departure warning, traffic sign assist, and intelligent headlight control. One or more cameras (e.g., all cameras) may simultaneously record and provide image data (e.g., video).

[0330] One or more cameras can be mounted in mounting assemblies, such as, but not limited to, custom-designed (3D-printed) assemblies, to eliminate stray light and reflections inside the vehicle 2600 (e.g., dashboard reflections in the windshield mirror) that could interfere with the camera's image data capture capabilities. Regarding rearview mirror mounting assemblies, these assemblies can be custom-3D printed to match the shape of the camera mounting plate to the rearview mirror. The camera can be integrated into the rearview mirror. For side-view cameras, the cameras can also be integrated into the four pillars in each corner of the cab.

[0331] A field of view including a portion of the environment in front of the vehicle 2600 (e.g., a front-facing camera) can be used for surround view to help identify the path and obstacles ahead, and, with the assistance of one or more controllers 2636 and / or control SoCs, provide information crucial for generating an occupancy grid and / or determining a preferred vehicle path. The front-facing camera can be used to perform many ADAS functions similar to LiDAR, including emergency braking, pedestrian detection, and collision avoidance. The front-facing camera can also be used in ADAS functions and systems, including lane departure warning (“LDW”), adaptive cruise control (“ACC”), and / or other functions such as, but not limited to, traffic sign recognition.

[0332] Various front-mounted cameras can be used, including, for example, monocular camera platforms including CMOS (“complementary metal-oxide-semiconductor”) color imagers. Wide-angle camera 2670 can be used to perceive objects entering the field of view from the periphery (e.g., pedestrians, crossing vehicles, or bicycles). Any number (including zero) of wide-angle cameras 2670 can be mounted on vehicle 2600. Any number of remote cameras 2698 (e.g., long-view stereo camera pairs) can be used for depth-based object detection, particularly for objects on which neural networks have not yet been trained. Remote cameras 2698 can also be used for object detection and classification, as well as basic object tracking.

[0333] It may also include any number of front-mounted stereo cameras 2668. One or more stereo cameras 2668 may include an integrated control unit that includes a scalable processing unit that can provide programmable logic (“FPGA”) and a multi-core microprocessor with an integrated controller area network (“CAN”) or Ethernet interface on a single chip. Such a unit can be used to generate a 3D map of the environment of the vehicle 2600, including distance estimates for all points in the image. One or more stereo cameras 2668 may include a compact stereo vision sensor that may include two camera lenses (one on each side) and an image processing chip that measures the distance from the vehicle 2600 to a target object and uses the generated information (e.g., metadata) to activate autonomous emergency braking and lane departure warning functions. Other types of stereo cameras 2668 may be used in addition to, or as an alternative to, the stereo cameras described herein.

[0334] Cameras (e.g., side-view cameras) that include portions of the environment along the sides of vehicle 2600 can be used in the surround view to provide information for creating and updating occupancy grids and generating side collision warnings. For example, surround view cameras 2674 (e.g., four surround view cameras) can be mounted on vehicle 2600. Surround view cameras 2674 can include any number and combination of wide-angle cameras, fisheye cameras, 360-degree cameras, and / or similar cameras. For example, four fisheye cameras can be mounted at the front, rear, and sides of vehicle 2600, respectively. Vehicle 2600 can use three surround view cameras 2674 (e.g., left, right, and rear) and can utilize one or more other cameras (e.g., front-facing cameras) as a fourth surround view camera.

[0335] A field of view including a portion of the environment behind vehicle 2600 (e.g., a rear-view camera) can be used for parking assistance, surround view, rear collision warning, and creating and updating occupancy grids. A wide variety of cameras can be used, including but not limited to cameras that are also suitable as front-facing cameras (e.g., long-range camera 2698 and / or mid-range camera 2676, stereo camera 2668, infrared camera 2672, etc.), as described herein.

[0336] Vehicle 2600 may include any number of SoCs 2604 or other processors described elsewhere herein, such as, but not limited to, the processors and / or components shown and described in Figures 10-22B. Each SoC 2604 may include a central processing unit (“CPU”) 2606, a graphics processing unit (“GPU”) 2608, a processor 2610, a cache 2612, an accelerator 2614, a data storage 2616, and / or other components and features not shown. SoC 2604 can be used to control vehicle 2600 in various platforms and systems. For example, SoC 2604 may be combined with a high-definition (“HD”) map 2622 in a system (e.g., the system of vehicle 2600), which may obtain map refreshes and / or updates from one or more servers (not shown) via network interface 2624.

[0337] The CPU 2606 may include a CPU cluster or CPU complex (referred to herein as “CCPLEX”). The CPU 2606 may include multiple cores and / or a secondary (“L2”) cache. For example, the CPU 2606 may include eight cores in a consistent multiprocessor configuration. The CPU 2606 may include four dual-core clusters, each with a dedicated L2 cache (e.g., 2 megabytes (MB) of L2 cache). The CPU 2606 (e.g., CCPLEX) may be configured to support simultaneous cluster operation, allowing any combination of CPU 2606 clusters to be active at any given time.

[0338] One or more of the CPU 2606 can implement power management capabilities, including one or more of the following features: individual hardware blocks can automatically clock-gated when idle to save dynamic power; clock gating can be applied to cores that are not actively executing instructions due to executing Wait for Interrupt (“WFI”) / Wait for Event (“WFE”) instructions; each core can independently perform power gating; each core cluster can clock-gated independently when all cores can perform clock or power gating; and / or each core cluster can independently perform power gating when all cores can perform power gating. The CPU 2606 can also implement enhanced algorithms for managing power states, where allowed power states and expected wake-up times can be specified, and the hardware / microcode determines which optimal power state the core, cluster, and CCPLEX should enter. Processing cores can enter a simplified power state sequence via software support, offloading work to the microcode.

[0339] The GPU 2608 may include an integrated GPU (referred to herein as an "iGPU"). The GPU 2608 may be programmable and capable of efficiently handling parallel workloads. The GPU 2608 may use an enhanced tensor instruction set. The GPU 2608 may include one or more streaming microprocessors, each of which may include a Level 1 ("L1") cache (e.g., an L1 cache with a storage capacity of at least 96KB), and two or more streaming microprocessors may share an L2 cache (e.g., an L2 cache with a storage capacity of 512KB). The GPU 2608 may include at least eight streaming microprocessors. The GPU 2608 may use a computation application programming interface (API). The GPU 2608 may use one or more parallel computing platforms and / or programming models (e.g., NVIDIA's CUDA model). A streaming microprocessor may be referred to as a streaming multiprocessor (“SM”), a streaming processor (“SP”), a streaming processing unit (“SPU”), a compute unit (“CU”), an execution unit (“EU”), and / or a slice, wherein, in this context, a slice may refer to a portion of the processing resources within a processing unit (e.g., 16 cores, a ray tracing unit, a thread bootstrap, or a scheduler).

[0340] One or more of the GPU 2608 can be power-optimized for optimal performance in automotive and embedded use cases. For example, the GPU 2608 can be fabricated on FinFET (“FinFET”) circuitry. Each streaming microprocessor can include multiple mixed-precision processing cores partitioned into blocks. For example, 64 PF32 cores and 32 FP64 cores can be partitioned into four processing blocks. Each processing block can be allocated 16 FP32 cores, 8 FP64 cores, 16 INT32 cores, two mixed-precision NVIDIA Tensor cores for deep learning matrix operations, a level-zero (“L0”) instruction cache, a scheduler (e.g., a warp scheduler) or sequencer, dispatch units, and / or a 64KB register file. Streaming microprocessors can include independent parallel integer and floating-point data paths for efficiently performing workloads involving mixed computation and addressing computation. Streaming microprocessors can include independent thread scheduling capabilities to enable finer-grained synchronization and cooperation between parallel threads. Streaming microprocessors may include a combination of L1 data cache and shared memory units to improve performance while simplifying programming.

[0341] One or more GPUs 2608 may include high-bandwidth memory (“HBM”) and / or a 16GB HBM2 memory subsystem, used in some examples to provide a peak memory bandwidth of approximately 900GB / s. In addition to HBM memory, or as an alternative to HBM memory, synchronous graphics random access memory (“SGRAM”) may be used, such as, but not limited to, graphics double data rate type 5 synchronous random access memory (“GDDR5”).

[0342] The GPU 2608 may include unified memory technology. An address translation service (“ATS”) supports page tables that allow the GPU 2608 to directly access the CPU 2606. When a memory miss occurs in one or more GPU 2608 memory management units (“MMUs”), an address translation request can be sent to the CPU 2606. In response, one or more CPUs in the CPU 2606 can look up the virtual-to-physical mapping of the address in their page tables and transfer the translation back to the GPU 2608. Unified memory technology can provide a single, unified virtual address space for the memory of both the CPU 2606 and GPU 2608, thereby simplifying GPU 2608 programming and application porting to the GPU 2608.

[0343] The GPU 2608 may include any number of access counters that can track the frequency of the GPU 2608's access to the memory of other processors. These access counters help ensure that memory pages are moved to the physical memory of the processor that accesses them most frequently, thereby improving the efficiency of shared memory ranges between processors.

[0344] One or more SoCs 2604 may include any number of caches 2612, including the caches described herein. For example, cache 2612 may include a Level 3 (“L3”) cache available for use by both CPU 2606 and GPU 2608 (e.g., connected to both CPU 2606 and GPU 2608). Cache 2612 may include a write-back cache that can track the state of rows, for example, but not limited to, by using a cache coherence protocol (e.g., MEI, ME...

Claims

1. A processor, comprising: One or more circuits, the one or more circuits being configured to execute an application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations.

2. The processor of claim 1, wherein the two or more first storage locations are not contiguous, or the two or more second storage locations are not contiguous.

3. The processor of claim 1, wherein one or more parameters indicate the size of the information to be copied.

4. The processor of claim 1, wherein one or more parameters indicate the order in which the information will be copied.

5. The processor of claim 1, wherein the two or more second storage locations comprise one or more memory addresses, the size of the one or more memory addresses being at least partially based on the size of the two or more first storage locations.

6. The processor of claim 1, wherein the two or more first storage locations and the two or more second storage locations are associated with different storage devices.

7. The processor of claim 1, wherein the information is used to execute one or more instructions based at least in part on accessing one or more of the two or more second storage locations.

8. A system comprising: One or more circuits, the one or more circuits being configured to execute an application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API indicating two or more first storage locations and two or more second storage locations.

9. The system of claim 1, wherein the two or more first storage locations are not contiguous, or the two or more second storage locations are not contiguous.

10. The system of claim 1, wherein one or more parameters indicate the size of the information to be copied.

11. The system of claim 1, wherein one or more parameters indicate the order in which the information will be copied.

12. The system of claim 1, wherein the two or more second storage locations comprise one or more memory addresses, the size of the one or more memory addresses being at least partially based on the size of the two or more first storage locations.

13. The system of claim 1, wherein the two or more first storage locations and the two or more second storage locations are associated with different storage devices.

14. The system of claim 1, wherein the information is used to execute one or more instructions based at least in part on accessing one or more of the two or more second storage locations.

15. A method comprising: This enables the execution of an application programming interface (API) such that information is copied from the two or more first storage locations to the two or more second storage locations, at least in part, based on one or more parameters of the API that indicate two or more first storage locations and two or more second storage locations.

16. The method of claim 1, wherein the two or more first storage locations are not contiguous, or the two or more second storage locations are not contiguous.

17. The method of claim 1, wherein one or more parameters indicate the size of the information to be copied.

18. The method of claim 1, wherein one or more parameters indicate the order in which the information will be copied.

19. The method of claim 1, wherein the two or more second storage locations comprise one or more memory addresses, the size of the one or more memory addresses being at least partially based on the size of the two or more first storage locations.

20. The method of claim 1, wherein the information is used to execute one or more instructions based at least in part on accessing the two or more second storage locations.

Images