DYNAMIC PERFORMANCE OF SECTIONS OF NEURAL NETWORKS

Abstract

Devices, systems, and methods for dynamically executing sections of neural networks are disclosed. In at least one embodiment, one or more sections of one or more neural networks are dynamically executed by two or more processors, for example, based on one or more estimated performance metrics of the one or more sections of the one or more neural networks.

Description

TECHNICAL AREA

[0001] At least one embodiment relates to processing resources used to perform and facilitate artificial intelligence for various tasks. For example, at least one embodiment relates to processors or computing systems that dynamically execute sections of neural networks on different processing devices. BACKGROUND

[0002] Artificial intelligence methods are used to implement various tasks. The architecture and size of parameters in neural networks, for example, can consume significant memory, time, or computing resources to perform different tasks. The amount of memory, time, or computing resources used to perform different tasks can be improved during training or inference using a neural network. BRIEF DESCRIPTION OF THE DRAWINGS Fig. Figure 1A illustrates an example of a block diagram of a scheduler implementing dynamic performance of neural network sections, according to at least one embodiment; Fig. Figure 1B illustrates an example of a scheduling hierarchy that can be used by a scheduler implementing dynamic performance of neural network sections, according to at least one embodiment; Fig. Figures 2A-2B illustrate examples of different scheduling decisions according to at least one embodiment; Fig. Figure 3 illustrates an example of a block diagram of a global scheduler implementing a dynamic performance of sections of neural networks of a distributed LLM (Large Language Model) system, according to at least one embodiment; Fig. Figure 4 illustrates an example of a block diagram of interactions between inferencing processing devices of a distributed LLM system according to at least one embodiment; Fig. Figure 5 illustrates an example of a block diagram of a frontend of a distributed LLM system according to at least one embodiment; Fig. Figure 6 illustrates an example of a block diagram of a planner performing various actions to adjust the performance of a distributed LLM system, according to at least one embodiment; Fig. Figure 7 illustrates an example of a flowchart of a method implementing the dynamic performance of sections of neural networks, according to at least one embodiment; Fig. Figure 8 illustrates an exemplary data center system according to at least one embodiment; Fig. Figure 9 illustrates a system-on-a-chip (SOC) according to at least one embodiment; Fig. 10A illustrates a parallel processor according to at least one embodiment; Fig. Figure 10B illustrates a processing cluster according to at least one embodiment; Fig. 10C illustrates a graphics multiprocessor according to at least one embodiment; Fig. Figure 11 illustrates an accelerator processor according to at least one embodiment; Fig. Figure 12A illustrates a central processing unit and a core of the central processing unit according to at least one embodiment; Fig. Figure 12B illustrates a core of the central processing unit in Fig. 12A according to at least one embodiment; Fig. Figure 13 illustrates another accelerator processor according to at least one embodiment; Fig. Figure 14 illustrates a neuromorphic processor according to at least one embodiment; Fig. 15 illustrates a supercomputer according to at least one embodiment; Fig. Figure 16 illustrates another accelerator processor according to at least one embodiment; Fig. 17 illustrates another processor according to at least one embodiment; Fig. Figure 18 illustrates another accelerator processor according to at least one embodiment; Fig. Figure 19 illustrates a tensor processing unit according to at least one embodiment; Fig. Figure 20 illustrates a RISC-V compatible processor according to at least one embodiment; Fig. 21A and Fig. 21B illustrates a speech processing unit according to at least one embodiment; Fig. 22 illustrates a software stack of a programming platform according to at least one embodiment; Fig. 23 illustrates software supported by a programming platform according to at least one embodiment; Fig. 24 illustrates compilation code for execution on programming platforms of Fig. 21 according to at least one embodiment; Fig. Figure 25 illustrates an example of an autonomous vehicle and its system architecture according to at least one embodiment; Fig. 26A illustrates an inference and / or training logic according to at least one embodiment; Fig. 26B illustrates an inference and / or training logic according to at least one embodiment; Fig. 26C illustrates a training and deployment of a neural network according to at least one embodiment. DETAILED DESCRIPTION

[0003] Fig. Figure 1A illustrates an example of a block diagram of a scheduler implementing dynamic performance of neural network sections according to at least one embodiment. In at least one embodiment, one or more schedulers 110 can implement neural network performance metric estimation to dynamically execute sections of neural network(s) 140 by scheduling inference across two or more processing devices 130a and 130n for an inference request 102, as specified in Figure 116. In at least one embodiment, the neural network(s) 140 can be trained for neural network performance metric estimation to perform various tasks and can include different stages, subtasks, or other sections, such as sections 140a and 140b, which may correspond, for example, to different layers or groups of neurons of the neural network.In at least one embodiment, different sections of one or more neural networks can encode, distribute, decode or perform various other operations and can correspond to one or more of transformer blocks, CNN (folding neural network) blocks, RNN (recurrent neural network) blocks, pooling blocks, tightly connected blocks and / or other types of blocks.

[0004] In at least one embodiment, the performance metric estimation of 120 neural networks can use machine learning models (e.g., neural networks) or other prediction methods (e.g., support vector machines, random forests, decision trees, naive Bayes, k-nearest neighbor, linear regression, logistic regression, and / or Bayesian networks). In at least one embodiment, performance metrics can be predicted using one or more heuristics and / or a cost function to model performance (e.g., as detailed below). In at least one embodiment, the estimated performance metrics can include the inference speed per inference request (e.g.,Latency or time spent completing inference requests), the inference throughput for a processing device, a group of processing devices, or an inference system as a whole, such as a distributed LLM system, referred to below. Fig. 2 is discussed, which includes the power consumed by individual processing devices (or a group of processing devices) and / or the power drawn from a power grid by a data center. In at least one embodiment, the estimation methods for performance metrics may be modified or based on other information 134, including workload predictions, fault predictions, or other information generated by a scheduler or other system management component that uses a different and longer prediction horizon than a scheduler or other component that uses estimated performance metrics to dynamically execute sections of one or more neural networks, as described below with reference to Fig. 6 is discussed in detail.

[0005] In at least one embodiment, the estimated performance metrics can be based on a set of data used for inferencing, such as a number of input tokens in inference request 102, and / or an expected number of output tokens as inputs for or information used by performance metric estimation methods (e.g., a number of tokens in a cost function as described above with reference to Fig. 1 is discussed). In at least one embodiment, the estimated performance metrics can be determined based on the resource utilization 132a and 132b of various processing devices 130a and 130n (e.g., the current utilization of a processing device performing other inference requests, GPU, CPU, memory, I / O, network, or other compute resource utilization information). In at least one embodiment, KV cache information, including the number of KV cache values ​​or other intermediate state information applicable to an inference request (e.g., as a result of an initial filling phase or stage performed with input tokens), can be used as input or a basis for determining the estimated performance metrics.In at least one embodiment, resource utilization, KV cache information, or other information used to estimate performance metrics can be captured or determined using various components and shared as events in an event stream transmitted using an event data layer, similar to the methods described below with reference to . Fig. 2-6 will be discussed.

[0006] In at least one embodiment, the neural network(s) 140, for which a dynamic performance is implemented by two or more processors, can be an LLM or another neural network(s) that uses an initial padding and decoding stage, as described below with reference to the Fig. 2A-4 is discussed. In at least one embodiment, an initial padding phase or section of the neural network(s) can process input data (e.g., tokenized into a sequence of input tokens) to compute intermediate states, keys, and values, which are then used to generate a new token based on previous tokens using the intermediate states. In at least one embodiment, (initial) padding can be used to compute intermediate states for large contexts (e.g., documents, images, text, audio, or various other data) that are used by neural networks to generate an inference as one or more output tokens. In at least one embodiment, the initial padding phase can be a long-lasting section of the inference generation process, since a large input data size may require a correspondingly longer time to compute intermediate states.In at least one embodiment, a decoding phase can be a new token generation phase, wherein one or more new tokens can be an inference output from an inference system. In at least one embodiment, decoding phase sections of a neural network can autoregressively generate output tokens one after the other until stop criteria are met, using previous iterations of output states, keys, and values. In at least one embodiment, key-value caching (KV caching) methods can be used to improve the speed of decoding phase processing. In at least one embodiment, a KV cache can store tensors or other data structures containing intermediate states that are used to generate new tokens, so that when a new token is generated, intermediate states can be retrieved and results added to the KV cache.

[0007] In at least one embodiment, the Scheduler(s) can use the procedures described below in pseudocode:

[0008] In at least one embodiment, the scheduler(s) 110 can schedule across different locations or levels of the organization's processing devices within a hierarchy and can be considered a hierarchical scheduler that independently schedules and / or scales resources across a hierarchy. Fig. Figure 1B illustrates an example of a scheduling hierarchy that can be used by a scheduler implementing dynamic performance of neural network segments, according to at least one embodiment. In at least one embodiment, the scheduling hierarchy 170 can be accessed via one or more processor(s) 172a, one or more nodes 172b (each containing, for example, one or more processors), one or more rack(s) 172c (each similar to the racks discussed below and containing one or more nodes), one or more clusters 172d (each containing, for example, one or more racks), one or more data center rooms 172e (each containing, for example, one or more clusters), one or more data centers 172f (each connected, for example, via dedicated high-speed communication links), and / or one or more data center regions 172g (each containing, for example,Each containing one or more data centers that are isolated from data centers in other regions (except for communication using public communication links, such as the Internet or another wide area network). In at least one embodiment, the scheduler 110 can schedule the performance of inference requests 102 taking into account different levels or locations in the hierarchy 170. In at least one embodiment, for example, an inference request can consist of using a decoding phase to generate new data, and performance metrics for the processor(s) 172, the node(s) 172b, the rack(s) 172c, the cluster(s) 172d, the data center room(s) 172e, the data center(s) 172f, and / or the data center region(s) 172f can be estimated to schedule performance.

[0009] In at least one embodiment, the Scheduler 110 can switch between scheduling levels or planning levels of the hierarchy 170 according to one or more scheduling policy(ies) or planning policy(ies) 160. In at least one embodiment, the scheduling policy(ies) can specify parameters, thresholds, criteria, or other information that can instruct or configure the performance of the Scheduler 110. In at least one embodiment, the scheduling policy(ies) 160 can be implemented by a planner system (e.g., the planner 610, referred to below). Fig. (discussed in detail in Section 6) are generated or updated. In at least one embodiment, the scheduling policy(ies) 160 can instruct the scheduler(s) 110 to schedule an inference request on one of a group of rack(s) that may have capacity, rather than on one or more nodes. In at least one embodiment, the hierarchical scheduling performed by the scheduler(s) 110 can support dynamic scheduling across different levels of the scheduling hierarchy 170, so that throughput, latency, performance, and / or various other performance metrics can be optimized for a system (e.g., the distributed system 301) implementing the scheduler(s) 110.

[0010] In at least one embodiment, the scheduler 110 can schedule processing for many different neural networks 150 (e.g., different types of LLMs, language models (VLMs), diffusion models, autoregressive or non-autoregressive models, among others). In at least one embodiment, the neural network(s) 150 can be arranged on or co-hosted by the same processing devices (or instances, nodes, or another level of the hierarchy 170), which, for example, makes it possible to execute decoding phases of several different LLMs on the same group of processing devices, for example, to optimize power consumption. In at least one embodiment, the scheduler 110 can schedule processing tasks of the neural network(s) 150 (e.g.,parallelize phases of (initial) filling or decoding) and assign or otherwise schedule inference requests to a group of processors at the same level of the scheduling hierarchy 170.

[0011] Fig. Figures 2A-2B illustrate examples of various planning and scheduling decisions according to at least one embodiment. At least one embodiment includes an example of a scheduling decision based on estimated performance metrics. Fig. Figure 2A shows that in at least one embodiment, the processing device(s) 210a can perform both the (initial) filling 220a and the decoding 220b for the input tokens 202 using the KV cache 230. In at least one embodiment, a wait or other pause can be specified for an inference request, so that the processing devices 210a can perform both the (initial) filling 220a and the decoding 220b, as discussed above.

[0012] In at least one embodiment, various processing devices, such as processing device(s) 240a and 280a, can be used to perform the (initial) padding and decoding. In at least one embodiment, the input tokens 204 are used to perform the (initial) padding 250a at the processing device(s) 240a, with intermediate states being stored in the KV cache 260 and moved to the KV cache 270 to perform the decoding 250b at the processing device(s) 280a. In at least one embodiment, further input tokens 206 (e.g., an instruction or a prompt) can also be provided to perform a small portion of the (initial) padding 205a before the decoding 250b.

[0013] Fig. Figure 3 illustrates an example of a block diagram of a global scheduler implementing dynamic performance of neural network sections of a distributed Large Language Model (LLM) system, according to at least one embodiment. In at least one embodiment, a distributed LLM system 301 can include various components to provide inferencing services to a client 310, including a frontend 320, a global scheduler 330, various inferencing processing devices, such as inferencing processing devices for decoding 350 and (initial) population 340, an inferencing data layer 360, an event data layer, and a tracking system 380 for global KV caches. In at least one embodiment, the frontend 320 can implement various types or styles of interfaces for the client(s) 310, including command lines, graphical user interfaces (e.g.,Chat windows) and / or program interfaces, such as application programming interfaces (APIs). In at least one embodiment, the frontend can host 320 different types or styles of interactions with LLMs. In at least one embodiment, an agent / fetch application can be hosted, as described below in relation to . Fig. 5 is discussed in detail.

[0014] In at least one embodiment, the scheduler(s) 110 can be implemented hierarchically as a hierarchical scheduler, wherein one (or a part of) the scheduler 110 can be used to schedule across different levels in a hierarchy (e.g., similar to the one in Fig. (Hierarchy shown in 1B). In at least one embodiment, a hierarchical scheduler can independently scale any part of the scheduling decisions up or down at any level in a hierarchy (e.g., based on guidelines or other information provided by the scheduler 610, referred to below). Fig. 6 is discussed). In at least one embodiment, the global scheduler 330 can be a hierarchical scheduler, as discussed above.

[0015] In at least one embodiment, the distributed LLM system 301 can implement the global scheduler 330. In at least one embodiment, the global scheduler 330 can implement scheduling procedures as described above with reference to Fig. 1-2B and below with reference to Fig. 7 will be discussed. In at least one embodiment, the global scheduler 330 can receive inference requests from the frontend 320 and determine which sections of a neural network (e.g., an LLM) are to be executed by different processors. In at least one embodiment, primary assignments can be made, such as for the inferencing processing device(s) 350 to decode and for the inferencing processing device(s) 340 to (initially) populate. In at least one embodiment, primary assignments can be modified / overridden by the global scheduler 330 using instructions, as described below with reference to Fig. 4 is discussed in detail.

[0016] In at least one embodiment, the distributed LLM system 301 can implement the global scheduler 330 to schedule and use multiple different LLMs (e.g., on different processing devices). In at least one embodiment, the distributed LLM system 301 can provide or assign the (initial) population and decoding phases of multiple different LLMs on the same processor, node, rack, cluster (or other hierarchy level), or compute instance, or on different identical processors, nodes, racks, clusters (or other hierarchy levels), or compute instances, according to available resource information.

[0017] In at least one embodiment, the global scheduler 330 can obtain the state of KV caches via the inferencing devices 340 and 350 through events transmitted to the event data layer 370. In at least one embodiment, the event data layer 370 can communicate events via various components of the distributed LLM system 301, including global KV cache tracking, which can provide occupancy and / or location information of the KV cache instance in the respective processing device(s) 350 and 340.In at least one embodiment, the event data layer 370 can be a data stream, bus, link, or other data processing, storage, and / or transmission system that exchanges data as events, enabling, for example, components that register for or wait for events to receive events in an event stream for use in various purposes (as discussed below). In at least one embodiment, the global scheduler 330 can use KV cache information obtained from tracking 380 global KV caches to estimate performance metrics for controlling the performance of inference requests.

[0018] In at least one embodiment, the Inferencing Data Layer 360 can be implemented to exchange KV cache data and / or other intermediate state across inferencing processing devices, and can, for example, include cross-device communications (e.g., NVlink), cross-server communications, cross-data center communications, and / or cross-data center room communications. In at least one embodiment, the Inferencing Data Layer 360 can be used to move some or all of the KV cache data from one inferencing processing device to another in order to dynamically execute different sections of the neural network(s) (as discussed above).

[0019] In at least one embodiment, components similar to those illustrated for the distributed LLM system 301 can be implemented for other neural networks in addition to or instead of LLMs. In at least one embodiment, a distributed VLM system, a distributed diffusion system, or various other distributed inferencing systems for neural networks can be implemented using components similar to the frontend 320, the global scheduler 330, various inferencing processing devices such as the inferencing processing devices for decoding 350 and (initial) padding 340, an inferencing data layer 360, an event data layer, and a global KV cache tracking system 380.

[0020] Fig. Figure 4 illustrates an example of a block diagram of interactions between inferencing processing devices of a distributed LLM system according to at least one embodiment. In at least one embodiment, the inferencing processing devices 350 for decoding and the inferencing processing devices 340 for (initial) population can receive inferencing requests from the global scheduler 330. In at least one embodiment, the inferencing processing devices 350 for decoding and the inferencing processing devices 340 for (initial) population can implement respective input queues 402a and 402b, from which inferencing requests are retrieved for processing using respective LLM runtime engines 410a and 420b.In at least one embodiment, an LLM runtime engine can be an executable application that instructs the execution of a neural network, specifying when to load weights, which operations to perform, which outputs to use as inputs for subsequent operations, and so on, to execute a neural network. In at least one embodiment, an LLM runtime engine can have a primary assignment (e.g., initial population or decoding). In at least one embodiment, a primary assignment does not preclude an LLM runtime agent from executing another section of a neural network to perform a different phase or task of the neural network (e.g., both initial population and decoding, as described above). Fig. 2A and Fig. 2B discussed).

[0021] In at least one embodiment, as above with reference to Fig. 1A-3 discusses scheduling procedures, such as the global scheduler 330, and other features of inferencing processing devices, such as a KV cache movement 422a and 422b, to stream, send, or otherwise communicate KV cache values ​​and / or other intermediate states from one KV cache instance to another over the inferencing data plane 360 ​​(e.g., from 430b to 430a or vice versa). In at least one embodiment, scheduling instructions or other information to initiate KV cache movements can be specified via an interface (e.g., an API) to identify a destination of KV cache values ​​and to identify which KV cache values ​​are to be sent.

[0022] In at least one embodiment, the Inferencing Data Layer 360 can implement or support one or more types of interfaces to obtain KV cache values, other intermediate states, or any other information that can be used for inferencing. In at least one embodiment, the Inferencing Data Layer 360 can implement one or more APIs to support various interactions. In at least one embodiment, the Inferencing Data Layer 360 can include an API to obtain specified data, read data, or receive another request (e.g., KV cache values, a file, an object, or other information, including documents or application state information, to perform an inferencing task).In at least one embodiment, the Inferencing Data Layer 360 may include an API for searching, querying, or otherwise requesting the retrieval of data according to parameters, predicates, or other criteria (e.g., a task identifier, application identifiers, or other information that links various pieces of information, values ​​of desired information, and / or metadata describing information). In at least one embodiment, the Inferencing Data Layer 360 may include an API for writing, updating, modifying, or otherwise altering specified data (e.g., KV cache values, a file, an object, or other information, including documents or application state information as a result of performing an inferencing task). In at least one embodiment, the Inferencing Data Layer 360 may include an API for deleting or otherwise requesting the removal of specific data (e.g.,KV cache values, a file, an object or other information, including documents or application state information in external storage 440).

[0023] In at least one embodiment, the inferencing processing devices 350 for decoding and the inferencing processing devices 340 for (initial) filling can each implement a KV cache management system, such as KV cache management 424a and 424b. In at least one embodiment, the KV cache management system can apply one or more cache management and / or cache eviction policies to determine whether KV cache values ​​are to be retained in the KV cache instance 430a or 430b, respectively. In at least one embodiment, the cache management policies can be modified, as described below with reference to Fig. 6 is discussed. In at least one embodiment, the inferencing data layer 360 can access external storage 440, such as an external database, data store, file system, or other storage devices that are separate from directly accessible on-chip or other on-processing device storage (e.g., GPU memory, CPU memory, and / or persistent storage devices of the host system). In at least one embodiment, the inferencing data layer 360 can implement an interface and / or protocol to support the exchange of KV cache values ​​and / or other intermediate state information with external sources, which (for example) allows large contexts that have already been generated and rarely used to be streamed via the inferencing data layer 360 or otherwise provided to processing devices without being regenerated.In at least one embodiment, the Inferencing Data Layer 360 can support requests to obtain additional data (e.g., raw data such as documents, images, or other data for which no (initial) padding or other encoding has been performed) in order to complete or perform an inferencing task.

[0024] In at least one embodiment, processing devices of the event data layer 370 can provide various types of information, such as metrics or events from respective event / metric generators 424a and 424b, input queues 402a and 402b, and / or from various other components of processing devices. In at least one embodiment, this can include information indicating whether (initial) population is complete, information on the occupancy or other use of a KV cache instance, acknowledgments of requests from the global scheduler 330, or other components. In at least one embodiment, events can correspond to state change events for KV cache instances, such as when a new block allocation is performed for a KV cache instance, when a usage counter changes, block eviction occurs, or the storage layer changes (e.g.,from CPU to SSD or another persistent storage device), and they can include information in event descriptions such as a block identifier, block metadata (e.g., token, parameter-efficient fine-tuning model (PEFT) ID), storage information such as pointer, size, device and storage type, and process information such as a runtime instance identifier and / or parallel configuration.

[0025] Fig. Figure 5 illustrates an example of a block diagram of a frontend of a distributed LLM system according to at least one embodiment. In at least one embodiment, a distributed LLM system 301 can be implemented with different types of frontends that host different types of applications. In at least one embodiment, the frontend 320 can implement an agent / retrieval augmented generation (RAG) hosting 510. In at least one embodiment, the agent / retrieval augmented generation (RAG) hosting 510 can implement an execution engine 520. In at least one embodiment, the execution engine 520 can perform one or more workflows to handle one or more loops, states, or sequences of interactions between a client and the agent / retrieval augmented generation (RAG) hosting 510.In at least one embodiment, the execution engine 520 can, for example, receive a workflow from the workflow database 530, which can specify a sequence of steps, stages, or instructions for executing a workflow using one or more loops of an inference request to an agent 502 (e.g., requesting a new agent or an existing agent). In at least one embodiment, the execution engine 520 can submit a request (e.g., with various prompt formatting information or other information) to a global queue(s) 550, which can distribute or assign the submission of inference requests 506 to the global scheduler 550. In at least one embodiment, the execution engine 520 can receive additional information 562 from external memory 560 before (or after) an inference 508 is returned for a workflow.In at least one embodiment, the external storage 560 can store various types of text, documents, images, video, audio, and / or other information that may be included in inference requests identified by the inference 508 and / or otherwise used to augment the inference 508 generated by a distributed LLM system. In at least one embodiment, the external storage 560 can include different types of data storage (including more than one). In at least one embodiment, the external storage 560 can be an external storage 440 used for KV cache values. In at least one embodiment, the agent / retrieval augmented generation (RAG) hosting 510 can implement an event / metric generator 540 that can capture, generate, and / or send events / metrics 542 to the event data layer 370 (e.g.,for the use of other components, such as the planner 610 in . Fig. 6).

[0026] In at least one embodiment, the frontend 320 can support more than one interface or type of inference request. For example, in addition to the agent inference request 502, at least one embodiment can support the inference request 504 (e.g., a query or other direct LLM request). In at least one embodiment, the direct inference request 504 can be supported by the frontend 320, but not other requests, such as inference requests for the agent / retrieval augmented generation (RAG) hosting 510.

[0027] In at least one embodiment, the global queue(s) 550 can coordinate, balance, or otherwise ensure fair or other specified distributions (e.g., by retrieving inference requests from respective queues of respective lengths that store inference requests of different types). In at least one embodiment, the global queue(s) 550 can transmit inference request(s) 506 to the global scheduler for inferencing, as discussed above.

[0028] In at least one embodiment, the global scheduler 330 can include planning or scheduling the performance (e.g., migrating from one system to another or creating a new replica) of application tasks that are to be scheduled together with inferencing devices. In at least one embodiment, the global scheduler can, for example, schedule the performance of agents that execute one or more loops of a defined set of interactions as part of an application that is to be scheduled together with a corresponding phase of an inferencing task (e.g., hosting and running agent application software on the same system as a processing device that generates inferences and returns them to that agent application software) to further improve application performance. In at least one embodiment, various sections of the execution of neural networks that are performed dynamically (e.g.,According to the global scheduler 330, these applications can also correspond to different agent applications, enabling dynamic performance of agent applications across two or more processors. In at least one embodiment, the scheduling of agent applications with inference processing tasks can implement a Program Agent in the Loop (PAL) system, which reduces communication between inferencing and client applications (e.g., agents) and results in other performance improvements.

[0029] Fig. Figure 6 illustrates an example of a block diagram of a scheduler that performs various actions to adjust or set the performance of a distributed LLM system according to at least one embodiment. In at least one embodiment, the scheduler 610 can be implemented as part of the distributed LLM system. In at least one embodiment, the scheduler 610 can operate with a different and longer time horizon for predictions than the global scheduler 330 (e.g., on an hourly, daily, weekly, or other time basis instead of per inference request or batch of inference requests for the next minute). In at least one embodiment, the scheduler 610 can receive information from various sources, including the frontend 320, the global KV cache tracking 380, and the event data layer 370.In at least one embodiment, the Frontend 320 can provide various types of information according to the type of LLM application hosted by the Frontend 320 (e.g., RAG, Agent, etc.). In at least one embodiment, the Frontend 320 can, for example, provide various types of received inferencing requests (e.g., according to inferencing types supported by the LLM application at the Frontend 320). In at least one embodiment, performance metrics, such as the total latency for responding to a query or request from a client, can be provided to the Scheduler 610. In at least one embodiment, the Frontend 320 can provide or indicate counts, frequencies, or other workload information to the Scheduler 610.

[0030] In at least one embodiment, the tracking of global KV caches can provide information to the planner 610. In at least one embodiment, the tracking of global KV caches can, for example, indicate the occupancy (e.g., usage) of KV caches across different processing devices as an indicator of utilization or other processing tasks distributed across the processing devices of a distributed LLM system.

[0031] In at least one embodiment, various other events or information collected or reported at event data level 370 can be provided to the scheduler 610, as discussed in detail above. For example, in at least one embodiment, processing device-specific task assignments (e.g., number of initial refills, number of decodes, number of initial refills and decodes), number of KV cache movements, or other reported information can be transmitted via event data level 370. In at least one embodiment, one or more preprocessing systems or components can combine, aggregate, filter, summarize, format, and / or analyze events to provide them to the scheduler 610.

[0032] In at least one embodiment, the planner 610 can implement a utilization forecast 612. In at least one embodiment, the utilization forecast 612 can use machine learning methods with a neural network or non-neural machine learning methods to predict future utilization according to a time horizon (e.g., the next 24 hours, the next 72 hours, the next week, the next month, etc.) using information obtained from various sources (e.g., the frontend 320, the tracking 380 of global KV caches, and / or the event data layer 370).In at least one embodiment, the utilization prediction 612 can predict the utilization as a number of inferencing requests, a number and / or use of computing resources and / or any other information that can be used to modify a distributed LLM system using various response measures, such as resource scaling 620, resource and / or data movement 630, resource modification 640 and processing reprioritization 650.

[0033] In at least one embodiment, the planner 610 can implement a failure prediction 614. In at least one embodiment, the failure prediction 614 can use machine learning methods with a neural network or non-neural machine learning methods to predict future failures or errors according to a time horizon (e.g., the next 24 hours, the next 72 hours, the next week, the next month, etc.) using information obtained from various sources (e.g., the frontend 320, the tracking 380 of global KV caches, and / or the event data layer 370).In at least one embodiment, the failure prediction 614 can predict the failure as a percentage of the available computing resources from an allocated set of computing resources, uptime ratios, failure interruptions during inferencing requests, any other failure scenarios that can be prevented or mitigated, and / or any other information that can be used to modify a distributed LLM system using various response measures, such as resource scaling 620, resource and / or data movement 630, resource modification 640, and processing reprioritization 650.

[0034] In at least one embodiment, the planner 610 can implement an engine simulation 616. In at least one embodiment, the engine simulation can be used to determine performance information for different runtimes of a neural network on different types of computing resources (e.g., different types of GPUs). In at least one embodiment, the engine simulation 616 can be implemented using machine learning techniques with a neural network or non-neural machine learning techniques to predict the performance of different inferencing engines (e.g., LLM runtimes) that are used or could be used in a distributed system.In at least one embodiment, the engine simulation 616 can provide information that can be used to perform various response actions, such as resource scaling 620, resource and / or data movement 630, resource modification 640 and processing reprioritization 650.

[0035] In at least one embodiment, the planner 610 can implement the power prediction 618, which can predict, model, or otherwise use methods for modeling power input and consumed power. In at least one embodiment, the power consumption of processing devices can be tracked (e.g., via the event data level 370) and used to predict the future power consumption of the processing devices at different levels of a hierarchy 170. Fig. 1B to predict. In at least one embodiment, the predicted future power consumption can be used to perform various response actions, such as resource scaling 620, resource and / or data movement 630, resource modification 640, and processing reprioritization 650. In at least one embodiment, the power consumption from a power grid (e.g., a public, private, or specific power generation system) can be tracked by (a) data center(s) or (a) data center region(s) (e.g., via the event data layer 370) and used to predict future power consumption from the power grid at various levels of a hierarchy 170 in Fig. 1B to predict. In at least one embodiment, the predicted power consumption can be used to perform various response measures, such as resource scaling 620, resource and / or data movement 630, resource modification 640 and processing reprioritization 650.

[0036] In at least one embodiment, various response measures, such as resource scaling 620, resource and / or data movement 630, resource modification 640, and processing reprioritization 650, can enable a distributed LLM system (or another distributed neural network processing system) to adapt to changing workloads, failure scenarios, external influences, or other events, such as the development of new or additional neural networks, during execution. In at least one embodiment, resource scaling 620 can add or remove computational resources that implement various components of a distributed system, such as the global scheduler 330, the frontend 320, the inferencing processing device(s) 350 for decoding, the inferencing processing device(s) for (initial) population, data layers, and / or the tracking 380 of global KV caches.In at least one embodiment, the resource scaling can apply 620 different criteria or thresholds to different components of a distributed system to determine when and how many computing resources to add or remove.

[0037] In at least one embodiment, the resource / data mover 630 can perform various procedures to move data or resources. For example, in at least one embodiment, the resource / data mover 630 can move a neural network to another processing device to make room for another neural network on a processing device. In at least one embodiment, the resource / data mover 630 can modify an eviction policy for KV cache instances on processing devices (e.g., according to the higher or lower demand expected for KV caches).

[0038] In at least one embodiment, resource modification 640 can perform various changes to applications or components of a distributed LLM system. For example, in at least one embodiment, resource modification 640 can change or modify a neural network runtime engine (e.g., by using an optimization tool to generate a new, modified version of the runtime engine, or by accessing and inserting a previously generated version of the runtime engine). In at least one embodiment, resource modification 460 can change settings or parameters (e.g., increasing or decreasing the size of queues, buffers, or other data structures used by various components of a distributed LLM system).

[0039] In at least one embodiment, the processing reprioritization 650 can be implemented to adjust weights, policies, or other parameters used for scheduling, routing, or making processing decisions for inference requests within a distributed LLM system. For example, in at least one embodiment, the processing reprioritization 650 can adjust various predefined values ​​or thresholds used by the global scheduler 330, such as a scheduling policy 160, to select which sections of a neural network are to be executed on a processing device (e.g., according to the above with reference to Fig. (Scenarios discussed in 1A-3). In at least one embodiment, the processing reprioritization 650 can adjust the processing device queue processing or the inference request processing at the frontend 320 using other control parameters to change the applied prioritization.

[0040] Fig. Figure 7 illustrates an example flowchart of a method implementing the dynamic performance of sections of neural networks according to at least one embodiment. In at least one embodiment, one or more performance metrics can be estimated for one or more sections of one or more neural networks, as described in Figure 710. In at least one embodiment, one or more neural networks can be trained to perform different tasks and can include different stages, subtasks, or other sections that may correspond, for example, to different layers or groups of neurons of the neural networks. In at least one embodiment, as described above with respect to Fig. As discussed in 1A-1B, different sections of one or more neural networks encode, distribute, decode, or perform various other operations and may correspond to one or more of transformer blocks, CNN (folding neural network) blocks, RNN (recurrent neural network) blocks, pooling blocks, tightly connected blocks, and / or other types of blocks.

[0041] In at least one embodiment, as above with regard to Fig. As discussed in 1A-1B, one or more performance metrics are estimated using machine learning models (e.g., neural networks) or other prediction methods (e.g., support vector machines, random forests, decision trees, naive Bayes, k-nearest neighbor, linear regression, logistic regression, and / or Bayesian networks). In at least one embodiment, the performance metrics can be predicted using one or more heuristics and / or a cost function to model performance (e.g., as above with respect to Fig. 1A discussed in detail). In at least one embodiment, the estimated performance metrics can be the inferencing rate per inferencing request (e.g., latency or time spent completing inferencing requests), the inferencing throughput for a processing device, a group of processing devices, or an inferencing system as a whole, such as a distributed LLM system discussed above with respect to Fig. 2 discussed above. In at least one embodiment, the estimation methods for performance metrics can be modified according to workload predictions, fault or failure predictions, or other information generated by a scheduler or other system management component that uses a different and longer prediction horizon than a scheduler or other component that uses estimated performance metrics to dynamically execute sections of one or more neural networks, as discussed above with respect to Fig. 6 is discussed in detail.

[0042] In at least one embodiment, as described above with regard to Fig. As discussed in 1A-1B, estimated performance metrics are based on a data set used for inferencing, such as a number of input tokens and / or an expected number of output tokens, as inputs for or information used by performance metric estimation methods (e.g., a number of tokens in a cost function, as discussed above in relation to Fig. 1A discussed). In at least one embodiment, the estimated performance metrics can be determined based on the resource utilization of various processing devices (e.g., the current utilization of a processing device performing other inference requests, GPU, CPU, memory, I / O, network, or other compute resource utilization information). In at least one embodiment, KV cache information, including the number of KV cache values ​​or other intermediate state information applicable to an inferencing request (e.g., as a result of a (initial) population phase or stage performed with input tokens), can be used as input or a basis for determining the estimated performance metrics.In at least one embodiment, resource utilization, KV cache information, or other information used to estimate performance metrics can be captured or determined using various components and shared as events in an event stream that is transmitted using an event data layer, similar to the methods described above in relation to . Fig. 2-6 are discussed.

[0043] In at least one embodiment, as described above with regard to Fig. As discussed in 1A-1B, one or more sections of one or more neural networks can be dynamically executed by two or more processors, at least partially, based on estimated performance metrics of one or more sections of the one or more neural networks, as described in 720. In at least one embodiment, the dynamic section performance can include various assignments, such as assignments described above with respect to Fig. 2A and Fig. 2B are discussed, wherein a processor can execute multiple sections of a neural network for one inference and only one section for another inference, so that a second section is executed on a different processor. In at least one embodiment, information can be shared between the processors as part of the dynamically executing section(s) of the one or more neural networks, such as by using high-speed communication links or protocols like NVlink, which provide direct high-speed communication between GPUs on a server. In at least one embodiment, KV cache values ​​and / or another intermediate state can be shared (e.g., via an inferencing data plane 360, as discussed above in relation to Fig. 3-5 is discussed). In at least one embodiment, the dynamic execution of sections of neural networks by two or more processors can provide faster inference speeds. In at least one embodiment, hierarchical scheduling methods can take into account two or more processors for scheduling or planning according to different levels of a hierarchy, such as hierarchy 170 in Fig. 1B. DATA CENTER

[0044] Fig. Figure 8 illustrates an exemplary data center 800 in which at least one embodiment may be used. The data center 800 may include one or more rooms with racks 802 and ancillary equipment used to house one or more racks 802 and one or more base plates 804. A rack 802 may contain one or more base plates 804. A rack 802 may include an enclosure that accommodates and supports individual base plates 804. Operational aspects of the rack 802 may be controlled, among other options, at a rack level corresponding to a group of base plates 804 or at a base plate level corresponding to individual base plates 804. The rack 802 or the base plates 804 may, in particular, have selected maximum operating parameters, such as, but not limited to, power consumption, operating frequencies, and others.The 800 data center can be supported by various cooling systems, such as, but not limited to, cooling towers, cooling circuits, pumps, and other support systems. Cooling systems can include sensors and controllers to monitor and manage the cooling characteristics of the 802 racks. Base plates 804 within the 802 racks can receive power from one or more power distribution units (PDUs; not shown). PDUs can be located within the 802 racks, for example, between the 802 racks that contain the 804 base plates, or within the 802 racks that also house the 804 base plates.

[0045] The 802 racks and 804 base plates can contain subsystems, modules, add-in cards, and other semiconductor components. The 804 base plates can contain one or more 806 arithmetic units, which can contain one or more 808 processors, one or more 810 memories, and an 812 interface controller. The 806 arithmetic units can contain any number of processors, such as, but not limited to, central processing units (“CPUs”), graphics processing units (“GPUs”), or other processors (including accelerators, field-programmable gate arrays (FPGAs), graphics processors, etc.), including any of the processors described herein, such as, but not limited to, the processors in the Fig. 9-21. The 806 Computing Units can include one or more 810 Storage Devices (e.g., dynamic read-only memory, solid-state memory, or hard disk drives), as well as network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power supply modules, cooling modules, etc. One or more 806 Computing Units can constitute a server with one or more of the aforementioned computing resources.

[0046] The 806 compute units can consist of separate groupings of compute units located within one or more racks (not shown), or many racks located in data centers at different geographic locations (also not shown). Separate groupings of compute units can include grouped compute, network, storage, or memory resources that can be configured or allocated to support one or more workloads. Multiple compute units (including, for example, CPUs and / or other processors) can be grouped within one or more racks to provide compute resources to support one or more workloads. An 814 resource orchestrator can configure or otherwise control one or more 806 compute units or groups of compute units.The Resource Orchestrator 814 can include a Software Design Infrastructure (“SDI”) management entity for the Data Center 800. The Resource Orchestrator 814 can include hardware, software, or a combination thereof.

[0047] The data center 800 can include any one or any combination of a framework layer 820, a software layer 830, and an application layer 840. As in Fig. As shown in Figure 8, the framework layer 820 includes a job scheduler 822, a configuration manager 824, a resource manager 826, and a distributed file system 828. The framework layer 820 can include a framework to support software 832 of software layer 830 and / or one or more applications 842 of application layer 840. Software 832 or application 842 can each include web-based service software or applications, such as those provided by Amazon Web Services, Google Cloud, and Microsoft Azure. The framework layer 820 can be a type of free and open-source software web application framework, such as Apache Spark™ (hereinafter "Spark"), which can use the distributed file system 828 for data processing at scale (e.g., for "big data").The Job Scheduler 822 can include a Spark driver to facilitate the scheduling of workloads supported by various layers of the Data Center 800. The Configuration Manager 824 can configure various layers, including the Software Layer 830 and the Framework Layer 820, which includes Spark and the Distributed File System 828, to support data processing at scale. The Resource Manager 826 can manage clustered or grouped compute units 806 that are allocated or assigned to support the Distributed File System 828 and the Job Scheduler 822. The Resource Manager 826 can coordinate with the Resource Orchestrator 814 to manage these allocated or assigned compute resources.

[0048] Software 832 can be contained in software layer 830 and can include software used by at least sections of a computing unit 806, one or more computing units 806, groups of computing units 806, and / or the distributed file system 828 of framework layer 820. One or more types of software can include, among others, internet website search software, email virus scanning software, database software, and streaming video content software.

[0049] Application(s) 842 may be contained in the application layer 840 and may include one or more types of applications used by at least sections of a compute unit 806, one or more compute units 806, groups of compute units 806, and / or the distributed file system 828 of the framework layer 820. One or more types of applications may include, among others, any number of a genomic application, a cognitive computing application, and a machine learning application, 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.

[0050] Any configuration manager 824, resource manager 826, or resource orchestrator 814 can implement any number and type of self-modifying actions based on any amount and type of data collected in any technically feasible way. Self-modifying actions can relieve a data center operator of data center 800 from potentially making poor configuration decisions and potentially avoiding underutilized and / or underperforming sections of a data center.

[0051] The Data Center 800 may include tools, services, software, or other resources for training one or more machine learning models or for predicting or deriving information using one or more machine learning models according to one or more embodiments described herein. For example, a machine learning model may be trained by calculating weighting parameters according to a neural network architecture using software and computing resources described above in relation to the Data Center 800. Trained machine learning models corresponding to one or more neural networks may be used to derive or predict information using resources described above in relation to the Data Center 800 by using weighting parameters calculated by one or more training methods described herein.

[0052] The Data Center 800 can support CPUs, application-specific integrated circuits (ASICs), GPUs, FPGAs, or other hardware (e.g., implementations in the Fig. 9-21) to perform some or all of the processes and procedures described elsewhere herein, such as, but not limited to, training and / or inferencing using the resources described above. In addition, one or more of the software and / or hardware resources described above may be configured as a service to enable users to train or perform information inferencing, such as, but not limited to, image recognition, speech recognition, or other artificial intelligence services.

[0053] In at least one embodiment, the processor 808 may comprise one of the processors listed below and / or include one or more circuits for dynamically executing sections of neural networks or for otherwise performing one of the operations described above or elsewhere herein. In at least one embodiment, the processor 808 is configured by the software 832 to dynamically execute the sections of the neural networks or otherwise perform one of the operations described above or elsewhere herein. The data center 800 may include logic, CPUs, application-specific integrated circuits (ASICs), GPUs, FPGAs, or other hardware (e.g., embodiments in the Fig. 9-21) to perform any of the operations described above or elsewhere herein. PROCESSORS

[0054] The following figures show, without limitation, examples of processors and processing systems that can be used to dynamically execute sections of neural networks or to otherwise perform some or all of the processes, operations, and / or procedures described herein. Exemplary processors and processing systems can be configured by software to dynamically execute the neural network sections or otherwise perform any of the operations described above or elsewhere herein. Processors and processing systems can be logic, central processing units (CPUs), application-specific integrated circuits (ASICs), graphics processing units (GPUs), field-programmable arrays (FPGAs), XPUs (i.e., any computer architecture that best meets the requirements of an application), or other hardware (e.g., embodiments in the Fig. 9-21) to perform one of the operations described above, below, or elsewhere herein. The processors and / or processing systems described herein may include one or more circuits that can be used to dynamically execute the neural network sections or to perform one of the operations described above or elsewhere herein. As used herein, one or more circuits may be configured by software to dynamically execute the neural network sections or to perform one of the operations described above or elsewhere herein. Fig. 26A and Fig. Figure 26B illustrates a logic 2615 which, as described elsewhere herein, can be used in one or more devices to perform operations such as, but not limited to, those discussed herein according to at least one embodiment. The logic may, for example, refer to any combination of software logic, hardware logic, and / or firmware logic to provide the functionality and / or operations described herein, wherein the logic, collectively or individually, may be designed as a circuit that forms part of a larger system, such as an integrated circuit (IC), an application-specific integrated circuit (ASIC), a field-programmable array (FPGA), a system-on-a-chip (SoC), or one or more processors (e.g., CPU, GPU).

[0055] Fig. Figure 9 illustrates a processor that is a System-on-a-Chip (SOC) 900 (which may be referred to as a System-on-Chip, a Superchip, or by another name) according to at least one embodiment. The SOC 900 may include a Processor Complex 910 and a Processor Complex 940. The SOC 900 may include any number of Processor Complexes 910 and / or Processor Complexes 940, which may include any number of processors described herein, such as, but not limited to, those in the Fig. 9-21, in any combination. For example, the 910 processor may include a central processing unit (CPU), and the 940 processor may include a graphics processor. Alternatively, the 910 processor may include a graphics processor, and the 940 processor may include a graphics processor. The 900 system-on-a-chip (SOC) may include any number of 992 display controllers, any number of 994 multimedia engines, any number of 970 I / O interfaces, any number of 980 memory controllers, and any number of 960 fabrics in any combination. For explanatory purposes, multiple instances of the same object are referred to herein by reference symbols that identify the object and by numbers in parentheses that identify the instance, if necessary. The 900 SOC may include a Broadcom processor in Palo Alto, CA.

[0056] The 910 processor complex can include a CPU, the 940 processor complex can include a GPU, and the 900 SOC can be a processing unit that integrates the 910 and 940 onto a single chip. Some tasks can be assigned to the 910 processor complex, and other tasks can be assigned to the 940 processor complex. The 910 processor complex can be configured to execute main control software associated with the 900 SOC, such as an operating system. The 910 processor complex can be the main processor of the 900 SOC, controlling and coordinating operations of other processors. The 910 processor complex can issue instructions that control the operation of the 940 processor complex to perform some or all of the operations described herein. The 910 processor complex can be configured to execute host-executable code generated by CUDA or other source code (e.g.,HIP source code) is derived, and the 940 processor complex can be configured to execute device-executable code derived from CUDA or other source code to perform any of the operations described herein.

[0057] The 910 processor complex can include 920(1)-920(4) cores and a 930 cache (e.g., L3 cache) to store information for performing the operations described herein. The 910 processor complex can include any number of 920 cores and any number and type of caches in any combination. The 920 cores can be configured to execute instructions of a specific instruction set architecture (ISA) to perform some or all of the operations described herein. Each 920 core can contain one CPU core. The 920(1)-920(4) cores can be referred to as compute units. The 900 system-on-a-chip (SOC) can include any number of 910 processor complexes, 960 fabrics, 970 I / O interfaces, and 980 memory controllers.

[0058] Each Core 920 can include a Fetch / Decode Unit 922, an Integer Execution Engine 924, a Floating-Point Execution Engine 926, and an L2 Cache 928. The Fetch / Decode Unit 922 can fetch instructions to perform some or all of the operations described herein (such as, but not limited to, an API compiled to instructions), decode such instructions, generate micro-operations, and output separate micro-instructions to the Integer Execution Engine 924 and / or the Floating-Point Execution Engine 926. The Fetch / Decode Unit 922 can simultaneously output one micro-instruction to the Integer Execution Engine 924 and another micro-instruction to the Floating-Point Execution Engine 926. The Integer Execution Engine 924 can perform integer and memory operations. The 926 floating-point execution engine can perform floating-point and vector operations.The 922 retrieval / decoding unit can output micro-instructions to one or more execution engines, replacing both the 924 integer execution engine and the 926 floating-point execution engine.

[0059] Each kernel 920(i), where i is an integer representing a specific instance of kernel 920, can access the L2 cache 928(i) contained within kernel 920(i). Each kernel 920 contained within kernel complex 910(j), where j is an integer representing a specific instance of kernel complex 910, can be connected to other kernels 920 contained within kernel complex 910(j) via the L3 cache 930(j) contained within kernel complex 910(j). All kernels 920 contained within kernel complex 910(j), where j is an integer representing a specific instance of kernel complex 910, can access the entire L3 cache 930(j) contained within kernel complex 910(j). The L3 cache 930 can contain any number of slices.

[0060] The Processor 940 can be a graphics complex configured to perform computational operations (e.g., computational operations involved in the operations described herein) in a highly parallel manner. The Processor 940 can be configured to execute graphics pipeline operations, such as, but not limited to, drawing commands, pixel operations, geometric calculations, and other operations associated with rendering an image on a display. The Processor 940 can be configured to perform non-graphics-related operations, such as, but not limited to, training and / or simulating a neural network. The Processor 940 can be configured to perform both graphics-related and non-graphics-related operations.

[0061] The 940 processor complex can include any number of 950(1)-950(N) arithmetic units, where N is any integer greater than 1, and an L2 cache 942. The 950 arithmetic units can use the L2 cache 942, which can store information to be used to perform some or all of the operations described herein. The L2 cache 942 can be partitioned. The 940 processor complex can include any number of 950 arithmetic units and any number (including zero) and type of caches. The 940 processor complex can include any amount of dedicated graphics hardware.

[0062] Each compute unit 950 can contain any number of SIMD units 952(1)-952(N), where N is any integer greater than 1, and a shared memory 954. Each SIMD unit 952 can implement a SIMD architecture and can be configured in parallel for some or all of the operations described herein. Each compute unit 950 can execute any number of thread blocks, but each thread block can be executed on a single compute unit 950, although in some embodiments a thread block can be executed on multiple compute units. A thread block can contain any number of execution threads. A workgroup can be a thread block. Each SIMD unit 952 can execute a group of threads. A group of threads (e.g., 16 threads), which can also be called a warp, subgroup, or wavefront (e.g.,(as used by AMD and Intel), where each thread in the warp, wave, subgroup, or wavefront can belong to a single thread block and is configured to process a different data set based on a single set of instructions. Prediction can be used to disable one or more threads in a warp, subgroup, or wavefront. A lane can be a thread. A worker can be a thread, as with OpenCL, but is not limited to this. Different warps, subgroups, or wavefronts in a thread block can synchronize with each other and communicate via shared memory. Each compute unit can contain one or more thread block clusters, where a thread block cluster can allow programmatic control of locality at a granularity greater than that of a single thread block on a single streaming multiprocessor (SM).Thread block clusters (also referred to as "clusters") enable multiple thread blocks running concurrently across streaming multiprocessors to synchronize and collaboratively retrieve, exchange, or otherwise use data. In at least one embodiment, streaming multiprocessors ("SMs") can be referred to as streaming microprocessors, stream processors ("SPs"), stream processing units ("SPUs"), compute units ("CUs"), execution units ("EUs"), and / or slices, where a slice in this context can refer to a portion of processing resources within a processing unit (e.g., 16 cores, a ray tracing unit, a thread director, or a scheduler).

[0063] Fabric 960 can be a system interconnect that enables data and control transfers across the 910 processor complex, the 940 processor complex, the 970 I / O interfaces, the 980 memory controllers, the 992 display controller, and the 994 multimedia engine to perform, for example, some or all of the operations described herein. The SOC 900 can include any number and type of system interconnects in addition to or instead of Fabric 960, enabling data and control transfers across any number and type of components directly or indirectly connected to the SOC 900, which may be internal or external to the SOC 900. The 970 I / O interfaces can represent any number and type of I / O interface (e.g., PCI, PCI-Extended (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). Various types of peripheral devices can be connected to the 970 I / O interfaces.Peripheral devices that can be connected to the I / O interfaces of the 970 can include keyboards, mice, printers, scanners, joysticks or other types of game controllers, media recording devices, external storage devices, network interface cards, and so on.

[0064] The display controller 992 can display images on one or more display devices, such as, but not limited to, a liquid crystal display ("LCD") device. The multimedia engine 994 can include any quantity and type of multimedia-related circuitry, such as, but not limited to, a video decoder, a video encoder, an image signal processor, etc. The memory controllers 980 can facilitate data transfers between the SOC 900 and a unified system memory 990. The processor complex 910 and the processor complex 940 can share the unified system memory 990. The unified system memory 990 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 double-rate graphics memory (GDDR).The Unified System Memory 990 can include 3D stacked memory, including, but not limited to, High Bandwidth Memory (HBM), HBM2e, or HDM3.

[0065] The SOC 900 can implement a memory subsystem that includes any number and type of memory controllers 980 and memory devices (e.g., shared memory 954) that can be allocated to a component or shared by multiple components to perform any of the operations described herein. The SOC 900 can implement a cache subsystem that includes one or more cache memories (e.g., L2 caches 928, L3 cache 930, and L2 cache 942), each of which can be private or shared by any number of components (e.g., cores 920, core complex 910, SIMD units 952, arithmetic units 950, and processor complex 940).

[0066] In at least one embodiment, the SOC 900 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0067] Fig. Figure 10A illustrates a parallel processor 1000 according to at least one embodiment. The parallel processor 1000 can be implemented using one or more circuits and can be a programmable processor (e.g., a CPU and / or GPU), logic, an application-specific integrated circuit (ASIC), a field-programmable gate array (FPGA), or other hardware (e.g., embodiments in the Fig. 9-21) to perform any of the operations described above or elsewhere herein.

[0068] The parallel processor 1000 may include a parallel processing unit 1002 to perform any of the operations described above or elsewhere herein. The parallel processing unit 1002 may include an I / O unit 1004, which enables communication with other devices, including other instances of the parallel processing unit 1002. The I / O unit 1004 may be directly connected to other devices. The I / O unit 1004 may be connected to other devices via a hub or switch interface, such as, but not limited to, a memory hub 1005. Connections between the memory hub 1005 and the I / O unit 1004 may form a communication link 1013.The I / O unit 1004 can be connected to a host interface 1006 and a memory crossbar 1016, with the host interface 1006 receiving commands to perform processing operations and the memory crossbar 1016 receiving commands to perform memory operations.

[0069] When the host interface 1006 receives a command buffer via the I / O unit 1004, the host interface 1006 can forward work operations to a frontend 1008 to execute these commands. The frontend 1008 can be coupled to a scheduler 1010 (which can be referred to as a sequencer) configured to distribute commands or other work items to a processing cluster array 1012. The scheduler 1010 can ensure that the processing cluster array 1012 is correctly configured and in a valid state before tasks can be distributed to a cluster of the processing cluster array 1012. The scheduler 1010 can be implemented using firmware logic running on a microcontroller.The microcontroller-implemented Scheduler 1010 can be configured to perform complex scheduling and workload distribution operations with coarse and fine granularity, enabling fast pre-allocation and context switching of threads running on the Processing Array 1012. Host software can discover workloads for scheduling on the Processing Array 1012 via one of several graphics processing paths. Workloads can then be automatically distributed across the Processing Array 1012 by the logic of Scheduler 1010 within a microcontroller that incorporates Scheduler 1010.

[0070] The processing cluster array 1012 can perform any of the operations described above or elsewhere herein and can contain up to "N" processing clusters (e.g., cluster 75A, cluster 75B to cluster 75N), where "N" is a positive integer (which may be a different integer "N" than used in other figures). Each cluster 75A-75N of the processing cluster array 1012 can execute a large number of concurrently running threads. The scheduler 1010 can allocate work to clusters 1014A-1014N of the processing cluster array 1012 using various scheduling and / or workload distribution algorithms, which may vary depending on the workload generated by each type of program or computation.Scheduling can be handled dynamically by the Scheduler 1010 or can be partially supported by compiler logic during the compilation of program logic designed for execution by the Processing Cluster Array 1012. Different clusters 1014A–1014N of the Processing Cluster Array 1012 can be assigned to process different types of programs or to perform different types of calculations.

[0071] The Processing Cluster Array 1012 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 1012 can be configured to perform parallel general-purpose computing operations. For example, the Processing Cluster Array 1012 can include logic for performing processing tasks, including filtering video and / or audio data, performing modeling operations, including physical operations, and performing data transformations.

[0072] The processing cluster array 1012 can be configured to perform parallel graphics processing operations. The processing cluster array 1012 can include additional logic to support 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. The processing cluster array 1012 can be configured to execute graphics processing-related shader programs, such as, but not limited to, vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. The parallel processing unit 1002 can transfer data from system memory via the I / O unit 1004 for processing. During processing, the transferred data can be stored in on-chip memory (e.g.,The data is stored in parallel processor memory (1022) and then written back to system memory.

[0073] When the parallel processing unit 1002 is used to perform graphics processing, the scheduler 1010 can be configured to divide a workload into approximately equal tasks to better distribute graphics processing operations across multiple clusters 1014A–1014N of the processing cluster array 1012. Sections of the processing cluster array 1012 can be configured to perform different types of processing. For example, a first section can be configured to perform vertex shading and topology generation, a second section can be configured to perform tessellation and geometry shading, and a third section 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 of the clusters 1014A - 1014N can be stored in buffers to allow intermediate data to be transferred between clusters 1014A - 1014N for further processing.

[0074] The processing cluster array 1012 can receive processing tasks to be executed via the scheduler 1010, which receives commands defining these tasks from the frontend 1008. The processing tasks can include indices 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 the data is to be processed (e.g., which program is to be executed). The scheduler 1010 can be configured to retrieve indices corresponding to tasks or to receive indices from the frontend 1008. The frontend 1008 can be designed to ensure that the processing cluster array 1012 is configured in a valid state before initiating a workload specified by incoming command buffers (e.g., batch buffer, push buffer, etc.).

[0075] Each of one or more instances of the parallel processing unit 1002 can be coupled to a parallel processor memory 1022 to perform one of the operations described above or elsewhere herein. The parallel processor memory 1022 can be accessed via the memory crossbar 1016, which can receive memory requests from the processing cluster array 1012 and the I / O unit 1004. The memory crossbar 1016 can access the parallel processor memory 1022 via a memory interface 1018. The memory interface 1018 can contain multiple partition units (e.g., partition unit 1020A, partition unit 1020B through partition unit 1020N), each of which can be coupled to a section (e.g., a memory unit) of the parallel processor memory 1022.A number of partition units 1020A–1020N can correspond to a number of storage units, such that a first partition unit 1020A has a corresponding first storage unit 1024A, a second partition unit 1020B has a corresponding storage unit 1024B, and an Nth partition unit 1020N has a corresponding Nth storage unit 1024N. A number of partition units 1020A–1020N cannot correspond to a number of storage units.

[0076] The 1024A-1024N memory units can contain various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM), including double data rate graphics memory (GDDR). The 1024A-1024N memory units can also contain 3D stacked memory, including high-bandwidth memory (HBM), HBM2e, or HDM3. Render targets, such as frame buffers or texture maps, can be stored across the 1024A-1024N memory units, allowing the 1020A-1020N partition units to write portions of each render target in parallel, thus efficiently utilizing the available bandwidth of the 1022 parallel processor memory.A local instance of the parallel processor memory 1022 can be excluded in favor of a unified memory design that uses system memory in conjunction with local cache memory.

[0077] Each of the clusters 1014A–1014N of the processing cluster array 1012 can process data written to any of the storage units 1024A–1024N within the parallel processor memory 1022. The memory crossbar 1016 can be configured to transfer an output from each cluster 1014A–1014N to any partition unit 1020A–1020N or to another cluster 1014A–1014N, which can perform additional processing operations on the output. Each cluster 1014A–1014N can communicate with the memory interface 1018 via the memory crossbar 1016 to read from or write to various external storage devices.The memory crossbar 1016 can have a connection to the memory interface 1018 to communicate with the I / O unit 1004, as well as a connection to a local instance of the parallel processor memory 1022, which allows processing units within different processing clusters 1014A–1014N to communicate with system memory or other memory that is not local to the parallel processing unit 1002. The memory crossbar 1016 can use virtual channels to separate traffic flows between the clusters 1014A–1014N and the partition units 1020A–1020N.

[0078] Multiple instances of the Parallel Processing Unit 1002 can be deployed on a single add-on card, or multiple add-on cards can be interconnected. Different instances of the Parallel Processing Unit 1002 can be configured to work together, even if they have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences. For example, some instances of the Parallel Processing Unit 1002 can include higher-precision floating-point units relative to other instances.Systems incorporating one or more instances of the Parallel Processing Unit 1002 or the Parallel Processor 1000 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.

[0079] Fig. 10A further includes a block diagram of a partitioning unit 1020 according to at least one embodiment. The partitioning unit 1020 is an instance of one of the partitioning units 1020A-1020N of Fig. 10A. The partitioning unit 1020 can include an L2 cache 1021, an image buffer interface 1025, and a ROP 1026 (raster operation unit). The L2 cache 1021 can be a read / write cache configured to perform load and store operations received from the memory crossbar 1016 and the ROP 1026. Read failures and urgent write-back requests can be passed from the L2 cache 1021 to the image buffer interface 1025 for processing. Updates can also be sent to an image buffer via the image buffer interface 1025 for processing. The image buffer interface 1025 can be connected to one of the memory units in the parallel processor memory, such as, among others, memory units 1024A–1024N. Fig. 10A (e.g., within the parallel processor memory 1022).

[0080] The ROP 1026 can be a processing unit that performs raster operations, such as stenciling, z-testing, blending, and others. The ROP 1026 can then output processed graphics data, which is stored in graphics memory. The ROP 1026 can include compression logic to compress depth or color data written to memory and decompress depth or color data read from memory. This compression logic can be lossless and employ one or more of several compression algorithms. The type of compression performed by the ROP 1026 can vary based on the statistical properties of the data being compressed. For example, delta color compression is performed on depth and color data on a per-tile basis.

[0081] The ROP 1026 can be used within any processing cluster (e.g., clusters 1014A-1014N of Fig. 10A) instead of being located within the partitioning unit 1020. Read and write requests for pixel data can be transmitted via the memory crossbar 1016 instead of pixel fragment data. Processed graphics data can be displayed on a screen, which can then be routed for further processing by the processor(s) 1802 or by one of the processing entities within the parallel processor 1000. Fig. 10A can be routed.

[0082] In at least one embodiment, the parallel processor 1000 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0083] Fig. Figure 10B includes a block diagram of a processing cluster 1014 within a parallel processing unit according to at least one embodiment. A processing cluster can be an instance of one of the processing clusters 1014A to 1014N of Fig. 10A, which can be used to perform any of the operations described above or elsewhere herein. Processing cluster 1014 can be configured to execute many threads in parallel, where a "thread" refers to an instance of a particular program running on a particular set of input data. Single-instruction, multiple-data (SIMD) instruction output methods can be used to support the parallel execution of a large number of threads without providing multiple independent instruction units. Single-instruction, multiple-thread (SIMT) methods can be used to support the parallel execution of a large number of generally synchronized threads, using a common instruction unit configured to issue instructions to a set of processing engines within each of the processing clusters.

[0084] The operation of the processing cluster 1014 can be controlled via a pipeline manager 1032, which distributes processing tasks to parallel SIMT processors. The pipeline manager 1032 can receive instructions from the scheduler 1010. Fig. 10A receives and manages the execution of these instructions via a graphics multiprocessor 1034 and / or a texture unit 1036. The graphics multiprocessor 1034 can be an example instance of a parallel SIMT processor. However, various types of parallel SIMT processors of different architectures can be contained within the processing cluster 1014. One or more instances of the graphics multiprocessor 1034 can be contained within a processing cluster 1014. The graphics multiprocessor 1034 can process data, and a data crossbar 1040 can be used to distribute processed data to one of several possible destinations, including other shader units. The pipeline manager 1032 can facilitate the distribution of processed data by specifying destinations for processed data to be distributed via the data crossbar 1040.

[0085] Each 1034 graphics multiprocessor within the 1014 processing cluster can contain an identical set of function execution logic (e.g., arithmetic logic units, load-store units, etc.) to perform computations for any of the operations described above or elsewhere herein. The function execution logic can be pipelined, allowing new instructions to be issued before previous instructions have completed. The function execution logic can support a wide variety of operations, including integer and floating-point arithmetic, comparison operations, Boolean operations, bit shifting, and the computation of various algebraic functions. The same function unit hardware can be used to perform different operations, and any combination of function units is permitted.

[0086] Instructions passed to the 1014 processing cluster can form threads, which can also be referred to as warps, subgroups, waves, or wavefronts. A set of threads executed across a set of parallel processing engines can be called a thread group. A thread group can execute a common program with different input data. Each thread within a thread group can be assigned to a different processing engine within a 1034 graphics multiprocessor. A thread group can contain fewer threads than the number of processing engines within the 1034 graphics multiprocessor. If a thread group contains fewer threads than the number of processing engines, one or more processing engines may be idle during cycles in which that thread group is being processed.A thread group can contain more threads than the number of processing engines within the 1034 graphics multiprocessor. If a thread group contains more threads than the number of processing engines within the 1034 graphics multiprocessor, processing can be performed over successive clock cycles. Multiple thread groups can run concurrently on a single 1034 graphics multiprocessor.

[0087] The 1034 graphics multiprocessor includes an internal cache for performing load and store operations, such as, but not limited to, any of the operations described above or elsewhere herein. The 1034 graphics multiprocessor may forgo an internal cache and use a cache (e.g., the 1048 L1 cache) within the 1014 processing cluster. Each 1034 graphics multiprocessor may also access L2 caches within partitioning units (e.g., partitioning units 1020A–1020N of the 1014). Fig. 10A) which can be shared by all processing clusters 1014 and used to transfer data between threads. The graphics multiprocessor 1034 can also access global off-chip memory, which may include one or more instances of local parallel processor memory and / or system memory. Any memory outside the parallel processing unit 1002 can be used as global memory. The processing cluster 1014 can contain multiple instances of the graphics multiprocessor 1034 and can use shared instructions and data that can be stored in the L1 cache 1048.

[0088] Each processing cluster 1014 can contain an MMU 1045 (memory management unit), which can be configured to map virtual addresses to physical addresses. One or more instances of the MMU 1045 can be located within the memory interface 1018 of Fig. The MMU 1045 can include a set of page table entries (PTEs) used to map a virtual address to a physical address of a tile and, optionally, a cache row index. The MMU 1045 can include address translation lookaside buffers (TLBs) or caches, which may reside within the graphics multiprocessor 1034, the L1 cache 1048, or the processing cluster 1014. A physical address can be processed to distribute surface data access locally, enabling efficient request nesting between partitioning units. A cache row index can be used to determine whether a request for a cache row is a hit or a failure.

[0089] A processing cluster 1014 can be configured such that each graphics multiprocessor 1034 is coupled to a texture unit 1036 for performing texture mapping operations, such as determining texture sample positions, reading texture data, and filtering texture data. Texture data can be read from an internal texture L1 cache (not shown) or from an L1 cache within the graphics multiprocessor 1034 and can be retrieved from an L2 cache, local parallel processor memory, or system memory as needed. Each graphics multiprocessor 1034 can output processed tasks to the data crossbar 1040 to provide a processed task to another processing cluster 1014 for further processing or to store a processed task in an L2 cache, local parallel processor memory, or system memory via the memory crossbar 1016.A Pre-ROP 1042 (Pre-Raster Operation Unit) can be configured to receive data from the Graphics Multiprocessor 1034 and to route data to ROP units which can be arranged with partitioning units as described herein (e.g. the partitioning units 1020A -1020N of . Fig. 10A). The Pre-ROP 1042 unit can perform optimizations for color mixing, organizing pixel color data, and performing address translations.

[0090] In at least one embodiment, the processing cluster 1014 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0091] Fig. Figure 10C shows a graphics multiprocessor 1034 according to at least one embodiment for performing, for example, any of the operations described above or elsewhere herein. The graphics multiprocessor 1034 can be coupled to the pipeline manager 1032 of the processing cluster 1014. The graphics multiprocessor 1034 can include an execution pipeline, including, but not limited to, an instruction cache 1052 (which can, for example, store instructions such as, but not limited to, compiled API instructions), an instruction unit 1054, an address allocation unit 1056, a register file 1058, one or more cores 1062 of a general-purpose graphics processing unit (GPGPU), and one or more load / store units 1066, wherein one or more load / store units 1066 can perform load / store operations to load / store instructions according to the execution of an operation.The GPGPU cores 1062 and the load / store units 1066 can be coupled to the cache memory 1072 and the shared memory 1070 via a memory and cache intermediate 1068. The GPGPU cores 1062 can be part of a SoC, such as, but not limited to, the integrated circuit 900 in [reference missing]. Fig. 9.

[0092] The instruction cache 1052 can receive a stream of instructions (e.g., to perform any of the operations described above or elsewhere herein) for execution from the pipeline manager 1032. Instructions can be cached in the instruction cache 1052 and issued for execution by an instruction unit 1054. The instruction unit 1054 can issue instructions as thread groups (e.g., warps, subgroups, wavefronts, or waves), with each thread in the thread group being assigned to a different execution unit within the GPGPU cores 1062. An instruction can access any address from a local, shared, or global address space by specifying an address within a unified address space.The address mapping unit 1056 can be used to translate addresses in a uniform address memory area into a unique memory address that can be accessed by the load / store units 1066.

[0093] Register file 1058 can provide a set of registers for functional units of the graphics multiprocessor 1034. Register file 1058 can provide temporary storage for operands associated with data paths of functional units (e.g., GPGPU cores 1062, load / store units 1066) of the graphics multiprocessor 1034. Register file 1058 can be subdivided among each of the functional units, so that each functional unit is assigned a dedicated section of register file 1058. Register file 1058 can be subdivided among different warps (which can be referred to as wavefronts, subgroups, and / or waves or threads) executed by the graphics multiprocessor 1034.

[0094] The GPGPU 1062 cores can each contain floating-point units (FPUs) and / or integer arithmetic logic units (ALUs) that can be used to execute instructions from the 1034 graphics multiprocessor. The GPGPU 1062 cores can be similar or different in architecture. A first section of the GPGPU 1062 cores can contain a single-precision FPU and an integer ALU, while a second section of the GPGPU cores can contain a double-precision FPU. FPUs can implement standard IEEE 754-2008 floating-point arithmetic or enable variable-precision floating-point arithmetic. The 1034 graphics multiprocessor can additionally include one or more fixed-function or special-function units to perform specific functions, such as, but not limited to, copy-rectangle or pixel-mixing operations.One or more of the GPGPU 1062 cores can also include logic with a fixed or special function.

[0095] The GPGPU Cores 1062 can include SIMD logic capable of executing a single instruction on multiple data sets. The GPGPU Cores 1062 can physically execute SIMD4, SIMD8, and SIMD16 instructions and logically execute SIMD1, SIMD2, and SIMD32 instructions. SIMD instructions for GPGPU Cores can be generated at compile time by a shader compiler or automatically when running programs written and compiled for Single Program Multiple Data (SPMD) or SIMT architectures. Multiple threads of a 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 over a single SIMD8 logic unit.

[0096] The memory and cache intermediary 1068 can include a link network that connects each functional unit of the graphics multiprocessor 1034 to the register file 1058 and the shared memory 1070. The memory and cache intermediary 1068 can be a crossbar intermediary, allowing the load / store unit 1066 to perform load and store operations between the shared memory 1070 and the register file 1058. The shared memory 1070 can be used to enable communication between threads running on functional units within the graphics multiprocessor 1034. For example, the cache memory 1072 can be used as a data cache to temporarily store texture data between the functional units and the texture unit 1036. The shared memory 1070 can also be used as a program-managed cache.Threads running on the GPGPU cores 1062 can programmatically store data within the shared memory in addition to automatically cached data stored within the cache memory 1072.

[0097] A parallel processor or GPGPU, as described herein, can be communicatively coupled to host / processor cores to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general-purpose GPU (GPGPU) functions. A GPU can be communicatively coupled to host processor / cores via a bus or other intermediary connection (e.g., a high-speed intermediary such as PCIe or NVLink, but not limited to). A system-on-a-chip (SoC) can incorporate a parallel processor or GPGPU, as described herein, with the parallel processor or GPGPU running on the SoC. A GPU can be integrated as cores within a package or chip and communicatively coupled to cores via an internal processor bus / intermediary within the package or chip.Regardless of how a GPU is connected, processor cores can assign work to such a GPU in the form of sequences of instructions contained in a work descriptor. The GPU can then use dedicated circuitry / logic to efficiently process these instructions to perform any of the operations described above or elsewhere herein.

[0098] In at least one embodiment, the graphics multiprocessor 1034 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0099] Fig. Figure 11 shows a Processor 1100 according to at least one embodiment. The Processor 1100 may include a hybrid-architecture processor (e.g., Lunar Lake or Meteor Lake) from Intel Corporation in Santa Clara, CA, or another processor that uses at least some of the components described herein. The Processor 900 may include one or more central processing units (CPUs), one or more graphics processing units (GPUs), and / or one or more neural processing units (NPUs), which may, for example, be a dedicated AI accelerator to which artificial intelligence (AI) workloads are offloaded from the CPU and GPU. The Processor 900 may use instructions that, when executed, cause the Processor 900 and / or any of its components to perform some or all of the processes and procedures described elsewhere herein.The 1100 processor can include any number of memory and cache units 1110 to facilitate processing among the various components. Memory and cache 1110 on the 1100 processor can 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. With respect to the 1100 processor and any of its components described above or elsewhere herein, one or more of the APIs described herein can, for example, be compiled into instructions that are retrieved by instruction-fetching logic or equivalent, decoded by a processor decoder or equivalent, and scheduled for execution by a scheduler or equivalent (e.g., a program).APIs (in order or out of order) can be executed by execution logic or equivalent, reordered, and then terminated by retirement logic or equivalent. APIs (and / or compiled instructions containing APIs) can be stored in any memory outside or inside the processor (e.g., cache and / or memory). A result of APIs can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache, or other memory. One or more of the APIs described herein can contain a call.

[0100] The 1100 processor can incorporate compute engines as CPUs 1102 and can include any number of cores, such as up to 16 cores / 22 threads. Cores in the CPU 1102 can be P cores (Performance), E cores (Efficient), and LP-E cores (Low-Power Efficient). Performance cores can be used for single-threaded, computationally intensive workloads with low latency, while Efficient cores can be used for multi-threaded, less computationally intensive workloads. Low-Power Efficient cores can be used for scalable multi-threaded performance and offloading background tasks. P cores can be used for single- and limited-threaded performance, while E and LP-E cores can be used for multi-threaded throughput and power efficiency.

[0101] The GPU 1106 can include any number of graphics engines, such as Intel® Arc™ Graphics Engines (Xe LPG) with 8 Xe cores (up to 128 execution units or EUs). As in Fig. As shown in Figure 11, the GPU 1106 can include vector engines 1110 and matrix engines 1112, which can, for example, execute FP, INT, and matrix operations all simultaneously, separately, or in batches. The GPU 1106 can include a load / store unit 1114 as well as other memory, such as, but not limited to, an instruction cache (I$) 1116 and an L1 cache / subsystem local memory (SLM) 1118, which can, for example, store instructions to perform any of the operations described above or elsewhere herein.

[0102] The NPU 1104 can contain one or more integrated neural processing units (NPUs) with Intel® AI Boost. The NPU 1104 can be mapped to the host processor as an integrated PCIe device. The NPU 1104 can include one or more (e.g., two) NCE tiles 1130 (NCE = Neural Compute Engine). Each tile can be configured with any combination of, but not limited to (e.g., 2300), MAC engines 1134 (MAC = Multiply Accumulate), a post-processing engine (not shown), an AI DSP processor (not shown), and memory (2 MB dedicated SRAM) per tile, as shown in Fig. Figure 11 shows that for general computing needs, the Neural Compute Engines 1130 can include SHAVE engines 1128 (SHAVE = Streaming Hybrid Architecture Vector Engines) 1128 for high-performance parallel computing, and DMA engines 1124 (DMA = Direct Memory Access) to transfer data between system memory DRAM (Dynamic Random Access Memory) 1126 and a software-managed cache. The integrated MMU (Memory Management Unit) 1122 plus IOMMU (Input-Output Memory Management Unit) (not shown) can support multiple hardware contexts simultaneously and provide security isolation between execution contexts according to the MCDM architecture (MCDM = Microsoft Compute Driver Model).The Processor 1100 may also include a media unit (not shown) that is on top of or separate from the XCDs or other components of the processor to enable video playback and video processing of compressed or uncompressed data, such as using HEVC, AV1, VP9 and AVC-HW accelerated decoding support and HEVC, VP9 and AVC-HW accelerated encoding support.

[0103] An Intel® Thread Director, which includes firmware integrated into the processor, can prioritize and manage workload distribution, sending tasks to optimized cores. For example, the Thread Director can link P cores, E cores, and / or LP-E cores (described above), along with task scheduling capabilities and the ability to send less demanding tasks to the E cores or LP-E cores. Intel® Deep Learning Boost (Intel® DL Boost) (not shown) can provide built-in AI acceleration for training and inference workloads and can include VNNI instruction set support (for a CPU) and DP4a instruction set support (for a GPU). This instruction set can be optimized with the OpenVINO™ toolkit and oneAPI to accelerate INT8 inferencing. A software stack, such as...As described elsewhere herein, it can be used to enable AI inference using the OpenVINO™ toolkit. The 1100 processor can be configured to run an application program, such as a CUDA program.

[0104] In at least one embodiment, the processor 1100 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0105] The Processor 1100 may alternatively include a processor based on an AI Engine Direct architecture from Qualcomm Corporation in Santa Clara, CA, or another processor that uses at least some of the components described herein, which may include any number of NPUs, GPUs, CPUs, and other associated components, such as, but not limited to, the NPU 1104 as a Hexagon NPU, the GPU 1106 as an Adreno GPU, the CPU 1102 as a Kryo or Qualcomm Oryon CPU, and a Qualcomm Sensing Hub (not shown) and a Memory Subsystem 1110 in any combination. The Hexagon NPU 1104 may include a power rail, a microtile inference unit, a hardware acceleration unit, a tensor unit, a scalar unit, and a vector unit (all not shown) that have dedicated memory or memory (e.g.,Cache or memory (such as HBM3) can be used to store instructions for performing any of the operations described above or elsewhere herein. The Adreno GPU 1106 can provide graphics and parallel processing for AI in formats such as, but not limited to, 32-bit floating point (FP32), 16-bit floating point (FP16), and 8-bit integer (INT8). Kryo or Qualcomm Oryon CPUs 1102 can perform AI workloads and can handle contextualization for ubiquitous generative AI applications. The CPU 1102 can also include an instruction fetch unit, a rename and retract unit, a memory management unit, a vector execution unit, an integer execution unit, and a load and store unit for processing and instruction management.With respect to the 1100 processor and any of its components described above or elsewhere herein, one or more of the APIs described herein may, for example, be compiled into instructions that may be fetched by the instruction fetch unit, decoded by a processor decoder or equivalent, scheduled for execution by a scheduler or equivalent (e.g., in sequence or out of sequence), executed by an execution logic or equivalent, reordered, and then completed by the rename and retirement unit. API(s) (and / or compiled instructions containing API(s)) may be stored in any memory outside or inside the processor (e.g., in a cache and / or memory).Any number of CPU 1102 cores can be contained in any number of CPU clusters, which can be coupled with memory and / or cache, such as a shared L2 cache. Memory can be segregated or shared; for example, CPU clusters of CPU 1102 cores can be coupled with the Memory Subsystem 1110, which can include fabric, system-level cache, and any number of memory management units capable of reading and writing memory (e.g., DRAM). The Qualcomm Sensing Hub (not shown) includes micro-NPUs, a power rail, and conventional sensors (a gyroscope, accelerometer, even a barometer) carrying voice and data streams. The 1110 memory subsystem can include memory and cache on the 1100 processor, which 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, for example, to store 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 1110 can be shared or used individually by any or all combinations of components (e.g., GPU 1106, NPU 1104, and CPU 1102) on the Processor 1100.

[0106] The Qualcomm AI Engine 1100 can be programmed and controlled with a software stack to perform some or all of the operations described herein. This stack includes, for example, a Qualcomm® Neural Processing SDK for inference, with versions for Android, Linux, and Windows. Developer libraries and services support the latest programming languages, virtual platforms, and compilers. At a lower level of the software stack, the system software includes the 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 infrastructure such as Prometheus, Kubernetes, and Docker. OpenCL and DirectML are supported for direct cross-platform access to the GPU. For the CPU, LLVM compiler infrastructure optimizations enable accelerated and efficient AI inference.With respect to the Qualcomm AI Engine 1100 and any of its components described above or elsewhere herein, one or more of the APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetch logic or equivalent, decoded by a processor decoder or equivalent, scheduled for execution by a scheduler or equivalent (e.g., in sequence or out of sequence), executed by execution logic or equivalent, reordered, and then terminated by retirement logic or equivalent. API(s) (and / or compiled instructions containing API(s)) may be stored in any memory outside or inside the processor (e.g., in cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache or other memory.

[0107] In at least one embodiment, the Processor 1100 or the Qualcomm AI Engine 1100 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0108] Fig. Figure 12A illustrates a Processor 1200 according to at least one embodiment. The Processor 1200 may include a processor from Intel Corporation's Scalable Family of processors in Santa Clara, CA, or any other processor that uses at least some of the components described herein. The Processor 1200 may include one or more Cores 1212(1)-1212(N), where N is any integer greater than 1, capable of performing the operations described elsewhere herein. The Cores 1212(1)-1212(N) may be interconnected using ring and / or mesh interconnects. The mesh interconnect architecture allows an array of vertical and horizontal communication paths to enable traversal from one Core to another 1212(1)-1212(N) via a shortest path (jumping on the vertical path to correct the row and jumping across the horizontal path to correct the column).For mesh interconnects, a die can have cores 1212(1)-1212(N) and can contain a grid of converged mesh stops (CMS) that can be associated with cores 1212(1)-1212(N) (e.g., 1:1). Each core can be associated with a lower-level cache (LLC) slice 1214(1)-1214(N), or cores 1212(1)-1212(N) can share a cache, e.g., a lower-level cache. The LLCs 1214(1)-1214(N) can be inclusive, by incorporating blocks in the higher-level cache (e.g., L2 cache), or non-inclusive (with blocks that may not be present in the higher-level cache). Each core and each LLC slice can include a Caching and Home Agent (CHA) (not shown) that can maintain cache coherence by providing resource scalability over mesh interconnects for the Intel® Ultra Path Interconnect (Intel® UPI 1216) cache coherence functionality.The UPI 1216 can provide a coherent intermediate connection for scalable systems and can allow multiple processors to share a single shared address memory area through connections, such as, but not limited to, two or three UPI connections per processor.

[0109] The 1200 processor may also include the 1210 system agent, which can have and / or perform various functionalities, such as, but not limited to, memory management, display functions, and / or input / output (I / O) functions. For example, the 1200 processor may include one or more 1208 integrated memory controllers (IMCs). The 1208 IMC can control and manage memory, such as, but not limited to, various memory types, e.g., DDR RAM, such as DDR4, or others described elsewhere herein. The 1210 system agent may include a display controller (not shown) to support display(s). The 1210 system agent may also include PCIe 1204 (e.g., up to 20 PCIe lanes), which can be used for... B. can be connected to an external dedicated graphics port via the DMI bus (e.g., Intel's DMI 3.0 bus) 1206.The System Agent 1210 can include an image processing unit (IPU) (not shown) which incorporates an image signal processor (ISP) on the die. The Fabric 1002 can provide scalability for connecting.

[0110] Fig. Figure 12B illustrates components within the core 1212 according to at least one embodiment. The core 1212 may include a frontend 1218, a backend or execution engine 1232, and a memory subsystem 1242. The frontend 1218 may provide operations (e.g., operations described elsewhere herein) to the execution engine 1232 by decoding instructions stored in memory. For example, the frontend 1218 may include a micro-operation (µOps) cache path and / or a legacy path, along with the branch prediction unit 1220, which can determine path instructions. A legacy path for instructions may involve retrieving variable-length instructions (e.g., x86) from the L1 instruction cache, queuing the instructions into the instruction queue 1224, and decoding instructions using the decoder 1226 into µOps, which can then be provided to the assignment queue 1228.Alternatively, a µOp cache path can include a cache containing pre-decoded µOps (µOps 1230) that can be sent to the allocation queue 1228. The allocation queue 1228 can act as an interface between the frontend 1218 and the execution engine 1232 and can provide commands to the execution engine 1232. For example, one or more of the API(s) described herein can be compiled into commands that can be stored, processed, and executed by the frontend 1218 and the execution engine 1232, and stored in the storage subsystem 1242.

[0111] The execution engine 1232 can receive micro-operations into the reorder buffer 1234, which can register the assignment, renaming, and retraction of micro-operations. From the reorder buffer, micro-operations can be sent to the scheduler 1236, which can be connected to one or more different execution units 1238. The execution units 1238 can perform, for example, basic arithmetic logic unit (ALU) operations such as multiplication, division, and / or more complex operations, such as various vector operations, but are not limited to these. The scheduler 1236 can manage the queuing of micro-operations for one or more of the execution units 1238, depending on factors such as the operations that need to be performed.

[0112] The 1242 memory subsystem can handle load and store requests as well as order operations. For example, micro-operations (µOps) can relate to memory access (e.g., load and store), and these can be sent to dedicated scheduler ports capable of performing these memory operations. Store and load operations can, for example, be sent to one or more load and store buffers 1244. The 1242 memory subsystem can also include shared or separate L1 data and instruction caches 1246, as well as an L2 cache 1248 that can be used and shared by the L1 data and instruction caches 1246. As above for Fig. As described in 12A, each core 1212 can be connected to a slice of a third cache level (e.g. LLC 1214) which can be shared by all cores 1212.

[0113] In at least one embodiment, the processor 1200 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0114] Fig. Figure 13 illustrates an AI Accelerator 1300 according to at least one embodiment. The Processor 1100 may include a processor with AI Accelerator architecture from Intel Corporation in Santa Clara, CA, or any other processor that uses at least some of the components described herein. The AI ​​Accelerator 1300 may use instructions which, when executed by the AI ​​Accelerator 1300, cause the AI ​​Accelerator 1300 to perform some or all of the processes and procedures described elsewhere herein. For example, with respect to the AI ​​Accelerator 1300 and any of its components described above or elsewhere herein, one or more of the APIs described herein may be compiled into instructions that are retrieved by instruction fetch logic or equivalent, decoded by a processor decoder or equivalent, scheduled for execution by a scheduler or equivalent (e.g.,APIs (in order or out of order) can be executed by execution logic or equivalent, reordered, and then completed by retirement logic or equivalent. APIs (and / or compiled instructions containing APIs) can be stored in any memory outside or inside the processor (e.g., cache and / or memory). A result of the API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache, or other memory. The AI ​​Accelerator 1300 can include one or more compute dies, which can contain homogeneous or heterogeneous processors. Compute dies can contain one or more central processing units (CPUs), one or more graphics processing units (GPUs), or combinations of both.

[0115] In at least one embodiment, compute dies can include compute engines for performing AI computations. In at least one embodiment, compute dies of the AI ​​accelerator 1300 can be divided into any number (e.g., four) clusters, which can be designated as DCORE (Deep Learning Core) 1306 and contain any number of matrix multiplication engines (MMEs) 1308, tensor processor cores (TPCs) 1310, and L2 cache 1314 in any combination. The MME(s) 1308 can perform operations that use matrix multiplication, such as fully connected layers, convolutions, and bundled general matrix multiplications (GEMMs). The MMEs 1308 can be equipped with MACs (multiply-accumulate units) (not shown) that can perform, for example, GEMM (general matrix multiplication) operations, such as...An AxB multiplication, which involves generating a tensor C[NxM] from two input tensors A[NxK] and B[KxN], is one of the possible functions, but not limited to this. The MME(s) 1308 can be programmed with array dimensions, memory locations, data types, and various execution operands. The MME(s) 1308 can retrieve the tensors A and B from memory and pull them into their streaming buffers for matrix multiplication, which is to be performed in parallel by the MACs. The MME(s) 1308 can push the tensor C back into memory upon completion. The TPC(s) 1310 can 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., a local vector memory), and load and store components for instructions that use memory or cache (e.g., a local vector memory).The TPCs can be coupled with HBM, L3 cache, and / or L2 cache (all not shown). The TPCs can support various types of parallel processing, such as VLIW (Very Long Instruction Word), SIMD (Single-Instruction Multiple-Data), and support data types such as FP32, BF16, FP16 & FP8 (both E4M3 and E5M2), UINT32, INT32, UINT16, INT16, UINT8, and INT8, but are not limited to these. Any number of compute dies can be connected via an interposer bridge. This interposer bridge can be transparent to software.

[0116] Memory on the AI ​​Accelerator 1300 can 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. Memory and / or cache systems can be unified or separate. Compute dies of the AI ​​Accelerator 1300 can include on-die memory containing one or more levels (e.g., two levels) of cache. On-die SRAM, or other memory described elsewhere herein, can be used as a unified, last-level (L3) cache or divided into slices of L2 cache accessible to groups of MMEs 1308 and TPCs 1310. The use of on-die memory as L2 or L3 cache can be fully configurable by software, which can dynamically determine its optimal cache allocation per I / O tensor.The AI ​​Accelerator 1300 can include one or more Memory Management Units (MMUs) 1322 for managing memory, such as enabling the AI ​​Accelerator 1300's memory subsystem to operate in a virtual memory area when accessing VRAM.

[0117] The AI ​​Accelerator 1300 can include a communication port (e.g., a PCIe Gen5 x16 port) 1302 for communicating with a host and a scheduling and synchronization unit 1304. The AI ​​Accelerator 1300 can include a media unit 1316, which can contain any number or combination of media decoder engines (DECs) 1320 and rotation engines (ROTs) 1318. The AI ​​Accelerator 1300 can include a network unit 1324, which can contain any number or combination of network ports 1326 and the associated RDMA engine(s) 1328, L2 cache, and memory stacks (e.g., HBM2e or HBM3). The AI ​​Accelerator 1300 can include a programmable control path entity (not shown) to manage the parallel and efficient execution of different engines.The control path can include submission queues (SQs) that can be issued by the runtime system, completion queues (CQs) 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 "Sync Manager (SM)" that can be used for hardware synchronization, and a programmable interrupt service mechanism or "Interrupt Manager (INTR)" that can enable the forwarding of asynchronous events to drivers.

[0118] The AI ​​Accelerator 1300 can include media decoding units that support video formats such as HEVC, Progressive H.264, SVC Base Layer, MVC, VP9, ​​JPEG, and Progressive JPEG. The AI ​​Accelerator 1300 can support post-processing of decoded media streams, including image downscaling (resizing the image), vertical and horizontal scaling with varying scaling ratios, image upscaling, image cropping, bilinear scaling, and Lancos scaling. The AI ​​Accelerator 1300 can implement two post-processing channels per decoding unit: one with scalar (up and down) and one for outputting the original image only.The AI ​​Accelerator 1300 can include a hardware rotation engine that performs the following transformations of an input image: 2D rotation, 3D rotation, projection, distortion and distortion correction of images, recalculation of input data at user-defined coordinates, and rescaling.

[0119] RDMA 1328 over Converged Ethernet on the AI ​​Accelerator 1300 enables scaling from a single node (i.e., a single AI Accelerator 1300) to hundreds or thousands of nodes or AI Accelerators 1300. The NW subsystem 1324 can include an Intel® Gaudi® Communication Library (IGCL), a master conductor that orchestrates data movement, and a programmable scheduling mechanism that enables smooth engine activation while maintaining task dependencies. An accelerator network subsystem can include Gigabit Ethernet NIC ports 1326, a Layer 2 MAC (not shown), and RDMA engines 1328. The AI ​​Accelerator 1300 can include aggregation engines for performing summation activities. All engines in the processor 1300 can operate in parallel, e.g., B. the MME(s) 1308, TPC(s) 1310 and NIC(s) 1326 can all work simultaneously.There can be dependencies between operations running on different engines; for example, the output of one engine can be used as the input of another, and / or it can be scheduled for MME, TPC, and NIC to run in parallel. Once one engine has completed its execution operation, another engine can be scheduled to begin work on the next operation (as soon as its inputs are available).

[0120] The AI ​​Accelerator 1300 can be operated and controlled using software layer 1328, which can include low-level components such as a graph compiler, an automatic kernel fuser, and a library of precompiled kernels, as well as integration with AI ecosystems such as PyTorch, DeepSpeed, Hugging Face, vLLM, Ray, and others, or as described elsewhere herein in relation to software and programming platforms. Software layer 1328 can include implementations of algorithms such as Paged Attention, Flash Attention, and others. Software layer 1328 can generate optimized binary code that implements the given model topology, including, but not limited to, operator fusion, data layout management, parallelization, pipelining and memory management, and graph-level optimizations.

[0121] In at least one embodiment, the AI ​​Accelerator 1300 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0122] A neuromorphic computing system is described that employs a multi-core architecture, with each core housing the computing elements including neurons, synapses with on-chip learning capability, and local memory for storing synaptic weights and routing tables. Fig. Figure 14 is a simplified block diagram 1400 illustrating an example of at least one section of such a neuromorphic computing device 1405 according to at least one embodiment. The neuromorphic computing device 1405 may include a neuromorphic processor from Intel Corporation in Santa Clara, CA, or another processor that uses at least some of the components described herein. As shown in this example, a device 1405 may be provided with a network 1410 of multiple cores for one or more neural networks interconnected by an in-device network, such that several different connections can potentially be defined between the cores.For example, a network 1410 of spiking cores for one or more neural networks can be provided in the device 1405 and can each communicate via spike messages packaged in short packets, which are sent from core to core over the network channels. Each core (e.g., 1415) can possess processing and memory resources and logic to implement a certain number of primitive nonlinear temporal computational elements, such as, but not limited to, multiple (e.g., 1300+) unique artificial neurons (referred to herein as "neurons"). For example, each core can be capable of implementing multiple neurons simultaneously, so the collection of neuromorphic cores can implement many multiples of neurons using the device.With respect to the neuromorphic computing device 1405 and any of its components described above or elsewhere herein, one or more of the APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled for execution by a scheduler or equivalents (e.g., in sequence or out of sequence), executed by execution logic or equivalents, reordered, and then terminated by retirement logic or equivalents. API(s) (and / or compiled instructions containing API(s)) may be stored in any memory outside or inside the processor (e.g., in cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache or other memory equivalent.

[0123] Continuing with the example of Fig. A neuromorphic computing device 1405 can additionally include a processor 1420 and system memory 1425 to implement one or more components for managing and providing the functionality of the device. For example, a system manager 1430 can be provided to manage global attributes and operations of the device (e.g., attributes affecting the network of cores 1410, multiple cores in the network, connections of the device 1405 with other devices, managing access to the global system memory 1425, among other possible examples). In one example, a system manager 1430 can define and provide a specific routing table for the various routers in the network 1410, orchestrate a network definition and attributes (e.g., weights, decay rates, etc.).), which are to be applied in the network, manage core synchronization and time-division multiplexing, the routing of inputs to the appropriate cores, among other possible functions.

[0124] As another example, a neuromorphic computing device 1405 can additionally include a programming interface 1435 through which a user or system can specify a definition for a neural network to be applied (e.g., by means of a routing table and individual neuron properties) and implemented by the network 1410 of neuromorphic cores. A software-based programming tool can be provided with or separate from the neuromorphic computing device 1405, through which a user can provide a definition for a specific neural network to be implemented using the network 1410 of neuromorphic cores. The programming interface 1435 can take the programmer's input and then generate appropriate routing tables and access the local memory of individual neuromorphic cores (e.g.,1415) to fill with the specified parameters in order to implement a corresponding customer-specific network of artificial neurons, which is implemented by the neuromorphic cores.

[0125] In some cases, a neuromorphic computing device 1405 can advantageously be connected to and cooperate with other devices, including general-purpose computing devices, to realize specific applications and use cases. Accordingly, in some cases, an external interface logic 1440 can be provided to communicate with one or more other devices (e.g., via one or more defined communication protocols). An external interface 1440 can be used to accept input data from another device or an external memory controller that serves as the source of the input data. An external interface 1440 can additionally or alternatively be used to allow results or outputs of computations of a neural network implemented using the neuromorphic computing device 1405 to be transmitted to another device (e.g., a computer ...(another general-purpose processor implementing a machine learning algorithm) to realize additional applications and improvements, among other examples.

[0126] As in Fig. Figure 14 shows a network 1410 of multiple cores for one or more neural networks interconnected by an in-device network, illustrating a section of a network fabric connecting multiple neuromorphic cores (e.g., 1415a-d). For example, a number of neuromorphic cores (e.g., 1415a-d) can be provided in a network, with each core connected by a network that includes a number of routers (e.g., 1450). In one implementation, each neuromorphic core (e.g., 1415a-d) can be connected to a single router (e.g., 1450), and each router can be connected to at least one other router (as in Figure 1410). Fig. (as shown in Figure 14). As an example, in a given implementation, four neuromorphic cores (e.g., 1415a-d) may be connected to a single router (e.g., 1450), and each of the routers may be connected to two or more other routers to form a multi-core network that allows each of the neuromorphic cores to connect to every other neuromorphic core in the device. Furthermore, since each neuromorphic core may be configured to implement multiple different neurons, the router network of the device may similarly allow connections, or artificial synapses (or simply “synapses”), to be defined between any two of the potentially many (e.g., 30,000+) neurons defined using the network of neuromorphic cores provided in a neuromorphic computing device.

[0127] Fig. Figure 14 shows a block diagram illustrating internal components of an exemplary implementation of a neuromorphic kernel 1415. In one example, a single neuromorphic kernel can implement a certain number of neurons (e.g., 1024) that share the neuromorphic kernel's architectural resources in a time-multiplexed manner. In one example, each neuromorphic kernel 1415 can include a processor block 1455 capable of performing arithmetic functions and routing in connection with the realization of a digitally implemented artificial neuron, as explained herein, but is not limited to this.Each neuromorphic core 1415 can additionally provide local memory in which a routing table for a neural network can be stored and retrieved, the accumulated potential of each soma of each neuron implemented with the core can be tracked, parameters of each neuron implemented with the core can be recorded, among other data and uses. Components or architectural resources of a neuromorphic core 1415 can further include an input interface 1465 for accepting input spike messages generated by other neurons on other neuromorphic cores, and an output interface 1470 for sending spike messages to other neuromorphic cores via the mesh network or net network. In some cases, the routing logic for the neuromorphic core 1415 can be implemented, at least partially, using the output interface 1470. Furthermore, in some cases, a core (e.g.,1415) implement several neurons within an exemplary SNN, and some of these neurons may be interconnected. In such cases, spike messages sent between the neurons hosted on the particular core can bypass communication via the routing structure of the neuromorphic computing device and can instead be managed locally on the particular neuromorphic core.

[0128] Each neuromorphic core can additionally include logic to implement an artificial dendrite 1480 and an artificial soma 1485 (referred to herein simply as "dendrite" and "soma," respectively) for each neuron 1475. The dendrite 1480 can be a hardware-implemented process that receives spikes from the network. The soma 1485 can be a hardware-implemented process that receives the accumulated neurotransmitter amounts of each dendrite for the current time and further develops the potential state of each dendrite and soma to generate outgoing spike messages at appropriate times. A dendrite 1480 can be defined for each connection that receives input from another source (e.g., another neuron). In one implementation, the dendrite 1480 process can receive and handle spike messages as they arrive serially from the network in a time-division multiplexed manner.When spikes are received, the activation of the neuron (tracked using Soma 1485 and local memory 1460) can increase. If the neuron's activation exceeds a threshold set for neuron 1475, the neuron can generate a spike message, which is passed on to a fixed set of fanout neurons via output interface 1470. The network distributes the spike messages to all target neurons, and in response, these neurons can update their activations in a temporary, time-dependent manner, and so on. This may cause the activation of some of these target neurons to also exceed corresponding thresholds and trigger further spike messages, as in real biological neural networks.

[0129] As noted above, a neuromorphic computing device can reliably implement a spike-based model of neural computation. Such models can also be referred to as spiking neural networks (SNNs). In addition to neuronal and synaptic state, SNNs also incorporate the concept of time. For example, in an SNN, communication occurs via event-driven action potentials, or spikes, which convey no explicit information other than the timing of the spike and an implicit source-target neuron pair corresponding to the transmission of the spike. Computation takes place in each neuron as a result of the dynamic, nonlinear integration of the weighted spike input. In some implementations, iterations and dynamic feedback can be integrated into an SNN computational model.Furthermore, a variety of network connectivity models can be used to model different real-world networks or relationships, including fully connected (all-to-all) networks, feedforward trees, completely random projections, and small-world networks, among other examples. A homogeneous, two-dimensional network of neuromorphic kernels, such as, but not limited to, the example of [missing example]. Fig. The device shown in Figure 14 can advantageously support all these network models. Since all cores of the device can be connected, all neurons defined in the cores can therefore also be fully connected via a number of router hops. Furthermore, the device can include fully configurable routing tables to define a variety of different neural networks by allowing the neurons of each core to distribute their spikes to any number of cores in the network, thus realizing completely arbitrary connectivity graphs.

[0130] In an improved implementation of a system capable of supporting SNNs, such as, but not limited to, those in the example of Fig. In the 12 illustrated very large integration (VLSI) hardware device, high-speed and reliable circuitry can be provided to implement SNNs to model the information processing algorithms used by the brain, but in a more programmable way. For example, while a biological brain can only execute a specific set of defined behaviors conditioned by years of development, a neuromorphic processor device can provide the ability to rapidly reprogram all neural parameters. Accordingly, a single neuromorphic processor can be used to realize a wider range of behaviors than those provided by a single slice of biological brain tissue.This distinction can be achieved by using a neuromorphic processor with neuromorphic design implementations that differ significantly from those of neural circuits found in nature.

[0131] As an example, a neuromorphic processor can use time-division multiplexing in both the spike communication network and the device's neural machinery to implement SNNs. Accordingly, the same physical circuitry of the processor device can be shared by many neurons to achieve a higher neuron density. With time-division multiplexing, the network can connect N cores with a total cabling length of O(N), whereas discrete point-to-point cabling would scale to O(N²), resulting in a significant reduction in cabling resources and enabling planar and non-plastic VLSI cabling technologies, among other applications.In neuromorphic nuclei, time-division multiplexing can be implemented through dense memory allocation, for example, using static random-access memory (SRAM) with shared buses, address decoding logic, and other multiplexed logic elements. The state of each neuron can be stored in the processor's memory, with data describing each neuron's state including, among other things, the state of each neuron's collective synapses, all currents and voltages across its membrane, and other information (such as, but not limited to, configuration and other information).

[0132] A neuromorphic processor can adopt a "digital" implementation, which differs from other processors that adopt more "analog" or "isomorphic" neuromorphic approaches. For example, a digital implementation might integrate synaptic current using digital adder and multiplier circuits, in contrast to the analog isomorphic neuromorphic approaches, which accumulate charge on capacitors in an electrically analogous way to how neurons accumulate synaptic charge on their lipid membranes. The accumulated synaptic charge could, for example, be stored for each neuron in the local memory of the corresponding nucleus.Furthermore, at the architectural level of an exemplary digital neuromorphic processor, reliable and deterministic operation can be achieved by synchronizing time across the network of cores, ensuring that any two implementations of the design, given the same initial conditions and configuration, produce identical results. Asynchronicity can be preserved at the circuit level, allowing individual cores to operate as quickly and freely as possible, while maintaining determinism at the system level. Accordingly, the concept of time as a temporal variable in neural computations can be abstracted, separating it from the "wall clock time" used by the hardware to perform the computation. Consequently, some implementations can provide a time synchronization mechanism that globally synchronizes the neuromorphic cores at discrete time intervals.The synchronization mechanism allows the system to perform a neural computation as quickly as the circuitry allows, with a divergence between the runtime and the biological time that models the neuromorphic system.

[0133] In operation, the neuromorphic network device can begin in an idle state, with all neuromorphic cores inactive. Since each core asynchronously traverses its neurons, it generates spike messages that route the network intermediary to the appropriate target cores containing all the target neurons. Because the implementation of multiple neurons on a single neuromorphic core can be time-division multiplexed, a time step can be defined in which all spikes involving the multiple neurons can be processed and accounted for using the shared resources of a corresponding core. Once each core has finished servicing its neurons for a given time step, in some implementations, the cores can communicate with neighboring cores using synchronization messages (e.g.,(using a handshake) to clear the network of all spike messages with the transmission, allowing the cores to reliably determine that all spikes have been served for the time step. At this point, all cores can be considered synchronized, allowing them to advance their time step, return to the initial state, and begin the next time step.

[0134] Given this context and as introduced above, a device (e.g., 1405) can be provided that implements a network 1410 of interconnected neuromorphic nuclei, with the nucleus potentially implementing multiple artificial neurons capable of being interconnected to implement a SNN. Each neuromorphic nucleus (e.g., 1415) can provide two loosely coupled asynchronous processes: an input dendrite process (e.g., 1480) that receives spikes from the network and applies them to the appropriate target dendrite compartments at suitable future time points, and an output soma process (e.g., 1485) that receives the accumulated neurotransmitter amounts of each dendrite compartment for the current time point and further develops the membrane potential state of each dendrite and soma, generating outgoing spike messages at the appropriate time points (e.g., when a threshold potential of the soma has been reached).It should be noted that from a biological point of view, the names dendrite and soma used here only approximate the role of these functions and should not be interpreted too literally.

[0135] In at least one embodiment, the neuromorphic computing device 1405 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0136] Fig. Figure 15 is a block diagram of an embodiment of a multi-node network in which remote storage computing can be implemented, according to any embodiment. System 1500 can represent a network of nodes described herein, which can be used, for example, to perform some or all of the operations described herein. System 1500 can represent a data center. System 1500 can represent a server farm. System 1500 can represent a data cloud or a processing cloud. System 1500 can represent a supercomputer. System 1500 can contain ten, hundreds, or thousands of nodes. The nodes of System 1500 can contain processors, such as, but not limited to, central processing units (CPUs), graphics processing units (GPUs), or any combination of processors described herein, such as, but not limited to, other processors in the Fig. 9-21. With respect to any of the processors in the System 1500 and any of its components described above or elsewhere herein, one or more of the APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetch logic or equivalent, decoded by a processor decoder or equivalent, scheduled for execution by a scheduler or equivalent (e.g., in sequence or out of sequence), executed by execution logic or equivalent, reordered, and then terminated by retirement logic or equivalent. API(s) (and / or compiled instructions incorporating API(s)) may be stored in any memory outside or inside the processor (e.g., in cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache, or other memory equivalent. The System 1500 can include over nine thousand nodes, with each node containing two Intel Xeon Max processors, six Intel Max series GPUs, and a unified memory architecture, such as, but not limited to, that used in Intel Corporation's Intel Aurora supercomputer in Santa Clara, CA, or any other supercomputer using at least some of the components described herein.

[0137] One or more clients 1502 send requests to system 1500 over network 1504. Network 1504 represents one or more local area networks, wide area networks, or a combination thereof. The clients 1502 can be human or machine clients that generate requests for system 1500 to perform operations. System 1500 executes applications or data processing tasks requested by the clients 1502.

[0138] The System 1500 can include one or more racks, which represent structural and interconnect resources to house and connect multiple compute nodes. The Rack 1510 can contain multiple nodes 1530. The Rack 1510 can host multiple blade components 1520. Hosting can refer to providing power, structural or mechanical support, and interconnection. Blades 1520 can refer to compute resources on printed circuit boards (PCBs), where a PCB contains the hardware components for one or more nodes 1530. Blades 1520 may or may not include a chassis or enclosure, or any other "box" besides that provided by the Rack 1510. Blades 1520 may include an enclosure with an exposed connector for connecting to the Rack 1510.The System 1500 may or may not include the Rack 1510, and each Blade 1520 may include a chassis or enclosure that can be stacked or otherwise located in close proximity to other Blades, enabling the connection of Nodes 1530. The System 1500 can contain 10,924 compute Blades, which include 63,104 Intel Max Series GPUs and 21,248 Intel Xeon Max CPUs across 166 Racks.

[0139] System 1500 can include Fabric 1570, which represents one or more connectors for nodes 1530. Fabric 1570 can include multiple switches 1572, routers, or other hardware to route signals between nodes 1530. Additionally, Fabric 1570 can connect System 1500 to network 1504 for access by clients 1502. Beyond routing equipment, Fabric 1570 can also include the cables, connectors, or other hardware to connect nodes 1530. Fabric 1570 can have one or more associated protocols to manage the routing of signals through System 1500. The protocol(s) depend, at least in part, on the hardware used in System 1500.

[0140] As illustrated, Rack 1510 can contain N Blades 1520. In addition to Rack 1510, System 1500 can contain Rack 1550. As illustrated, Rack 1550 can contain M Blades 1560. M is not necessarily equal to N; thus, it is understood that various different hardware equipment components can be used and coupled to each other in System 1500 via Fabric 1570. The Blades 1560 can be the same as or similar to the Blades 1520. The Nodes 1530 can be any type of node, as described herein, and need not all be the same type of node. System 1500 is neither restricted to being homogeneous nor restricted to being heterogeneous.

[0141] A node in Blade 1520(0) is illustrated in detail. However, other nodes in System 1500 may be the same or similar. At least some nodes 1530 may be compute nodes with a processor 1532 and memory 1540. A compute node refers to a node with processing resources (e.g., one or more processors) that runs an operating system and can receive and process one or more tasks. At least some nodes 1530 may include storage server nodes with a server as processing resources 1532 and memory 1540. A storage server refers to a node with more storage resources than a compute node, and instead of having processors to perform tasks, a storage server includes processing resources to manage access to the storage nodes within the storage server.

[0142] The node 1530 can include the interface controller 1534, which can represent logic to control access by the node 1530 to the fabric 1570. This logic can include hardware resources to connect to the physical linking hardware. It can also include software or firmware logic to manage the connection. The interface controller 1534 can be a host fabric interface or can include a fabric interface according to any embodiment described herein.

[0143] Node 1530 can contain the storage subsystem 1540. Storage 1540 can contain the memory computation resources (comp) 1542, which represent one or more capabilities through storage 1540 to perform memory computations. System 1500 enables remote memory operations, such as those described elsewhere herein. Thus, nodes 1530 can request memory computations from remote nodes, with data for the computation remaining local to the executing node instead of being sent across Fabric 1570 or from storage to the Fabric interface. In response to the execution of the memory computation, the executing node can provide a result to the requesting node.

[0144] The 1532 processor can include one or more separate processors. Each separate processor can include a single processing unit, a multi-core processing unit, or a combination thereof. The processing unit can be a primary processor, such as a CPU (central processing unit), a peripheral processor, such as a GPU (graphics processing unit), or a combination of both. The 1540 memory can be or include memory devices and a memory controller.

[0145] A reference to storage devices can apply to different types of memory. Storage devices generally refer to volatile memory technologies. Volatile memory is memory whose state (and therefore the data stored on it) is indeterminate when the power supply to the device is interrupted. Non-volatile memory refers to memory whose state is definite, even when the power supply to the device is interrupted. Dynamic volatile memory requires the data stored in the device to be updated in order to maintain its state. An example of dynamic volatile memory includes DRAM (Dynamic Random Access Memory) or a variant such as synchronous DRAM (SDRAM).A memory subsystem, as described herein, can be compatible with a number of memory technologies, including but not limited to DDR3 (version 3 with double the data rate, originally published by JEDEC (Joint Electronic Device Engineering Council) on 27.June 2007, currently on Publication 21), DDR4 (Version 4 with DDR, original specification, published in September 2012 by JEDEC), DDR4E (Version 4 with DDR, extended, currently under discussion by JEDEC), LPDDR3 (Version 3 with low-power DDR, JESD209-3B, Aug 2013 by JEDEC), LPDDR4 (Version 4 with double data rate low-power (LPDDR), JESD209-4, originally published by JEDEC in August 2014), WIO2 (Wide I / O 2 (WideI02), JESD229-2, originally published by JEDEC in August 2014), HBM (High Bandwidth Memory DRAM, JESD235, originally published by JEDEC in October 2013), DDR5 (Version 5 with DDR, currently under discussion by JEDEC), LPDDR5 (currently under discussion by JEDEC), HBM2 (HBM Version 2), currently under discussion by JEDEC), or other or combinations of memory technologies and technologies based on derivatives or extensions of such specifications.

[0146] In addition to or as an alternative to volatile memory, in one embodiment a reference to memory devices may refer to a non-volatile memory device whose state is determined even when the power supply to the device is interrupted. In one embodiment, the non-volatile memory device is a block-addressable memory device, such as, among others, NAND or NOR technologies. Thus, a memory device may also include future-generation non-volatile devices, such as, among others, a three-dimensional crosspoint memory (3DXP) device, other byte-addressable non-volatile memory devices, or memory devices that use chalcogenide phase-change material (e.g., chalcogenide glass).In one embodiment, the storage device may be or include a multi-threshold NAND flash memory, a NOR flash memory, a single- or multi-level phase-change memory (PCM) or a single-switch phase-change memory (PCMS), a resistive memory, a nanowire memory, a ferroelectric transistor random-access memory (FeTRAM), a magnetoresistive random-access memory (MRAM) incorporating memristor technology, or a spin-transfer-torque MRAM (STT-MRAM), or a combination of any of the above, or any other memory.

[0147] In at least one embodiment, the system 1500 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0148] Fig. Figure 16 illustrates an accelerated processing unit 1600 according to at least one embodiment. The accelerated processing unit 1600 may include a processor based on a CDNA architecture from AMD Corporation in Santa Clara, CA, or another processor that uses at least some of the components described herein. The accelerated processing unit 1600 may include one or more accelerator complex dies (XCDs) 1604 for performing operations described elsewhere herein, such as, but not limited to, graphics processing and / or parallel processing, as well as instruction-level parallel computations, including support for a wide range of precisions (INT8, FP8, BF16, FP16, TF32, FP32, and FP64) and sparsity matrix data. XCDs may, in some cases, be referred to as graphics compute dies (GCDs).The Accelerated Processing Unit 1600 can include one or more Complex Computing Dies (CCDs) 1606 for performing operations described elsewhere herein, such as, but not limited to, those operations performed by host processors. CCDs may in some cases be referred to as core complexes or CCXs, such as, but not limited to, CCXs used in AMD Ryzen processors. XCDs and CCDs can use any type of cache or memory (e.g., one or more Memory Units 1602) or have a cache or memory allocated to each XCD or CCD or group of XCDs or CCDs. For example, the on-package AMD Infinity Fabric connects XCDs and CCDs to the shared AMD Infinity Cache 1608 and, in some embodiments, high-bandwidth memory (e.g., HMB3).The Accelerated Processing Unit 1600 can be an AMD MI300a processor, which includes three CPU chiplets (or CCDs) and six accelerator chiplets (XCDs) on four input / output dies (IODs). These can be stacked on a single piece of silicon, which connects them (e.g., via an AMD Infinity Fabric) to eight stacks of high-bandwidth DRAM surrounding the superchip. An AMD MI300x processor replaces the CCDs with two additional XCDs for an accelerator-only system.

[0149] The Accelerated Processing Unit 1600 can include one or more input / output (I / O) interfaces. For example, the XCDs 1604 and CCDs 1606 can be located together on one or more input / output dies (IODs) 1610, which can include one or more I / O interfaces. The IODs 1610 can include any number and type of I / O interfaces (e.g., PCI, PCI-Extended ("PCI-X"), PCIe, Gigabit Ethernet ("GBE"), USB, etc.). Various types of peripheral devices can be connected to the I / O interfaces 970. I / O interfaces of IODs 1610 can also be used to connect one or more Accelerated Processing Units 1600, for example, in a server architecture.

[0150] The Accelerated Processing Unit 1600 can include one or more Memory Units 1602 for storing instructions and other information used to perform operations described elsewhere herein. The Memory Units 1602 can include any type of volatile memory, such as, but not limited to, memory types described elsewhere herein, and can include, for example, high-bandwidth memory (e.g., HMB3) or high-bandwidth DRAM. Memory allocated to the Accelerated Processing Unit 1600 (e.g., Memory Units 1602) can include system memory, which can be used, for example, for instructions, statements, constants, and inputs and outputs. The Memory Units 1602 can also include device memory, which can be used, for example, for instructions, statements, constants, inputs and outputs, as a return buffer, and for private data.The memory units 1602 can be connected to one or more IODs 1610. In at least one embodiment, a memory hierarchy begins with the L1 cache 1620 and includes the shared L2 cache 1628, for example, within the XCDs. AMD Infinity Cache™ is a last-level cache (LLC) located on an active I / O die (IOD). The CCDs 1606 and XCDs 1604 can have separate or shared memory. AMD Infinity Architecture and AMD Infinity Fabric™ technology enable coherent unification of GPU and CPU chiplet technologies (e.g., XCDs, CCDs, and / or CCXs) with memory (e.g., stacked HBM3 memory) in single devices and across multiple device platforms with high throughput.

[0151] As in Fig. As shown in Figure 16, an XCD 1604 can include a shared set of global resources 1630, which may include a hardware scheduler 1612 and asynchronous compute engines (ACEs) 1624 that send tasks (e.g., compute shader workgroups) to compute units (CUs or cores) 1630. ACEs 1624 (e.g., four) can each be associated with CUs 1630 (e.g., 40 CUs), and some of the CUs may be disabled for yield management. The CUs 1630 may have a dedicated cache or a shared cache (e.g., L2 cache) 1628 that can be used to aggregate all memory traffic for the die. The CUs 1630 can include thread and parallel processor cores, including instruction retrieval and scheduling with the scheduler (S) 1612, the matrix core unit (MCU) 1616 and the shader core (SC) 1618 (e.g.Execution units for scalar, vector, and matrix data types) as well as load / store pipelines with an L1 cache 1620 and local data sharing (LDS) 1614. Local data sharing can include, for example, a scratch RAM with integrated arithmetic capabilities that allow data to be shared between threads in a workgroup. An instruction cache 1640 (for example, for storing and providing the instructions for performing operations described elsewhere herein) can be associated with one or more CUs and can be shared by two CUs. The matrix cores 1616 can process a variety of data types, such as, but not limited to, INT8, FP8, FP16, BF16, and TF32 data types. The accelerated processing unit 1600 can include compute units 1630, which can be arranged in an array format, e.g., B. as a data parallel processor array (DPP array).The Ultra-Threaded Dispatch Processor 1642 can communicate with the Arithmetic Units 1630, and the Instruction Processor 1644 can read instructions written by the host to memory-mapped registers in a system memory address space (not shown). The Instruction Processor 1644 can send hardware-generated interrupts to a host processor (such as a CCD) when the instruction completes. The Memory Controller 1636 can also have direct access to all device memory and the host-specified areas of system memory. To satisfy read and write requests, the Memory Controller 1636 can perform Direct Memory Access (DMA) controller functions, including calculating memory address offsets based on the format of the requested data in memory.For example, one or more of the APIs described herein can be compiled into instructions that are stored in the 1640 instruction cache and then retrieved by instruction fetch logic in the 1640 processor, decoded by a processor decoder or equivalent, scheduled for execution by a scheduler or equivalent (e.g., in sequence or out of sequence), executed by execution logic or equivalent, reordered, and then terminated by retirement logic or equivalent. API(s) (and / or compiled instructions containing API(s)) can be stored in any memory outside or inside the 1640 processor (e.g., in the cache and / or memory). A result of API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache, or other memory equivalent.

[0152] An application can consist of a program running on a host processor (e.g., a CCD) and programs, called kernels, running on one or more XCDs. Programs can be controlled by host commands that set an internal base address and other configuration registers, specify a data domain on which the Accelerated Processing Unit 1600 can operate, invalidate and flush caches on the Accelerated Processing Unit 1600, and cause the Accelerated Processing Unit 1600 to begin executing a program. Kernels can be referred to as programs executed by the Accelerated Processing Unit 1600. A kernel can be executed independently on each work element or as groups of work elements, which can be called a wavefront, that can execute the kernel on all work elements in the group (e.g., 64) in a single pass.The 1630 compute units can include a scalar arithmetic logic unit (ALU) that can operate on one value per wavefront (common to all worker elements), a vector ALU that can operate on unique values ​​per worker element, a local data share 1614 that allows worker elements within a workgroup to communicate and share data, a scalar memory (not shown) that can transfer data between scalar general-purpose registers (SGPRs) and memory through a cache, and a vector memory that can transfer data between vector general-purpose registers (VGPRs) and memory, including sampling texture maps. Kernel control flow can be handled using instructions for a scalar ALU, which can include if-else branches and loops. Instructions for a scalar ALU (SALU) and memory can operate on an entire wavefront and on one or more SGPRs.Vector memory and ALU commands can operate on all work elements in the wavefront simultaneously.

[0153] In at least one embodiment, the accelerated processing unit 1600 can include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits can be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0154] Fig. Figure 17 illustrates a Processor 1700, such as, but not limited to, a processor based on a Zen architecture (such as Zen 1, 2, 3, 4, 5, or others) from AMD Corporation in Santa Clara, CA, or any other processor that uses at least some of the components described herein. The Processor 1700 includes one or more CPU dies 1702(1)-1702(N), where N is any integer greater than 1. The CPU die 1702 can include any number of processor cores 1716 (for example, to perform any of the operations described elsewhere herein) and any number of cache memories (for example, to store instructions and other information to perform any of the operations described elsewhere herein) in any combination.For example, L2 cache units 1718 can be coupled with one or more processor cores 1716, which can share the L2 cache units 1718 and / or be individually coupled to them. The processor cores 1716 can be individually coupled to the L3 cache 1722 and / or share the L3 cache, which can be a lowest-level cache (LLC) 1722 for accessing data and other information used by the processor cores 1716. One or more processor cores 1716 and one or more L2 cache units 1718 can be contained in a core complex (CCX) 1720, which can contain a shared cache (e.g., the L3 cache 1722) of, for example, 32 MB. The 1720 core complex can be manufactured on a 1702 die (CCD or CPU die). For example, up to 12 1720 core complexes can be combined with 8 1702 CPU dies to form a processor, providing up to 96 1716 processor cores.A "Zen 4c" core complex 1720, for example, can contain up to eight cores 1716 and a shared 16 MB L3 cache 1722. Two of these core complexes 1720 can be combined on a single CPU die 1702 for 16 cores per die and a total of 32 MB L3 cache 1722 per die. Up to eight of the CPU dies 1702 can be combined with an I / O unit 1704 to provide CPUs with up to 128 processor cores 1716. Up to four "Zen 4c" dies as described above can be combined to provide CPUs with up to 64 processor cores 1716.

[0155] The 1700 processor can incorporate a variety of input / output configurations, which are described in more detail herein. The 1704 I / O unit can include one or more 1706 memory controllers, which manage memory usage (e.g., DDR5 memory) for the 1700 processor. The 1704 I / O unit can include one or more 1712 SATA hard disk controllers for managing memory and one or more 1714 Compute Express Link (CXL™) 1.1+ memory controllers, which provide CPU-to-device and CPU-to-memory connections and can be flexibly assigned specific functions at server design time. The 1704 I / O unit can include the 1708 PCIe controller for connecting peripherals and other components connected to the 1700 processor. The I / O unit 1704 can include USB ports 1710 for connecting to other components that are separate from the processor 1700.The CPU dies 1702 can support any number of connections, such as one or two, to the I / O unit 1704. As shown, the I / O unit 1704 includes the components described in more detail herein, and the I / O unit 1704 can be an I / O die containing several different components. The memory controller 1706, the PCIe controller 1708, the USB ports 1710, the SATA controller 1712, and / or the CXL controller 1714 can be integrated at any location within the processor 1700, either separately or in any groups or combinations thereof.

[0156] The 1700 processor can include Infinity Fabric 1724 intermediate links (which may be similar to or based on PCIe architectures) that can provide connections between CPUs (e.g., 1702(1)-1702(N) CPU dies), 1726 graphics processor(s), 1732 inference engine(s), and other components in the multi-chip architecture, such as a 1728 security processor(s) and a 1704 I / O unit. One or more AMD Infinity Fabric™ 1710 intermediate links can be connected to the 1702(1)-1702(N) CPU dies and serve as a connection used between CPUs. One or more Infinity Fabric 1710 links can connect each 1702 CPU die to the 1710 I / O unit.

[0157] In at least one embodiment, the Processor 1700 may include central processing units (CPUs) and other associated hardware and software, as described above and further herein. The Processor 1700 may also include graphics processor(s) 1726. The graphics processor 1726 may be used for image generation and processing, as well as other computations and operations further described herein. The graphics processor 1726 may be based on an AMD RDNA 3 or 3.5 architecture in Santa Clara, CA. The graphics processor 1726 may include graphics compute dies (GCDs) and memory cache dies (MCDs). GCDs may include any number of compute units (CUs) for graphics or other processing, such as operations performed by arithmetic logic units (ALUs), as further described herein. The graphics processor 1726 may include an L2 cache that can be used by compute units.MCDs (not shown) can contain any number of memory units and can include cache, such as L3 cache, as well as memory interfaces for coupling to memory, such as Memory 1742(1)-(N), where N is an integer. Components within the GPU 1726 can be interconnected using various approaches, such as Infinity Fabric 1724 interleaves outside or inside the GPU 1726.

[0158] The inference engine 1732 can provide neural processing capabilities for the processor 1700 for computational procedures used for neural networks, deep learning, and other operations related to artificial intelligence, which are further described herein. The processor 1700 can include a security processor 1728 for managing the security of the processor, a display controller 1730 for controlling displays, a system management unit 1734 for managing and operating some or all of the components on the processor 1700, multimedia engines 1736 for audio and video operations, a fusion control hub 1738 for managing USB, SATA, and PCIe connections to the processor, and a sensor fusion hub 1740 for managing sensors, such as accelerometers. The processor 1700 can also include memory 1742(1)-(N), where N is any integer. The storage can include different types of storage, such as...LPDDR5 and / or DDR5 or others described elsewhere herein.

[0159] To perform the operations described herein, the 1700 processor can include an execution pipeline containing a front end, which may include a cache (e.g., an L1 cache) that stores instructions (not shown). The flow of instructions can be modified by a branch predictor. Instructions can be decoded by a decoder, output to a back end for execution, and renamed. Instruction fetch and decode pipes can, for example, output to integer or floating-point execution operations, which can be scheduled by a scheduler and transferred to vector and / or general-purpose registers. Floating-point multiplier and / or addition operations can be processed, and arithmetic logic units (ALUs) can also be used to perform calculations, such as arithmetic and logical operations.Outputs from the compute units can be coupled with a load / store queue, which may be connected to a cache, such as an L1 cache and / or L2 cache.

[0160] With respect to the 1700 processor and any of its components described above or elsewhere herein, one or more of the APIs or equivalents described herein may, for example, be compiled into instructions or equivalents (e.g., AVX-812 instructions based on a SIMD model) that may be fetched by instruction fetch logic or equivalent, decoded by a processor decoder or equivalent, scheduled for execution by a scheduler or equivalent (e.g., in sequence or out of sequence), executed by execution logic or equivalent, reordered, and then terminated by retirement logic or equivalent. API(s) (and / or compiled instructions containing API(s)) may be stored in any memory outside or inside the processor (e.g., in the cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache or other memory equivalent.

[0161] In at least one embodiment, the 1700 processor may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0162] Fig. Figure 18 illustrates an example of a 1800 processing core that can implement an Arm architecture (e.g., v9.0-A) or another processor that uses at least some of the components described herein. The Neoverse™ V2 1800 core can be implemented within a DynamIQ Shared Unit (DSU) cluster via the DSU-110 Intermediate 1854 for one or more interconnected cores, for example, for parallel processing. The Neoverse™ V2 core can be implemented as a single core in a DSU cluster configured for direct interconnection with or without L3 cache, snoop filtering, or snoop control unit (SCU) logic (not shown). The Neoverse™ V2 core can include a CPU bridge 1852, which connects the core 1800 to the DSU-110 intermediate link, which can also connect the core 1800 to an external storage system and the rest of a system-on-a-chip. The L1 instruction memory system 1802 can retrieve instructions from an instruction cache 1804 and execute the instructions (e.g., commands, commands, etc.).B. one or more APIs described herein, which can be compiled to instructions) to an instruction decoding unit 1810, e.g., to perform some or all of the operations described above or elsewhere herein. The L1 instruction storage system 1802 may include the L1 instruction cache 1804, e.g., with 64-byte cache lines; the L1 instruction translation lookaside buffer (TLB) 1806, e.g., with native support for 4KB, 16KB, 64KB, and 2MB page sizes; the macro operation cache (MOP) 1808 (e.g., 1836 entry, 4-way skewed associative L0 MOP cache), which may contain decoded and optimized instructions for higher performance. The 1810 instruction decoding unit can decode AArch64 instructions into an internal format. An 1812 register renaming unit can perform register renaming to facilitate out-of-order execution and sends decoded instructions to various output queues.The instruction output unit 1814 can control when decoded instructions can be issued to the execution pipelines, and it can include output queues to store instructions until they are issued to execution pipelines. The integer execution pipeline 1816 can be included in an execution pipeline and can include an integer execution unit 1818, which can perform arithmetic and logical data processing operations. The vector execution unit 1820 can be included in an execution pipeline and can perform Advanced SIMD and Floating Point Operations (FPU) 1822, execute Scalable Vector Extension (SVE) and Scalable Vector Extension 2 (SVE2) instructions 1824, and can optionally execute the cryptographic (Crypto) instructions 1826. Advanced-SIMD can include a media and signal processing architecture that adds commands primarily for audio, video, 3D graphics, image and speech processing.A floating-point architecture provides support for single- and double-precision floating-point operations. The L1 Data Storage System 1830 can execute load and store instructions as well as service memory coherence requests. The L1 Data Storage System 1830 can include an L1 Data Cache 1832 and a fully associative L1 Data TLB 1834 with native support for 4KB, 16KB, and 64KB page sizes and 2MB and 812MB block sizes. The Memory Management Unit (MMU) 1828 can provide fine-grained memory system control through a set of mappings from virtual to physical addresses and memory attributes, which can be maintained in translation tables that can be stored in the TLB 1834 when an address is translated. The L2 memory system 1836 can include the L2 cache 1838 and can be connected to the DSU-110 1854 by an asynchronous CPU bridge 1852.The Neoverse™ V2 Core 1800 can support a range of debugging, testing, and tracing options, including a Trace Unit 1842, a Trace Buffer 1840, and an Embedded Logic Analyzer (ELA) 1848. The Neoverse™ V2 Core 1800 can implement the Statistical Profiling Extension (SPE) 1844 to provide a statistical view of the performance characteristics of executed instructions, which software writers can use to optimize their code for better performance. The Performance Monitoring Unit (PMU) 1846 can provide performance monitoring, which can be configured to collect statistics on the operation of each core and the memory system. This information can be used for debugging and code profiling.The Generic Interrupt Controller (GIC) CPU interface 1850, when integrated with an external distributor component, can be a resource for supporting and managing interrupts in a cluster system. In a cluster, there can be a CPU bridge 1852 between each Neoverse™ V2 Core 1800 and the DSU-110 1854. The CPU bridge 1852 can control buffering and synchronization between the Core 1800 and the DSU-110 1854. The CPU bridge 1852 can be asynchronous to allow different frequency, power, and range implementation points for each Core 1800. The CPU bridge 1852 can run synchronously without affecting other interfaces, such as, but not limited to, debug and trace, which can be asynchronous.

[0163] In at least one embodiment, the Core 1800 can include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits can be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0164] Fig. Figure 19 illustrates one or more chips incorporating one or more tensor processing units (TPUs) 1900 according to at least one embodiment. The TPUs 1900 in Fig. 19 TPUs may include application-specific integrated circuits (ASICs) to perform, for example, some or all of the operations described above or elsewhere herein, such as, but not limited to, accelerating machine learning workloads that perform matrix operations. The TPUs 1900 may be ASICs from Alphabet Corporation in Mountain View, CA. The Cloud TPU includes a cloud service that makes TPUs available 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.

[0165] The Chip 1900 can contain any number of TPUs, which can contain Tensor Cores 1906. The Tensor Core 1906 can contain one or more Core Sequencers 1908, a Vector Processing Unit (VPU) 1910, a Matrix Multiplication Unit (MXU) 1912(A)-1914(N), where N is any integer greater than 1, and a Transpose Permutation Unit 1916. The Core Sequencer 1908 can retrieve instructions from the core's Instruction Memory (Imem) 1906 (e.g., VLIW (Very Long Instruction Word)), perform scalar operations using a scalar data memory (Smem) and scalar registers (Sregs) (not shown), and forward vector instructions to the Vector Processing Unit (VPU) 1910. The commands can, for example, initiate eight operations: two scalar ALUs, two vector ALUs, vector loading and storage, and a pair of slots that queue data to and from the matrix multiplication and transposition units.The VPU 1910 can perform vector operations using a large on-chip vector memory (Vmem) and vector registers (Vregs). The VPU 1910 can stream data to and from the MXU by decoupling FIFOs. The VPU 1910 can collect and distribute data to the Vmem via data-level parallelism (2D matrix and vector functional units) and instruction-level parallelism (8 operations per instruction). A large two-dimensional matrix multiplier unit (MXU) 1912(A)-1912(N), for example, can use a systolic array to reduce area and power consumption, plus large, software-controlled on-chip memories instead of caches. The transpose-reduction-permutation unit 1916 can perform matrix transpositions, reductions, and permutations of the VPU 1910 orbits (e.g., 128 x 128). The high-bandwidth 1904 memory can be used for on-chip applications. One or more 1900 chips can be interconnected for computation.For example, one or more Chip 1900s can be connected as a torus, e.g., a 2D torus. The Chip 1900 can also include any number (e.g., four) of Inter-Core Interconnect (ICI) links 1918, which can provide direct connections between chips to form a supercomputer.

[0166] With respect to any of the processors in the Chip 1900 and any of its components described above or elsewhere herein, one or more of the APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled for execution by a scheduler or equivalents (e.g., in sequence or out of sequence), executed by execution logic or equivalents, reordered, and then terminated by retirement logic or equivalents. API(s) (and / or compiled instructions containing API(s)) may be stored in any memory outside or inside the processor (e.g., in cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache or other memory equivalent.

[0167] In at least one embodiment, the Chip 1900 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0168] Fig. Figure 20 illustrates a vector processor according to at least one embodiment. The Vector Processor 2000 can support a RISC-V standard. The Vector Processor 2000 can include one or more Cores 2010 (e.g., scalar units) with one or more Vector Processing Units (VPUs) 2042 (e.g., vector units) that can perform, for example, some or all of the operations described above or elsewhere herein. The Core 2010 can include a Custom Extension (ACE) 2016 that can be used to communicate custom instructions to the Processor 2000. The Core 2010 can include a 1-cycle multiplier and a 1-cycle instruction / data local memory (ILM / DLM) for increased parallelism by enabling simultaneous instruction fetching and data access.The Memory Management Unit (MMU) 2024 can manage system memory and cache and provide branch execution, instruction pair output, L1 instruction / data caches, and local memory storage. The Core 2010 can include a Physical Memory Protection Unit and a Programmable Physical Memory Attribute Unit (PMP / PPMA) 2022. The Core 2010 can include a Digital Signal Processor (DSP) 2028 and a Floating Point Unit (FPU) 2026, as well as a Load-Storage Unit (LSU) 2032 to interface with the memory hierarchy (D$ 2034 and I$ 2030). The Core 2010 can include a Branch Prediction Unit 2018 and a Multiplier Unit 2020.

[0169] The Vector Processing Unit (VPU) 2042 can include one or more Vector Functional Units (FUs) 2046(A)-2046(N) which can be chained together for parallel processing, independent memory paths for loading / saving RISC-V vector (RVV) via ACE-RVV 2048 and Andes Streaming Port (ASP) 2044 and a Vector Load / Storage Unit (VLSU) 2050.

[0170] The Vector Processor 2000 can include bus interfaces such as the L2 cache memory port 2056 for cacheable access, an MMIO port 2054 for non-cacheable access, an input / output coherence port (IOCP) 2058 for a cacheless bus master, local memory access ports for ILM / DLM 2012 and access to high-bandwidth vector memory (HVM) 2036, and a shared peripheral port (SPP) 2052 for external peripherals. Other memory ports include the LM slave port AXI 2002 and the HVM sub-port AXI 2004.

[0171] With respect to any of the processors in the Processor 2000 and any of its components described above or elsewhere herein, one or more of the APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled for execution by a scheduler or equivalents (e.g., in sequence or out of sequence), executed by execution logic or equivalents, reordered, and then terminated by retirement logic or equivalents. API(s) (and / or compiled instructions containing API(s)) may be stored in any memory outside or inside the processor (e.g., in cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache or other memory equivalent.

[0172] In at least one embodiment, the Vector Processor 2000 may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0173] Fig. Figure 21A illustrates a diagram of an exemplary multi-core tile processor microarchitecture. A multi-core tile processor in Fig. 21A may include a speech processing processor. As in Fig. As illustrated in Figure 21A, each "tile" of the processor architecture is a processing element connected using a network-on-chip (NoC) that can be used, for example, to perform some or all of the operations described above or elsewhere herein. For example, each tile may include an instruction output 2104 and an integer (INT) 2106 and floating-point (FP) unit 2108, as well as a load-memory unit (LSU) 2112 to interface with the memory hierarchy (data cache (D$) 2110 and instruction cache (I$) 2114) and a network (NET) interface 2116 for communicating with other tiles of the architecture. Some tiles in the 2100 processor may include a memory controller 2102 for managing and controlling memory, as further described herein. The 2100 processor may have a functional slice architecture.The 2100 processor can be located on an application-specific integrated circuit (ASIC), and . Fig. 21A can represent the layout of the ASIC. The 2100 processor can include a coprocessor designed to execute instructions for a predictive model. The predictive model is any model designed to make a prediction based on input data. The predictive model can use a classifier to make a classification prediction. The predictive model can be a machine learning model, such as, but not limited to, a tensor flow model, and the 2100 processor is a tensor streaming processor.

[0174] The 2100 processor can use different microarchitectures, which are in Fig. 21B divide the functional units shown in each tile. Instead, the function tiles of the 2100 processor can be divided into a variety of functional process units (hereafter referred to as "slices") 2104, each corresponding to a specific function type (e.g., FP / INT, NET, MEM). For example, as shown in Fig. Figure 21B illustrates that each slice corresponds to a column of function tiles extending in a north-south direction. Additionally, the processor includes communication lanes to transport data between tiles of different slices, each running horizontally in an east-west direction. Each communication lane can be connected to any of the slices 2104 of the processor 2100.

[0175] The 2104 slices of the 2100 processor can each correspond to a different function and can include slices for arithmetic logic (e.g., FP / INT), path-switching slices (e.g., NET), and memory slices (e.g., MEM). The arithmetic logic units perform one or more arithmetic and / or logical operations on the data received via the communication paths to produce output data. Examples of arithmetic logic units include matrix multiplication units and vector multiplication units. The memory slices contain memory cells that store data. The memory slices can provide data to other slices via the communication paths. The memory slices can also receive data from other slices via the communication paths. The path-switching slices can configurably forward data from one communication path to any other.For example, data from a first lane can be provided to a second lane via a lane-switching slice. In some embodiments, the lane-switching slice can be implemented as a crossbar switch. Each 2104 slice also includes its own instruction queue (not shown), which stores instructions, and an instruction control unit (ICU) to control the execution of the instructions. The instructions in a given instruction queue can only be executed by tiles in their associated functional slice and cannot be executed by the other slice of the processor.

[0176] By arranging the tiles of the 2100 processor into different functional slices 2104, the on-chip instruction and control flow of the 2100 processor can be decoupled from the data flow. For example, an arrow in Fig. 21B the flow of instructions within the processor architecture according to some embodiments. Another arrow in Fig. Figure 21B illustrates a data flow within the processor architecture according to at least one embodiment. As illustrated, the instructions and control flow in a first direction across the tiles of the 2100 processor (e.g., north-south along the length of the functional slices, as shown by the first arrow), while the data flows in a second direction across the tiles of the 2100 processor (e.g., east-west across the functional slices, as shown by the second arrow), which is perpendicular to the first direction.

[0177] Different functional slices of the processor can correspond to MEM (memory), VXM (vector execution module), MXM (matrix execution module), NIM (numeric interpretation module), and SXM (switching and permutation module). Each slice can contain N tiles, all of which can be controlled by the same instruction control unit (ICU) (not shown). Each of the slices can operate completely independently and can only be coordinated using barrier-like synchronization primitives or by the compiler exploiting tractable determinism. Each tile of the processor can correspond to an execution unit organized as an xM SIMD tile. For example, each tile of the processor's on-chip memory can be organized to atomically store an L-element vector. Thus, a MEM slice with N tiles can work together to store or process a large vector (e.g., with a total of N x M elements).

[0178] Tiles within the same slice can execute instructions in a staggered manner, with instructions being issued tile by tile within the slice over a period of N cycles. Functional slices can be physically arranged on the chip to enable efficient data flow for pipelined execution over hundreds of cycles for common patterns. Data flows can perform a single "U-turn" (change of direction) corresponding to a single matrix operation before being written back to memory; in some embodiments, a given data flow can change direction multiple times (due to multiple matrix and vector operations) before the resulting data is written back to memory.

[0179] To achieve good single-threaded performance, a conventional multi-core processor design (e.g., as in Fig. Figure 21A illustrates this. Typically, a significant portion of the silicon area is reserved for exposing and exploiting instruction-level parallelism (ILP). This usually involves register renaming schemes and large instruction windows, through which the instructions have no explicit understanding of the hardware on which they are executed, while maintaining the illusion of in-order program execution. In contrast, when using a processor (e.g., TSP) with a functional slice architecture, the TSP compiler generates an explicit plan for how the processor will execute the microprogram. The compiler specifies when each operation will be performed, which functional slices will carry out the work, and which STREAM registers will hold the operands. The compiler maintains a highly reliable (cycle-accurate) model of the TSP's hardware state, allowing the microprogram to orchestrate the data flow.

[0180] The Processor 2100 (e.g., a TSP) can use a web-hosted compiler that takes a model (e.g., a machine learning model, such as, but not limited to, a TensorFlow model) as its input and outputs a proprietary instruction stream targeting the processor's TSP hardware. The compiler is responsible for coordinating the program's control and data flow, specifying any instruction-level parallelism by explicitly bundling instructions that can and should be executed concurrently so they can be output together. The primary hardware structure is the architecturally visible streaming register file (STREAMs), described in more detail below, which serves as the conduit through which operands flow from memory slices (e.g., SRAM) to functional slices and vice versa.

[0181] The processor's MEM unit serves as: (1) memory for model parameters, microprograms, and the data on which they operate, and (2) network-on-chip (NoC) for communicating data operands from the MEM to the functional slices and computed results back to the MEM. In some embodiments, the on-chip memory occupies approximately 75% of the processor's chip area. In some embodiments, the on-chip memory of the MEM tiles may be SRAM rather than DRAM due to the processor's bandwidth requirements. The processor's on-chip memory capacity determines (i) the number of machine learning models that can reside on the chip simultaneously, (ii) the size of any given model, and (iii) the partitioning of large models to fit into multi-chip systems. In some embodiments, the processor's MEM system provides a plurality of memory slices arranged in two distinct hemispheres (referred to as "MEM WEST" and "MEM 0").“MEM EAST” is the name given to these organizations.

[0182] The memory slices of each hemisphere can be mirrored, so that the slices can be physically numbered {0,..., L} in the east hemisphere (410) and {L,..., 0} in the west hemisphere (405), such that memory slice 0 for each hemisphere corresponds to the slice closest to the VXM slices (415) between the hemispheres, with each hemisphere containing L slices. The direction of data transmission toward the center of the chip can be described as inward, while data transmission toward the outermost (eastern or westernmost) edge of the chip can be described as outward. Although the processor's memory hemispheres can be referred to as east and west, it is understood that in other embodiments, other names may be used to refer to the different memory hemispheres.

[0183] In some embodiments, a streaming register file, referred to as STREAMS, transfers operands and results between the SRAM of the MEM slices and the functional slices of the processor. In some embodiments, a large number of MEM slices (e.g., between 2 and 10 adjacent MEM slices) can be physically organized as a set. Each set of slices can be located between a pair of STREAM register files, so that each slice is able to read from or write to the STREAM registers in both directions. Placing STREAM register files between sets of MEM slices reduces the number of cycles required for data operands to be transferred across one hemisphere (e.g., by a factor equal to the number of slices per set). The number of slices per set can be configured based on a distance over which data can be transferred in a single clock cycle.

[0184] With regard to any of the processors in Fig. 21 and any of the components described above or elsewhere herein can, for example, compile one or more of the APIs or equivalents described herein into instructions or equivalents that can be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled for execution by a scheduler or equivalents (e.g., in sequence or out of sequence), executed by execution logic or equivalents, reordered, and then terminated by retirement logic or equivalents. API(s) (and / or compiled instructions containing API(s)) can be stored in any memory outside or inside the processor (e.g., in cache and / or memory).A result from API(s) can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache or other memory equivalent.

[0185] In at least one embodiment, the 2100 processor may include one or more circuits for dynamically executing sections of neural networks or otherwise performing one of the operations described above or elsewhere herein. One or more circuits may be configured by software to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein. Software designs

[0186] The following figures present, without limitation, examples of software constructs for implementing at least one embodiment.

[0187] Fig. Figure 22 illustrates a software stack of a programming platform according to at least one embodiment. A programming platform may include a platform for utilizing hardware in a computing system to accelerate computational tasks. A programming platform may be accessible to software developers through libraries, compiler guidelines, and / or extensions of programming languages ​​in at least one embodiment. A programming platform may be CUDA, Radeon Open Compute Platform (“ROCm”), OpenCL (OpenCL™ is developed by the Khronos group), SYCL, or Intel oneAPI.

[0188] A software stack 2200 of a programming platform can provide an execution environment for an application 2201. The application 2201 can include any computer software capable of running on the software stack 2200. The application 2201 can include an artificial intelligence (“AI”) / machine learning (“ML”) application, a high-performance computing (“HPC”) application, a virtual desktop infrastructure (“VDI”), or a data center workload.

[0189] The 2201 application and the 2200 software stack run on 2208 hardware. The 2208 hardware can include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices that support a programming platform. The 2200 software stack can be vendor-specific and compatible only with devices from a particular vendor (or vendors), such as CUDA, ROCm, OneAPI, OpenCL, or other implementations. The 2208 hardware can include a host connected to one or more devices that can be accessed to perform computing tasks via calls to application programming interfaces (APIs).In at least one embodiment, a device within the hardware 2208 can include a GPU, an FPGA, an AI engine, or another computing device (but it can also include a CPU) and its memory, in contrast to a host within the hardware 2208, which can include a CPU (but it can also include a computing device) and its memory. With respect to any of the hardware 2208 described above or elsewhere herein, one or more of the APIs described herein can, for example, be compiled into instructions that can be retrieved by an instruction fetch logic, decoded by a processor decoder, scheduled for execution by a scheduler (e.g., in sequence or out of sequence), executed by an execution logic, reordered, and then terminated by a retirement logic.APIs (and / or compiled instructions containing APIs) can be stored in any memory outside or inside the processor (e.g., cache and / or memory). A result of APIs can then be stored in memory inside or outside the processor, including registers, DRAM, flash, SRAM, cache, or other memory. One or more of the APIs described herein can include a call. One or more of the APIs described herein can include a library or part of a library to perform a function described by the call. One or more of the APIs described herein can include a call and a library or part of a library to perform a function described by the call.

[0190] The software stack 2200 of a programming platform can include a number of libraries 2203, a runtime environment 2205, an optional driver / interface 2207, and a device kernel driver 2208. Each of the libraries 2203 can contain data and programming code that can be used by computer programs and employed during software development. The libraries 2203 can include predefined code and subroutines, classes, values, type specifications, configuration data, documentation, auxiliary data, and / or message templates. The libraries 2203 can include functions that can be optimized for execution on one or more device types. The libraries 2203 can include functions for performing mathematical, deep learning, and / or other types of operations on devices.The libraries 2203 can be mapped to corresponding APIs 2202, which may include one or more APIs that expose functions implemented in the libraries 2203. A processor (e.g., CPU, GPU) can execute, call, or otherwise use one or more APIs to prioritize kernels. For example, a first kernel (e.g., parent kernel) can start a second kernel (e.g., child kernel), and the second kernel can be used by a processor to start additional kernels (e.g., grandchildren kernels) independently of the first kernel. A processor can execute an API or call an API from memory to support dynamic stream priority (e.g., updating the priority while a stream is being used to perform operations).For example, when a processor executes the API, it allows a programmer to copy the stream priority from one stream to one or more other streams.

[0191] Software Stack 2200 can include an API to support dynamic stream priority (e.g., updating the priority while a stream is being used to perform operations), allowing a programmer to set a stream's priority at any time after its creation. Software Stack 2200 can also include an API to support dynamic stream priority (e.g., updating the priority while the stream is being used to perform operations), allowing a programmer to retrieve a stream's current priority, where the priority is one of several attributes of the stream. Finally, Software Stack 2200 can include an API to support dynamic stream priority (e.g., updating the priority while the stream is being used to perform operations), allowing a programmer to retrieve a stream's current priority as a single attribute.Software Stack 2200 can include an API to support dynamic stream priority (e.g., updating the priority while the stream is being used to perform operations), allowing a programmer to start a kernel to perform operations on a stream with a set priority that may differ from the stream priority. Software Stack 2200 can also include an API to indicate whether an object (e.g., a thread synchronization object, such as, but not limited to, a barrier) is tracking whether all data movement operations for a set of threads running on a GPU have completed after a certain period of time and are in a specific state. This specific state could indicate that data has been moved and is ready for use, and is specified as input to the API using an expected parity value.

[0192] The Software Stack 2200 can include one or more APIs for updated kernels. A processor can execute an API or call an API from memory to update an existing API and support context-free kernels, allowing a programmer to add a kernel node to a graph without a graphics context, so that a graphics context can be dynamically associated with a kernel at runtime. The Software Stack 2200 can include one or more APIs to allow a programmer to retrieve a kernel identifier and a graphics context as separate parameters from a kernel node, so that parameters can be retrieved from kernels and from context-free kernels. The Software Stack 2200 can include one or more APIs to use parallel processor(s), such as, but not limited to, one or more graphics processing units (GPUs), to process task graphs (e.g., graphics).to start task graphs) and execute one or more task graphs (e.g., including one or more programs).

[0193] The Software Stack 2200 can include one or more APIs to associate one or more commands with one or more memory order operations, such as, but not limited to, a fence or membar operation. Commands can be associated with one or more domains, so that a memory order operation is performed in conjunction with one or more specific domains without interfering with commands in other domains. An API can indicate that a thread has arrived (for example, at a thread synchronization barrier) or has completed a stage of work with respect to asynchronous data movement operations on a GPU.The Software Stack 2200 can include one or more features to allow programmers to manually specify an expected transaction count when a thread has completed a work stage, which can be used to update an object that tracks whether all data movement operations have been completed for a set of threads.

[0194] Application 2201 can be written as source code that is compiled into executable code, as shown below in conjunction with the Fig. 23 and Fig. 24 is discussed in more detail. The executable code of application 2201 can run, at least partially, on an execution environment provided by software stack 2200. During the execution of application 2201, code may be required to run on a device rather than a host. In such a case, runtime environment 2205 can be invoked to load and start the necessary code on the device. Runtime environment 2205 can include any technically feasible runtime system capable of supporting the execution of application 2201.

[0195] The 2205 runtime environment can be implemented as one or more runtime libraries associated with corresponding APIs, shown as API(s) 2204. One or more such runtime libraries may include, among other things, memory management, execution control, device management, error handling, and / or synchronization functions. Memory management functions may include functions for allocating, releasing, and copying device memory, as well as for transferring data between host memory and device memory. Execution control functions may include functions for starting a function (sometimes called a "kernel" if the function is a global function that can be called by a host) on a device and for setting attribute values ​​in a buffer maintained by a runtime library for a given function to be executed on a device.

[0196] Runtime libraries and their corresponding APIs can be implemented in any technically feasible way. One (or any number of) APIs can expose a set of low-level functions for fine-grained control of a device, while another (or any number of) APIs can expose a set of such functions at a higher level. A high-level runtime API can be built upon a low-level API. One or more runtime APIs can be language-specific APIs that may be built upon a language-independent runtime API.

[0197] An optional driver or interface 2207 can be implemented, for example, for CUDA and ROCm implementations, which are described below. The optional driver / interface 2207 can be associated with an optional driver or interface API(s), such as, but not limited to, CUDA and / or ROCm API(s).

[0198] One or more processors disclosed in "Processing Systems" can execute, access, or otherwise use Software Stack 2200. For example, System-on-a-Chip 900, Parallel Processor 1000, Graphics Multiprocessor 1034, Processor 1100, Processor 1200, Accelerator 1300, Neuromorphic Processor 1405, Supercomputer 1500, Acceleration Processing Unit 1600, Processor 1700, Processor 1800, Tensor Processing Unit 1900, Processor 2000, and Speech Processing Unit 2100 can execute, use, call, or otherwise implement (e.g., by accessing memory) one or more APIs included in Software Stack 2200.

[0199] The Device Kernel Driver 2208 can be configured to facilitate communication with an underlying device. The Device Kernel Driver 2208 can provide low-level functionality upon which APIs, such as, but not limited to, the API(s) 2204 and / or other software, depend. The Device Kernel Driver 2208 can be configured to compile intermediate representation code (“IR” code) into binary code at runtime. For CUDA or other implementations, such as, but not limited to, ROCm, OneAPI, or OpenCL, the Device Kernel Driver 2208 can compile non-hardware-specific parallel thread execution (“PTX”) IR code at runtime into binary code for a specific target device (with caching of the compiled binary code), sometimes referred to as “completion code.”This can allow completed code to run on a target device that may not have existed when the source code was originally compiled into PTX code. Alternatively, the device source code can be compiled offline into binary code without the need for the 2208 device kernel driver to compile IR code at runtime.

[0200] Processors described elsewhere herein, such as, but not limited to, processors in Fig. 9-21, may include one or more circuits to dynamically execute sections of neural networks or otherwise perform any of the operations described above or elsewhere herein. One or more circuits may be configured by software, e.g., software stack 2200, to dynamically execute the sections of neural networks or otherwise perform any of the operations described above or elsewhere herein.

[0201] According to at least one embodiment, the software stack 2200 can be used by Fig. 22 in a CUDA implementation. A CUDA software stack 2200, on which an application 2201 can be started, can include CUDA libraries 2203, a CUDA runtime environment 2205, a CUDA driver 2207, and a device kernel driver 2208. The CUDA software stack 2200 can be run on hardware 2609, which can include a GPU that supports CUDA and is developed by NVIDIA Corporation in Santa Clara, CA.

[0202] The application 2201, the CUDA runtime environment 2205, and the device kernel driver 2208 can perform functionalities described above and elsewhere herein. The CUDA driver 2207 can include a library (libcuda.so) that can implement a CUDA driver API 2206. Similar to a CUDA runtime API 2204 implemented by a CUDA runtime library (cudart), the CUDA driver API 2206 can expose, among other things, memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability functions. The CUDA driver API 2206 may differ from the CUDA runtime API 2204 in that the CUDA runtime API 2204 simplifies device code management by providing implicit initialization, context management (analogous to a process), and module management (analogous to dynamically loaded libraries).In contrast to the high-level CUDA Runtime API 2204, the CUDA Driver API 2206 can be a low-level API that provides finer control over the device, particularly regarding contexts and module loading. The CUDA Driver API 2206 can expose context management features that might not be exposed by the CUDA Runtime API 2204. In addition to being language-independent and supportive, the CUDA Driver API 2206 can also be language-independent and supportive, for example, OpenCL. Furthermore, development libraries, including the CUDA Runtime Environment 2205, can be considered separate from driver components, including the user-mode CUDA Driver 2207 and the kernel-mode Device Driver 2208 (sometimes referred to as the "display" driver).

[0203] The CUDA libraries 2203 can include mathematical libraries, deep learning libraries, libraries for parallel algorithms, and / or signal / image / video processing libraries that can be used by parallel computing applications, such as, but not limited to, the application 2201. The CUDA libraries 2203 can include mathematical libraries such as, but not limited to, a cuBLAS library, which is an implementation of Basic Linear Algebra Subprograms (“BLAS”) for performing linear algebra operations, a cuFFT library for calculating fast Fourier transforms (“FFTs”), and a cuRAND library for generating random numbers, among others. The CUDA libraries 2203 can include deep learning libraries such as, but not limited to, a cuDNN library of primitives for deep neural networks and a TensorRT platform for high-performance deep learning inference, among others.

[0204] In at least one embodiment, processors described elsewhere herein, such as, but not limited to, processors in Fig. 9-21, have one or more circuits to dynamically execute sections of neural networks or otherwise perform one of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., software stack 2200, to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0205] According to at least one embodiment, the software stack 2200 can be used by Fig. 22 in an ROCm implementation. An ROCm software stack 2200, on which an application 2201 can be started, includes a language runtime environment 2203, a system runtime environment 2205, a thunk 2207, and an ROCm kernel driver 2208. The ROCm software stack 2200 runs on hardware 2209, which may include a GPU that supports ROCm and is developed by AMD Corporation in Santa Clara, CA.

[0206] Application 2201 can perform similar functionalities as described above in conjunction with Fig. 22 discussed. Furthermore, the Language Runtime Environment 2203 and the System Runtime Environment 2205 can perform similar functionalities to those mentioned above in conjunction with Fig. 22 discusses runtime environment 2205. The language runtime environment 2203 and the system runtime environment 2205 can differ in that the system runtime environment 2205 is a language-independent runtime environment that implements a ROCr system runtime API 2204 and uses a heterogeneous system architecture runtime environment (“HSA” runtime environment). The HSA runtime API can include, among other things, a lean, user-mode API that exposes interfaces for accessing and interacting with an AMD GPU, including memory management functions, execution control via an architecture-based distribution of kernels, error handling, system and agent information, and runtime initialization and shutdown. In contrast to the system runtime environment 2205, the language runtime environment 2203 can be an implementation of a language-specific runtime API 2202 built on top of the ROCr system runtime API 2204.The language runtime API can include, among other things, a Heterogeneous Compute Interface for Portability (“HIP”) language runtime API, a Heterogeneous Compute Compiler (“HCC”) language runtime API, or an OpenCL API. The HIP language, in particular, is an extension of the C++ programming language with functionally similar versions of CUDA mechanisms, and an HIP language runtime API can include, among other things, functions similar to those of the aforementioned APIs. Fig. 22 discussed CUDA runtime APIs may be similar, such as, but not limited to, memory management, execution control, device management, error handling and synchronization functions.

[0207] The Thunk (ROCt) 2207 can be an interface 2206 that can be used to interact with the underlying ROCm driver 2208. The ROCm driver 2208 can be a ROCK driver that is a combination of an AMDGPU driver and an HSA kernel driver (AMDKFD). The AMDGPU driver can be a device kernel driver for AMD-designed GPUs that performs similar functionalities to the one described above in conjunction with Fig. 22 discussed device kernel drivers 2209. The HSA kernel driver can be a driver that allows different types of processors to share system resources more effectively via hardware features.

[0208] Various libraries (not shown) may be included in the ROCm software stack 2200 above the language runtime environment 2203 and provide functionality that complements the above in conjunction with Fig. 22 discussed CUDA libraries are similar to 2203. Various libraries may include mathematical, deep learning, and / or other libraries, such as, but not limited to, a cuBLAS library implementing functions similar to those of CUDA cuBLAS, a rocFFT library for calculating FFTs similar to CUDA cuFFT, among others.

[0209] Processors described elsewhere herein, such as, but not limited to, processors in Fig. 9-21, may include one or more circuits to dynamically execute sections of neural networks or otherwise perform any of the operations described above or elsewhere herein. One or more circuits may be configured by software, e.g., software stack 2200, to dynamically execute the sections of neural networks or otherwise perform any of the operations described above or elsewhere herein.

[0210] According to at least one embodiment, the software stack 2200 can be used by Fig. 22 in an OpenCL implementation. An OpenCL software stack 2200, on which an application 2201 can be started, may include an OpenCL framework 2203, an OpenCL runtime environment 2205, and a driver 2208. The OpenCL software stack 2200 can run on hardware 2209 that is not vendor-specific. Because OpenCL is supported by devices developed by various vendors, it may be necessary for specific OpenCL drivers to work with hardware from such vendors.

[0211] The application 2201, the OpenCL runtime environment 2205, the device kernel driver 2208, and the hardware 2209 can perform similar functionalities to other implementations of the application 2201, the runtime environment 2205, the device kernel driver 2208, and the hardware 2209, respectively, described above in conjunction with Fig. 22 are discussed. Application 2201 may further include an OpenCL kernel (not shown) with code to be executed on a device.

[0212] OpenCL can define a "platform" that allows a host to control devices connected to it. An OpenCL framework can provide a platform layer API and a runtime API, shown as Platform API 2202 and Runtime API 2204, respectively. Runtime API 2204 can use contexts to manage kernel execution on devices. Each identified device can be associated with a specific context, which Runtime API 2204 can use to manage command queues, program objects, and kernel objects that, among other things, share memory objects for that device. Platform API 2202 can, among other things, expose functionality that allows the use of device contexts to select and initialize devices, submit work to devices via command queues, and enable data transfer to and from devices.In addition, the OpenCL framework can provide various integrated functions (not shown), including mathematical functions, relational functions, and image processing functions.

[0213] A compiler (not shown) may also be included in the OpenCL framework 2203. Source code can be compiled offline before running an application or online during application execution. Unlike CUDA and ROCm, OpenCL applications can be compiled online by a compiler representative of any number of compilers that can be used to compile source code and / or IR code, such as, but not limited to, Standard Portable Intermediate Representation (SPIR-V) code, into binary code. Alternatively, OpenCL applications can be compiled offline before such applications are executed.

[0214] In at least one embodiment, processors described elsewhere herein, such as, but not limited to, processors in Fig. 9-21, have one or more circuits to dynamically execute sections of neural networks or otherwise perform one of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., software stack 2200, to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

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

[0216] The programming platform can be one of a CUDA, ROCm, or OpenCL platform, as described above in conjunction with Fig. The programming platform can support multiple programming models, which can be abstractions of an underlying computing system that enable the representation of algorithms and data structures. The programming models can expose features of the underlying hardware to improve performance. The programming models can include CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++AMP”), Open Multi-Processing (“OpenMP”), Open Accelerators (“OpenACC”), and / or Vulkan Compute.

[0217] Libraries and / or middleware can provide implementations of programming model abstractions. Such libraries can include data and code that can be used by computer programs and employed during software development. Middleware can also include software that provides services for applications beyond those available from the programming platform. Libraries and / or middleware can include cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Additionally, libraries and / or middleware can include NCCL and ROCm Communication Collective Library ("RCCL") libraries, which provide communication routines for GPUs, an MIOpen library for deep learning acceleration, and / or an Eigen library for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0218] Application frameworks can depend on libraries and / or middleware. Each application framework can be a software framework used to implement a standard application software structure. Referring again to the AI / ML example discussed above, an AI / ML application can be implemented using a framework such as, but not limited to, Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet deep learning frameworks.

[0219] In at least one embodiment, processors described elsewhere herein, such as, but not limited to, processors in Fig. 9-21, include one or more circuits to dynamically execute sections of neural networks or otherwise perform one of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., programming platforms described herein, to dynamically execute the sections of neural networks or otherwise perform one of the operations described above or elsewhere herein.

[0220] Fig. Figure 23 illustrates compiling code for execution on one of the programming platforms described above. Fig. 22 according to at least one embodiment. A compiler 2301 is configured to receive source code 2300, compile source code 2300, and output an executable file 2310. The compiler 2301 can be configured to convert source code 2300 into executable host code 2307 for execution on a host and executable device code 2308 for execution on a device. The source code 2300 can be compiled either offline before the execution of an application or online during the execution of an application. The source code 2300 can contain code in any programming language supported by the compiler 2301, such as, but not limited to, C++, C, Fortran, etc. The source code 2300 can be contained in a single source file that has a mixture of host code and device code, specifying memory locations for the device code therein. A single source file can be a .cu file containing CUDA code, or a .hip file.The source code 2300 may consist of a .cpp file containing HIP code, or a file in another format containing both host code and device code. Alternatively, the source code 2300 may consist of multiple source code files instead of a single source file, allowing host code and device code to be separated. The compiler 2301 includes or has access to one or more libraries to detect a sequence of API calls that execute a single fused API, where a single fused API is a combined API for two or more APIs. In at least one embodiment, the compiler 2301 may be an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code into .cu files, an HCC compiler for compiling HIP code into .hip.cpp files, or other compilers.

[0221] The 2301 compiler can be configured to compile 2300 source code into 2307 host executable code for execution on a host and 2308 device executable code for execution on a device. The 2301 compiler introduces operations that include parsing 2300 source code into an abstract system tree (AST), performing optimizations, and generating executable code. If the 2300 source code consists of a single source file, the 2301 compiler can separate device code from host code in such a single source file, compile the device code and host code into 2308 device executable code and 2307 host executable code, respectively, and then join the device executable code and host executable code into a single file.

[0222] The compiler 2301 can include a compiler frontend 2302, a host compiler 2305, a device compiler 2306, and a linker 2309. The compiler frontend 2302 can be configured to separate device code 2304 from host code 2303 in the source code 2300. In at least one embodiment, device code 2304 can be compiled by the device compiler 2306 into executable device code 2308, which, as described, can contain binary code or IR code. Separately, host code 2303 can be compiled by the host compiler 2305 into executable host code 2307. For other NVCC compilers, such as, but not limited to, those for oneAPI, ROCm and OpenCL, the host compiler 2305 can be a general-purpose C / C++ compiler that outputs native object code, while the device compiler 2306 can be a low-level virtual machine (“LLVM”)-based compiler that branches an LLVM compiler infrastructure and outputs PTX code or binary code.For HCC, both the host compiler 2305 and the device compiler 2306 can be LLVM-based compilers that output binary target code.

[0223] After compiling source code 2300 into executable host code 2307 and executable device code 2308, the linker 2309 can link the executable host code and device code 2307 and 2308 together in an executable file 2310. Native object code for a host and PTX or binary code for a device can be linked together in an (“ELF”) file in an executable and linkable format, which is a container format used to store object code. The executable host code 2307 and the executable device code 2308 can be in any suitable format, such as, but not limited to, binary code and / or IR code. In the case of CUDA, the executable host code 2307 can, in at least one embodiment, contain native object code, and the executable device code 2308 can contain code in a PTX intermediate representation.In the case of ROCm, both the executable host code 2307 and the executable device code 2308 can include binary target code in at least one embodiment. Other implementations, such as, but not limited to, oneAPI and OpenCL, are considered and can be executed similarly to the CUDA and ROCm implementations described above.

[0224] The source code 2300 can be translated before compiling the source code. The source code is passed through a translation tool (not shown) that translates the source code 2300 into the translated source code. A compiler 2301 can be used to compile the translated source code into the executable host code 2307 and the executable device code 2308 in a procedure similar to the compilation of the source code 2300 by the compiler 2301 into the executable host code 2307 and the executable device code 2308, as described above in conjunction with Fig. 23 discussed.

[0225] A translation performed by the translation tool can be used to port the 2300 source code for execution in a different environment than the one in which it was originally intended to run. The translation tool may include a HIP translator, which is used to "HIPify" CUDA code intended for a CUDA platform into HIP code that can be compiled and executed on a ROCm platform. The translation of the 2300 source code may involve parsing the 2300 source code 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 described below in conjunction with Fig. 24 is discussed in more detail. Referring again to the example of HIP-ifying CUDA code, calls to the CUDA runtime API, the CUDA driver API, and / or the CUDA libraries can be converted into corresponding HIP API calls. Automated translations performed by the translation tool 2301 can sometimes be incomplete, requiring additional manual effort to fully port the source code 2300.

[0226] One or more of the methods described herein may use other methods for converting one type of code to another type of code to enable interchangeability between different device architectures. In at least one embodiment, an application for one platform (e.g., a CUDA application) may be compiled into code for implementation on another platform (e.g., an AMD processor, an Intel processor, or another processor). For example, the source code 2300 may contain source code for one platform (e.g., CUDA). The compiler 2301 may compile the source code 2300 into an executable file 2310 that can be used by another platform (e.g., AMD or Intel). Programming toolkits may enable applications for one platform (e.g., CUDA) to be compiled (e.g., natively) for another platform (e.g., AMD or Intel).For example, a GPGPU programming toolkit can enable CUDA applications to be compiled natively for AMD GPUs. Programs (e.g., CUDA programs) or their build system do not need to be modified or translated into another language before being compiled to code for a different platform. A compiler can accept the same command-line options and programming dialect (e.g., CUDA dialect) as another compiler (e.g., nvcc for CUDA), thus replacing the need to install a toolkit (e.g., NVIDIA CUDA Toolkit), allowing existing build tools and scripts (e.g., such as CMake) to function without further modification. In at least one embodiment, an nvcc-compatible compiler can be used to compile nvcc-dialect CUDA for AMD GPUs, including PTX-ASM. Implementations of CUDA runtime and driver APIs for AMD GPUs can be used. Libraries (e.g.,Open-source wrapper libraries can provide APIs, such as CUDA-X APIs, by delegating to the corresponding ROCm libraries. One exemplary implementation is SCALE from Spectral Compute in London, England. Rather than providing a new way to write GPGPU software, SCALE allows programs written using the widely adopted CUDA language to be compiled directly for AMD GPUs. Additional implementations may include a Clang compiler, which provides a language frontend and tooling infrastructure for languages ​​in the C language family (C, C++, Objective C / C++, OpenCL, CUDA, and RenderScript). In at least one embodiment, the compilers described herein, such as, but not limited to, the 2301, 2305, and / or 2306 compilers, may include one or more circuits for compiling code (e.g.,CUDA, HIP, OpenCL, OneAPI or others) for the dynamic execution of sections of the neural network and / or for performing any of the operations described above or elsewhere herein.

[0227] Fig. Figure 24 illustrates a System 2400 configured to compile and execute CUDA source code 2410 using various types of processing units according to at least one embodiment. The System 2400 includes CUDA source code 2410, a CUDA compiler 2450, host executable code 2470(1), host executable code 2470(2), CUDA device executable code 2484, a CPU 2490, a CUDA-enabled GPU 2494, a GPU 2492, a CUDA-to-HIP translation tool 2420, HIP source code 2430, a HIP compiler driver 2440, an HCC 2460, and HCC device executable code 2482.

[0228] The CUDA source code 2410 can be a collection of human-readable code written in a CUDA programming language. A CUDA programming language can be an extension of the C++ programming language that includes mechanisms to define device code and distinguish 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 the CUDA-enabled GPU 2490, the GPU 2492, or another GPGPU, etc. Host code is source code that, after compilation, can be executed on a host. A host is a processor optimized for sequential instruction processing, such as the CPU 2490, among others.

[0229] The CUDA source code 2410 can contain any number (including zero) of global functions 2412, any number (including zero) of device functions 2414, any number (including zero) of host functions 2416, and any number (including zero) of host / device functions 2418. The global functions 2412, device functions 2414, host functions 2416, and host / device functions 2418 can be mixed in the CUDA source code 2410. Each of the global functions 2412 can be executable on a device and callable by a host. One or more of the global functions 2412 can therefore act as entry points for a device. Each of the global functions 2412 can be a kernel. In a technique known as dynamic parallelism, one or more of the global functions 2412 can define a kernel that is executable on and invoked by a device.A kernel can be executed N times in parallel by N different threads on a device during execution (where N is any positive integer).

[0230] Each of the device functions 2414 can only be executed on and called by one device. Each of the host functions 2416 can only be executed on and called by one host. Each of the host / device functions 2416 can define both a host version of a function, which is executable on and called by one host, and a device version of the function, which is executable on and called by one device.

[0231] CUDA source code 2410 can also include any number of calls to any number of functions that can be defined via a CUDA runtime API 2402. The CUDA runtime API 2402 can include any number of functions that run on a host to allocate and release device memory, transfer data between host memory and device memory, manage systems with multiple devices, and so on. CUDA source code 2410 can also include any number of calls to any number of functions that can be specified in any number of other CUDA APIs. A CUDA API can be any API designed for use by CUDA code. CUDA APIs can include the CUDA runtime API 2402, a CUDA driver API, APIs for any number of CUDA libraries, and so on, including any API(s) described elsewhere herein.With regard to the CUDA runtime API 2402, a CUDA driver API can be a lower-level API, but it can provide finer-grained control of a device or instrument. Examples of CUDA libraries include cuBLAS, cuFFT, cuRAND, cuDNN, etc.

[0232] The CUDA compiler 2450 can compile input CUDA code (e.g., CUDA source code 2410) to produce executable host code 2470(1) and executable CUDA device code 2484. The CUDA compiler 2450 can be NVCC, among others. The executable host code 2470(1) can be a compiled version of host code contained in the input source code, executable on the CPU 2490. The CPU 2490 can be any processor optimized for sequential instruction processing.

[0233] The CUDA device code 2484 executable can be a compiled version of device code contained in the input source code, executable on the CUDA-enabled GPU 2494. The CUDA device code 2484 executable can include binary code. The CUDA device code 2484 executable can include IR code, such as, but not limited to, PTX code, which is further compiled at runtime by a device driver into binary code for a specific target device (e.g., the CUDA-enabled GPU 2494). The CUDA-enabled GPU 2494 can include any processor optimized for parallel instruction processing that supports CUDA. The CUDA-enabled GPU 2494 may have been developed by NVIDIA Corporation in Santa Clara, CA.

[0234] The CUDA-to-HIP translation tool 2420 can be designed to translate CUDA source code 2410 into functionally similar HIP source code 2430. The HIP source code 2430 can include a collection of human-readable code in a HIP programming language. The HIP code can include human-readable code in a HIP programming language. A HIP programming language can include an extension of the C++ programming language that incorporates functionally similar versions of CUDA mechanisms for defining device code and distinguishing between device code and host code. A HIP programming language can include a subset of the functionality of a CUDA programming language.For example, a HIP programming language includes a mechanism to define global functions 2412, but such a HIP programming language may lack support for dynamic parallelism, and therefore global functions 2412 defined in the HIP code may only be callable by one host.

[0235] The HIP source code 2430 can include any number (including zero) of global functions 2412, any number (including zero) of device functions 2414, any number (including zero) of host functions 2416, and any number (including zero) of host / device functions 2418. The HIP source code 2430 can also include any number of calls to any number of functions that may be specified in a HIP runtime API 2432. The HIP runtime API 2432 can include functionally similar versions of a subset of functions contained in the CUDA runtime API 2402. The HIP source code 2430 can also include any number of calls to any number of functions that may be specified in any number of other HIP APIs. A HIP API can be any API designed for use by HIP code and / or ROCm.HIP APIs can include the HIP runtime API 2432, a HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, etc.

[0236] The CUDA-to-HIP translation tool 2420 can convert any kernel call in CUDA code from CUDA syntax to HIP syntax and can convert any number of other CUDA calls in CUDA code to any number of other functionally similar HIP calls. A CUDA call can involve a call to a function specified in a CUDA API, and a HIP call can involve a call to a function specified in a HIP API. The CUDA-to-HIP translation tool 2420 can convert any number of calls to functions specified in the CUDA runtime API 2402 to any number of calls to functions specified in the HIP runtime API 2432.

[0237] The CUDA-to-HIP translation tool 2420 may include a tool known as hipify-perl, which performs a text-based translation process. Alternatively, the CUDA-to-HIP translation tool 2420 may include a tool known as hipify-clang, which performs a more complex and robust translation process compared to hipify-perl. This process involves parsing CUDA code using clang (a compiler frontend) and then translating the resulting symbols. Converting CUDA code to HIP code may involve modifications (e.g., manual edits) in addition to those performed by the CUDA-to-HIP translation tool 2420.

[0238] The HIP compiler driver 2440 can include a front end that determines a target device 2446 and then configures a compiler compatible with the target device 2446 to compile the HIP source code 2430. The target device 2446 can include a processor optimized for parallel instruction processing. The HIP compiler driver 2440 can determine the target device 2446 in any technically feasible way.

[0239] If the target device 2446 is compatible with CUDA (e.g., the CUDA-enabled GPU 2494), then the HIP compiler driver 2440 can generate a HIP / NVCC compilation instruction 2442. The HIP / NVCC compilation instruction 2442 can configure the CUDA compiler 2450 to compile the HIP source code 2430 using a HIP-to-CUDA translation header and a CUDA runtime library. In response to the HIP / NVCC compilation instruction 2442, the CUDA compiler 2450 can generate host executable code 2470(1) and CUDA device executable code 2484.

[0240] If the target device 2446 is not CUDA-compatible, the HIP compiler driver 2440 can generate a HIP / HCC compilation command 2444. The HIP / HCC compilation command 2444 can configure the HCC 2460 to compile the HIP source code 2430 using an HCC header and a HIP / HCC runtime library. In response to the HIP / HCC compilation command 2444, the HCC 2460 can generate host executable code 2470(2) and HCC device executable code 2482. The HCC device executable code 2482 can be a compiled version of device code contained in the HIP source code 2430, which is executable on the GPU 2492. The GPU 2492 can be any processor optimized for parallel instruction processing, not CUDA-compatible, and HCC-compatible. The GPU 2492 may have been developed by AMD Corporation in Santa Clara, CA. The GPU 2492 may include a non-CUDA-enabled GPU 2492.

[0241] The following are for illustrative purposes only: Fig. Figure 24 shows three different processes that can be implemented in at least one embodiment to compile CUDA source code 2410 for execution on the CPU 2490 and various devices. A direct CUDA process can compile CUDA source code 2410 for execution on the CPU 2490 and the CUDA-enabled GPU 2494 without translating the CUDA source code 2410 into the HIP source code 2430. An indirect CUDA flow can translate the CUDA source code 2410 into the HIP source code 2430 and then compile the HIP source code 2430 for execution on the CPU 2490 and the CUDA-enabled GPU 2494. A CUDA / HCC flow can translate the CUDA source code 2410 into the HIP source code 2430 and then compile the HIP source code 2430 for execution on the CPU 2490 and the GPU 2492.

[0242] A direct CUDA flow that can be implemented is represented by dashed lines and a series of bubbles labeled A1-A3. As shown by bubble A1, the CUDA compiler 2450 can receive the CUDA source code 2410 and a CUDA compilation instruction 2448, which can configure the CUDA compiler 2450 to compile the CUDA source code 2410. The CUDA source code 2410, which can be used in a direct CUDA flow, 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 the CUDA compilation instruction 2448, the CUDA compiler 2450 can generate executable host code 2470(1) and executable CUDA device code 2484 (shown in bubble A2). As shown in bubble A3, the executable host code 2470(1) and the executable CUDA device code 2484 can be executed on the CPU 2490 and the CUDA-enabled GPU 2494, respectively.The executable CUDA device code 2484 can contain binary code. The executable CUDA device code 2484 can contain PTX code and can be further compiled at runtime into binary code for a specific target device.

[0243] An indirect CUDA flow that can be implemented is represented by dashed lines and a series of bubbles labeled B1-B6. As shown by bubble B1, the CUDA-to-HIP translation tool 2420 can receive the CUDA source code 2410. As shown by bubble B2, the CUDA-to-HIP translation tool 2420 can translate the CUDA source code 2410 into the HIP source code 2430. As shown by bubble B3, the HIP compiler driver 2440 can receive the HIP source code 2430 and determine that the target device 2446 is CUDA-enabled.

[0244] As illustrated in bubble B4, the HIP compiler driver 2440 can generate the HIP / NVCC compilation instruction 2442 and transfer both the HIP / NVCC compilation instruction 2442 and the HIP source code 2430 to the CUDA compiler 2450. The HIP / NVCC compilation instruction 2442 can configure the CUDA compiler 2450 to compile the HIP source code 2430 using a HIP-to-CUDA translation header and a CUDA runtime library. The HIP-to-CUDA translation header can translate any number of mechanisms (e.g., functions) specified in any number of HIP APIs into any number of mechanisms specified in any number of CUDA APIs. The CUDA compiler 2450 can use the HIP-to-CUDA translation header in conjunction with a CUDA runtime library that conforms to the CUDA runtime API 2402 to generate executable host code 2470(1) and executable CUDA device code 2484.In response to the HIP / NVCC compilation instruction 2442, the CUDA compiler 2450 can generate host executable code 2470(1) and CUDA device executable code 2484 (shown in bubble B5). As shown in bubble B6, host executable code 2470(1) and CUDA device executable code 2484 can be executed on the CPU 2490 and the CUDA-enabled GPU 2494, respectively. CUDA device executable code 2484 can contain binary code. CUDA device executable code 2484 can contain PTX code and can be further compiled at runtime into binary code for a specific target device.

[0245] A CUDA / HCC workflow that can be implemented is represented by solid lines and a series of bubbles labeled C1-C6. As shown by bubble C1, the CUDA-to-HIP translation tool 2420 can receive the CUDA source code 2410. As shown by bubble C2, the CUDA-to-HIP translation tool 2420 can translate the CUDA source code 2410 into the HIP source code 2430. As shown by bubble C3, the HIP compiler driver 2440 can receive the HIP source code 2430 and determine that the target device 2446 is not CUDA-enabled.

[0246] The HIP compiler driver 2440 can generate the HIP / HCC compilation command 2444 and transfer both the HIP / HCC compilation command 2444 and the HIP source code 2430 to the HCC 2460 (represented by bubble C4). The HIP / HCC compilation command 2444 can configure the HCC 2460 to compile the HIP source code 2430 using an HCC header and a HIP / HCC runtime library. The HIP / HCC runtime library can correspond to the HIP runtime API 2432. The HCC header can include any number and type of interoperability mechanisms for HIP and HCC. In response to the HIP / HCC compilation instruction 2444, the HCC 2460 can generate host executable code 2470(2) and HCC device executable code 2482 (shown in bubble C5). As shown in bubble C6, host executable code 2470(2) and HCC device executable code 2482 can be executed on the CPU 2490 and GPU 2492, respectively.

[0247] After the CUDA source code 2410 has been translated into the HIP source code 2430, the HIP compiler driver 2440 can then be used to generate executable code for either the CUDA-enabled GPU 2494 or the GPU 2492 without re-executing the CUDA-to-HIP translation tool 2420. The CUDA-to-HIP translation tool 2420 can translate the CUDA source code 2410 into the HIP source code 2430, which is then stored in memory. The HIP compiler driver 2440 can then configure the HCC 2460 to generate host executable code 2470(2) and HCC device executable code 2482 based on the HIP source code 2430. In at least one embodiment, the HIP compiler driver 2440 then configures the CUDA compiler 2450 to generate executable host code 2470(1) and executable CUDA device code 2484 based on the stored HIP source code 2430.

[0248] An example kernel can be generated by the CUDA-to-HIP translation tool 2420 from Fig. 24 can be translated according to at least one embodiment. The CUDA source code 2410 partitions an overall problem, which a given kernel is to solve, into relatively coarse subproblems that can be solved independently using thread blocks. Each thread block contains an arbitrary number of threads. Each subproblem can be partitioned into relatively small pieces that can be solved cooperatively in parallel by threads within a thread block. Threads within a thread block can cooperate by exchanging data over shared memory and synchronizing execution to coordinate memory accesses.

[0249] The CUDA source code 2410 can organize thread blocks associated with a specific kernel into a one-dimensional, two-dimensional, or three-dimensional grid of thread blocks. Each thread block can contain any number of threads, and a grid can contain any number of thread blocks.

[0250] A kernel can be a function in device code defined using a declaration specifier "__global__". The dimension of a grid that executes a kernel for a specific kernel call and associated streams can be specified using a CUDA kernel start syntax. The CUDA kernel start syntax is denoted as "KernelName<<<GridSize, BlockSize, SharedMemorySize, Stream> >>(KernelArguments);” specifies. An execution configuration syntax can include a construct “<<<...>>>” inserted between a kernel name (“KernelName”) and a list of kernel arguments (“KernelArguments”) enclosed in parentheses. The CUDA kernel startup syntax can include a CUDA startup function syntax instead of an execution configuration syntax.

[0251] `GridSize` can be of type `dim3` and specify the dimension and size of a 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 one. If `y` is not specified, it defaults to one. The number of thread blocks in a grid can be the product of `GridSize.x`, `GridSize.y`, and `GridSize.z`. `BlockSize` can be of type `dim3` and specify the dimension and size of each thread block. The number of threads per thread block can be the product of `BlockSize.x`, `BlockSize.y`, and `BlockSize.z`. Each thread executing a kernel can be given a unique thread ID, accessible within the kernel via a built-in variable (e.g., `threadIdx`).

[0252] Regarding the CUDA kernel startup syntax, `SharedMemorySize` can be an optional argument that specifies the number of bytes in shared memory that are dynamically allocated per thread block for a given kernel call, in addition to statically allocated memory. By default, `SharedMemorySize` can be zero. Similarly, `Stream` can be an optional argument that specifies an associated stream and is zero by default to specify a standard stream. A stream can be a sequence of commands (possibly issued by different host threads) that are executed sequentially. Different streams can execute commands in any order or concurrently.

[0253] The CUDA source code 2410 can contain a kernel definition for an example kernel "MatAdd" and a main function. The main function can be host code that runs on a host and includes a kernel call that causes the MatAdd kernel to run on a device. The MatAdd kernel can add two matrices A and B of size NxN, where N is a positive integer, and store the result in a matrix C. The main function can define a ThreadsPerBlock variable as 16x16 and a numBlocks variable as N / 16xN / 16. The main function can then make the kernel call "MatAdd<<<numBlocks, threadsPerBlock> >>(A, B, C);“ specify. According to the CUDA kernel startup syntax, the kernel MatAdd can be executed using a grid of thread blocks with a dimension of N / 16 by N / 16, where each thread block has a dimension of 16 by 16.Each thread block can contain 256 threads, a grid can be created with enough blocks to have one thread per matrix element, and each thread in such a grid can execute the kernel MatAdd to perform pairwise addition.

[0254] During the translation of the CUDA source code 2410 to the HIP source code 2430, the CUDA-to-HIP translation tool 2420 can translate each kernel call in the CUDA source code 2410 from the CUDA kernel launch syntax to a HIP kernel launch syntax and can convert any number of other CUDA calls in the source code 2410 to any number of other functionally similar HIP calls. The HIP kernel launch syntax can be specified as "hipLaunchKernelGGL(KernelName, GridSize, BlockSize, SharedMemorySize, Stream, KernelArguments);". Each of the following elements—KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments—can have the same meaning in the HIP kernel startup syntax as in the CUDA kernel startup syntax (described earlier herein). The SharedMemorySize and Stream arguments may be required in the HIP kernel startup syntax and may be optional in the CUDA kernel startup syntax.

[0255] A section of the HIP source code 2430 can be identical to a section of the depicted CUDA source code 2410, except for a kernel call that causes the MatAdd kernel to run on a device. The MatAdd kernel can be defined in the HIP source code 2430 with the same declaration specifier "__global__" as the MatAdd kernel in the CUDA source code 2410. A kernel call in the HIP source code 2430 can be "hipLaunchKernelGGL(MatAdd, numBlocks, threadsPerBlock, 0, 0, A, B, C);", while a corresponding kernel call in the CUDA source code 2410 might be "MatAdd<<numBlocks, threadsPerBlock> >>(A, B, C);“ is.

[0256] Other implementations are considered and can be executed similarly to the CUDA and HIP implementations described above, such as oneAPI, OpenCL, and other programming platforms. Code can be translated in either direction. For example, CUDA can be translated to HIP, and CUDA can be translated to OpenCL. SnuCL-Tr and CUCL can be used to translate OpenCL to CUDA and CUDA to OpenCL, respectively. Compiled code or intermediate representations (e.g., CUDA PTX code) can also be translated to run on other processor platforms (e.g., AMD or Intel). For example, PTX code can be translated using a translation tool such as ZLUDA to run on Intel or AMD processors.

[0257] One or more of the methods described herein may use a oneAPI programming model. A oneAPI programming model may refer to a programming model for interacting with different compute accelerator architectures. OneAPI may refer to an application programming interface (API) designed to interact with different compute accelerator architectures. A oneAPI programming model may use a DPC++ programming language. A DPC++ programming language may refer to a high-level language for the productivity of data-parallel programming. A DPC++ programming language may be based, at least in part, on C and / or C++ programming languages. A oneAPI programming model may be a programming model such as, among others, the one developed by Intel Corporation in Santa Clara, CA.

[0258] OneAPI and / or oneAPI programming models can be used to interact with various accelerators, GPUs, processors, and / or variants thereof, as well as associated architectures. OneAPI can include a set of libraries that implement various functionalities. OneAPI can include at least one oneAPI DPC++ library, oneAPI library for a mathematical kernel, oneAPI data analysis library, oneAPI library for a deep neural network, oneAPI collection communication library, oneAPI threading component library, oneAPI video processing library, and / or variants thereof.

[0259] A oneAPI-DPC++ 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 variants thereof. OneDPL can implement one or more classes and / or functions from a C++ standard library. OneDPL can implement one or more random number generator functions.

[0260] A oneAPI library for a mathematical kernel, also known as oneMKL, can be a library that implements various optimized and parallelized routines for different mathematical functions and / or operations. OneMKL can implement one or more Basic Linear Algebra Subprograms (BLAS) and / or Linear Algebra Packages (LAPACK) for 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 mathematics (VM) routines for mathematical operations on vectors. OneMKL can implement one or more Fast Fourier Transform (FFT) functions.

[0261] A oneAPI data analysis library, also known as oneDAL, can include a library that implements various data analysis applications and distributed computations. OneDAL can implement various algorithms for preprocessing, transformation, analysis, modeling, validation, and decision-making for data analysis in batch, online, and distributed computing processing modes. OneDAL can implement various C++ and / or Java APIs and various connectors to one or more data sources. OneDAL can implement DPC++ API extensions to a standard C++ interface and enables GPU utilization for various algorithms.

[0262] A oneAPI library for a deep neural network, also known as oneDNN, can include a library that implements various deep learning functions. OneDNN can implement different functions, algorithms, and / or variations thereof for neural networks, machine learning, and deep learning.

[0263] A oneAPI collective communication library, also known as oneCCL, can include a library that implements various applications for deep learning and machine learning workloads. OneCCL can be built on top of lower-level communication middleware, such as, but not limited to, a Message Passing Interface (MPI) and libfabrics. OneCCL can enable a set of specific deep learning optimizations, such as, but not limited to, prioritization, persistent operations, out-of-order execution, and / or variations thereof. OneCCL can implement various CPU and GPU functions.

[0264] A oneAPI Threading Building Blocks library, also known as oneTBB, can include a library that implements various parallelized processes for different applications. OneTBB can be used for task-based, shared parallel programming on a single host. OneTBB can implement generic parallel algorithms. OneTBB can implement concurrently running containers. OneTBB can implement a scalable memory allotment. OneTBB can implement a work-stealing task scheduler. OneTBB can implement low-level synchronization primitives. OneTBB can be compiler-independent and usable on various processors, such as, but not limited to, GPUs, PPUs, CPUs, and / or variants thereof.

[0265] A oneAPI video processing library, also known as oneVPL, can include a library used to accelerate video processing in one or more applications. OneVPL can implement various video decoding, encoding, and processing functions. OneVPL can implement various functions for media pipelines on CPUs, GPUs, and other accelerators. OneVPL can implement device detection and selection in media-oriented and video analytics workloads. OneVPL can implement API primitives for zero-copy buffer sharing.

[0266] A oneAPI programming model can use a DPC++ programming language. A DPC++ programming language can include functionally similar versions of CUDA mechanisms to define device code and distinguish between device code and host code. A DPC++ programming language can include a subset of the functionality of a CUDA programming language. One or more CUDA programming model operations can be executed using a oneAPI programming model with a DPC++ programming language.

[0267] Each application programming interface (API) described here can be compiled by a compiler, interpreter, or other software tool into one or more instructions, operations, or any other signal. Compilation can involve generating one or more machine-executable instructions, operations, or other signals from the source code. An API compiled into one or more instructions, operations, or other signals, when executed, can cause one or more processors to execute, such as processors found, for example, in the Fig. 9-21, or any other logic circuit further described herein, to perform one or more arithmetic operations.

[0268] In at least one embodiment, for example, the translation tools described elsewhere herein may include one or more circuits to translate CUDA code for the dynamic execution of neural network sections into HIP, oneAPI, OpenCL, or any other language used to perform any of the operations described above or elsewhere herein. One or more circuits may be configured by software to translate the CUDA code for the dynamic execution of neural network sections into HIP, oneAPI, OpenCL, or any other language used to perform any of the operations described above or elsewhere in this document. AUTONOMOUS VEHICLE

[0269] Fig. Figure 25 illustrates an example of an autonomous vehicle 2500 according to at least one embodiment. The autonomous vehicle 2500 (hereinafter referred to alternatively as "vehicle 2500") can be a passenger vehicle, such as, but not limited to, a car, a truck, a bus, and / or another type of vehicle that carries one or more passengers. In at least one embodiment, the vehicle 2500 can be a semi-trailer truck used for transporting cargo. The vehicle 2500 can be an aircraft, a robotic vehicle, or another type of vehicle.

[0270] Autonomous vehicles can be described in terms of automation levels 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”) “Taxonomy and Definitions for Terms Related to Driving Automation Systems for On-Road Motor Vehicles” (e.g., Standard No. J3016-201806, published June 15, 2018, Standard No. J3016-201609, published September 30, 2016, and previous and future versions of this standard). In at least one embodiment, the Vehicle 2500 may be capable of operating according to one or more of Levels 1 through 5 of autonomous driving levels. For example, depending on the embodiment, the Vehicle 2500 may be capable of conditional automation (Level 3), high automation (Level 4), and / or full automation (Level 5) in at least one embodiment.

[0271] The vehicle 2500 can include components such as, but not limited to, a chassis, a vehicle body, wheels (e.g., 2, 4, 6, 8, 18, etc.), tires, axles, and other vehicle components. The vehicle 2500 can include a propulsion system 2550, such as an internal combustion engine, a hybrid-electric drive, a fully electric motor, and / or another type of propulsion system. The propulsion system 2550 can be connected to a drivetrain of the vehicle 2500, which may include a transmission to enable the propulsion of the vehicle 2500. The propulsion system 2550 can be controlled in response to receiving signals from a throttle / accelerator pedal(s) 2552.

[0272] A steering system 2554, which may include a steering wheel, is used to steer the vehicle 2500 (e.g., along a desired path or route) when the propulsion system 2550 is in operation (e.g., when the vehicle 2500 is in motion). The steering system 2554 can receive signals from a steering actuator 2556. A steering wheel may be optional for full automation (Level 5) functionality. A brake sensor system 2546 can be used to apply vehicle brakes in response to receiving signals from a brake actuator 2548 and / or brake sensors.

[0273] The controller(s) 2536, which may include one or more system-on-chips (“SoCs”) and / or graphics processing units (“GPU(s)”), can provide signals (e.g., representative of commands) to one or more components and / or systems of the vehicle 2500. For example, the controller(s) 2536 can send signals to actuate vehicle brakes via brake actuator(s) 2548, to actuate the steering system 2554 via steering actuator(s) 2556, and to actuate the propulsion system 2550 via a throttle / accelerator pedal(s) 2552. The control(s) 2536 may include one or more onboard (e.g., integrated) computing devices that process sensor signals and issue operating commands (e.g., signals representing commands) to enable autonomous driving and / or to assist a human driver in driving the vehicle 2500.The controller(s) 2536 can include a first controller for autonomous driving functions, a second controller for functional safety functions, a third controller for artificial intelligence functionality (e.g., machine vision), a fourth controller for infotainment functionality, a fifth controller for redundancy in emergency conditions, and / or other controllers. A single controller can handle two or more of the above functionalities; two or more controllers can handle a single functionality and / or any combination thereof.

[0274] The controller(s) 2536 can provide signals to control one or more components and / or systems of the vehicle 2500 in response to sensor data received from one or more sensors (e.g. sensor inputs). Sensor data can be obtained, for example, from one or more sensors 2558 of the global navigation satellite system (“GNSS”) (e.g., global positioning system sensor(s)), one or more radar sensors 2560, one or more ultrasonic sensors 2562, one or more lidar sensors 2564, one or more sensors 2566 of an inertial measurement unit (“IMU”) (e.g., accelerometers, gyroscopes, a magnetic compass or compasses, magnetometer, etc.), one or more microphones 2596, one or more stereo cameras 2568, one or more wide-angle cameras 2570 (e.g., fisheye cameras), one or more infrared cameras 2572, one or more surround cameras 2574 (e.g.,360-degree cameras), long-range cameras 2598, one or more medium-range cameras 2576, one or more speed sensors 2544 (e.g., to measure the speed of the vehicle 2500), one or more vibration sensors 2542, one or more steering sensors 2540, one or more brake sensors (e.g., as part of the brake sensor system 2546), and / or other sensor types.

[0275] One or more of the controller(s) 2536 can receive inputs (e.g., represented by input data) from an instrument cluster 2532 of the vehicle 2500 and provide outputs (e.g., represented by output data, display data, etc.) via a display 2534, a human-machine interface (“HMI”), an acoustic signal generator, a loudspeaker, and / or other components of the vehicle 2500. Outputs can include information such as, but are not limited to, vehicle speed, velocity, time, map data (e.g., a high-resolution map (not shown)), position data (e.g., the position of the vehicle 2500, as shown on a map), direction, the position of other vehicles (e.g., an occupancy grid), information about objects, and the status of objects as perceived by the controller(s) 2536, etc. For example, the HMI display 2534 can provide information about the presence of one or more objects (e.g.,a road sign, warning sign, traffic light change, etc.) and / or display information about driving maneuvers that the vehicle has performed, is performing, or will perform (e.g., changing lanes now, taking exit 34B in two miles, etc.).

[0276] Each of the components, features and systems of the vehicle 2500 in Fig. 25 can be connected via a bus 2502. The bus 2502 can include a CAN data interface (referred to herein alternatively as a "CAN bus"). A CAN can be a network within the vehicle 2500 that is used to assist in controlling various features and functionalities of the vehicle 2500, such as, but not limited to, actuation of brakes, acceleration, braking, steering, windshield wipers, etc. The bus 2502 can be configured to have dozens or even hundreds of nodes, each with its own unique identifier (e.g., a CAN ID). Data can be read from the bus 2502 to obtain steering wheel angle, ground speed, engine revolutions per minute ("RPMs"), button positions, and / or other vehicle status indicators. The bus 2502 can be a CAN bus that is ASIL B compliant.

[0277] In addition to or as an alternative to CAN, FlexRay and / or Ethernet protocols can be used. There can be any number of buses that comprise the 2502 bus, which can 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 can be used to perform different functions and / or for redundancy. For example, a first bus can be used for collision avoidance functionality, and a second bus can be used for actuation control. Each bus of the 2502 bus can communicate with any of the components of the 2500 vehicle, and two or more buses of the 2502 bus can communicate with corresponding components.Each of any number of system(s) on chip(s) (“SoC(s)”) 2504 (such as, but not limited to, SoC 2504(A) and SoC 2504(B)), each of the controller(s) 2536 and / or each computer within the vehicle can have access to the same input data (e.g. inputs from sensors of the vehicle 2500) and can be connected to a common bus, such as a CAN bus.

[0278] Any number of cameras can be installed according to at least one embodiment at any selection of camera positions and fields of view for the autonomous vehicle 2500 of Fig. 25. Cameras and their respective fields of view can be an exemplary embodiment and are not intended to be limiting. For example, additional and / or alternative cameras may be included and / or cameras may be located in different positions on the vehicle 2500.

[0279] Camera types may include digital cameras that can be adapted for use with components and / or systems of the vehicle 2500. The camera(s) may operate at Automotive Safety Integrity Level (ASIL) B and / or another ASIL. Depending on the configuration, camera types may be capable of any frame rate, such as, but not limited to, 60 frames per second (fps), 1220 fps, 240 fps, etc. Cameras may be equipped with rolling shutter, global shutter, or other shutter types, or a combination thereof. In at least one embodiment, the color filter array may include a red-clear-clear-clear color filter array (“RCCC”), a red-clear-clear-blue color filter array (“RCCB”), a red-blue-green-clear color filter array (“RBGC”), a Foveon X3 color filter array, a Bayer sensor color filter array (“RGGB”), a monochrome sensor color filter array and / or another type of color filter array.Clear pixel cameras, such as cameras with an RCCC, RCCB and / or RBGC color filter array, can be used to increase light sensitivity.

[0280] One or more of the camera(s) can be used to perform functions of advanced driver assistance systems (ADAS) (e.g., as part of a redundant or fail-safe design). For example, a multi-function mono camera can be installed to provide functions including lane departure warning, traffic sign recognition, and intelligent headlight control. One or more of the camera(s) (e.g., all cameras) can simultaneously record and provide image data (e.g., video).

[0281] One or more cameras may be mounted in a bracket, such as, but not limited to, a custom-designed (three-dimensional ("3D") printed) bracket to avoid stray light and reflections from inside the vehicle (e.g., reflections from the dashboard reflected in the windshield mirrors) that could impair the camera's image data acquisition capabilities. With reference to exterior mirror brackets, exterior mirror brackets may be custom 3D printed so that a camera mounting plate conforms to the shape of an exterior mirror. Camera(s) may be integrated into exterior mirrors. For side-view cameras, camera(s) may also be integrated into four pillars at each corner of a cabin.

[0282] Cameras with a field of view that includes portions of the environment in front of the vehicle (e.g., forward-facing cameras) can be used for surround view to help identify forward paths and obstacles, and to help provide, with the help of one or more controllers and / or control SoCs, information critical for generating an occupancy grid and / or determining preferred vehicle paths. Forward-facing cameras can be used to perform many similar ADAS functions as LiDAR, including emergency braking, pedestrian detection, and collision avoidance. Forward-facing cameras can also be used for 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.

[0283] A variety of cameras can be used in a forward-facing configuration, including, for example, a monocular camera platform incorporating a CMOS (Complementary Metal Oxide Semiconductor) color image sensor. A wide-angle camera 2570 can be used to detect objects approaching from the periphery (e.g., pedestrians, crossing traffic, or bicycles). Any number (including zero) of wide-angle cameras 2570 can be mounted on the vehicle 2500. Any number of long-range camera(s) 2598 (e.g., a pair of long-range stereo cameras) can be used for depth-based object detection, particularly for objects for which a neural network has not yet been trained. Long-range camera(s) 2598 can also be used for object detection and classification, as well as basic object tracking.

[0284] Any number of stereo camera(s) 2568 can also be present in a forward-facing configuration. One or more of the stereo camera(s) 2568 can include an integrated control unit comprising 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 vehicle 2500's environment, including a distance estimate for all points in an image. One or more of the stereo camera(s) 2568 can include one or more compact stereo vision sensors, which can include two camera lenses (one each on the left and right) and an image processing chip that measures the distance from the vehicle 2500 to the target object and generates information (e.g.,Metadata) can be used to activate autonomous emergency braking and lane departure warning functions. Other types of stereo camera(s) 2568 can be used in addition to or as an alternative to those described here.

[0285] Cameras with a field of view that includes sections of the environment to the sides of the vehicle 2500 (e.g., side-view cameras) can be used for surround view, providing information used to create and update an occupancy grid and to generate side-impact collision warnings. For example, one or more surround-view cameras 2574 (e.g., four surround-view cameras) could be positioned on the vehicle 2500. The surround-view camera(s) 2574 can include any number and combination of wide-angle cameras, fisheye cameras, 360-degree cameras, and / or similar cameras. For example, four fisheye cameras could be positioned on the front, rear, and sides of the vehicle 2500. The vehicle 2500 can use three surround-view camera(s) 2574 (e.g. left, right and rear) and can use one or more other camera(s) (e.g. a forward-facing camera) as a fourth surround-view camera.

[0286] Cameras with a field of view that includes sections of the environment behind the vehicle 2500 (e.g., rear-view cameras) can be used for parking assistance, surround view, rear collision warnings, and creating and updating an occupancy grid. A wide variety of cameras can be used, including, but not limited to, cameras that can also be used as forward-facing cameras (e.g., long-range cameras 2598 and / or medium-range cameras 2576, stereo cameras 2568, infrared cameras 2572, etc.), as described herein.

[0287] The Vehicle 2500 can include any number of SoCs 2504 or other processors described elsewhere herein, such as, but not limited to, processors and / or components that are compatible with the Fig. Figure 9-21 illustrates and describes the SoC(s) 2504. Each SoC can include central processing units (“CPU(s)”) 2506, graphics processing units (“GPU(s)”) 2508, processor(s) 2510, cache(s) 2512, accelerator(s) 2514, data storage(s) 2516, and / or other components and features not shown. SoC(s) 2504 can be used to control the vehicle 2500 in a variety of platforms and systems. For example, the SoC(s) 2504 can be combined in a system (e.g., the vehicle 2500 system) with a high-resolution (“HD”) map 2522, which can receive map updates and / or updates via the network interface 2524 from one or more servers (not shown).

[0288] The CPU(s) 2506 can include a CPU cluster or CPU complex (hereafter referred to as "CCPLEX"). The CPU(s) 2506 can include multiple cores and / or Level Two ("L2") caches. For example, the CPU(s) 2506 can include eight cores in a coherent multiprocessor configuration. The CPU(s) 2506 can include four dual-core clusters, each cluster having a dedicated L2 cache (e.g., a 2-megabyte (MB) L2 cache). The CPU(s) 2506 (e.g., CCPLEX) can be configured to support concurrent cluster operations, allowing any combination of CPU(s) 2506 clusters to be active at any given time.

[0289] One or more of the CPU(s) 2506 can implement power management capabilities that include one or more of the following features: Individual hardware blocks can be automatically clocked when idle to save dynamic power; each core can be clocked when such core is not actively executing instructions due to the execution of Wait for Interrupt (WFI) / Wait for Event (WFE) instructions; each core can be clocked independently; each core cluster can be clocked independently if all cores can be clocked or power-locked; and / or each core cluster can be power-locked independently if all cores can be power-locked.The CPU(s) 2506 can further implement an extended algorithm for managing power states, where allowed power states and expected wake-up times can be defined, and the hardware / microcode determines which power state is most appropriate for the core, cluster, and CCPLEX. Processing cores can support simplified sequences for entering the power-saving state in software, offloading the work to microcode.

[0290] The GPU(s) 2508 may include an integrated GPU (referred to herein as the "iGPU"). The GPU(s) 2508 may be programmable and efficient for parallel workloads. The GPU(s) 2508 may use an extended Tensor instruction set. The GPU(s) 2508 may include one or more streaming microprocessors, each of which may include a Level One ("L1") cache (e.g., an L1 cache with a minimum storage capacity of 96 KB), and two or more streaming microprocessors may share an L2 cache (e.g., an L2 cache with a storage capacity of 512 KB). The GPU(s) 2508 may include at least eight streaming microprocessors. The GPU(s) 2508 can use one or more application programming interfaces (APIs). The GPU(s) 2508 can use one or more parallel computing platforms and / or programming models (e.g.,(using NVIDIA's CUDA model). Streaming microprocessors can be referred to as streaming multiprocessors (“SMs”), stream processors (“SPs”), stream processing units (“SPUs”), compute units (“CUs”), execution units (“EUs”) and / or slices, where a slice in this context can refer to a portion of processing resources within a processing unit (e.g., 16 cores, a ray tracing unit, a thread director, or a scheduler).

[0291] One or more of the GPU(s) 2508 can be performance-optimized for best performance in automotive and embedded applications. For example, the GPU(s) 2508 could be manufactured on a FinFET (Fin-Fin) field-effect transistor circuit. Each streaming microprocessor can include a number of mixed-precision processing cores partitioned into multiple blocks. For example, 64 PF32 cores and 32 FP64 cores could be partitioned into four processing blocks. Each processing block could be allocated 16 FP32 cores, 8 FP64 cores, 16 INT32 cores, two NVIDIA mixed-precision Tensor cores for deep learning matrix arithmetic, a level-zero instruction cache (“L0”), a scheduler (e.g., warp scheduler) or sequencer, an output unit, and / or a 64 KB register file.Streaming microprocessors can include independent parallel integer and floating-point data paths to provide efficient execution of workloads with a mix of computational and addressing operations. Streaming microprocessors can include independent thread scheduling capabilities to enable finer synchronization and cooperation between parallel threads. Streaming microprocessors can include a combined L1 data cache and shared memory to improve performance while simplifying programming.

[0292] One or more of the GPU(s) 2508 can include high-bandwidth memory (“HBM”) and / or a 16 GB HBM2 memory subsystem to provide a peak memory bandwidth of approximately 900 GB / second in some examples. In addition to or as an alternative to HBM memory, synchronous graphics random-access memory (“SGRAM”) can be used, such as, but not limited to, Graphics Double Data Rate Type Five (“GDDR5”) synchronous random-access memory.

[0293] The GPU(s) 2508 can incorporate a unified memory technology. Support for Address Translation Services (“ATS”) can be used to allow the GPU(s) 2508 to directly access the page tables of the CPU(s) 2506. If a GPU of the memory management unit (“MMU”) of the GPU(s) 2508 fails, an address translation request can be sent to the CPU(s) 2506. In response, two CPUs of the CPU(s) 2506 can search their page tables for a virtual-to-physical mapping for an address and send the translation back to the GPU(s) 2508. The unified memory technology can enable a single, unified virtual address space for the memory of both the CPU(s) 2506 and the GPU(s) 2508, thereby simplifying the programming of the GPU(s) 2508 and the porting of applications to the GPU(s) 2508.

[0294] The GPU(s) 2508 can include any number of access counters that track the frequency of GPU(s) accessing the memory of other processors. The access counter(s) can help ensure that memory pages are moved to the physical memory of the processor that accesses pages most frequently, thereby improving the efficiency of memory areas shared between processors.

[0295] One or more of the SoC(s) 2504 can include any number of caches 2512, including those described herein. For example, the cache(s) 2512 could include a level three (“L3”) cache that is available to both the CPU(s) 2506 and the GPU(s) 2508 (e.g., one connected to both the CPU(s) 2506 and the GPU(s) 2508). The cache(s) 2512 can include a write-back cache capable of tracking the states of rows, such as by using a cache coherence protocol (e.g., MEI, MESI, MSI, etc.). An L3 cache can include 4 MB of memory or more, depending on the implementation, although smaller cache sizes are possible.

[0296] One or more of the SoC(s) 2504 can include one or more Accelerators 2514 (e.g., hardware accelerators, software accelerators, or a combination thereof). The SoC(s) 2504 can include a hardware acceleration cluster, which may contain optimized hardware accelerators and / or large on-chip memory. Large on-chip memory (e.g., 4 MB SRAM) can enable a hardware acceleration cluster to accelerate neural networks and other computations. A hardware acceleration cluster can be used to complement the GPU(s) 2508 and offload some tasks from the GPU(s) 2508 (e.g., to free up more GPU cycles for other tasks). The Accelerator 2514 could be used for targeted workloads (e.g., perception, convolutional neural networks (“CNNs”), reciprocal neural networks (“RNNs”), etc.) that may be stable enough to be amenable to acceleration.A CNN can include a region-based or regional convolutional neural network (“RCNNs”) and fast RCNNs (e.g., as used for object detection) or another type of CNN.

[0297] The Accelerator 2514 (e.g., Hardware Acceleration Cluster) can include one or more Deep Learning Accelerators (“DLAs”). The DLA(s) can include one or more Tensor Processing Units (“TPUs”), which can be configured to provide an additional ten trillion operations per second for deep learning applications and inference, such as the TPU(s) in Fig. 19. The TPUs can be accelerators designed and optimized to perform image processing functions (e.g., for CNNs, RCNNs, etc.). The DLA(s) can further be optimized for a specific set of neural network types and floating-point operations, as well as inference. The design of the DLA(s) can provide more performance per millimeter than a typical general-purpose GPU and typically far exceeds the performance of a CPU. The TPU(s) can perform multiple functions, including a single-instance convolution function that supports, for example, INT8, INT16, and FP16 data types for both features and weights, as well as post-processing functions.The DLA(s) can quickly and efficiently execute neural networks, especially CNNs, on processed or unprocessed data for any of a wide variety of functions, including, for example: a CNN for object identification and recognition using data from camera sensors; a CNN for distance estimation using data from camera sensors; a CNN for emergency vehicle detection and identification using data from microphones; a CNN for facial recognition and vehicle owner identification using data from camera sensors; and / or a CNN for safety and / or security-related events.

[0298] The DLA(s) can perform any function of the GPU(s) 2508, and by using an inference accelerator, a designer can, for example, target either the DLA(s) or the GPU(s) 2508 for any given function. For example, a designer can concentrate the processing of CNNs and floating-point operations on the DLA(s) and leave other functions to the GPU(s) 2508 and / or the accelerator(s) 2514.

[0299] The accelerator(s) 2514 may include a programmable image processing accelerator (“PVA”), which may alternatively be referred to herein as a computer vision accelerator. The PVA may be designed and configured to accelerate computer vision algorithms for an advanced driver assistance system (“ADAS”) 2538, autonomous driving, augmented reality (“AR”) applications, and / or virtual reality (“VR”) applications. The PVA may provide a balance between performance and flexibility. For example, each PVA may include any number of reduced instruction set computer (“RISC”) cores, direct memory access (“DMA”), and / or any number of vector processors.

[0300] RISC cores can interact with image sensors (e.g., image sensors of any camera described herein), image signal processor(s), etc. Each RISC core can include any amount of memory. Depending on the embodiment, RISC cores can use any of a number of protocols. RISC cores can run a real-time operating system (RTOS). RISC cores can be implemented using one or more integrated circuit devices, application-specific integrated circuits (ASICs), and / or memory devices. For example, RISC cores could include an instruction cache and / or tightly coupled RAM.

[0301] A DMA can allow components of the PVA to access system memory independently of the CPU(s) 2506. The DMA can support any number of features used to provide optimization for a PVA, including support for multidimensional addressing and / or circular addressing. The DMA can support up to six or more dimensions of addressing, which may include block width, block height, block depth, horizontal block step, vertical block step, and / or depth step.

[0302] Vector processors can be programmable processors designed to efficiently and flexibly execute programming for computer vision algorithms and provide signal processing capabilities. A PVA can include a PVA core and two partitions of a vector processing subsystem. A PVA core can include a processor subsystem, one or more DMA engines (e.g., two DMA engines), and / or other peripherals. A vector processing subsystem can act as the primary processing engine of a PVA and may include a vector processing unit (VPU), an instruction cache, and / or vector memory (e.g., VMEM). The VPU core may include, but is not limited to, a digital signal processor, such as a single-instruction, multiple-data (SIMD), or very-long instruction word (VLIW) digital signal processor.A combination of SIMD and VLIW can improve throughput and speed.

[0303] Each vector processor can include an instruction cache and can be coupled to dedicated memory. Consequently, each vector processor can be configured to run independently of other vector processors. Vector processors included in a given PVA can be configured to utilize data parallelism. For example, multiple vector processors included in a single PVA can execute a common computer vision algorithm, but on different regions of an image. Vector processors included in a given PVA can simultaneously execute different computer vision algorithms on an image, or even execute different algorithms on sequential images or sections of an image.Among other things, any number of PVAs can be included in a hardware acceleration cluster, and any number of vector processors can be included in each PVA. The PVA can include additional memory with error correction code (ECC) to improve overall system security.

[0304] The Accelerator 2514 can include an on-chip computer vision network and static random-access memory (“SRAM”) to provide high-bandwidth, low-latency SRAM for the Accelerator 2514. The on-chip memory can include a minimum of 4 MB of SRAM, including, for example, eight field-configurable memory blocks accessible to both a PVA and a DLA. Each pair of memory blocks can include an interface for an advanced peripheral bus (“APB”), a configuration circuit, a controller, and a multiplexer. Any type of memory can be used. A PVA and a DLA can access the memory via a backbone that provides high-speed memory access for both. A backbone can include an on-chip computer vision network that connects a PVA and a DLA to the memory (e.g., using APB).

[0305] An on-chip computer vision network can include an interface that, prior to the transmission of any control signal, address, or data, determines that both a PVA and a DLA provide available and valid signals. An interface can provide separate phases and channels for transmitting control signals, addresses, and data, as well as burst-type communication for continuous data transmission. An interface can conform to the standards of the International Organization for Standardization (ISO) 26262 or the International Electrotechnical Commission (IEC) 61508, although other standards and protocols may be used.

[0306] One or more of the SoC(s) 2504 can include a real-time ray-tracing hardware accelerator. The real-time ray-tracing hardware accelerator can be used to quickly and efficiently determine the positions and extents of objects (e.g., within a world model), to generate real-time visualization simulations for radar signal interpretation, sound propagation synthesis and / or analysis, SONAR system simulation, general wave propagation simulation, comparison with LiDAR data for localization purposes, and / or other functions, and / or for other uses.

[0307] The Accelerator 2514 can have a wide range of uses for autonomous driving. A PVA can be used for key processing stages in ADAS and autonomous vehicles. The capabilities of a PVA may be well-suited for algorithmic areas that require predictable processing with low power consumption and low latency. In other words, a PVA can perform well in semi-dense or dense regular computation, even with small datasets that require predictable runtimes with low latency and low power consumption. In the Vehicle 2500, PVAs could be configured to run classic computer vision algorithms, as they can be efficient at object detection and working with integer mathematics. For example, a PVA is used to perform computer stereo vision.In some examples, a semi-global, matching-based algorithm can be used, although this is not a limitation. Level 3-5 autonomous driving applications use real-time motion estimation / stereo matching (e.g., structure from motion, pedestrian detection, lane detection, etc.). A PVA can perform computer stereo vision functions with input from two monocular cameras. A PVA can be used to perform dense optical flow processing. For example, a PVA could process raw radar data (e.g., using a fast 4D Fourier transform) to provide processed radar data. A PVA is used for runtime depth processing, for example, by processing raw data at runtime to provide processed runtime data.

[0308] A DLA (Dynamic Load Alert) can be used to run any type of network to improve control and driving safety, including, for example, a neural network that outputs a confidence measure for each object detection. Confidence can be represented or interpreted as a probability, or as providing a relative "weight" for each detection compared to other detections. A confidence measure allows a system to make further decisions about which detections to consider as true positives rather than false positives. A system can set a confidence threshold and consider only detections that exceed the threshold as true positives. If an automatic emergency braking (AEB) system is used, false positives can cause the vehicle to automatically perform emergency braking, which is obviously undesirable.High-reliability detections can be considered triggers for AEB. A DLA can execute a neural network to reduce the confidence level. A neural network can take as its input at least a subset of parameters, such as, but not limited to, boundary box dimensions, a ground plane estimate obtained (e.g., from another subsystem), an output from IMU sensor(s) 2566 correlated with the vehicle's orientation 2500, a distance, and 3D position estimates of the object obtained from the neural network and / or other sensors (e.g., LIDAR sensor(s) 2564 or RADAR sensor(s) 2560).

[0309] One or more of the SoC(s) 2504 can include data storage (e.g., memory) 2516. The data storage (e.g., memory) 2516 can be on-chip memory of the SoC(s) 2504 capable of storing neural networks to be executed on the GPU(s) 2508 and / or a DLA. The data storage (e.g., memory) 2516 can be large enough to store multiple instances of neural networks for redundancy and security. The data storage (e.g., memory) 2516 can include L2 or L3 cache(s).

[0310] One or more of the SoC(s) 2504 can include any number of processor(s) 2510 (e.g., embedded processors). The processor(s) 2510 can include a boot and power management processor, which may be a dedicated processor and subsystem to handle boot power and management functions and the associated security measures. A boot and power management processor can be part of a boot sequence of SoC(s) 2504 and can provide runtime power management services. A boot power and management processor can provide clock and voltage programming, support for transitions to system power-saving modes, management of thermal and temperature sensors of the SoC(s) 2504, and / or management of power-saving modes of the SoC(s) 2504.Each temperature sensor can be implemented as a ring oscillator whose output frequency is proportional to the temperature, and the SoC(s) 2504 can use ring oscillators to detect temperatures of the CPU(s) 2506, the GPU(s) 2508, and / or the accelerator(s) 2514. If temperatures are detected exceeding a threshold, a boot and power management processor can invoke a temperature fault routine and put the SoC(s) 2504 into a power-saving mode and / or put the vehicle 2500 into a chauffeur-to-safety stop mode (e.g., bring the vehicle 2500 to a safe stop).

[0311] The 2510 processor(s) can further include a set of embedded processors that can serve as an audio processing engine, which can be an audio subsystem providing full hardware support for multi-channel audio across multiple interfaces and a wide and flexible range of audio I / O interfaces. An audio processing engine is a dedicated processor core with a digital signal processor and dedicated RAM.

[0312] The 2510 processor(s) may also include a continuously active processor engine, which can provide the necessary hardware features to support low-power use cases for sensor management and wake-up. A continuously active processor engine may include a processor core, tightly coupled RAM, supporting peripherals (such as timers and interrupt controllers), various I / O control peripherals, and routing logic.

[0313] The 2510 processor(s) can also include a security cluster engine, which may contain a dedicated processor subsystem to handle security management for automotive applications. A security cluster engine can include two or more processor cores, tightly coupled RAM, supporting peripherals (e.g., timers, interrupt control, etc.), and / or routing logic. In a security mode, two or more cores can operate in lockstep mode, functioning as a single core with comparison logic to detect differences between their operations. The 2510 processor(s) can also include a real-time camera engine, which may contain a dedicated processor subsystem to handle real-time camera management.The 2510 processor(s) may also include a high dynamic range signal processor, which may include an image signal processor, which is a hardware engine that is part of a camera processing pipeline.

[0314] The 2510 processor(s) may include a video image compositor, which may be a processing block (e.g., implemented on a microprocessor) that implements video post-processing functions required by a video playback application to produce a final image for a player window. A video image compositor may perform lens distortion correction on the 2570 wide-angle camera(s), the 2574 omnidirectional camera(s), and / or the indoor surveillance camera sensor(s). The indoor surveillance camera sensor(s) may preferably be monitored by a neural network running on a separate instance of the 2504 SoC, configured to identify indoor events and respond accordingly.An in-vehicle system can lip-read to activate cellular services and make phone calls, dictate emails, change a vehicle's destination, activate or change the vehicle's infotainment system and settings, or provide voice-activated web browsing. Certain functions may be available to a driver when a vehicle is operating in autonomous mode and may be disabled otherwise.

[0315] A video image compositor can include enhanced temporal noise reduction for both spatial and temporal noise reduction. For example, if there is motion in a video, the noise reduction appropriately weights spatial information, thereby reducing the weight of information provided by neighboring frames. If a frame or portion of a frame does not contain motion, the temporal noise reduction performed by the video image compositor can use information from a previous frame to reduce noise in the current frame.

[0316] A video image compositor can also be designed to perform stereo correction on input stereo lens images. Furthermore, a video image compositor can be used for user interface composition when an operating system desktop is used, and GPU(s) 2508 may not be required to continuously render new surfaces. If GPU(s) 2508 are powered on and active for 3D rendering, a video image compositor can be used to offload the GPU(s) 2508, thus improving performance and responsiveness.

[0317] One or more SoCs of the SoC(s) 2504 may further include a serial camera interface of a mobile industrial processor interface (“MIPI”) for receiving video and camera inputs, a high-speed interface, and / or a video input block, which can be used for a camera and associated pixel input functions. One or more of the SoC(s) 2504 may further include input / output controller(s) that can be controlled by software and used to receive I / O signals that may not be assigned to a specific function.

[0318] One or more SoCs of the SoC(s) 2504 can further include a wide range of peripheral interfaces to enable communication with peripheral devices, audio encoders / decoders ("codecs"), power management, and / or other devices. The SoC(s) 2504 can be used to process data from cameras (e.g., connected via Gigabit Multimedia Serial Link and Ethernet channels), sensors (e.g., LiDAR sensor(s) 2564, RADAR sensor(s) 2560, etc., which may be connected via Ethernet channels), data from the bus 2502 (e.g., vehicle speed 2500, steering wheel position, etc.), data from GNSS sensor(s) 2558 (e.g., connected via an Ethernet bus or a CAN bus), etc. One or more SoCs of the SoC(s) 2504 can further include dedicated high-performance mass storage controllers, which may include their own DMA engines and which can be used to offload routine data management tasks from the CPU(s) 2506.

[0319] The 2504 SoC(s) can be an end-to-end platform with a flexible architecture spanning automation levels 3-5, providing a comprehensive functional safety architecture that employs and efficiently utilizes machine vision and ADAS techniques for diversity and redundancy, and provides a platform for a flexible, reliable driving software stack along with deep learning tools. The 2504 SoC(s) can be faster, more reliable, and even more energy-efficient and compact than conventional systems. For example, the 2514 accelerator(s), when combined with the 2506 CPU(s), 2508 GPU(s), and 2516 data storage(s), can provide a fast, efficient platform for autonomous vehicles at levels 3-5.

[0320] Computer vision algorithms can be run on CPUs, which can be designed using a high-level programming language, such as, but not limited to, C, to execute a wide variety of processing algorithms over a wide variety of visual data. However, CPUs are often unable to meet the performance requirements of many computer vision applications, such as, but not limited to, those related to execution time and power consumption. Many CPUs are unable to execute complex object recognition algorithms in real time, which is used in in-vehicle ADAS applications and in practical Level 3-5 autonomous vehicles.

[0321] The embodiments described herein allow multiple neural networks to run simultaneously and / or sequentially and to combine their results to enable autonomous driving functionality at levels 3-5. For example, a CNN running on a DLA or a discrete GPU (e.g., GPU(s) 2520) can include text and word recognition, enabling the reading and understanding of traffic signs, including signs for which a neural network has not been specifically trained. Furthermore, a DLA can include a neural network capable of identifying, interpreting, and providing a semantic understanding of a sign and passing this semantic understanding to path planning modules running on a CPU complex.

[0322] Multiple neural networks can run simultaneously, such as when driving at levels 3, 4, or 5. For example, a warning sign stating "Warning: Flashing lights indicate icy conditions" can be interpreted independently or joi...

Claims

Processor comprising: one or more circuits to cause one or more sections of one or more neural networks to be dynamically executed by two or more processors, at least partially, based on one or more estimated performance metrics of the one or more sections. Processor according to claim 1, wherein the one or more neural networks are an LLM, Large Language Model, and the one or more sections correspond to different agent application actions. Processor according to claim 1 or 2, wherein one of the one or more sections is a padding section that generates respective intermediate states of one or more input tokens, and another of the one or more sections is a decoding section that generates one or more output tokens. Processor according to one of the preceding claims, wherein the one or more estimated performance metrics are determined at least partially based on a transfer time of one or more key-value cache values, KV cache values, between different of the two or more processors. Processor according to one of the preceding claims, wherein the two or more processors are located at different locations in a hierarchy of computing resources used to execute the one or more neural networks, and wherein the different locations in the hierarchy of computing resources are selected to estimate the one or more performance metrics according to a dynamically determined scheduling policy. Processor according to one of the preceding claims, wherein the one or more estimated performance metrics are determined at least partially based on utilization information of the two or more processors. Processor according to one of the preceding claims, wherein a first and a second section of the one or more sections are executed on one of the two or more processors to generate a first inference, and the first and the second section are executed on different of the two or more processors to generate a second inference. Method encompassing: Causing one or more sections of one or more neural networks to be dynamically executed by two or more processors, at least partially, based on one or more estimated performance metrics of the one or more sections. Method according to claim 8, wherein the one or more neural networks are an LLM, Large Language Model, and the one or more sections correspond to different agent application actions. Method according to claim 8 or 9, wherein one of the one or more sections is a padding section that generates respective intermediate states of one or more input tokens, and another of the one or more sections is a decoding section that generates one or more output tokens. Method according to one of claims 8 to 10, wherein the one or more estimated performance metrics are determined at least partially based on a transmission time of one or more key-value cache values, KV cache values, between different of the two or more processors. Method according to any one of claims 8 to 11, wherein the two or more processors are located at different locations in a hierarchy of computing resources used to execute the one or more neural networks, and wherein the different locations in the hierarchy of computing resources are selected to estimate the one or more performance metrics according to a dynamically determined scheduling policy. Method according to one of claims 8 to 12, wherein the one or more estimated performance metrics are determined at least partially based on utilization information of the two or more processors. Method according to any one of claims 8 to 13, wherein a first and a second section of the one or more sections are executed on one of the two or more processors to generate a first inference, and the first and the second section are executed on different of the two or more processors to generate a second inference. System comprising: one or more processors to cause one or more sections of one or more neural networks to be dynamically executed by two or more processors, at least partially, based on one or more estimated performance metrics of the one or more sections; and one or more memories to store weights of the one or more neural networks. System according to claim 15, wherein the one or more neural networks are an LLM, Large Language Model, and the one or more sections correspond to different agent application actions. System according to claim 15 or 16, wherein one of the one or more sections is a padding section that generates respective intermediate states of one or more input tokens, and another of the one or more sections is a decoding section that generates one or more output tokens. System according to one of claims 15 to 17, wherein the one or more estimated performance metrics are determined at least partially based on a transfer time of one or more key-value cache values, KV cache values, between different of the two or more processors. System according to one of claims 15 to 18, wherein the two or more processors are located at different locations in a hierarchy of computing resources used to execute the one or more neural networks, and wherein the different locations in the hierarchy of computing resources are selected to estimate the one or more performance metrics according to a dynamically determined scheduling policy. System according to one of claims 15 to 19, wherein the one or more estimated performance metrics are determined at least partially based on utilization information of the two or more processors.

Images