Tile based programming model
Patent Information
- Application Number
- PCT/CN2025/082799
- Authority / Receiving Office
- WO · WO
- Patent Type
- Applications
- Current Assignee / Owner
- Filing Date
- 2025-03-16
- Publication Date
- 2026-09-24
Smart Images

Figure CN2025082799_24092026_PF_FP_ABST
Abstract
Description
TILE BASED PROGRAMMING MODELTECHNICAL FIELD
[0001] At least one embodiment pertains to systems and methods for compiling and / or performing GPU code using a tile-based programming model.BACKGROUND
[0002] Code portability is important to ensure that code written for an earlier version of hardware is usable by a newer version of hardware. Current versions of portable GPU code may include first programming code in a general programming language (e.g., C++) and then converting the general code into GPU-compatible code to use the GPU.BRIEF DESCRIPTION OF DRAWINGS
[0003] FIG. 1 illustrates an example of the architecture of a Tile IR compiler and toolchain components, according to at least one embodiment;
[0004] FIG. 2 illustrates an example including a processor and modules, in accordance with at least one embodiment, according to at least one embodiment;
[0005] FIG. 3 is a block diagram illustrating a driver and / or runtime including one or more libraries to provide one or more application programming interfaces (APIs) , according to at least one embodiment;
[0006] FIG. 4 illustrates an example data center system, in accordance with at least one embodiment;
[0007] FIG. 5 illustrates an system-on-a-chip (SOC) , in accordance with at least one embodiment;
[0008] FIG. 6A illustrates a parallel processor, in accordance with at least one embodiment;
[0009] FIG. 6B illustrates a processing cluster, in accordance with at least one embodiment;
[0010] FIG. 6C illustrates a graphics multiprocessor, in accordance with at least one embodiment;
[0011] FIG. 7 illustrates an accelerator processor, in accordance with at least one embodiment;
[0012] FIG. 8A illustrate a central processing unit, in accordance with at least one embodiment;
[0013] FIG. 8B illustrates a core of central processing unit in FIG. 8A, in accordance with at least one embodiment;
[0014] FIG. 9 illustrates another accelerator processor, in accordance with at least one embodiment;
[0015] FIG. 10 illustrates a neuromorphic processor, in accordance with at least one embodiment;
[0016] FIG. 11 illustrates a supercomputer, in accordance with at least one embodiment;
[0017] FIG. 12 illustrates another accelerator processor, in accordance with at least one embodiment;
[0018] FIG. 13 illustrates another processor, in accordance with at least one embodiment;
[0019] FIG. 14 illustrates another accelerator processor, in accordance with at least one embodiment;
[0020] FIG. 15 illustrates a tensor processing unit, in accordance with at least one embodiment;
[0021] FIG. 16 illustrates a RISC-V-compatible processor, in accordance with at least one embodiment;
[0022] FIGs. 17A and 17B illustrate a language processing unit, in accordance with at least one embodiment;
[0023] FIG. 18 illustrates a software stack of a programming platform, in accordance with at least one embodiment;
[0024] FIG. 19 illustrates software that is supported by a programming platform, in accordance with at least one embodiment;
[0025] FIG. 20 illustrates compiling code to execute on programming platforms of FIG. 19, in accordance with at least one embodiment;
[0026] FIG. 21 illustrates an example of an autonomous vehicle and its system architecture, in accordance with at least one embodiment;
[0027] FIG. 22A illustrates inference and / or training logic, in accordance with at least one embodiment;
[0028] FIG. 22B illustrates inference and / or training logic, in accordance with at least one embodiment; and
[0029] FIG. 22C illustrates training and deployment of a neural network, in accordance with at least one embodiment.DETAILED DESCRIPTION
[0030] In the following description, numerous specific details are set forth to provide a more thorough understanding of at least one embodiment. However, it will be apparent to one skilled in the art that the inventive concepts may be practiced without one or more of these specific details.
[0031] In an embodiment, a compiler supports a programming model in which application data is divided into tiles and that maps those tiles to threads to be performed by the GPU, taking into account the specific architectural features of the GPU. Because this mapping occurs at a higher level, new GPU hardware can be used with the same code by modifying how the compiler maps to the new hardware instead of having to modify application code to be aware of the new GPU features or use low-level GPU instructions specific to that GPU. In at least one embodiment, a compiler maps the application data directly to threads usable by the GPU driver, without the need of converting to lower-level languages.
[0032] In some embodiments, SIMT (Single Instruction Multiple Thread) programming model is used GPU programming, and within the CUDA ecosystem. SIMT is designed to leverage the parallel processing capabilities of GPUs by executing the same instruction across multiple threads simultaneously.
[0033] The SIMT model begins with an array of data that needs processing. This data is divided into smaller chunks, known as tiles, which are then organized into a grid of thread blocks. Each thread block contains multiple threads, and the data within each block is mapped onto these threads programmatically. This mapping determines how the data is distributed across the threads for parallel execution.
[0034] Once the data is divided and mapped onto threads, the SIMT model executes the same program across all threads in parallel. This parallelism is a property of the SIMT model, allowing for efficient processing of large datasets. Each thread operates on individual elements of data, typically indexed by thread and block identifiers. This indexing allows threads to access specific data elements and perform operations independently.
[0035] In the SIMT model, each thread performs operations on individual data elements. For example, a thread might calculate the greyscale value of a pixel by accessing its RGB components and applying a formula. These operations are performed independently by each thread, contributing to the overall parallel processing capability of the GPU. The model's design ensures that threads can efficiently handle tasks that require element-wise operations.
[0036] A CUDA Tile Programming Model is designed to address the complexities and limitations of previous models, such as the SIMT (Single Instruction Multiple Thread) model. The CUDA Tile Programming Model simplifies the process of programming tensor cores by introducing higher-level abstractions. Unlike the SIMT model, which requires programmers to map data onto threads manually, the Tile model operates on whole arrays or tensors at a time. This approach reduces the complexity involved in programming, making it more accessible to developers who may not be familiar with the intricacies of GPU architecture. In another aspect, the CUDA Tile programming model may rely on dynamically determined features of GPU hardware to map data onto threads in a way that improves or maximizes utilization of the determined features.
[0037] One of the challenges addressed by the CUDA Tile model is the breaking of GPU generational compatibility due to changes in tensor core hardware design. The Tile model restores this compatibility by abstracting the architecture-specific details, allowing programs to run on new hardware without modification. This is achieved through a compiler that maps data tensors onto the threads of the tensor cores, ensuring that old code remains functional on new GPU architectures.
[0038] The CUDA Tile model is designed to be targetable from various programming languages, with a particular emphasis on Python. Python's array-based programming style, exemplified by libraries like NumPy, aligns well with the Tile model's approach to data handling. These techniques further extend into other dynamically-typed languages, such as JavaScript, Ruby, and PHP.
[0039] The Tile model offers a simpler yet high-performing interface for tensor core programming. It defines a stable abstraction for tensor cores and related hardware, allowing developers to achieve high performance without delving into low-level architecture-specific details. This is particularly important as precision scheduling and data management become increasingly critical with each GPU generation. By simplifying these aspects, CUDA Tile makes tensor cores more accessible to all GPU programmers, reducing the burden associated with using tools like CUTLASS and PTX.
[0040] One of the motivations for developing the CUDA Tile Programming Model was to address the issue of GPU generational compatibility, particularly with the evolving design of tensor cores. As new GPU architectures emerged, the hardware design for tensor cores changed significantly, breaking the compatibility of programs targeting these cores. This posed a challenge to the core precept of CUDA, which is that old code should run on new hardware. By providing a higher-level abstraction for tensor core programming, CUDA Tile aims to restore this compatibility guarantee, allowing programs to use accelerated tensor core operations without needing to describe architecture-specific mappings.
[0041] Another motivation is to simplify the programming interface for tensor cores, making it more accessible to a broader range of developers. CUDA Tile an interface that operates on data tensors instead of individual data elements. This approach reduces the complexity of programming tensor cores, making them more accessible to all GPU programmers, and aligns with array-based models familiar to Python developers.
[0042] The CUDA Tile Programming Model allows for a better target for compiler writers and support a wider range of programming languages. By defining CUDA Tile as an MLIR dialect, it is a viable target for both compilers and application developers in different languages. This allows for the integration of CUDA Tile with various front-ends, including C++ and Python.
[0043] The CUDA Tile Programming Model introduces a higher-level abstraction that operates on whole arrays or tensors, rather than individual data elements. This shift from the traditional SIMT (Single Instruction Multiple Thread) model reduces the complexity of programming, making it more accessible to developers. By focusing on tiles of data, the model allows programmers to work with larger data structures in a more intuitive manner, aligning with array-based programming paradigms familiar to many developers, especially those using Python.
[0044] The CUDA Tile Programming Model addresses the issue of GPU generational compatibility, particularly with the evolving design of tensor cores through the use of a compiler that maps data tensors onto the threads of the tensor cores, ensuring that old code remains functional on new GPU architectures. The Tile model restores this compatibility by abstracting architecture-specific details, allowing programs to run on new hardware without modification.
[0045] The Tile model offers a simpler yet high-performing interface for tensor core programming. It defines a stable abstraction for tensor cores and related hardware, allowing developers to achieve high performance without delving into low-level architecture-specific details. This is particularly important as precision scheduling and data management become increasingly critical with each GPU generation. By simplifying these aspects, CUDA Tile makes tensor cores more accessible to all GPU programmers, reducing the burden associated with using tools like CUTLASS and PTX.
[0046] The CUDA Tile Programming Model simplifies the programming process by dropping the second step of mapping data onto threads used in the SIMT model, leaving this task to the compiler. This change addresses several challenges and opportunities in GPU programming. In the SIMT model, programmers are required to manually map data onto threads within each block, a process that can be complex and error-prone. The CUDA Tile model simplifies this by allowing programmers to focus on dividing data into tiles, while the compiler automatically handles the mapping of these tiles onto internal threads.
[0047] By abstracting the thread mapping process, the CUDA Tile model restores inter-generational compatibility. Programs written using the Tile model can run on new GPU architectures without modification, as the compiler adapts the thread mapping to suit the specific hardware design, which ensures that old code remains functional on new hardware.
[0048] The compiler's role in mapping data onto threads allows for more sophisticated performance optimizations that are tailored to the specific GPU architecture. This is particularly important as precision scheduling and data management become increasingly critical with each GPU generation. The compiler can make informed decisions about how to best utilize the available resources, leading to improved performance without requiring developers to have deep knowledge of the hardware.
[0049] The CUDA Tile Programming Model performs GPU programming by mapping execution blocks to data tiles, which offers a more efficient and flexible way to handle data parallelism. In the Tile model, the application defines the mapping of execution blocks to data tiles, which eliminates the need for developers to manually map data onto threads, leaving this task to the compiler.
[0050] In the CUDA Tile model, the application defines how data is divided into tiles, similar to the SIMT model where a block can operate on one or multiple tiles. This flexibility allows developers to choose tile sizes based on execution resources, ensuring that they fit within shared memory or cache and align with tensor core MMA requirements. The tile size is a characteristic size in the problem and may vary between GPU architectures, providing adaptability across different hardware configurations.
[0051] The mapping of execution blocks to data tiles in the CUDA Tile model is designed to be architecture-agnostic, allowing for integration with various GPU architectures. This is achieved by abstracting the hardware-specific details and focusing on the logical representation of data tiles. The compiler automatically maps data onto internal threads based on the defined tile sizes. This automation reduces the complexity of programming and ensures that the application can leverage the full potential of the GPU's parallel processing capabilities.
[0052] The hierarchy of execution models in the CUDA Tile Programming Model is designed to streamline and optimize GPU programming by organizing operations at different levels: thread-level, block-level, and grid-level. This hierarchical approach allows developers to manage data parallelism more effectively, leveraging the strengths of each level to achieve high performance and flexibility.
[0053] At the thread-level, the execution model focuses on individual threads operating on specific elements of data. This is akin to the traditional SIMT (Single Instruction Multiple Thread) model, where each thread is responsible for a particular piece of data. The thread-level model provides fine-grained control over data operations, allowing developers to optimize performance by managing thread-specific tasks. However, this level of control can be complex and requires detailed knowledge of the GPU architecture to achieve optimal results.
[0054] The block-level model introduces a higher level of abstraction by grouping threads into blocks. In this model, the application maps data onto blocks, which are then managed by the compiler to optimize thread execution. This approach simplifies the programming process by reducing the need for manual thread management, while still allowing for efficient data processing. Block-level models are particularly useful for operations that require synchronization and shared memory access, as they provide a structured way to manage these resources within a block.
[0055] At the grid-level, the execution model operates on a larger scale, managing multiple blocks across the GPU. Grid-level models leave all parallelism to the compiler, allowing it to map data onto both blocks and threads automatically. This level of abstraction is ideal for domain-specific languages where the best data decomposition strategies are known, enabling developers to focus on high-level programming without worrying about the intricacies of thread and block management. Grid-level models are particularly effective for applications that perform per-element operations, as they can leverage the full power of the GPU's parallel processing capabilities.
[0056] The Tile IR Stack is a sophisticated compilation framework designed to enhance the CUDA programming model by introducing a tile-based approach to GPU programming. This stack is composed of several components, each playing a crucial role in transforming high-level tile-based programs into executable code optimized for NVIDIA GPUs. The stack begins with the Public CUDA Tile dialect, which serves as the stable and documented interface for developers. This dialect is designed to be targetable from various programming languages, including Python, making it accessible to a wide range of developers. It abstracts the complexities of tensor core programming, allowing developers to focus on high-level operations without worrying about the underlying hardware specifics.
[0057] At the heart of the Tile IR Stack are the TileAA and TileAS dialects, which are responsible for the critical task of compiling tile-based programs into thread-level representations. These dialects leverage NVIDIA's proprietary technologies, such as CuTe and NVVM, to ensure that the compiled code is both stable and portable across different GPU architectures. The compilation process involves mapping data tiles onto internal threads, optimizing the execution for the specific hardware configuration.
[0058] The Tile IR Stack also integrates with NVIDIA's existing toolchain, including JIT compilation, debugging, profiling, and autotuning capabilities. This integration ensures that tile-based applications can benefit from the full suite of development and deployment tools available in the CUDA ecosystem. The stack supports both offline and JIT compilation paths, allowing developers to choose the most suitable approach for their applications. Additionally, the stack is designed to accommodate third-party compilers and domain-specific languages, enabling a diverse range of applications to leverage the power of NVIDIA GPUs.
[0059] Finally, the Tile IR Stack is built on top of the LLVM / PTX Compiler Stack, which provides the foundational infrastructure for code generation and optimization. This integration allows the Tile IR Stack to target any dialect, ensuring that it can adapt to future hardware developments and maintain its relevance in the rapidly evolving landscape of GPU technology.
[0060] The CUDA Tile Programming Model is integrated into the existing CUDA toolchain through the use of MLIR (Multi-Level Intermediate Representation) layers. This integration allows for a smooth transition from traditional SIMT programming models to the tile-based approach, without disrupting the existing ecosystem. The integration with MLIR layers provides robust compiler support for the CUDA Tile Programming Model. The MLIR framework allows for the definition of custom dialects, which can be used to represent the tile-based operations at a higher level of abstraction. This enables the compiler to optimize the mapping of data tiles onto internal threads, ensuring efficient execution on tensor cores. The use of MLIR also facilitates the development of domain-specific languages (DSLs) that can target the CUDA Tile model, making it accessible to a wider range of programming languages and environments.
[0061] One of the benefits of integrating the CUDA Tile model with MLIR layers is the ability to provide a stable and portable representation of tile-based programs. The MLIR framework allows for the definition of a stable bytecode representation, which can be used to ensure compatibility across different GPU architectures. By abstracting away hardware-specific details, the CUDA Tile model can adapt to future changes in GPU architecture without requiring modifications to existing code.
[0062] The integration with MLIR layers also enables support for third-party DSLs, allowing developers to leverage the CUDA Tile model in conjunction with other programming frameworks. This is achieved through the use of adaptors that can translate third-party IRs into the CUDA Tile dialect, ensuring compatibility and interoperability
[0063] The compilation process from Tile to Thread begins with the high-level representation of data as tiles, which are essentially arrays or tensors that encapsulate the data to be processed. The Tile model abstracts away the complexities of thread-level programming by allowing developers to focus on the logical structure of their data, rather than the intricate details of thread management.
[0064] Once the data is represented as tiles, the compilation process involves mapping these tiles onto the internal threads of the GPU. TileAA and TileAS dialects serve as the intermediary layers that translate the high-level tile operations into thread-level instructions. These dialects are built on NVIDIA's CuTe and NVVM technologies, which provide the necessary infrastructure to ensure that the compiled code is both stable and portable across different GPU architectures. The compiler automatically handles the mapping of tiles onto threads, optimizing the execution based on the specific hardware configuration. This automation not only simplifies the programming model but also enhances performance by leveraging the full capabilities of the GPU.
[0065] The next stage in the compilation process involves the integration with NVIDIA's existing toolchain, including JIT compilation, debugging, profiling, and autotuning capabilities. This integration ensures that tile-based applications can benefit from the full suite of development and deployment tools available in the CUDA ecosystem. The compiled code is packaged into a format that can be easily loaded and executed on the GPU, with the compiler making informed decisions about resource allocation and scheduling. This integration is a advantage of the CUDA Tile model, as it allows developers to focus on high-level programming while the compiler takes care of the low-level details.
[0066] Finally, the compiled code is executed on the GPU, with the tile operations being mapped onto the tensor cores for optimal performance. The compiler ensures that the data is efficiently loaded into shared memory or cache, aligned with the tensor core MMA requirements. This alignment is crucial for achieving high performance, as it allows the GPU to process large amounts of data in parallel, leveraging the power of the tensor cores.
[0067] The End-to-End Tile IR Compilation Stack is a comprehensive framework designed to facilitate the efficient compilation and execution of tile-based programs on NVIDIA GPUs. This stack is built upon several components, each contributing to the integration of tile-based programming models with the existing CUDA ecosystem. The Public CUDA Tile dialect, serves as the stable and documented interface for developers. This dialect is designed to be targetable from various programming languages, including Python, making it accessible to a wide range of developers. It abstracts the complexities of tensor core programming, allowing developers to focus on high-level operations without worrying about the underlying hardware specifics.
[0068] The Tile IR Stack also integrates with NVIDIA's existing toolchain, including JIT compilation, debugging, profiling, and autotuning capabilities. This integration ensures that tile-based applications can benefit from the full suite of development and deployment tools available in the CUDA ecosystem. The stack supports both offline and JIT compilation paths, allowing developers to choose the most suitable approach for their applications. Additionally, the stack is designed to accommodate third-party compilers and domain-specific languages, enabling a diverse range of applications to leverage the power of NVIDIA GPUs.
[0069] The Tile IR Stack is built on top of the LLVM / PTX Compiler Stack, which provides the foundational infrastructure for code generation and optimization. This integration allows the Tile IR Stack to target any dialect, ensuring that it can adapt to future hardware developments and maintain its relevance in the rapidly evolving landscape of GPU technology.
[0070] Finally, the Tile IR Stack incorporates technologies, such as CuTe and NVVM, to ensure that the compiled code is both stable and portable across different GPU architectures. The compilation process involves mapping data tiles onto internal threads, optimizing the execution for the specific hardware configuration.
[0071] The integration of third-party compilers and domain-specific languages (DSLs) with the CUDA Tile Programming Model is facilitated through the use of the Multi-Level Intermediate Representation (MLIR) framework, which serves as a bridge between various programming models and the CUDA Tile IR stack. By leveraging MLIR, third-party compilers can target the CUDA Tile dialect, enabling them to utilize NVIDIA's tensor cores and other GPU features without being constrained by architecture-specific limitations. This approach not only enhances the portability of third-party applications but also ensures that they can benefit from the performance optimizations inherent in the CUDA Tile model.
[0072] One of the components of this integration is the development of adaptors that translate third-party IRs into the CUDA Tile dialect. These adaptors are designed to be open-source, allowing developers to incrementally port their applications to the CUDA Tile model while maintaining compatibility with existing codebases. This incremental porting process is beneficial for applications that have been built using popular frameworks like PyTorch or TensorFlow, as it allows them to gradually adopt the CUDA Tile model without requiring a complete overhaul of their existing infrastructure.
[0073] Furthermore, the integration with third-party compilers is supported by NVIDIA's toolchain, which includes JIT compilation, debugging, profiling, and autotuning capabilities. The integration of these tools with third-party compilers allows developers to optimize their applications for NVIDIA GPUs, taking advantage of features like tensor core acceleration and warp specialization.
[0074] The adaptation of Triton IR to CUDA Tile IR involves a transformation that leverages NVIDIA's technologies to enhance the programming model for tensor cores. One of the aspects of this adaptation is the definition of CUDA-Tile as an MLIR dialect, which serves as an extension to the widely used LLVM intermediate representation. This choice enables CUDA-Tile to be a target for compiler writers, as it includes front-ends in C++, Python, and other languages, making tensor core programming more accessible to developers across different platforms. The integration of CUDA-Tile into the CUDA ecosystem allows for interoperation with existing CUDA SIMT code, enabling more general parallel programs to be written. This is particularly important for applications that require both tensor-based and SIMT-based operations, as it allows for a fusion of different programming models within a single kernel.
[0075] The adaptation process also involves the development of a stable bytecode representation for CUDA-Tile programs, which is crucial for maintaining the portability guarantee of the CUDA ecosystem. Unlike other MLIR compilation stacks such as Triton, which do not define a stable bytecode, CUDA-Tile ensures that its programs can be executed on all future NVIDIA GPUs. This stability is achieved through the use of technologies such as CuTe and NVVM, which provide a robust compilation path that can adapt to future hardware developments. The integration of these technologies into the CUDA Tile IR stack allows for efficient mapping of data tiles onto internal threads, optimizing execution for specific hardware configurations.
[0076] Finally, the adaptation of Triton IR to CUDA Tile IR is designed to support third-party compilers and domain-specific languages, enabling a diverse range of applications to leverage the power of NVIDIA GPUs. This flexibility is achieved through the use of adaptors that can translate third-party IRs into the CUDA Tile dialect, ensuring compatibility and interoperability.
[0077] The End-to-End Tile IR Compilation Stack is a comprehensive framework designed to facilitate the efficient compilation and execution of tile-based programs on NVIDIA GPUs. This stack is built upon several components, each contributing to the integration of tile-based programming models with the existing CUDA ecosystem. The Public CUDA Tile dialect serves as the stable and documented interface for developers. This dialect is designed to be targetable from various programming languages, including Python, making it accessible to a wide range of developers. It abstracts the complexities of tensor core programming, allowing developers to focus on high-level operations without worrying about the underlying hardware specifics.
[0078] The Tile IR Stack integrates with NVIDIA's existing toolchain, including JIT compilation, debugging, profiling, and autotuning capabilities. This integration ensures that tile-based applications can benefit from the full suite of development and deployment tools available in the CUDA ecosystem. The stack supports both offline and JIT compilation paths, allowing developers to choose the most suitable approach for their applications. Additionally, the stack is designed to accommodate third-party compilers and domain-specific languages, enabling a diverse range of applications to leverage the power of NVIDIA GPUs.
[0079] The Tile IR Stack is built on top of the LLVM / PTX Compiler Stack, which provides the foundational infrastructure for code generation and optimization. This integration allows the Tile IR Stack to target any dialect, ensuring that it can adapt to future hardware developments and maintain its relevance in the rapidly evolving landscape of GPU technology.
[0080] Finally, the Tile IR Stack incorporates NVIDIA's proprietary technologies, such as CuTe and NVVM, to ensure that the compiled code is both stable and portable across different GPU architectures. The compilation process involves mapping data tiles onto internal threads, optimizing the execution for the specific hardware configuration. This approach not only simplifies the programming model but also restores the inter-generational compatibility guarantee that is a hallmark of the CUDA ecosystem.
[0081] The public CUDA Tile is designed to provide a stable abstraction for tensor core programming, allowing developers to write code that remains compatible across different generations of NVIDIA GPUs. The stability is achieved by abstracting away architecture-specific details that may change with each new GPU release, such as thread mapping and memory management. By focusing on data tensors rather than individual data elements, the CUDA Tile dialect enables developers to write high-level code without being constrained by the underlying hardware specifics.
[0082] The dialect is documented to ensure that developers can understand and utilize its features effectively. This includes detailed explanations of the operations and constructs available within the dialect, as well as guidelines on how to integrate it with existing CUDA tools and libraries. The documentation also covers the integration of the dialect with various programming languages, including Python, which is a major area of growth for CUDA.
[0083] The public CUDA Tile dialect is designed to be targetable from multiple programming languages, making it a versatile tool for developers across different domains. This is facilitated by defining the dialect as an MLIR (Multi-Level Intermediate Representation) dialect, which is an extension to the widely used LLVM intermediate representation. This choice allows the dialect to be easily integrated with third-party compilers and domain-specific languages, enabling a diverse range of applications to benefit from NVIDIA's GPU architecture.
[0084] Finally, the stability and documentation of the CUDA Tile dialect are complemented by its integration with the broader CUDA ecosystem. As a formal extension of CUDA, the dialect offers all of CUDA's "forever" portability guarantees and interoperation with the rest of the CUDA stack, including JIT compilers, developer tools, linkers, and libraries. This integration ensures that programs written using the CUDA Tile dialect can benefit from the full suite of development and deployment tools available in the CUDA ecosystem.
[0085] The use of lower-level dialects for internal code and NDA partners in the context of the Tile-based Programming Model Extension to CUDA is an approach to enhance the flexibility and functionality of NVIDIA's programming model. These dialects are designed to provide specialized operations and optimizations that are not exposed to the public, allowing NVIDIA and its partners to leverage advanced features of the GPU architecture. The lower-level dialects, such as CUTLASS IR and CuTe IR, offer bindings that enable efficient mapping of data tiles onto tensor cores, optimizing performance for specific hardware configurations. This approach ensures that internal applications and libraries, like cuDNN and Myelin, can achieve high performance by utilizing architecture-specific instructions and optimizations.
[0086] The integration of lower-level dialects into the CUDA ecosystem is facilitated by the use of MLIR (Multi-Level Intermediate Representation) passes, which provide a flexible and efficient compilation path for tile-based programs. These passes enable the transformation of high-level operations into architecture-specific instructions, optimizing execution for specific hardware configurations. The MLIR framework allows NVIDIA to define custom passes that implement advanced optimizations, such as warp specialization and tensor memory management, ensuring that internal applications can achieve high performance on NVIDIA GPUs.
[0087] Finally, the use of lower-level dialects for internal code and NDA partners is complemented by the public CUDA Tile dialect, which provides a stable and documented interface for developers. This dialect serves as the formal extension of the CUDA ecosystem, offering all of CUDA's "forever" portability guarantees and interoperation with the rest of the CUDA stack. By defining a stable bytecode representation of a program, the public CUDA Tile dialect ensures that programs can be executed on all future NVIDIA GPUs. This stability is achieved through the integration of proprietary technologies, such as NVVM and CuTe, which provide a robust compilation path that can adapt to future hardware developments.
[0088] The Cutlass IR plays a role in the tile programming model by providing a specialized intermediate representation that facilitates efficient mapping of data tiles onto tensor cores. This IR is designed to handle the complexities associated with tensor core programming, offering a more streamlined and high-level approach compared to traditional methods. By abstracting away the intricate details of thread management and memory allocation, Cutlass IR allows developers to focus on the logical structure of their programs, making it easier to achieve high performance across different GPU architectures. This abstraction is particularly beneficial in scenarios where tensor cores are used for matrix multiplication and other computationally intensive operations.
[0089] One of the features of Cutlass IR is its ability to define operations at a tile level, rather than at the individual thread level. This tile-based approach aligns well with the architecture of tensor cores, which are optimized for handling large blocks of data simultaneously. By operating on tiles, Cutlass IR can efficiently utilize the parallel processing capabilities of tensor cores, leading to significant performance improvements. The IR also includes dialects that define specific operations and transformations, such as the pipeline dialect for managing producer-consumer relationships between different stages of computation. This ensures that data movement and computation are synchronized effectively, further enhancing performance.
[0090] Cutlass IR is integrated into the broader CUDA ecosystem, allowing it to interoperate with other components such as cuDNN and Myelin. This integration is crucial for maintaining the portability and stability of programs across different GPU generations. As part of the CUDA toolchain, Cutlass IR benefits from the extensive suite of development tools available, including JIT compilers, debuggers, and profilers. This makes it easier for developers to optimize their programs and ensure compatibility with future hardware releases. The IR's design also supports the use of lower-level dialects for internal code and NDA partners, providing additional flexibility for specialized applications.
[0091] Finally, the role of Cutlass IR in the tile programming model is complemented by its support for various programming languages, including Python. This is achieved through the use of MLIR (Multi-Level Intermediate Representation) passes, which provide a flexible compilation path for tile-based programs. By defining a stable bytecode representation, Cutlass IR ensures that programs can be executed reliably on all future NVIDIA GPUs. This stability is an advantage over other solutions, such as OpenAI Triton, which may not offer the same level of portability and interoperation with the CUDA stack.
[0092] The producer-consumer relationships in the pipeline dialect are an aspect of the tile-based programming model extension to CUDA. This concept is designed to manage the synchronization and data flow between different stages of computation within a GPU kernel. In the context of CUDA Tile IR, the pipeline dialect defines how data is transferred and processed between various computational units, such as tensor cores and memory units. The producer-consumer model ensures that data is loaded, processed, and stored efficiently, minimizing idle time and maximizing throughput.
[0093] In the tile-based programming model, the producer-consumer relationships are implemented through a series of operations that define the dependencies between different computational tasks. For example, a producer operation might involve loading data from global memory into shared memory, while a consumer operation could be a matrix multiplication using tensor cores. The pipeline dialect specifies the synchronization barriers and signals required to coordinate these operations, ensuring that data is available when needed and that computational resources are utilized effectively.
[0094] The pipeline dialect's ability to manage producer-consumer relationships is particularly important in the context of warp specialization. Warp specialization involves dividing the workload among different warps, each responsible for a specific task, such as data loading, computation, or storing results. The pipeline dialect facilitates this by defining the data flow and synchronization between warps, allowing them to operate concurrently without conflicts. This approach enhances the performance of tile-based programs by optimizing the use of GPU resources and reducing latency.
[0095] Overall, the producer-consumer relationships in the pipeline dialect are an innovation in the tile-based programming model extension to CUDA. They provide a structured framework for managing data dependencies and synchronization, enabling efficient execution of complex GPU kernels. By abstracting these relationships into a dialect, NVIDIA can offer a more flexible and powerful programming model that adapts to different hardware architectures and computational requirements.
[0096] At least one embodiment of the disclosure can be described in view of the following clauses:
[0097] 1. A processor comprising: circuits to cause a compiler to compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data, wherein the compiler: identifies one or more hardware features of a GPU, the one or more hardware features comprising one or more capabilities to perform blocks of threads in parallel; and generates code to perform operations on tiles of the input data using the one or more hardware features, including the capabilities to perform blocks of thread in parallel.
[0098] 2. The processor of clause 1, wherein the compiler further optimizes the code for execution on a specific version of the GPU by adjusting the mapping of tiles to threads based on the GPU's memory architecture.
[0099] 3. The processor of clauses 1 or 2, wherein the compiler includes a module to analyze the input data's structure to determine optimal tile sizes for parallel processing.
[0100] 4. The processor of any one of clauses 1-3, wherein the compiler generates code that utilizes the GPU's tensor cores for enhanced performance in matrix operations.
[0101] 5. The processor of any one of clauses 1-4, wherein the compiler supports multiple programming languages, including Python and C++, for generating the code.
[0102] 6. The processor of any one of clauses 1-5, wherein the compiler is configured to automatically update the mapping of tiles to threads when new GPU hardware is detected.
[0103] 7. The processor of any one of clauses 1-6, wherein the compiler includes error-checking mechanisms to ensure that the generated code does not exceed the GPU's resource limits.
[0104] 8. The processor of any one of clauses 1-7, wherein the compiler can generate code that is compatible with both CUDA and OpenCL environments.
[0105] 9. The processor of any one of clauses 1-8, wherein the compiler provides a user interface for developers to manually adjust the mapping of tiles to threads if desired.
[0106] 10. The processor of any one of clauses 1-9, wherein the compiler includes a profiling tool to analyze the performance of the generated code on different GPU architectures.
[0107] 11. The processor of any one of clauses 1-10, wherein the compiler supports integration with third-party development tools for enhanced debugging and optimization capabilities.
[0108] 12. The processor of any one of clauses 1-11, wherein the processor is comprised in at least one of: a control system for an autonomous or semi-autonomous machine; a perception system for an autonomous or semi-autonomous machine; a system for performing one or more simulation operations; a system for performing one or more digital twin operations; a system for performing one or more light transport simulation; a system for performing collaborative content creation for 3D assets; a system for performing one or more wireless cellular transmissions using a wireless cellular network; a system that provides one or more cloud gaming applications; a system for performing one or more deep learning operations; a system implemented using an edge device; a system implemented using a robot; a system for performing one or more generative AI operations; a system for performing one or more conversational AI operations; a system for performing operations using one or more large language models (LLMs) ; a system for performing operations using one or more vision language models (VLMs) ; a system for performing operations using one or more multi-modal language models (MMLMs) ; a system for performing operations using one or more vision-language-action (VLA) models; a system for performing one or more conversational AI operations; a system for performing one or more synthetic data generation operations; a system for presenting at least one of virtual reality content, augmented reality content, or mixed reality content; systems using or deploying one or more inference microservices; systems that incorporate deploy one or more machine learning models in a service or microservice along with an OS-level virtualization package (e.g., a container) ; a system incorporating one or more virtual machines (VMs) ; a system implemented at least partially in a data center; or a system implemented at least partially using cloud computing resources.
[0109] 13. A method of compiling source code for GPU execution, comprising: receiving source code with indications of a mapping between input data and tiles; identifying hardware features of a GPU, including capabilities for parallel thread execution; and generating executable code that maps tiles to threads based on the identified hardware features.
[0110] 14. A method for enhancing GPU code compatibility, comprising: analyzing source code with tile mappings for GPU execution; detecting changes in GPU hardware architecture; adjusting the mapping of tiles to threads to maintain code compatibility across different GPU generations.
[0111] 15. One or more processors, comprising: circuitry to provide a higher-level abstraction for tensor core programming, ensuring compatibility across GPU generations by abstracting architecture-specific mappings, simplifying the programming process, and aligning with array-based models.
[0112] 16. One or more processors, comprising: circuitry to provide a stable bytecode representation of programs having interoperation with the CUDA stack, and integration of tile-based tensor core programming with general parallel programs.
[0113] 17. One or more processors, comprising: circuitry to organize operations at thread-level, block-level, and grid-level models, and manage data parallelism by dividing an array of data into tiles and mapping the tiles onto a grid of thread blocks for parallel execution.
[0114] 18. One or more processors, comprising: circuitry to provide a programming interface compatible with array-based programming paradigms.
[0115] At least one embodiment of the disclosure can be understood in view of the following clauses:1. One or more processors, comprising:circuitry to cause a compiler to compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data, wherein the compiler:identifies one or more hardware features of a GPU, the one or more hardware features comprising one or more capabilities to perform blocks of threads in parallel; andgenerates code to perform operations on tiles of the input data using the one or more hardware features, including the capabilities to perform blocks of thread in parallel.2. The one or more processors of claim 1, wherein the circuits are further to:translate one or more tile-oriented intermediate representation instructions to one or more intermediate representation instructions to cause the operations on tiles of the input data to be performed using threads performed in parallel.3. The one or more processors of claim 1, wherein the source code is written in a domain specific language ( “DSL” ) that maps between the input data and the tiles.4. The one or more processors of claim 3, wherein the circuits are further to:compile the source code written the DSL to one or more tile-oriented intermediate representation instructions.5. The one or more processors of claim 4, wherein the tile-oriented intermediate representation instructions comprise one or more arrays of tiles.6. The one or more processors of claim 1, wherein the source code comprises one or more expressions written in the Python programming language.7. The one or more processors of claim 1, wherein the source code comprises one or more expressions written in a dynamically-typed language.8. The one or more processors of claim 1, wherein the compiler generates code that integrates tile-oriented intermediate representation instructions with thread-oriented intermediate representation instructions.9. The one or more processors of claim 7, wherein the thread-oriented intermediate representation instructions correspond to parallel thread execution ( “PTX” ) instructions.10. The one or more processors of claim 1, wherein the compilation is performed within at least one of a runtime or a driver of the GPU.11. The one or more processors of claim 1, wherein the circuits are further to:compile the source code to first binary instructions on the GPU, wherein the GPU has the first hardware features; andcompile the source code to second binary instructions for second hardware features, wherein the first binary instructions and second binary instructions are functionally equivalent, and wherein the second binary instructions use one or more of the second hardware features not included in the first hardware features
[0116] At least one embodiment of the disclosure can be understood in view of the following clauses:1. One or more processors, comprising:circuitry to:compile one or more intermediate representation instructions to instructions compatible with a graphics processing unit ( “GPU” ) , wherein the intermediate representation instructions have operands that correspond to one or more tiles of input data, wherein the compilation comprises mapping the one or more tiles of input data to operations to be performed on a plurality of threads in parallel.2. The one or more processors of clause 1, wherein the intermediate representation instructions are generated by compilation of source code comprising expressions written in a domain-specific language.3. The one or more processors of clauses 1 or 2, wherein the domain specific language comprises one or more expressions to indicate one or more strategies for mapping the one or more tiles of input data to the operations to be performed on the plurality of threads in parallel.4. The one or more processors of any of clauses 1-3, wherein the intermediate representation instructions are generated by compilation of source code comprising expressions written in Python.5. The one or more processors of any of clauses 1-4, wherein the intermediate representation instructions are generated by compilation of source code comprising expressions written in an interpreted language.6. The one or more processors of any of clauses 1-5, wherein the intermediate representation instructions comprise one or more expressions to map input data to the one or more tiles of input data.7. The one or more processors of any of clauses 1-6, wherein the intermediate representation instructions are tile-oriented intermediate representation instructions.8. The one or more processors of any of clauses 1-7, the circuitry further to:integrate one or more tile-oriented intermediate representation instructions with one or more second intermediate representation instructions, wherein the second intermediate representation instructions are thread-oriented intermediate representation instructions.9. The one or more processors of claim 8, wherein the second intermediate representation instructions are parallel thread execution ( “PTX” ) instructions.10. The one or more processors of any of clauses 1-9, wherein the compilation is performed within at least one of a runtime or a driver of the GPU.11. The one or more processors of any of clauses 1-10, wherein the mapping is dependent on hardware features obtained by at least one of a runtime or driver of the GPU.12. The one or more processors of any of clauses 1-11, wherein the mapping is based, at least in part, on one or more dynamically discovered hardware features of the GPU.
[0117] At least one embodiment of the disclosure can be understood in view of the following clauses:1. A method, comprising:performing a first compilation of one or more source code expressions indicating a mapping between input data and a plurality of tiles, wherein output of the first compilation comprises one or more tile-oriented intermediate representation instructions;performing a second compilation of the one or more tile-oriented intermediate representation instructions, wherein output of the second compilation comprises one or more thread-oriented intermediate representation instructions; andperforming a third compilation, wherein output of the third compilation comprises instructions natively executable by a GPU.2. The method of clause 1, wherein the source code expressions comprise expressions in a domain-specific language.3. The method of clauses 1 or 2, wherein the source code expressions comprises expressions in a dynamically typed language.4. The method of clause 3, wherein the dynamically typed language is the Python programming language.5. The method of any of clauses 1-4, wherein the source code expressions comprise a tile program.6. The method of any of clauses 1-5, wherein performing the second compilation comprises mapping tiles of the plurality of tiles to blocks of threads to be performed in parallel by the GPU.7. The method of clause 6, wherein the mapping is based, at least in part, on one or more dynamically discovered hardware features of the GPU.8. The method of any of clauses 1-7, wherein the second compilation is based, at least in part, on one or more dynamically discovered hardware features of the GPU.9. The method of any of clauses 1-8, wherein the thread-oriented intermediate representation instructions are parallel thread execution ( “PTX” ) instructions.
[0118] At least one embodiment of the disclosure can be understood in view of the following clauses:1. A method, comprising:compiling one or more domain-specific language expressions to one or more tile-based intermediate representation instructions, wherein the tile-based intermediate representation instructions are to be compiled, at runtime, to thread-oriented intermediate representation instructions based, at least in part, on one or more hardware features determined at runtime.
[0119] At least one embodiment of the disclosure can be understood in view of the following clauses:1. One or more processors, comprising:circuitry to operate a graphics processing unit ( “GPU” ) , wherein the circuitry at least:compiles one or more tile-oriented intermediate representation instructions to other instructions, wherein the tile-oriented intermediate representation instructions indicate operations on tiles, and wherein the other instructions indicate corresponding operations to be performed by threads of the GPU; andcauses the other instructions to be performed by the GPU.2. The one or more processors of clause 1, wherein the other instructions comprise parallel thread execution ( “PTX” ) instructions.3. The one or more processors of clauses 1 or 2, wherein the PTX instructions are compiled to instructions natively performable by the GPU in order to cause the other instructions to be performed by the GPU.
[0120] At least one embodiment of the disclosure can be understood in view of the following clause:1. One or more processors, comprising:circuitry to operate a graphics processing unit ( “GPU” ) , wherein the circuitry at least:compile one or more intermediate representation instructions to instructions compatible with a graphics processing unit ( “GPU” ) , wherein the intermediate representation instructions have operands that correspond to one or more tiles of input data, wherein the compilation comprises mapping the one or more tiles of input data to operations to be performed on a plurality of threads in parallel.
[0121] At least one embodiment of the disclosure can be understood in view of the following clauses:1. One or more processors, comprising:circuitry to operate a graphics processing unit ( “GPU” ) , wherein the circuitry at least:compile instructions compatible with a graphics processing unit ( “GPU” ) , wherein the instructions have operands that correspond to one or more tiles of input data, wherein the compilation comprises mapping the one or more tiles of input data to operations to be performed on a plurality of threads in parallel, and wherein the mapping is based, at least in part, on one or more dynamically discovered capabilities of the GPU.2. The one or more processor of clause 1, wherein dynamically discovering capabilities of the GPU comprises at least one of a driver or runtime library of the GPU determining one or features of the GPU.
[0122] FIG. 1 illustrates an example of the architecture of the Tile IR Stack Compiler and Toolchain Components, illustrating the layers and components involved in compiling and deploying tile-based applications, from domain-specific representations to thread representation and development deployment, within a CUDA toolchain. FIG. 1 depicts the following components in the compiler stack.
[0123] Tile-Based Application: The tile-based application component represents the high-level domain-specific representation of the programming model. It encompasses various domains such as deep learning, graphics, and high-performance computing (HPC) . This component defines the specific use cases and applications that will benefit from the tile-based programming model. By focusing on domain-specific needs, the tile-based application ensures that the programming model is tailored to efficiently handle the computational requirements of each domain, providing optimized performance and resource utilization.
[0124] Tile Code (DSLs) : Tile code, represented by domain-specific languages (DSLs) , is the layer where developers write their programs using high-level abstractions. These DSLs are designed to simplify the programming process by providing constructs that are closer to the problem domain rather than the hardware specifics. This component enables developers to express complex algorithms and computations in a more intuitive and readable manner, reducing the cognitive load and potential for errors associated with low-level programming.
[0125] Tile Language Compilers: Tile language compilers, such as nvcc, Numba-Tile, and Triton, are responsible for translating the high-level tile code into an intermediate representation that can be further processed by the CUDA toolchain. These compilers bridge the gap between the high-level abstractions provided by the DSLs and the low-level machine instructions executed on the GPU. By optimizing the translation process, tile language compilers ensure that the resulting code is both efficient and portable across different GPU architectures.
[0126] CUDA Toolchain MLIR Layers: The CUDA toolchain MLIR (Multi-Level Intermediate Representation) layers provide a structured framework for transforming the intermediate representation generated by the tile language compilers into a form that can be executed on the GPU. This component includes various dialects, such as the public CUDA Tile dialect, which define the operations and transformations specific to the tile-based programming model. The MLIR layers enable advanced optimizations and ensuring that the tile-based programs can be efficiently mapped onto the underlying hardware.
[0127] Public CUDA Tile Dialect: The public CUDA Tile dialect is a key component of the MLIR layers, providing a stable and well-defined interface for expressing tile-based computations. This dialect serves as the formal extension of the CUDA ecosystem, offering all of CUDA's portability guarantees and interoperation with the rest of the CUDA stack. By defining a stable bytecode representation, the public CUDA Tile dialect ensures that tile-based programs can be reliably executed on all future NVIDIA GPUs, maintaining compatibility and performance across generations.
[0128] Tile IR Compilation (Tile->Thread) : The Tile IR compilation process involves transforming the tile-based intermediate representation into a thread-based representation that can be executed on the GPU. This component includes the TileAA and TileAS dialects, which handle the complexities of mapping tiles onto threads and scheduling the computations. By automating this process, the Tile IR compilation ensures that developers can focus on high-level programming without worrying about the intricate details of thread management and synchronization.
[0129] MMA Code Generation : MMA (Matrix Multiply-Accumulate) code generation is a specialized component within the CUDA toolchain that focuses on optimizing matrix multiplication operations, which are common in many computational workloads. This component leverages the capabilities of NVIDIA's tensor cores to accelerate these operations, providing significant performance improvements for applications such as deep learning and scientific computing. By integrating MMA code generation into the tile-based programming model, developers can achieve high performance with minimal effort.
[0130] CUTLASS / CuTe IR Dialects: The CUTLASS and CuTe IR dialects are part of the MLIR layers, providing additional abstractions and optimizations for tensor core programming. These dialects define operations and transformations specific to tensor core computations, enabling developers to take full advantage of the hardware's capabilities. By incorporating these dialects into the tile-based programming model, the CUDA toolchain ensures that tensor core operations are both efficient and easy to use, further enhancing the performance of tile-based applications.
[0131] NVVM Dialect: The NVVM (NVIDIA Virtual Machine) dialect is a lower-level intermediate representation that serves as a bridge between the high-level tile-based abstractions and the final machine code executed on the GPU. This dialect provides a stable and portable interface for expressing GPU computations, ensuring that the resulting code can be executed on a wide range of NVIDIA hardware. By leveraging the NVVM dialect, the CUDA toolchain can optimize the execution of tile-based programs, providing both performance and portability.
[0132] Compiler Stack: The compiler stack is responsible for transforming the intermediate representations into executable machine code. This component includes various stages, such as NVVM SOLID, ToT LLVM, PTX, Mercury, and SASS, each of which performs specific optimizations and transformations.
[0133] CUDA Tile Toolchain Integration: The CUDA Tile toolchain integration encompasses various development and deployment tools, such as JIT compilation, debugging, profiling, autotuning, and fatbinaries. These tools provide developers with the necessary capabilities to optimize, test, and deploy their tile-based applications, ensuring that they can achieve the desired performance and functionality.
[0134] FIG. 2 illustrates an example of a system 200 that can include software and hardware to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described herein, according to at least one embodiment. System 200 can include storage 202 and processor (s) 208. Storage 202 can include, for example, memory, cache, or other storage described further herein. Storage 202 can be separate from processor (s) 208, or storage 202 can be included in processor (s) 208 (e.g., in storage 212) . In at least one embodiment, software program 204 and / or software libraries (or instructions) 206 can be stored in memory, cache, or other storage and provided to processor (s) 208 to cause one or more circuits of processor (s) 208 to perform operations described herein. In at least one embodiment, software program 204 and / or software libraries (or instructions) 206 can be integrated into one or more circuits of processor (s) 208. Software program 204, which can be used to perform any of the operations described herein, may be stored on storage 202.
[0135] In at least one embodiment, software program 204 can include one or more software modules. In at least one embodiment, as used in any implementation described herein, unless otherwise clear from context or stated explicitly to contrary, a module refers to any combination of software logic, firmware logic, hardware logic, and / or circuitry configured to provide functionality described herein. In at least one embodiment, software is embodied as a software package, code and / or instruction set or instructions, and “hardware, ” as used in any implementation described herein, includes, for example, singly or in any combination, hardwired circuitry, programmable circuitry, state machine circuitry, fixed function circuitry, execution unit circuitry, and / or firmware that stores instructions performed by programmable circuitry. In at least one embodiment, modules are, collectively or individually, embodied as circuitry that forms part of a larger system, for example, an integrated circuit (IC) , system on-chip (SoC) , and so forth. In at least one embodiment, a module performs one or more processes in connection with any suitable processing unit and / or combination of processing units, such as one or more CPUs, GPUs, GPGPUs, PPUs, and / or variations thereof including those further described herein.
[0136] In at least one embodiment, software program 204 can include a collection of software code, commands, instructions, or other sequences of text to instruct a computing device to perform one or more computational operations and / or invoke one or more other sets of instructions, such as API (s) or API function (s) or Instruction Set Architecture (ISA) level instructions, to be executed or otherwise performed. Instructions (e.g., hardware instructions) or microcode can involve ISA level instructions, which can include native ISA instructions or non-native ISA instructions. Software program 204 and / or software libraries (or instructions) 206 (e.g., one or more modules) can be distributed among multiple processors that communicate over a bus, network, by writing to shared memory, and / or any suitable communication process such as those described herein.
[0137] In at least one embodiment, system 200 can include one or more software libraries 206 that can, for example, provide one or more APIs and / or ISA instructions. In at least one embodiment, one or more APIs and / or ISA instructions can be used to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data. In at least one embodiment, one or more software libraries 206 can be included in drivers and / or runtimes. In at least one embodiment, software libraries 206 (e.g., including one or more APIs and / or ISA instructions) can include sets of software instructions that, if executed or otherwise performed, cause processor (s) 208 to perform one or more computational operations, such as any of the operations described herein. In at least one embodiment, one or more APIs and / or ISA instructions can be distributed or otherwise provided as a part of one or more software libraries 206, runtimes, drivers, and / or any other grouping of software and / or executable code further described herein. In at least one embodiment, one or more APIs and / or ISA instructions can perform one or more computational operations in response to invocation by software program 204.
[0138] Processor (s) 208 may include any number of processors and any suitable processing unit and / or combination of processing units, 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, parallel processors, GPGPUs, DPUs, and / or variations thereof including those further described herein) , including any processors described herein, such as, but not limited to, processors in FIGS. 4-22C. In at least one embodiment, processor (s) 208 can retrieve or fetch instructions (e.g., one or more APIs and / or ISA instructions) from storage 202 using, for example, instruction fetch 216 (e.g., for an Instruction Fetch stage) . Instructions can include instructions to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data. In at least one embodiment, processor (s) 208 can include storage 212 and instruction queue 210 to store and queue instructions fetched from storage 202. In at least one embodiment, fetched instructions can be decoded by decode 218 to determine what operation should be performed by processor (s) 208 (e.g., in an Instruction Decode stage) . In at least one embodiment, processor (s) 208 can fetch additional operands (data) that may be used for instructions, and operands can be stored, e.g., in registers or storage 212. In at least one embodiment, micro-operations 220 can perform operations on data stored in one or more registers or storage 212. For example, each step of instructions fetched by processor (s) 208 can be decomposed during execution so processor (s) 208 can execute instructions in steps through a series of micro-operations 220. In at least one embodiment, program counter (PC) 214 can hold an address for a next instruction and can be updated to point to the next instruction to be executed by processor (s) 208.
[0139] In at least one embodiment, processor (s) 208 can perform instructions (e.g., in an Execution stage) . For example, processor (s) 208 can perform an operation specified by the instructions, such as an arithmetic operation, a logical operation, or a data transfer. In at least one embodiment, compute unit (s) 222 can execute instructions to perform any of the operations described herein. In at least one embodiment, compute unit (s) can include ALU (s) 224 (Arithmetic Logic Units) , which may be used for performing arithmetic and logical operations. In at least one embodiment, compute unit (s) can include FPU (s) (Floating Point Units) 226, which may be used for performing floating-point calculations. In at least one embodiment, other circuits 228 can be used to perform other operations, such as vector and / or scalar operations. In at least one embodiment, accelerator (s) 230 can include one or more matrix multiplication accelerators, one or more parallel processing units (PPUs) , such as GPUs, or any other accelerator or processor further described herein. In at least one embodiment, software program 204 can utilize one or more APIs and / or ISA instructions to perform various computing operations with accelerator (s) 230, such as matrix multiplication, arithmetic operations, or any other computing operation further described herein. In at least one embodiment, one or more computing operations using accelerator (s) 230 can include at least one or more groups of computing operations to be accelerated by execution at least in part by accelerator (s) 230, including to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data.
[0140] In at least one embodiment, system 200 can be used to perform one or more instructions that include functions or operations, such as those described in connection with FIGS. 1-6. In at least one embodiment, system 200 comprising one or more processors causes one or more circuits Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data and / or otherwise perform operations described herein. In at least one embodiment, system 200 is included in and / or otherwise includes systems illustrated in FIG. 1 to cause one or more circuits Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data and / or otherwise perform operations described herein. In at least one embodiment, system 200 includes one or more hardware illustrated in FIGS. 4-22C such as Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data and / or otherwise perform operations described herein.
[0141] FIG. 3 is a block diagram 300 illustrating a driver and / or runtime including one or more libraries to provide one or more application programming interfaces (APIs) , according to at least one embodiment. In at least one embodiment, a software program 302 is a software module. In at least one embodiment, a software program 302 includes one or more software modules. In at least one embodiment, one or more APIs 310 are sets of software instructions that, if executed, cause one or more processors to perform one or more computational operations. In at least one embodiment, one or more APIs 310 are distributed or otherwise provided as a part of one or more libraries 306, runtimes 304, drivers 304, and / or any other grouping of software and / or executable code further described herein. In at least one embodiment, one or more APIs 310 perform one or more computational operations in response to invocation by software programs 302. In at least one embodiment, a software program 302 is a collection of software code, commands, instructions, or other sequences of text to instruct a computing device to perform one or more computational operations and / or invoke one or more other sets of instructions, such as APIs 310 or API functions 312, to be executed.
[0142] In at least one embodiment, API functions 312 included but are not limited function to verify whether objects indicated in a description of an image are depicted in said image, functions to generate a textual description of visual content, functions to accept a natural language prompt to parse, edit, modify, and / or alter a description of an image, functions to identify whether objects descripted in a caption are depicted in an image sought to be described, and functions to generate an evaluation metric of a degree of similarity between an input image sought to be captioned and a generated caption. In at least one embodiment, functionality provided by one or more APIs 310 include software functions 312, such as those usable to accelerate one or more portions of software programs 302 using one or more parallel processing units (PPUs) , such as graphics processing units (GPUs) . In at least one embodiment, a software program is a compiler.
[0143] In at least one embodiment, APIs 310 are hardware interfaces to one or more circuits to perform one or more computational operations. In at least one embodiment, one or more software APIs 310 described herein are implemented as one or more circuits to perform one or more techniques described in conjunction with FIG. 1. In at least one embodiment, one or more software programs 302 includes instructions that, if executed, cause one or more hardware devices and / or circuits to perform one or more techniques further described in conjunction with FIG. 1.
[0144] In at least one embodiment, software programs 302, such as user-implemented software programs, utilize one or more application programming interfaces (APIs) 310 to perform various computing operations, such as memory reservation, matrix multiplication, arithmetic operations, or any computing operation performed by parallel processing units (PPUs) , such as graphics processing units (GPUs) , as further described herein. In at least one embodiment, one or more APIs 310 provide a set of callable functions 312, referred to herein as APIs, API functions, and / or functions, that individually perform one or more computing operations, such as computing operations related to parallel computing. In at least one embodiment, one or more APIs 310 provide functions 312 to adjusting a description of an image. In at least one embodiment, one or more APIs 310 provide functions 312 to cause a neural network to perform one or more operations, such as by returning a called function to a processor where said processor invokes said neural network. In at least one embodiment, one or more APIs 310 provide functions 312 to execute an application programming interface to cause software to be corrected based on a previous version of said software 316.
[0145] In at least one embodiment, one or more software programs 302 interact or otherwise communicate with one or more APIs 310 to perform one or more computing operations using one or more PPUs, such as GPUs. In at least one embodiment, one or more computing operations using one or more PPUs include at least one or more groups of computing operations to be accelerated by execution at least in part by said one or more PPUs. In at least one embodiment, one or more software programs 302 interact with one or more APIs 310 to facilitate parallel computing using a remote or local interface.
[0146] In at least one embodiment, an interface is software instructions that, if executed, provide access to one or more functions 312 provided by one or more APIs 310. In at least one embodiment, a software program 302 uses a local interface when a software developer compiles one or more software programs 302 in conjunction with one or more libraries 306 including or otherwise providing access to one or more APIs 310. In at least one embodiment, one or more software programs 302 are compiled statically in conjunction with pre-compiled libraries 306 or uncompiled source code including instructions to perform one or more APIs 310. In at least one embodiment, one or more software programs 302 are compiled dynamically and said one or more software programs utilize a linker to link to one or more pre-compiled libraries 306 including one or more APIs 310.
[0147] In at least one embodiment, a software program 302 uses a remote interface when a software developer executes a software program that utilizes or otherwise communicates with a library 306 including one or more APIs 310 over a network or other remote communication medium. In at least one embodiment, one or more libraries 306 including one or more APIs 310 are to be performed by a remote computing service, such as a computing resource services provider. In another embodiment, one or more libraries 306 including one or more APIs 310 are to be performed by any other computing host providing said one or more APIs 310 to one or more software programs 302.
[0148] In at least one embodiment, a processor performing or using one or more software programs 302 calls, uses, performs, or otherwise implements one or more APIs 310 to allocate and otherwise manage memory to be used by said software programs 302. In at least one embodiment, one or more software programs 302 utilize one or more APIs 310 to allocate and otherwise manage memory to be used by one or more portions of said software programs 302 to be accelerated using one or more PPUs, such as GPUs or any other accelerator or processor further described herein. Those software programs 302 request a neural network to generate a modified bounding box based, at least in part, on one or more second bounding boxes.
[0149] In at least one embodiment, an API 310 is an API to facilitate parallel computing. In at least one embodiment, an API 310 is any other API further described herein. In at least one embodiment, an API 310 is provided by a driver and / or runtime 304. In at least one embodiment, an API 310 is provided by a CUDA user-mode driver. In at least one embodiment, an API 310 is provided by a CUDA runtime. In at least one embodiment, a driver 304 is data values and software instructions that, if executed, perform or otherwise facilitate operation of one or more functions 312 of an API 310 during load and execution of one or more portions of a software program 302. In at least one embodiment, a runtime 304 is data values and software instructions that, if executed, perform or otherwise facilitate operation of one or more functions 312 of an API 310 during execution of a software program 302. In at least one embodiment, one or more software programs 302 utilize one or more APIs 310 implemented or otherwise provided by a driver and / or runtime 304 to perform combined arithmetic operations by said one or more software programs 302 during execution by one or more PPUs, such as GPUs.
[0150] In at least one embodiment, one or more software programs 302 utilize one or more APIs 310 provided by a driver and / or runtime 304 to perform combined arithmetic operations of one or more PPUs, such as GPUs. In at least one embodiment, one or more APIs 310 provide combined arithmetic operations through a driver and / or runtime 304, as described above. In at least one embodiment, one or more software programs 302 utilize one or more APIs 310 provided by a driver and / or runtime 304 to allocate or otherwise reserve one or more blocks of memory 314 of one or more PPUs, such as GPUs. In at least one embodiment, one or more software programs 302 utilize one or more APIs 310 provided by a driver and / or runtime 304 to allocate or otherwise reserve blocks of memory. In at least one embodiment, one or more APIs 310 are to perform combined arithmetic operations, as described below in conjunction with any FIG. 1.
[0151] To improve software programs 302 usability and / or optimization of one or more portions of said software programs 302 to be accelerated by one or more PPUs, such as GPUs, in an embodiment, one or more APIs 310 provide one or more API functions 312 to perform a software correction system usable or used by one or more computing devices as described above and further described in conjunction with FIG. 1. In at least one embodiment, a block diagram 300 depicts a processor, including one or more circuits to perform one or more software programs to combine two or more application programming interfaces (APIs) into a single API. In at least one embodiment, a block diagram 300 depicts a system, including one or more processors to perform one or more software programs to combine two or more application programming interfaces (APIs) into a single API.
[0152] In at least one embodiment, at least a portion of block diagram 300 is implemented using at least a portion of any system (s) depicted in and / or described with respect to FIGS. 4-22C. In at least one embodiment, at least a portion of block diagram 300 is used to implement at least a portion of any system (s) depicted in and / or described with respect to FIGS. 4-22C.DATA CENTER
[0153] FIG. 4 illustrates an example data center 400, in accordance with at least one embodiment. Data center 400 may include one or more rooms having racks 402 and auxiliary equipment used to house one or more racks 402 and one or more baseboards 404. Rack 402 can include one or more baseboards 404. Rack 402 can include a housing that receives and supports individual baseboards 404. Operational aspects of rack 402 may be regulated at a rack level, corresponding to a group of baseboards 404, or at a baseboard level, corresponding to individual baseboards 404, among other options. Rack 402 or baseboards 404 can have particularly selected maximum operating parameters, such as, but not limited to, power consumption, operating frequencies, and others. Data center 400 can be supported by various cooling systems, such as, but not limited to, cooling towers, cooling loops, pumps, and other support systems. Cooling systems may include sensors and controllers to monitor and managing cooling properties for racks 402. Baseboards 404 within racks 402 can get operational power from one or more power distribution units (PDUs; not shown) . PDUs may be arranged within racks 402, for example between racks 402 including baseboards 404, or within racks 402 that also house baseboards 404.
[0154] Racks 402 and baseboards 404 can include sub-systems, modules, add-in cards, and other semiconductor components. Baseboards 404 can include one or more computing units 406 that can include one or more processors 408, one or more memory 410, and an interface controller 412. Computing units 406 may include 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 processors described herein, such as, but not limited to, processors in FIGs. 5-17. Computing units 406 can include one or more memory storage devices 410 (e.g., dynamic read-only memory, solid state storage or disk drives) , as well as network input / output ( “NW I / O” ) devices, network switches, virtual machines ( “VMs” ) , power modules, and cooling modules, etc. One or more computing units 406 may be a server having one or more of above-mentioned computing resources.
[0155] Computing units 406 can include separate groupings of computing units housed within one or more racks (not shown) , or many racks housed in data centers at various geographical locations (also not shown) . Separate groupings of computing units may include grouped compute, network, memory or storage resources that may be configured or allocated to support one or more workloads. Several computing units (e.g., including CPUs and / or other processors) may be grouped within one or more racks to provide compute resources to support one or more workloads. A resource orchestrator 414 may configure or otherwise control one or more computing units 406 or groups of computing units. Resource orchestrator 414 may include a software design infrastructure ( “SDI” ) management entity for data center 400. Resource orchestrator 414 may include hardware, software or some combination thereof.
[0156] Data center 400 can include any one of or any combination of a framework layer 420, a software layer 430 and an application layer 440. As shown in FIG. 4, framework layer 420 includes a job scheduler 422, a configuration manager 424, a resource manager 426 and a distributed file system 428. Framework layer 420 may include a framework to support software 432 of software layer 430 and / or one or more application (s) 442 of application layer 440. Software 432 or application (s) 442 may respectively include web-based service software or applications, such as, but not limited to, those provided by Amazon Web Services, Google Cloud and Microsoft Azure. Framework layer 420 may be a type of free and open-source software web application framework such as, but not limited to, Apache SparkTM (hereinafter “Spark” ) that may utilize distributed file system 428 for large-scale data processing (e.g., “big data” ) . Job scheduler 422 may include a Spark driver to facilitate scheduling of workloads supported by various layers of data center 400. Configuration manager 424 may be capable of configuring different layers such as, but not limited to, software layer 430 and framework layer 420 including Spark and distributed file system 428 for supporting large-scale data processing. Resource manager 426 may be capable of managing clustered or grouped computing units 406 mapped to or allocated for support of distributed file system 428 and job scheduler 422. Resource manager 426 may coordinate with resource orchestrator 414 to manage these mapped or allocated computing resources.
[0157] Software 432 can be included in software layer 430 and may include software used by at least portions of a computing unit 406, one or more computing units 406, groups of computing units 406, and / or distributed file system 428 of framework layer 420. One or more types of software may include, but are not limited to, Internet web page search software, e-mail virus scan software, database software, and streaming video content software.
[0158] Application (s) 442 can be included in application layer 440 and may include one or more types of applications used by at least portions of a computing unit 406, one or more computing units 406, groups of computing units 406, and / or distributed file system 428 of framework layer 420. One or more types of applications may include, but are not limited to, any number of a genomics application, a cognitive compute, application and a machine learning application, including training or inferencing software, machine learning framework software (e.g., PyTorch, TensorFlow, Caffe, etc. ) or other machine learning applications used in conjunction with one or more embodiments.
[0159] Any of configuration manager 424, resource manager 426, and resource orchestrator 414 may implement any number and type of self-modifying actions based on any amount and type of data acquired in any technically feasible fashion. Self-modifying actions may relieve a data center operator of data center 400 from making possibly bad configuration decisions and possibly avoiding underutilized and / or poor performing portions of a data center.
[0160] Data center 400 may include tools, services, software or other resources to train one or more machine learning models or predict or infer information using one or more machine learning models in accordance with one or more embodiments described herein. For example, a machine learning model may be trained by calculating weight parameters in accordance with a neural network architecture using software and computing resources described above with respect to data center 400. Trained machine learning models corresponding to one or more neural networks may be used to infer or predict information using resources described above with respect to data center 400 by using weight parameters calculated through one or more training techniques described herein.
[0161] Data center 400 may use CPUs, application-specific integrated circuits (ASICs) , GPUs, FPGAs, or other hardware (e.g., embodiments in FIGs. 5-17) to perform some or all of processes and techniques described elsewhere herein, such as, but not limited to, training and / or inferencing using above-described resources. Moreover, one or more software and / or hardware resources described above may be configured as a service to allow users to train or performing inferencing of information, such as, but not limited to, image recognition, speech recognition, or other artificial intelligence services.
[0162] In at least one embodiment, processor 408 can include one of the processors below and / or comprises one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. In at least one embodiment, processor 408 is configured by software 432 to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. Data center 400 may use logic, CPUs, application-specific integrated circuits (ASICs) , GPUs, FPGAs, or other hardware (e.g., embodiments in FIGs. 5-17) to perform any of the operations described above or elsewhere herein.PROCESSORS
[0163] The following figures set forth, without limitation, example processors and processing systems that can be used to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform some or all of processes, operations and / or and techniques described elsewhere herein. Example processors and processing systems can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. Processors and processing systems can include logic, central processing units (CPUs) , application-specific integrated circuits (ASICs) , graphics processing units (GPUs) , field programmable arrays (FPGAs) , XPUs (i.e., any compute architecture that best fits the need of an application) or other hardware (e.g., embodiments in FIGs. 5-17) to perform any of the operations described above, below, or elsewhere herein. Processors and / or processing systems described herein can include one or more circuits that can be used to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. As used herein, one or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. FIGs. 22A and 22B illustrate logic 2215 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 in accordance with at least one embodiment. Logic can refer, for example, to any combination of software logic, hardware logic, and / or firmware logic to provide functionality and / or operations described herein, wherein logic may be, collectively or individually, embodied as circuitry that forms part of a larger system, for example, an integrated circuit (IC) , an application-specific integrated circuit (ASIC) , a field programmable array (FPGA) , system-on-chip (SoC) , or one or processors (e.g., CPU, GPU) .
[0164] FIG. 5 illustrates a processor which is a system-on-a-chip (SOC) 500 (which may be referred to as system-on-chip, a superchip, or another name) , in accordance with at least one embodiment. SOC 500 can include processor complex 510 and processor complex 540. SOC 500 can include any number of processor complexes 510 and / or processor complexes 540 that may include any number of processors that are described herein, such as, but not limited to, those in FIGs. 5-17, in any combination. For example, processor 510 may include a central processing unit (CPU) , and processor 540 may include a graphics processor. Alternatively, processor 510 may include a graphics processor, and processor 540 may include a graphics processor. SOC 500 may include any number of display controllers 592, any number of multimedia engines 594, any number of I / O Interfaces 570, any number of memory controllers 580, and any number of fabrics 560 in any combination. For explanatory purposes, multiple instances of like objects are denoted herein with reference numbers identifying the object and parenthetical numbers identifying the instance where needed. SOC 500 can include a processor from Broadcom in Palo Alto, CA.
[0165] Processor complex 510 can include a CPU, processor complex 540 can include a GPU, and SOC 500 can include a processing unit that integrates 510 and 540 onto a single chip. Some tasks may be assigned to processor complex 510 and other tasks may be assigned to processor complex 540. Processor complex 510 can be configured to execute main control software associated with SOC 500, such as, but not limited to, an operating system. Processor complex 510 can be the master processor of SOC 500, controlling and coordinating operations of other processors. Processor complex 510 can issue commands that control the operation of processor complex 540 to perform some or all of the operations described herein. Processor complex 510 can be configured to execute host executable code derived from CUDA or other source code (e.g., HIP source code) , and processor complex 540 can be configured to execute device executable code derived from CUDA or other source code in order to perform any of the operations described herein.
[0166] Processor complex 510 can include cores 520 (1) -520 (4) and a cache (e.g., L3 cache) 530 to store information to perform operations described herein. Processor complex 510 may include any number of cores 520 and any number and type of caches in any combination. Cores 520 can be configured to execute instructions of a particular instruction set architecture ( “ISA” ) to perform some or all of the operations described herein. Each core 520 can include a CPU core. Core 520 (1) -520 (4) can be referred to as a computing units or compute units. SOC 500 can includes any number of processor complexes 510, fabric 560, I / O interfaces 570, and memory controllers 580.
[0167] Each core 520 can include a fetch / decode unit 522, an integer execution engine 524, a floating point execution engine 526, and an L2 cache 528. Fetch / decode unit 522 can fetch instructions to perform some or all of the operations described herein (such as, but not limited to, an API that is compiled into instructions) and decode such instructions, generate micro-operations, and dispatch separate micro-instructions to integer execution engine 524 and / or floating point execution engine 526. Fetch / decode unit 522 can concurrently dispatch one micro-instruction to integer execution engine 524 and another micro-instruction to floating point execution engine 526. Integer execution engine 524 can execute integer and memory operations. Floating point engine 526 can execute floating point and vector operations. Fetch-decode unit 522 can dispatch micro-instructions to one or more execution engines that replaces both integer execution engine 524 and floating point execution engine 526.
[0168] Each core 520 (i) , where i is an integer representing a particular instance of core 520, may access L2 cache 528 (i) included in core 520 (i) . Each core 520 included in core complex 510(j) , where j is an integer representing a particular instance of core complex 510, can be connected to other cores 520 included in core complex 510 (j) via L3 cache 530 (j) included in core complex 510 (j) . Cores 520 included in core complex 510 (j) , where j is an integer representing a particular instance of core complex 510, can access all of L3 cache 530 (j) included in core complex 510 (j) . L3 cache 530 may include any number of slices.
[0169] Processor complex 540 can be a graphics complex that can be configured to perform compute operations (e.g., compute operations involved in operations described herein) in a highly-parallel fashion. Processor complex 540 can be configured to execute graphics pipeline operations such as, but not limited to, draw commands, pixel operations, geometric computations, and other operations associated with rendering an image to a display. Processor complex 540 can be configured to execute operations unrelated to graphics, such as, but not limited to, neural network training and / or simulations. Processor complex 540 can be configured to execute both operations related to graphics and operations unrelated to graphics.
[0170] Processor complex 540 can include any number of compute units 550 (1) -550 (N) , where N is any integer greater than 1, and an L2 cache 542. Compute units 550 can share L2 cache 542, which may store information to be used to perform some or all of the operations described herein. L2 cache 542 can be partitioned. Processor complex 540 can include any number of compute units 550 and any number (including zero) and type of caches. Processor complex 540 can include any amount of dedicated graphics hardware.
[0171] Each compute unit 550 can include any number of SIMD units 552 (1) -552 (N) , where N is any integer greater than 1, and a shared memory 554. Each SIMD unit 552 can implement a SIMD architecture and can be configured to some or all of the operations described herein, in parallel. Each compute unit 550 may execute any number of thread blocks, but each thread block can execute on a single compute unit 550, although in some embodiments a thread block can execute on multiple compute units. A thread block can include any number of threads of execution. A workgroup can be a thread block. Each SIMD unit 552 can execute a group of threads. A group of threads (e.g., 16 threads) , which can also be referred to as a warp, or 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 set of data based on a single set of instructions. Predication can be used to disable one or more threads in a warp, subgroup, or wavefront. A lane can be a thread. A work item can be a thread, such as, but not limited to, e.g., with OpenCL. Different warps, subgroups, or wavefronts in a thread block may synchronize together and communicate via shared memory 554. Each compute unit 550 can include one or more thread block clusters, where a thread block cluster can enable programmatic control of locality at a granularity larger than a single thread block of a single streaming multiprocessor (SM) . Thread block clusters (also referred to as “clusters” ) can enable multiple thread blocks running concurrently across streaming multiprocessors to synchronize and collaboratively fetch, exchange, or otherwise use data. In at least one embodiment, streaming multiprocessors ( “SMs” ) can be referred to 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 in a processing unit (e.g., 16 cores, a ray tracing unit, a thread director or scheduler) .
[0172] Fabric 560 can be a system interconnect that facilitates data and control transmissions across processor complex 510, processor complex 540, I / O interfaces 570, memory controllers 580, display controller 592, and multimedia engine 594, e.g., to perform some or all of the operations described herein. SOC 500 may include any amount and type of system interconnect in addition to or instead of fabric 560 that facilitates data and control transmissions across any number and type of directly or indirectly linked components that may be internal or external to SOC 500. I / O interfaces 570 can be representative of 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 coupled to I / O interfaces 570. Peripheral devices that can be coupled to I / O interfaces 570 may include keyboards, mice, printers, scanners, joysticks or other types of game controllers, media recording devices, external storage devices, network interface cards, and so forth.
[0173] Display controller 592 may display images on one or more display device (s) , such as, but not limited to, a liquid crystal display ( “LCD” ) device. Multimedia engine 594 can include any amount and type of circuitry that is related to multimedia, such as, but not limited to, a video decoder, a video encoder, an image signal processor, etc. Memory controllers 580 may facilitate data transfers between SOC 500 and a unified system memory 590. Processor complex 510 and processor complex 540 may share unified system memory 590. Unified system memory 590 can include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as, but not limited to, synchronous graphics random access memory (SGRAM) , including graphics double data rate (GDDR) memory. Unified system memory 590 may include 3D stacked memory, including but not limited to high bandwidth memory (HBM) , HBM2e, or HDM3.
[0174] SOC 500 may implement a memory subsystem that includes any amount and type of memory controllers 580 and memory devices (e.g., shared memory 554) that may be dedicated to one component or shared among multiple components in order to perform any of the operations described herein. SOC 500 can implement a cache subsystem that includes one or more cache memories (e.g., L2 caches 528, L3 cache 530, and L2 cache 542) that may each be private to or shared between any number of components (e.g., cores 520, core complex 510, SIMD units 552, compute units 550, and processor complex 540) .
[0175] In at least one embodiment, SOC 500 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0176] FIG. 6A illustrates a parallel processor 600, in accordance with at least one embodiment. Parallel processor 600 may be implemented using one or more circuits and may be referred to as 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 FIGs. 5-17) to perform any of the operations described above or elsewhere herein.
[0177] Parallel processor 600 can include a parallel processing unit 602 to perform any of the operations described above or elsewhere herein. Parallel processing unit 602 can include an I / O unit 604 that enables communication with other devices, including other instances of parallel processing unit 602. I / O unit 604 may be directly connected to other devices. I / O unit 604 may connect with other devices via use of a hub or switch interface, such as, but not limited to, a memory hub 605. Connections between memory hub 605 and I / O unit 604 can form a communication link 613. I / O unit 604 may connect with a host interface 606 and a memory crossbar 616, where host interface 606 receives commands directed to performing processing operations and memory crossbar 616 receives commands directed to performing memory operations.
[0178] When host interface 606 receives a command buffer via I / O unit 604, host interface 606 can direct work operations to perform those commands to a front end 608. Front end 608 can couple with a scheduler 610 (which may be referred to as a sequencer) , which is configured to distribute commands or other work items to a processing cluster array 612. Scheduler 610 can ensure that processing cluster array 612 is properly configured and in a valid state before tasks may be distributed to a cluster of processing cluster array 612. Scheduler 610 may be implemented via firmware logic executing on a microcontroller. Microcontroller-implemented scheduler 610 can be configurable to perform complex scheduling and work distribution operations at coarse and fine granularity, enabling rapid preemption and context switching of threads executing on processing array 612. Host software can prove workloads for scheduling on processing cluster array 612 via one of multiple graphics processing paths. Workloads can then be automatically distributed across processing array cluster 612 by scheduler 610 logic within a microcontroller including scheduler 610.
[0179] Processing cluster array 612 can perform any of the operations described above or elsewhere herein and can include up to “N” processing clusters (e.g., cluster 614A, cluster 614B, through cluster 614N) , where “N” represents a positive integer (which may be a different integer “N” than used in other figures) . Each cluster 614A-614N of processing cluster array 612 can execute a large number of concurrent threads. Scheduler 610 can allocate work to clusters 614A-614N of processing cluster array 612 using various scheduling and / or work distribution algorithms, which may vary depending on workload arising for each type of program or computation. Scheduling can be handled dynamically by scheduler 610, or can be assisted in part by compiler logic during compilation of program logic configured for execution by processing cluster array 612. Different clusters 614A-614N of processing cluster array 612 can be allocated for processing different types of programs or for performing different types of computations.
[0180] Processing cluster array 612 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. Processing cluster array 612 can be configured to perform general-purpose parallel compute operations. For example, processing cluster array 612 can include logic to execute processing tasks including filtering of video and / or audio data, performing modeling operations, including physics operations, and performing data transformations.
[0181] Processing cluster array 612 can be configured to perform parallel graphics processing operations. Processing cluster array 612 can include additional logic to support execution of such graphics processing operations, including but not limited to, texture sampling logic to perform texture operations, as well as tessellation logic and other vertex processing logic. Processing cluster array 612 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. Parallel processing unit 602 can transfer data from system memory via I / O unit 604 for processing. During processing, transferred data can be stored to on-chip memory (e.g., parallel processor memory 622) during processing, then written back to system memory.
[0182] When parallel processing unit 602 is used to perform graphics processing, scheduler 610 can be configured to divide a processing workload into approximately equal sized tasks, to better enable distribution of graphics processing operations to multiple clusters 614A-614N of processing cluster array 612. Portions of processing cluster array 612 can be configured to perform different types of processing. For example, a first portion may be configured to perform vertex shading and topology generation, a second portion may be configured to perform tessellation and geometry shading, and a third portion may be configured to perform pixel shading or other screen space operations, to produce a rendered image for display. Intermediate data produced by one or more of clusters 614A-614N may be stored in buffers to allow intermediate data to be transmitted between clusters 614A-614N for further processing.
[0183] Processing cluster array 612 can receive processing tasks to be executed via scheduler 610, which receives commands defining processing tasks from front end 608. Processing tasks can include indices of data to be processed, e.g., surface (patch) data, primitive data, vertex data, and / or pixel data, as well as state parameters and commands defining how data is to be processed (e.g., what program is to be executed) . Scheduler 610 may be configured to fetch indices corresponding to tasks or may receive indices from front end 608. Front end 608 can be configured to ensure processing cluster array 612 is configured to a valid state before a workload specified by incoming command buffers (e.g., batch-buffers, push buffers, etc. ) is initiated.
[0184] Each of one or more instances of parallel processing unit 602 can couple with a parallel processor memory 622 to perform any of the operations described above or elsewhere herein. Parallel processor memory 622 can be accessed via memory crossbar 616, which can receive memory requests from processing cluster array 612 as well as I / O unit 604. Memory crossbar 616 can access parallel processor memory 622 via a memory interface 618. Memory interface 618 can include multiple partition units (e.g., partition unit 620A, partition unit 620B, through partition unit 620N) that can each couple to a portion (e.g., memory unit) of parallel processor memory 622. A number of partition units 620A-620N can be configured to be equal to a number of memory units, such that a first partition unit 620A has a corresponding first memory unit 624A, a second partition unit 620B has a corresponding memory unit 624B, and an N-th partition unit 620N has a corresponding N-th memory unit 624N. A number of partition units 620A-620N may not be equal to a number of memory units.
[0185] Memory units 624A-624N can include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as, but not limited to, synchronous graphics random access memory (SGRAM) , including graphics double data rate (GDDR) memory. Memory units 624A-624N may also include 3D stacked memory, including but not limited to high bandwidth memory (HBM) , HBM2e, or HDM3. Render targets, such as, but not limited to, frame buffers or texture maps may be stored across memory units 624A-624N, allowing partition units 620A-620N to write portions of each render target in parallel to efficiently use available bandwidth of parallel processor memory 622. A local instance of parallel processor memory 622 may be excluded in favor of a unified memory design that utilizes system memory in conjunction with local cache memory.
[0186] Any one of clusters 614A-614N of processing cluster array 612 can process data that will be written to any of memory units 624A-624N within parallel processor memory 622. Memory crossbar 616 can be configured to transfer an output of each cluster 614A-614N to any partition unit 620A-620N or to another cluster 614A-614N, which can perform additional processing operations on an output. Each cluster 614A-614N can communicate with memory interface 618 through memory crossbar 616 to read from or write to various external memory devices. Memory crossbar 616 can have a connection to memory interface 618 to communicate with I / O unit 604, as well as a connection to a local instance of parallel processor memory 622, enabling processing units within different processing clusters 614A-614N to communicate with system memory or other memory that is not local to parallel processing unit 602. Memory crossbar 616 can use virtual channels to separate traffic streams between clusters 614A-614N and partition units 620A-620N.
[0187] Multiple instances of parallel processing unit 602 can be provided on a single add-in card, or multiple add-in cards can be interconnected. Different instances of parallel processing unit 602 can be configured to interoperate even if different instances have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences. For example, some instances of parallel processing unit 602 can include higher precision floating point units relative to other instances. Systems incorporating one or more instances of parallel processing unit 602 or parallel processor 600 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.
[0188] FIG. 6A further includes a block diagram of a partition unit 620, in accordance with at least one embodiment. Partition unit 620 is an instance of one of partition units 620A-620N of FIG. 6A. Partition unit 620 can include an L2 cache 621, a frame buffer interface 625, and a ROP 626 (raster operations unit) . L2 cache 621 can be a read / write cache that is configured to perform load and store operations received from memory crossbar 616 and ROP 626. Read misses and urgent write-back requests can be output by L2 cache 621 to frame buffer interface 625 for processing. Updates can also be sent to a frame buffer via frame buffer interface 625 for processing. Frame buffer interface 625 may interface with one of memory units in parallel processor memory, such as, but not limited to, memory units 624A-624N (shown as 624) of FIG. 6A (e.g., within parallel processor memory 622) .
[0189] ROP 626 can be a processing unit that performs raster operations such as, but not limited to, stencil, z test, blending, etc. ROP 626 can then output processed graphics data that is stored in graphics memory. ROP 626 can include compression logic to compress depth or color data that is written to memory and decompress depth or color data that is read from memory. Compression logic can be lossless compression logic that makes use of one or more of multiple compression algorithms. A type of compression that is performed by ROP 626 can vary based on statistical characteristics of data to be compressed. For example, delta color compression is performed on depth and color data on a per-tile basis.
[0190] ROP 626 can be included within each processing cluster (e.g., cluster 614A-614N of FIG. 6A) instead of within partition unit 620. Read and write requests for pixel data may be transmitted over memory crossbar 616 instead of pixel fragment data. Processed graphics data may be displayed on a display routed for further processing by processor (s) , or routed for further processing by one of processing entities within parallel processor 600 of FIG. 6A.
[0191] In at least one embodiment, parallel processor 600 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0192] FIG. 6B includes a block diagram of a processing cluster 614 within a parallel processing unit, in accordance with at least one embodiment. A processing cluster can be an instance of one of processing clusters 614A-614N of FIG. 6A that can be used to perform any of the operations described above or elsewhere herein. Processing cluster 614 can be configured to execute many threads in parallel, where “thread” refers to an instance of a particular program executing on a particular set of input data. Single-instruction, multiple-data (SIMD) instruction issue techniques can be used to support parallel execution of a large number of threads without providing multiple independent instruction units. Single-instruction, multiple-thread (SIMT) techniques may be used to support 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 one of processing clusters.
[0193] Operation of processing cluster 614 can be controlled via a pipeline manager 632 that distributes processing tasks to SIMT parallel processors. Pipeline manager 632 can receive instructions from scheduler 610 of FIG. 6A and manages execution of those instructions via a graphics multiprocessor 634 and / or a texture unit 636. Graphics multiprocessor 634 may be an example instance of a SIMT parallel processor. However, various types of SIMT parallel processors of differing architectures may be included within processing cluster 614. One or more instances of graphics multiprocessor 634 can be included within a processing cluster 614. Graphics multiprocessor 634 can process data and a data crossbar 640 can be used to distribute processed data to one of multiple possible destinations, including other shader units. Pipeline manager 632 can facilitate distribution of processed data by specifying destinations for processed data to be distributed via data crossbar 640.
[0194] Each graphics multiprocessor 634 within processing cluster 614 can include an identical set of functional execution logic (e.g., arithmetic logic units, load-store units, etc. ) to perform computations for any of the operations described above or elsewhere herein. Functional execution logic can be configured in a pipelined manner in which new instructions can be issued before previous instructions may be complete. Functional execution logic can support a variety of operations including integer and floating point arithmetic, comparison operations, Boolean operations, bit-shifting, and computation of various algebraic functions. Same functional-unit hardware can be leveraged to perform different operations and any combination of functional units may be present.
[0195] Instructions transmitted to processing cluster 614 may constitute a thread, which can also be referred to as a warp, subgroup, wave, or a wavefront. A set of threads executing across a set of parallel processing engines can be referred to as a thread group. A thread group can execute a common program on different input data. Each thread within a thread group can be assigned to a different processing engine within a graphics multiprocessor 634. A thread group may include fewer threads than a number of processing engines within graphics multiprocessor 634. When a thread group includes fewer threads than a number of processing engines, one or more of processing engines may be idle during cycles in which that thread group is being processed. A thread group may also include more threads than a number of processing engines within graphics multiprocessor 634. When a thread group includes more threads than number of processing engines within graphics multiprocessor 634, processing can be performed over consecutive clock cycles. Multiple thread groups can be executed concurrently on a graphics multiprocessor 634.
[0196] Graphics multiprocessor 634 includes an internal cache memory to perform load and store operations, such as, but not limited to, any of the operations described above or elsewhere herein. Graphics multiprocessor 634 can forego an internal cache and use a cache memory (e.g., L1 cache 648) within processing cluster 614. Each graphics multiprocessor 634 may also have access to L2 caches within partition units (e.g., partition units 620A-620N of FIG. 6A) that can be shared among all processing clusters 614 and may be used to transfer data between threads. Graphics multiprocessor 634 may also access off-chip global memory, which can include one or more of local parallel processor memory and / or system memory. Any memory external to parallel processing unit 602 may be used as global memory. Processing cluster 614 can include multiple instances of graphics multiprocessor 634 and can share common instructions and data, which may be stored in L1 cache 648.
[0197] Each processing cluster 614 may include an MMU 645 (memory management unit) that can be configured to map virtual addresses into physical addresses. One or more instances of MMU 645 may reside within memory interface 618 of FIG. 6A. MMU 645 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 line index. MMU 645 may include address translation lookaside buffers (TLB) or caches that may reside within graphics multiprocessor 634 or L1 648 cache or processing cluster 614. A physical address can be processed to distribute surface data access locally to allow for efficient request interleaving among partition units. A cache line index may be used to determine whether a request for a cache line is a hit or miss.
[0198] A processing cluster 614 may be configured such that each graphics multiprocessor 634 is coupled to a texture unit 636 for performing texture mapping operations, e.g., 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 graphics multiprocessor 634 and can be fetched from an L2 cache, local parallel processor memory, or system memory, as needed. Each graphics multiprocessor 634 can output processed tasks to data crossbar 640 to provide processed task to another processing cluster 614 for further processing or to store processed task in an L2 cache, local parallel processor memory, or system memory via memory crossbar 616. A preROP 642 (pre-raster operations unit) can be configured to receive data from graphics multiprocessor 634, and direct data to ROP units, which may be located with partition units as described herein (e.g., partition units 620A-620N of FIG. 6A) . PreROP 642 unit can perform optimizations for color blending, organizing pixel color data, and performing address translations.
[0199] In at least one embodiment, processing cluster 614 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0200] FIG. 6C shows a graphics multiprocessor 634, in accordance with at least one embodiment, e.g., to perform any of the operations described above or elsewhere herein. Graphics multiprocessor 634 can couple with pipeline manager 632 of processing cluster 614. Graphics multiprocessor 634 can include an execution pipeline including but not limited to an instruction cache 652 (that, e.g., can store instructions, such as, not limited to compiled API instructions) , an instruction unit 654, an address mapping unit 656, a register file 658, one or more general purpose graphics processing unit (GPGPU) cores 662, and one or more load / store units 666, where one or more load / store units 666 can perform load / store operations to load / store instructions corresponding to performing an operation. GPGPU cores 662 and load / store units 666 can be coupled with cache memory 672 and shared memory 670 via a memory and cache interconnect 668. GPGPU cores 662 can be part of an SoC such as, but not limited to, part of integrated circuit 500 in FIG. 5.
[0201] Instruction cache 652 can receive a stream of instructions (e.g., to perform any of the operations described above or elsewhere herein) to execute from pipeline manager 632. Instructions can be cached in instruction cache 652 and dispatched for execution by an instruction unit 654. Instruction unit 654 can dispatch instructions as thread groups (e.g., warps, subgroups, wavefronts, or waves) , with each thread of thread group assigned to a different execution unit within GPGPU cores 662. An instruction can access any of a local, shared, or global address space by specifying an address within a unified address space. Address mapping unit 656 can be used to translate addresses in a unified address space into a distinct memory address that can be accessed by load / store units 666.
[0202] Register file 658 can provide a set of registers for functional units of graphics multiprocessor 634. Register file 658 may provide temporary storage for operands connected to data paths of functional units (e.g., GPGPU cores 662, load / store units 666) of graphics multiprocessor 634. Register file 658 may be divided between each of functional units such that each functional unit is allocated a dedicated portion of register file 658. Register file 658 can be divided between different warps (which may be referred to as wavefronts, subgroups, and / or waves or threads) being executed by graphics multiprocessor 634.
[0203] GPGPU cores 662 can each include floating point units (FPUs) and / or integer arithmetic logic units (ALUs) that can be used to execute instructions of graphics multiprocessor 634. GPGPU cores 662 can be similar in architecture or can differ in architecture. A first portion of GPGPU cores 662 can include a single precision FPU and an integer ALU while a second portion of GPGPU cores include a double precision FPU. FPUs can implement IEEE 754-2008 standard floating point arithmetic or enable variable precision floating point arithmetic. Graphics multiprocessor 634 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 blending operations. One or more of GPGPU cores 662 can also include fixed or special function logic.
[0204] GPGPU cores 662 can include SIMD logic capable of performing a single instruction on multiple sets of data. GPGPU cores 662 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 generated when executing programs written and compiled for single program multiple data (SPMD) or SIMT architectures. Multiple threads of a program can be configured for an SIMT execution model that can be executed via a single SIMD instruction. For example, eight SIMT threads that perform same or similar operations can be executed in parallel via a single SIMD8 logic unit.
[0205] Memory and cache interconnect 668 can include an interconnect network that connects each functional unit of graphics multiprocessor 634 to register file 658 and to shared memory 670. Memory and cache interconnect 668 may be a crossbar interconnect that allows load / store unit 666 to implement load and store operations between shared memory 670 and register file 658. register file 658 can operate at a same frequency as GPGPU cores 662, thus data transfer between GPGPU cores 662 and register file 658 can have very low latency. Shared memory 670 can be used to enable communication between threads that execute on functional units within graphics multiprocessor 634. Cache memory 672 can be used as a data cache for example, to cache texture data communicated between functional units and texture unit 636. Shared memory 670 can also be used as a program managed cache. Threads executing on GPGPU cores 662 can programmatically store data within shared memory in addition to automatically cached data that is stored within cache memory 672.
[0206] A parallel processor or GPGPU as described herein may 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 may be communicatively coupled to host processor / cores over a bus or other interconnect (e.g., a high-speed interconnect such as, but not limited to, PCIe or NVLink) . An SoC may include a parallel processor or GPGPU as described herein, where said parallel processor or said GPGPU is performed on said SoC. A GPU may be integrated on a package or chip as cores and communicatively coupled to cores over an internal processor bus / interconnect internal to a package or chip. Regardless a manner in which a GPU is connected, processor cores may allocate work to such GPU in a form of sequences of commands / instructions contained in a work descriptor. GPU then may use dedicated circuitry / logic for efficiently processing these commands / instructions to perform any of the operations described above or elsewhere herein.
[0207] In at least one embodiment, graphics multiprocessor 634 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0208] FIG. 7 shows a processor 700, in accordance with at least one embodiment. Processor 700 can include a processor with hybrid architecture (e.g., Lunar Lake or Meteor Lake) from Intel Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. Processor 700 can include one or more Central Processing Unit (s) (CPU 702) , one or more Graphics Processing Unit (s) (GPU 706) , and / or one or more Neural Processing Unit (s) (NPU 708) that can be, e.g., a dedicated AI accelerator that offloads artificial intelligence (AI) workloads from CPU 702 and GPU 706. Processor 700 can use instructions that, if executed cause processor 700 and / or any of its components to perform some or all of processes and techniques described elsewhere herein. Processor 700 may include any number of memory and cache units 710 to facilitate processing amongst different components of processor 700. Memory and cache 710 on processor 700 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. With respect to processor 700 and any of its components described above or elsewhere herein, one or more of APIs described herein can, for example, get compiled into instructions, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of processor 700 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 700, including registers, DRAM, flash, SRAM, cache, or other memory. One or more of APIs described herein can include a call.
[0209] Processor 700 can include compute engines as CPUs 702 and can include any number of cores, such as, but not limited to, up to 16 cores / 22 threads. Cores in CPU 702 can include P-cores (Performance) , E-cores (Efficient) &LP-E cores (Low-power Efficient) . Performance-cores can be used for low latency single-threaded, compute-intensive workloads, while Efficient-cores can be used for multi-threaded, less compute-intensive workloads. Low-power Efficient cores can be used for scalable multithreaded performance and offloading background tasks. P-cores can be used for single &limited threading performance, whereas E-and LP-E cores can be used for multi-threaded throughput and power efficiency.
[0210] GPU 706 can include any number of graphics engines, such as, but not limited to, ArcTM graphics engines (Xe LPG) with 8 Xe cores (up to 128 Execution Units or EUs) . As shown in FIG. 7, GPU 706 can include vector engines 710 and matrix engines 712, that, for example, can run FP, INT, and matrix operation tasks all at the same time or separately or in batches. GPU 706 can include a load / store unit 714, as well as other memory, such as, but not limited to, an instruction cache (I$) 716 and L1 cache / subsystem local memory (SLM) 718 that can, e.g., store instructions to perform any of the operations described above or elsewhere herein.
[0211] NPU 704 can include one or more AI Boost built-in neural processing unit (s) (NPUs) . NPU 704 can be enumerated to a host processor as an integrated PCIe device. NPU 704 can include one or more (e.g., two) Neural Compute Engine (NCE) tiles 730. Each tile can be configured with any combination of, but not limited to, (e.g., 2000) Multiply Accumulate (MAC) Engines 734, a Post Processing Engine (not shown) , a AI DSP Processor (not shown) , and memory (2 MB of dedicated SRAM) per tile as shown in FIG. 7. For general compute needs, Neural Compute Engines 730 can include interference pipeline 732, activation function (AF) 736, data conversion 738, load / store 740, and Streaming Hybrid Architecture Vector Engines (SHAVE) 728 for high performance parallel computing, which can include DMA (Direct Memory Access) engines 724 to shuttle data between system memory DRAM (Dynamic Random Access Memory) 726 and a software managed cache. Built-in device MMU (Memory Management Unit) 722 plus IOMMU (Input-Output Memory Management Unit) (not shown) can support multiple simultaneous hardware contexts and provide security isolation between execution contexts as per MCDM (Microsoft Compute Driver Model) architecture. Processor 700 can also include a media unit (not shown) that is included on or separately from XCDs or other components of processor 700 to enable video playback and video processing of compressed or non-compressed data, such using HEVC, AV1, VP9 and AVC HW accelerated decode support and HEVC, VP9 and AVC HW accelerated encode support.
[0212] A Thread Director, which includes firmware that is built into processor 700, can prioritize and manage distribution of workloads, sending tasks to optimized cores. For example, Thread Director can tie P-cores, E-cores and / or LP-E cores (described above) together with task-scheduling capabilities and ability to send less-demanding tasks to E-cores or LP-E cores. Deep Learning Boost ( DL Boost) (not shown) can provide built in AI acceleration for training and inference workloads, and may include VNNI (for CPU) and DP4a (for GPU) instruction set support. This instruction set may be optimized with OpenVINOTM Toolkit and oneAPI to accelerate INT8 inferencing. A software stack, e.g., as described elsewhere herein, can be used to enable AI inference using OpenVINOTM toolkit. Processor 700 can be configured to execute an application program, such as, but not limited to, a CUDA program.
[0213] In at least one embodiment, processor 700 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0214] Processor 700 can alternatively include a processor based on AI Engine Direct architecture from Qualcomm Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. that may include any number of NPUs, GPUs, CPUs and other related components, such as, but not limited to, NPU 704 as a Hexagon NPU, GPU 706 as a Adreno GPU, CPU 702 as a Kryo or Qualcomm Oryon CPU, as well as a Qualcomm Sensing Hub (not shown) and a memory subsystem 710, in any combination. Hexagon NPU 704 can include a power rail a micro-tile inferencing unit, a hardware acceleration unit, a tensor unit, a scalar unit, and a vector unit (all not shown) , which can have dedicated memory or share memory (e.g., cache or memory, such HBM3) for, e.g., storing instructions to perform any of the operations described above or elsewhere herein. Adreno GPU 706 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 702 can perform AI workloads, and can handle contextualization for pervasive generative AI applications. CPU 702 can also include an instruction fetch unit, a rename and retire 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 processor 700 and any of its components described above or elsewhere herein, one or more of APIs described herein can, for example, get compiled into instructions, which may be fetched by instruction fetch unit, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by rename and retire unit. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of processor 700 (e.g., in cache and / or memory) . Any number of CPU cores 702 may be included in any number of CPU cluster (s) that can be coupled to memory and / or cache, such as, but not limited to a shared L2 cache. Memory can be separate or shared, e.g., CPU clusters of CPU cores 702 can couple to memory subsystem 710 that can include fabric, system level cache and any number of memory management units that can, for example, read and write memory (e.g., DRAM) . Qualcomm Sensing Hub (not shown) includes micro NPUs, a power rail, and traditional sensors (agyrometer, accelerometer, even a barometer) with voice and data streams. Memory subsystem 710 can include memory and cache on processor 700, 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, e.g., for storing information and / or instructions to perform any of the operations described above or elsewhere herein. All or some of memory and / or cache in memory subsystem 710 can be shared or used individually by any one or combinations of components (e.g., GPU 706, NPU 704, and CPU 702) on processor 700.
[0215] Qualcomm AI Engine 700 may be programmed and controlled with an a software stack to perform some or all of the operations described herein, and include, e.g., a Neural Processing SDK for inferencing with versions for Android, Linux, and Windows. Developer libraries and services support programming languages, virtual platforms, and compilers. At a lower level of software stack, system software includes basic real-time operating system (RTOS) , system interfaces, and drivers. Software stack supports different operating systems, including Android, Windows, Linux, and QNX, and deployment and monitoring infrastructure like Prometheus, Kubernetes, and Docker. For direct cross-platform access to GPU 706, OpenCL and DirectML may be supported. For CPU 702, a LLVM compiler infrastructure optimizations enable accelerated and efficient AI inference. With respect to Qualcomm AI Engine 700 and any of its components described above or elsewhere herein, one or more of APIs described herein can, for example, get compiled into instructions, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of Qualcomm AI Engine 700 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of Qualcomm AI Engine 700, including registers, DRAM, flash, SRAM, cache, or other memory.
[0216] In at least one embodiment, processor 700 or Qualcomm AI Engine 700 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0217] FIG. 8A illustrates a processor 800, in accordance with at least one embodiment. Processor 800 can include an processor with scalable family from Intel Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. Processor 800 can include one or more cores 812 (1) -812 (N) , where N is any integer greater than 1 that can perform the operations described elsewhere herein. Cores 812 (1) -812 (N) can be interlinked together using ring and / or mesh interconnects. With a mesh interconnects architecture, an array of vertical and horizontal communication paths may allow traversal from one core to another 812 (1) -812 (N) through a shortest path (hop on vertical path to correct row, and hop across horizontal path to correct column) . For mesh interconnects, a die can house cores 812 (1) -812 (N) and can include a grid of converged mesh stops (CMS) that may be associated (e.g., 1: 1) with cores 812 (1) -812 (N) . Each core can be associated with one lower level cache (LLC) slice 814 (1) -814 (N) , or cores 812 (1) -812 (N) can share cache, e.g., lower level cache. LLCs 814 (1) -814 (N) can be inclusive by incorporating blocks in higher level cache (e.g., L2 cache) or non-inclusive (having blocks that may be not present in higher level cache) . Each core and LLC slice can include a Caching and Home Agent (CHA) (not shown) that can maintain cache coherency by providing scalability of resources across mesh interconnects for Ultra Path Interconnect ( UPI 816) cache coherency functionality. UPI 816 can provide a coherent interconnect for scalable systems and can allow for multiple processors to share a single shared address space through links, such as, but not limited to, two or three UPI links per processor.
[0218] Processor 800 can also include System Agent 810 that can house and / or perform various functionalities, such as, but not limited to, memory management, display functions, and / or input / output (I / O) functions. For example, processor 800 can include one or more integrated memory controller (s) (IMC) 808. IMC 808 can control and manage memory, such as, but not limited to, different memory types e.g., DDR ram, like DDR4 or others described elsewhere herein. System Agent 810 can include a display controller (not shown) to support display (s) . System Agent 810 can also incorporate PCIe 804 (e.g., up to 20 lanes of PCIe) , e.g., that can connect with an external dedicated graphics hookup over DMI bus (e.g., Intel’s DMI 3.0 bus) 806. System Agent 810 can include an Image Processing Unit (IPU) (not shown) which incorporates an image signal processor (ISP) on-die. Fabric 802 can provide scalability for connecting to other nodes (e.g., processors, such as processor 800) , and can, for example, be used with Cornelis Networks, an element of Scalable System Framework, that delivers the performance for high performance computing (HPC) workloads and the ability to scale to tens of thousands of nodes.
[0219] FIG. 8B illustrates components within core 812, in accordance with at least one embodiment. Core 812 can include front-end 818, back-end or execution engine 832, and memory subsystem 842. Front-end 818 can provide execution engine 832 with operations (e.g., operations described elsewhere herein) by decoding instructions stored in memory. For example, front-end 818 can include a micro-operations (μOps) cache path and / or a legacy path, along with branch prediction unit 821 that can determine paths instructions. A legacy path for instructions may include fetching variable-length (e.g., x86) instructions from L1 instruction cache 820 with instruction fetch and predecode 822, queuing the instructions in instruction queue 824, and decoding instructions using decoder 826 into μOps that can be provided to allocation queue 828. Alternatively, a μOPs cache path may include a cache containing already decoded μOps (μOps 830) that can be sent to allocation queue 828. Allocation queue 828 can perform as an interface between front-end 818 and execution engine 832, and can provide instructions to execution engine 832. One or more of API (s) described herein can, for example, get compiled into instructions that can be stored, processed, and executed by front-end 818, execution engine 832, and stored in memory subsystem 842.
[0220] Execution engine 832 can receive micro-operations into reorder buffer 834, which can register allocation, rename, and retire μOPs. From reorder buffer, μOPs can be sent to scheduler 836 that can be connected one or more different execution units 838, which can be connected to address generation unit (AGU) 840. Execution units 838 can perform, e.g., basic arithmetic logic unit (ALU) operations, multiplication, division, and / or more complex operations, such as, but not limited to, various vector operations. Scheduler 836 may manage queuing μOPs for one or more of execution units 838 depending, e.g., on operations needed to be performed.
[0221] Memory subsystem 842 can process load and store requests as well as ordering operations. For example, μOPs may relate to memory access (e.g. load and store) , and those can be sent on dedicated scheduler ports that can perform those memory operations. Store and load operations, for example, can be sent to load and store buffer (s) 844. Memory subsystem 842 can also include shared or separate L1 data and instruction cache 846, as well as L2 cache 848 that can be used and shared by L1 data and instruction cache 846. As described above for FIG. 8A, each core 812 can be connected to a slice of a third level of cache (e.g., LLC 814) that can be shared by all core 812.
[0222] In at least one embodiment, processor 800 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0223] FIG. 9 illustrates an AI accelerator 900, in accordance with at least one embodiment. Processor 900 can include a processor with AI accelerator architecture from Intel Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. AI accelerator 900 may use instructions that, if executed by AI accelerator 900, cause AI accelerator 900 to perform some or all of processes and techniques described elsewhere herein. For example, with respect to AI accelerator 900 and any of its components described above or elsewhere herein, one or more of APIs described herein can, for example, get compiled into instructions, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of AI accelerator 900 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of AI accelerator 900, including registers, DRAM, flash, SRAM, cache, or other memory. AI accelerator 900 may include one or more compute dies that can include homogeneous or heterogeneous processors. Compute dies may include one or more central processing units (CPU) , one or more graphics processing units (GPU) , or combinations of both.
[0224] In at least one embodiment, compute dies may include compute engines to perform AI computations. In at least one embodiment, AI accelerator 900 compute dies may be split into any number of (e.g., four) clusters that may be referred to as a DCORE (Deep Learning Core) 906 and contain any number of Matrix Multiplication Engines (MMEs) 908, Tensor Processor Cores (TPCs) 910, memory management unit 912, and L2 Cache 914, in any combination. MME (s) 908 can perform operations that use Matrix Multiplication, like fully connected layers, convolutions and batched-General Matrix Multiplications (GEMMs) . MMEs 908 may be equipped with Multiply-Accumulate Units (MACs) (not shown) that, for example, may perform General Matrix Multiplication (GEMM) operations, such as, but not limited to, an AxB multiplication that involves generating tensor C [NxM] from two input tensors, A [NxK] and B [KxN] . MME (s) 908 may be programmed with array dimensions, locations, data types, and various execution operands. MME (s) 908 can retrieve tensors A and B from memory, pulling them into its streaming buffers for matrix multiplication to be performed in parallel by MACs. MME (s) 908 may push tensor C back to memory upon completion. TPC (s) 910 may include any number of scalar units for performing scalar operations, any number of vector units for performing vector operations, any number of register files or local memory units (e.g., a vector local memory) , and load and store components for instructions, which can be coupled to memory or cache (e.g., HBM, L3 cache and / or L2 cache) (all not shown) . TPCs can support different types of parallel processing, e.g., Very Long Instruction Word (VLIW) Single-Instruction Multiple-Data (SIMD) that supports data types, such as, but not limited to, FP32, BF16, FP16 &FP8 (both E4M3 and E5M2) , UINT32, INT32, UINT16, INT16, UINT8 and INT8 datatypes. Any number of compute dies may be connected through an interconnect. An interconnect that can connect compute dies can be over an interposer bridge that, e.g., is transparent to software.
[0225] Memory on AI Accelerator 900 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. Memory and / or cache systems can be unified or separate. Compute dies of AI accelerator 900 may include on-die memory that includes one or more levels (e.g., two-levels) of cache. On-die SRAM or other memory described elsewhere herein can be used as a uniformly accessible last-level cache (L3) or split to slices of L2 cache that may be accessible to groups of MMEs 908 and TPCs 910. Using on-die memory as L2 or L3 cache can be fully configurable by software, which dynamically may decide per I / O tensor its optimal cache allocation. AI Accelerator 900 may include one or more Memory Management Units (MMUs) 922 for managing memory, such as allowing AI accelerator 900 memory subsystem to operate in a virtual space when accessing VRAM.
[0226] AI accelerator 900 may include a communications port (e.g., a PCIe Gen5 X16 port) 902 for communicating with a host and Scheduling and Synchronization Unit 904. AI accelerator 900 may include Media Unit 916 that may include any number or combinations of Media Decoder Engines (DECs) 920 and Rotator Engines (ROT) 918. AI accelerator 900 may include a network unit 924 that may include any number or combinations of network ports 926 and accompanied RDMA Engine (s) 928, L2 Cache, and memory (e.g., HBM2e or HBM3) stacks. AI accelerator 900 can incorporate a programmable Control Path entity (not shown) to manage parallel and efficient execution of various engines. Control Path can include Submission Queues (SQs) that may be issued by runtime system, Completion Queues (CQs) that may be used for job completion reporting, a Programmable Scheduling Mechanism that may be utilized for task scheduling, a Programmable Hardware Synchronization Mechanism or ‘Sync Manager (SM) ’ that may be used for hardware synchronization, a Programmable Interrupt Service Mechanism or ‘Interrupt Manager (INTR) ’ that can enable passing of asynchronous events to drivers.
[0227] AI accelerator 900 may include media decoding units that support Video Formats, such as, but not limited to, HEVC, Progressive H. 264, SVC base layer, MVC, VP9, JPEG, Progressive JPEG. AI accelerator 900 may support post processing of decoded media streams, such as, but not limited to, image down-scaling (resizing an image) , vertical and horizontal scaling at different scaling ratios, Image up-scaling, Image cropping, bilinear scaling, and Lancos scaling. AI accelerator 900 may implement two post processing channels per decoder unit, one with scalar (up and down) and one just to output the original image. AI accelerator 900 may include a hardware rotator engine that performs the following transformations of an input image: 2D rotation, 3D rotation, Projection, distorting and undistorting images, resampling input data at user-defined coordinates, and rescaling.
[0228] RDMA 928 over Converged Ethernet on AI accelerator 900 may enable scaling from a single node (i.e., a single AI Accelerator 900 to hundreds or thousands of nodes or AI Accelerators 900) . NW Subsystem 924 can include an Communication Library (IGCL) , a master conductor that orchestrates data movement, and a programable scheduling mechanism that can enable smooth activation of engines while maintaining task dependencies. A accelerator networking sub-system can include Gigabit Ethernet NIC ports 926, a Layer2 MAC (not shown) , and RDMA Engines 928. AI Accelerator 900 can include Aggregation Engines for performing summing activities. All engines in processor 900 can operate in parallel, e.g., MME (s) 908, TPC (s) 910 and NIC (s) 926 can all work at the same time. There can be dependency between operations running on different engines, e.g., output of one engine can be used as input of another engine, and / or MME, TPC and NIC can be scheduled to run in parallel. When one engine has completed its executing operation, another engine can be scheduled to start working on the next operation (immediately upon readiness of its inputs) .
[0229] AI Accelerator 900 can be operated and controlled using software layer 928 that may include low-level components, such as, but not limited to, a graph compiler, an automatic kernel fuser and a library of precompiled kernels, as well as integration to AI ecosystems, such as, but not limited to, PyTorch, DeepSpeed, Hugging Face, vLLM, Ray and more, or as described elsewhere herein with respect to software and programming platforms. Software layer 928 may include implementations of algorithms, such as, but not limited to, Paged Attention, Flash Attention and more. Software layer 928 may generate optimized binary code that implements a given model topology, such as, but not limited to, performing operator fusion, data layout management, parallelization, pipelining and memory management, and graph-level optimizations.
[0230] In at least one embodiment, AI accelerator 900 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0231] A neuromorphic computing system is described that adopts a multicore architecture where each core houses computing elements including neurons, synapses with on-chip learning capability, and local memory to store synaptic weights and routing tables. FIG. 10 is a simplified block diagram 1000 illustrating an example of at least a portion of such a neuromorphic computing device 1005, in accordance with at least one embodiment. Neuromorphic computing device 1005 can include a neuromorphic processor from Intel Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. As shown in this example, a device 1005 may be provided with a network 1010 of multiple neural network cores interconnected by an on-device network such that multiple different connections may be potentially defined between cores. For instance, a network 1010 of spiking neural network cores may be provided in device 1005 and may each communicate via short packetized spike messages sent from core to core over network channels. Each core (e.g., 1015) may possess processing and memory resources and logic to implement some number of primitive nonlinear temporal computing elements, such as, but not limited to, multiple (e.g., 1000+) distinct artificial neurons (referred to herein as “neurons” ) . For instance, each core may be capable of concurrently implementing multiple neurons such that neuromorphic cores may implement many multiples of neurons using device 1005. With respect to neuromorphic computing device 1005 and any of its components described above or elsewhere herein, one or more of APIs or equivalents described herein can, for example, get compiled into instructions or equivalents, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of neuromorphic computing device 1005 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of neuromorphic computing device 1005, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.
[0232] Continuing with the example of FIG. 10, neuromorphic computing device 1005 may additionally include processor 1020 and system memory 1025 to implement one or more components to manage and provide functionality of neuromorphic computing device 1005. For instance, system manager 1030 may be provided to manage global attributes and operations of neuromorphic computing device 1005 (e.g., attributes affecting network of cores 1010, multiple cores in network 1010, interconnections of neuromorphic computing device 1005 with other devices, manage access to global system memory 1025, among other potential examples) . In one example, system manager 1030 may manage the definition and provisioning of a specific routing tables to various routers in network 1010, orchestration of a network definition and attributes (e.g., weights, decay rates, etc. ) to be applied in network 1010, core synchronization and time multiplexing management, routing of inputs to appropriate cores, among other potential functions.
[0233] As another example, neuromorphic computing device 1005 may additionally include programming interface 1035 through which a user or system may specify a neural network definition to be applied (e.g., through a routing table and individual neuron properties) and implemented by mesh 1010 of neuromorphic cores. A software-based programming tool may be provided with or separate from neuromorphic computing device 1005 through which a user may provide a definition for a particular neural network to be implemented using network 1010 of neuromorphic cores. Programming interface 1035 may take an input of a programmer to then generate corresponding routing tables and populate local memory of individual neuromorphic cores (e.g., 1015) with specified parameters to implement a corresponding, customized network of artificial neurons implemented by neuromorphic cores 1015.
[0234] In some cases, neuromorphic computing device 1005 may advantageously interface with and interoperate with other devices, including general purpose computing devices, to realize certain applications and use cases. Accordingly, external interface logic 1040 may be provided in some cases to communicate (e.g., over one or more defined communication protocols) with one or more other devices. An external interface 1040 may be utilized to accept input data from another device or external memory controller acting as a source of input data. External interface 1040 may be additionally or alternatively utilized to allow results or output of computations of a neural network implemented using neuromorphic computing device 1005 to be provided to another device (e.g., another general purpose processor implementing a machine learning algorithm) to realize additional applications and enhancements, among other examples.
[0235] As shown in FIG. 10, network 1010 of multiple neural network cores interconnected by an on-device network is shown illustrating a portion of a network fabric interconnecting multiple neuromorphic cores (e.g., 1015 a-d) . For instance, a number of neuromorphic cores (e.g., 1015 a-d) may be provided in a mesh, with each core being interconnected by a network including a number of routers (e.g., 1050) . In one implementation, each neuromorphic core (e.g., 1015 a-d) may be connected to a single one of routers (e.g., 1050) and routers may be connected to at least one other router (as shown at 1010 in FIG. 10) . As an example, in one particular implementation, four neuromorphic cores (e.g., 1015 a-d) may be connected to a single router (e.g., 1050) and each of routers 1050 may be connected to two or more other routers to form a manycore mesh, allowing each neuromorphic core to interconnect with each other neuromorphic core in neuromorphic computing device 1005. Moreover, as each neuromorphic core may be configured to implement multiple distinct neurons, router network of neuromorphic computing device 1005 may similarly enable connections, or artificial synapses (or, simply, “synapses” ) , to be defined between any two of potentially many (e.g., 30,000+) neurons defined using network of neuromorphic cores 1010 provided in neuromorphic computing device 1005.
[0236] FIG. 10 shows a block diagram illustrating internal components of one example implementation of neuromorphic core 1015. In one example, a single neuromorphic core may implement some number of neurons (e.g. 1024) that share architectural resources of neuromorphic core 1015 in a time-multiplexed manner. In one example, each neuromorphic core 1015 may include processor block 1055 capable of performing arithmetic functions and routing in connection with the realization of a digitally implemented artificial neuron, such as, but not limited to, explained herein. Each neuromorphic core 1015 may additionally provide local memory in which a routing table may be stored and accessed for a neural network, accumulated potential of each soma of each neuron implemented using core 1015 may be tracked, parameters of each neuron implemented by core may 1015 be recorded, among other data and usage. Components, or architectural resources, of neuromorphic core 1015 may further include input interface 1065 to accept input spike messages generated by other neurons on other neuromorphic cores and output interface 1070 to send spike messages to other neuromorphic cores over mesh network 1010. In some instances, routing logic for neuromorphic core 1015 may be at least partially implemented using output interface 1070. Further, in some cases, core (e.g., 1015) may implement multiple neurons within an example SNN and some of these neurons may be interconnected. In such instances, spike messages sent between neurons hosted on core 1015 may forego communication over routing fabric of neuromorphic computing device 1005 and may instead by managed locally at particular neuromorphic core 1015.
[0237] Each neuromorphic core may additionally include logic to implement, for each neuron 1075, artificial dendrite 1080 and artificial soma 1085 (referred to herein, simply, as “dendrite” and “soma” respectively) . Dendrite 1080 may be a hardware-implemented process that receives spikes from network 1010. Soma 1085 may be a hardware-implemented process that receives each dendrite's accumulated neurotransmitter amounts for the current time and evolves each dendrite and soma's potential state to generate outgoing spike messages at the appropriate times. Dendrite 1080 may be defined for each connection receiving inputs from another source (e.g., another neuron) . In one implementation, dendrite process 1080 may receive and handle spike messages as they serially arrive in time-multiplexed fashion from network 1010. As spikes are received, neuron's activation (tracked using soma 1085 (and local memory 1060) ) may increase. When neuron's activation exceeds a threshold set for neuron 1075, neuron 1075 may generate a spike message that is propagated to a fixed set of fanout neurons via output interface 1070. Network distributes spike messages to all destination neurons, and in response those neurons, in turn, may update their activations in a transient, time-dependent manner, and so on, potentially causing the activation of some of these destination neurons to also surpass corresponding thresholds and trigger further spike messages, as in real biological neural networks.
[0238] As noted above, neuromorphic computing device 1005 may reliably implement a spike-based model of neural computation. Such models may also be referred to as Spiking Neural Networks (SNNs) . In addition to neuronal and synaptic state, SNNs also incorporate the concept of time. For instance, in an SNN, communication occurs over event-driven action potentials, or spikes, that convey no explicit information other than the spike time as well as an implicit source and destination neuron pair corresponding to the transmission of the spike. Computation occurs in each neuron as a result of the dynamic, nonlinear integration of weighted spike input. In some implementations, recurrence and dynamic feedback may be incorporated within an SNN computational model. Further, a variety of network connectivity models may be adopted to model various real world networks or relationships, including fully connected (all-to-all) networks, feed-forward trees, fully random projections, “small world” networks, among other examples. A homogeneous, two-dimensional network of neuromorphic cores, such as, but not limited to, shown in the example of FIG. 10 may advantageously supports all of these network models. As some or all cores of neuromorphic computing device 1005 may be connected, some or all neurons defined in cores may be therefore also fully connected through some number of router hops. Neuromorphic computing device 1005 may further include fully configurable routing tables to define a variety of different neural networks by allowing each core's neurons to distribute their spikes to any number of cores in mesh 1010 to realize fully arbitrary connectivity graphs.
[0239] In an improved implementation of a system capable of supporting SNNs, such as, but not limited to, a very large scale integration (VLSI) hardware device illustrated in the example of FIG. 10, high speed and reliable circuits may be provided to implement SNNs to model information processing algorithms as employed by a brain, but in a more programmable manner. For instance, while a biological brain can only implement a specific set of defined behaviors, as conditioned by years of development, a neuromorphic processor device may provide a capability to rapidly reprogram all neural parameters. Accordingly, a single neuromorphic processor may be utilized to realize a broader range of behaviors than those provided by a single slice of biological brain tissue. This distinction may be realized by adopting a neuromorphic processor with neuromorphic design realizations that differ markedly from those of neural circuits found in nature.
[0240] As an example, a neuromorphic processor may utilize time-multiplexed computation in both a spike communication network and neuron machinery of neuromorphic computing device 1005 to implement SNNs. Accordingly, physical circuitry of neuromorphic computing device 1005 may be shared among many neurons to realize higher neuron density. With time multiplexing, a network can connect N cores with O (N) total wiring length, whereas discrete point-to-point wiring would scale as O (N2) , realizing a significant reduction in wiring resources to accommodate planar and non-plastic VLSI wiring technologies, among other examples. In neuromorphic cores, time multiplexing may be implemented through dense memory allocation, for instance, using Static Random Access Memory (SRAM) , with shared buses, address decoding logic, and other multiplexed logic elements. State of each neuron may be stored in processor's memory, with data describing each neuron state including state of each neuron's collective synapses, all currents and voltages over its membrane, among other example information (such as, but not limited to, configuration and other information) .
[0241] A neuromorphic processor may adopt a “digital” implementation that diverts from other processors adopting more “analog” or “isomorphic” neuromorphic approaches. For instance, a digital implementation may implement integration of synaptic current using digital adder and multiplier circuits, as opposed to analog isomorphic neuromorphic approaches that accumulate charge on capacitors in an electrically analogous manner to how neurons accumulate synaptic charge on their lipid membranes. Accumulated synaptic charge may be stored, for instance, for each neuron in local memory of a corresponding core. Further, at an architectural level of an example digital neuromorphic processor, reliable and deterministic operation may be realized by synchronizing time across a network of cores such that any two executions of a design, given same initial conditions and configuration, will produce identical results. Asynchrony may be preserved at a circuit level to allow individual cores to operate as fast and freely as possible, while maintaining determinism at a system level. Accordingly, a notion of time as a temporal variable may be abstracted away in neural computations, separating it from a “wall clock” time that the hardware utilized to perform the computation. Accordingly, in some implementation, a time synchronization mechanism may be provided that globally synchronizes neuromorphic cores at discrete time intervals. A synchronization mechanism allows neural computation to complete as fast as circuitry allows, with a divergence between run time and biological time that a neuromorphic system models.
[0242] In operation, neuromorphic computing device 1005 may begin in an idle state with all neuromorphic cores inactive. As each core asynchronously cycles through its neurons, it generates spike messages that a mesh interconnect routes to appropriate destination cores containing all destination neurons. Implementation of multiple neurons on a single neuromorphic core may be time-multiplexed, and a time step may be defined in which all spikes involving multiple neurons may be processed and considered using shared resources of a corresponding core. As each core finishes servicing its neurons for a respective time step, cores may, in some implementations, communicate (e.g., using a handshake) with neighboring cores using synchronization messages to flush a mesh of all spike messages in flight, allowing cores to safely determine that all spikes have been serviced for a time step. At that point all cores may be considered synchronized, allowing them to advance their time step and return to an initial state and begin a next time step.
[0243] Given this context, and as introduced above, a device (e.g., 1005) implementing a mesh 1010 of interconnected neuromorphic cores may be provided, with core 1015 implementing potentially multiple artificial neurons capable of being interconnected to implement an SNN. Each neuromorphic core (e.g., 1015) may provide two loosely coupled asynchronous processes: an input dendrite process (e.g., 1080) that receives spikes from network 1010 and applies them to an appropriate destination dendrite compartments at the appropriate future times, and output soma process (e.g., 1085) that receives each dendrite compartment's accumulated neurotransmitter amounts for the current time and evolves each dendrite and soma's membrane potential state, generating outgoing spike messages at appropriate times (e.g., when a threshold potential of a soma has been reached) . Note that, from a biological perspective, dendrite and soma names used here only approximate a role of these functions and should not be interpreted too literally.
[0244] In at least one embodiment, neuromorphic computing device 1005 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0245] FIG. 11 is a block diagram of an embodiment of a multi-node network in which remote memory computation can be implemented, in accordance with any embodiment. System 1100 may represent a network of nodes described herein that can, e.g., be used to perform some or all of the operations described herein. System 1100 can represent a data center. System 1100 may represent a server farm. System 1100 may represent a data cloud or a processing cloud. System 1100 can represent a supercomputer. System 11 may include tens, hundreds, or thousands of nodes. Nodes of system 1100 may include 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 FIGs. 5-17. With respect to any of processors in system 1100 and any of its components described above or elsewhere herein, one or more of APIs or equivalents described herein can, for example, get compiled into instructions or equivalents, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of a processor or node (e.g., in cache and / or memory) . A result of API(s) can then be stored in storage within or outside of a processor or node, including registers, DRAM, flash, SRAM, cache, or other memory equivalents. System 1100 may include over nine thousand nodes, with each node including 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 Aurora Supercomputer from Intel Corporation in Santa Clara, CA or another supercomputer that shares at least some of the components described herein.
[0246] One or more clients 1102 make requests over network 1104 to system 1100. Network 1104 represents one or more local networks, or wide area networks, or a combination. Clients 1102 can be human or machine clients, which generate requests for execution of operations by system 1100. System 1100 executes applications or data computation tasks requested by clients 1102.
[0247] System 1100 can include one or more racks, which represent structural and interconnect resources to house and interconnect multiple computation nodes. Rack 1110 can include multiple nodes 1130. Rack 1110 may host multiple blade components 1120 (0) to 1120 (N-1) , where N is an integer greater than or equal to 2. Hosting can refer to providing power, structural or mechanical support, and interconnection. Blades 1120 (0) to 1120 (N-1) can refer to computing resources on printed circuit boards (PCBs) , where a PCB houses hardware components for one or more nodes 1130. Blades 1120 (0) to 1120 (N-1) may or may not include a chassis or housing or other "box" other than that provided by rack 1110. Blades 1120 (0) to 1120 (N-1) may include housing with exposed connector to connect into rack 1110. System 1100 may or may not include rack 1110, and each blade (e.g., 1120 (0) ) can include a chassis or housing that can stack or otherwise reside in close proximity to other blades and allow interconnection of nodes 1130. System 1100 may include 10, 624 compute blades, which include 63, 744 Intel Max Series GPUs and 21, 248 Intel Xeon Max CPUs across 166 racks.
[0248] System 1100 can include fabric 1170, which represents one or more interconnectors for nodes 1130. Fabric 1170 can include multiple switches 1172 or routers or other hardware to route signals among nodes 1130. Additionally, fabric 1170 can couple system 1100 to network 1104 for access by clients 1102. In addition to routing equipment, fabric 1170 can be considered to include cables or ports or other hardware equipment to couples nodes 1130 together. Fabric 1170 can have one or more associated protocols to manage routing of signals through system 1100. A protocol or protocols is at least partly dependent on hardware equipment used in system 1100.
[0249] As illustrated, rack 1110 can include N blades (e.g., 1120 (0) to 1120 (N-1) ) . In addition to rack 1110, system 1100 can include rack 1150. As illustrated, rack 1150 may include M blades (e.g., 1160 (0) to 1160 (M-1) ) . M is not necessarily the same as N; thus, it will be understood that various different hardware equipment components could be used, and coupled together into system 1100 over fabric 1170. Blades 1160 (0) to 1160 (M-1) can be the same or similar to blades 1120 (0) to 1120 (N-1) . Nodes 1130 can be any type of node as described herein, and may not be necessarily all the same type of node. System 1100 is not limited to being homogenous, nor is it limited to not being homogenous.
[0250] A node in blade 1120 (0) is illustrated in detail. However, other nodes in system 1100 can be the same or similar. At least some nodes 1130 may be computation nodes, with processor 1132 and memory 1140. A computation node refers to a node with processing resources (e.g., one or more processors) that executes an operating system and can receive and process one or more tasks. At least some nodes 1130 can include storage server nodes with a server as processing resources 1132 and memory 1140. A storage server refers to a node with more storage resources than a computation node, and rather than having processors for execution of tasks, a storage server includes processing resources to manage access to storage nodes within a storage server.
[0251] Node 1130 can include interface controller 1134, which can represent logic to control access by node 1130 to fabric 1170. Logic can include hardware resources to interconnect to physical interconnection hardware. Logic can include software or firmware logic to manage interconnection. Interface controller 1134 can include a host fabric interface, which can include a fabric interface in accordance with any embodiment described herein.
[0252] Node 1130 may include memory subsystem 1140. Memory 1140 can include memory computation resources (comp) 1142, which represent one or more capabilities by memory 1140 to perform memory computations. System 1100 enables remote memory operations, such as, but not limited to, the operations described elsewhere herein. Thus, nodes 1130 can request memory computations by remote nodes, where data for computation remains local to an executing node instead of being sent over fabric 1170 or instead of being sent from memory to a fabric interface. In response to execution of memory computation, executing node can provide a result to a requesting node.
[0253] Processor 1132 can include one or more separate processors. Each separate processor can include a single processing unit, a multicore processing unit, or a combination. A processing unit can include a primary processor such as, but not limited to, a CPU (central processing unit) , a peripheral processor such as, but not limited to, a GPU (graphics processing unit) , or a combination. Memory 1140 can be or include memory devices and a memory controller.
[0254] Reference to memory devices can apply to different memory types. Memory devices generally refer to volatile memory technologies. Volatile memory is memory whose state (and therefore data stored on it) is indeterminate if power is interrupted. Nonvolatile memory refers to memory whose state is determinate even if power is interrupted. Dynamic volatile memory can refresh data stored in a device to maintain state. One example of dynamic volatile memory includes DRAM (dynamic random access memory) , or some variant such as, but not limited to, synchronous DRAM (SDRAM) . A memory subsystem as described herein may be compatible with a number of memory technologies, such as, but not limited to, DDR3 (dual data rate version 3, original release by JEDEC (Joint Electronic Device Engineering Council) on June 27, 2007, currently on release 21) , DDR4 (DDR version 4, initial specification published in September 2012 by JEDEC) , DDR4E (DDR version 4, extended, currently in discussion by JEDEC) , LPDDR3 (low power DDR version 3, JESD209-3B, Aug 2013 by JEDEC) , LPDDR4 (LOW POWER DOUBLE DATA RATE (LPDDR) version 4, 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 (DDR version 5, currently in discussion by JEDEC) , LPDDR5 (currently in discussion by JEDEC) , HBM2 (HBM version 2) , currently in discussion by JEDEC) , or others or combinations of memory technologies, and technologies based on derivatives or extensions of such specifications.
[0255] In addition to, or alternatively to, volatile memory, in one embodiment, reference to memory devices can refer to a nonvolatile memory device whose state is determinate even if power is interrupted. In one embodiment, nonvolatile memory device is a block addressable memory device, such as, but not limited to, NAND or NOR technologies. Thus, a memory device can also include a future generation nonvolatile devices, such as, but not limited to, a three dimensional crosspoint (3DXP) memory device, other byte addressable nonvolatile memory devices, or memory devices that use chalcogenide phase change material (e.g., chalcogenide glass) . In one embodiment, a memory device can be or include multi-threshold level NAND flash memory, NOR flash memory, single or multi-level phase change memory (PCM) or phase change memory with a switch (PCMS) , a resistive memory, nanowire memory, ferroelectric transistor random access memory (FeTRAM) , magnetoresistive random access memory (MRAM) memory that incorporates memristor technology, or spin transfer torque (STT) -MRAM, or a combination of any of the above, or other memory.
[0256] In at least one embodiment, system 1100 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0257] FIG. 12 illustrates accelerated processing unit 1200, in accordance with at least one embodiment. Accelerated processing unit 1200 can include a processor based on CDNA architecture from AMD Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. Accelerated processing unit 1200 can include one or more accelerator complex dies (XCDs) 1204 for performing operations described elsewhere herein, such as, but not limited to, graphics processing and / or parallel processing as well as computations with instruction-level parallelism, including support for a broad range of precisions (INT8, FP8, BF16, FP16, TF32, FP32, and FP64) and sparse matrix data (i.e. sparsity) . XCDs may, in some instances, be referred to as Graphics Compute Dies (GCDs) . Accelerated processing unit 1200 can include one or more complex compute dies (CCDs) 1206 for performing operations described elsewhere herein, such as, but not limited to, those operations performed by host processors. CCDs may, in some instances, be referred to as core complexes or CCXs, such as, but not limited to, CCXs used in AMD Ryzen processors. XCDs and CCDs can share any type of cache or memory (e.g., one or more memory units 1202) , or have cache or memory allocated to each XCD or CCD or groups of XCDs or CCDs. For example, on-package AMD Infinity Fabric connects XCDs and CCD into shared AMD Infinity Cache 1208 and, in some embodiments, high-bandwidth memory (e.g., HMB3) . Accelerated processing unit 1200 can include an AMD MI300a processor that includes three CPU chiplets (or CCDs) and six accelerator chiplets (XCDs) on top of four input-output dies (IODs) that may be layered on a piece of silicon that links them together (e.g., via AMD Infinity Fabric) to eight stacks of high-bandwidth DRAM that ring a superchip. An AMD MI300x processor substitutes CCDs for two more XCDs, for an accelerator-only system.
[0258] Accelerated processing unit 1200 can include one or more input / output (I / O) interfaces. For example, XCDs 1204 and CCDs 1206 can be together on one or more input-output dies (IODs) 1210 that can include one or more I / O interfaces. IODs 1210 can include of 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 coupled to I / O interfaces 1270. I / O interfaces from IODs 1210 can also be used for connected one or more accelerated processing units 1200, e.g., in a server architecture.
[0259] Accelerated processing unit 1200 can include one or more memory units 1202 for storing instructions and other information used to perform operations described elsewhere herein. Memory units 1202 can include any volatile memory, such as, but not limited to, memory types described elsewhere herein and can include, e.g., high-bandwidth memory (e.g., HMB3) or high-bandwidth DRAM. Memory associated with accelerated processing unit 1200 (e.g., memory units 1202) can include system memory that can be used, for example, for commands, instructions and constants, and inputs and outputs. Memory units 1202 can also include device memory that can be used as storage and, for example, for commands, instructions and constants, and inputs and outputs, as return buffer (s) and for private data. Memory units 1202 can be linked to one or more IODs 1210. L1 cache 1220 starts a memory hierarchy that includes shared L2 cache 1228, e.g., within XCDs. AMD Infinity CacheTM, which is a last level cache (LLC) located on an active I / O die (IOD) . CCDs 1206 and XCDs 1204 may have separate or shared memory. AMD Infinity Architecture and AMD Infinity FabricTM technology can enable coherent, high-throughput 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 multi-device platforms.
[0260] As shown in FIG. 12, an XCD 1204 can include a shared set of global resources 1230, which can include hardware scheduler 1232 and Asynchronous Compute Engines (ACE) 1224 that send tasks (e.g., compute shader workgroups) to Compute Units (CUs or cores) 1234. ACEs 1224 (e.g., four) can be each associated with CUs 1234 (e.g., 40 CUs) , and some of CUs 1234 can be disabled for yield management. CUs 1234 can have dedicated cache or share cache (e.g., L2 cache) 1228 that may be used to coalesce all memory traffic for a die. CUs 1234 can include threaded and parallel processor cores including instruction fetching and scheduling with Scheduler (S) 1212, matrix core unit (MCU) 1216 and shader core (SC) 1218 (e.g., execution units for scalar, vector and matrix data types) , as well as load / store pipelines with an L1 cache 1220 and Local Data Share (LDS) 1214. Local data share can include, for example, a scratch RAM with built-in arithmetic capabilities that allow data to be shared between threads in a workgroup. An instruction cache 1240 (e.g., for storing and providing instructions for performing operations described elsewhere herein) and a constant cache 1238 can be connected to one or more CUs and can be shared between two CUs. Matrix cores 1216 can process a variety of data types, such as, but not limited to, INT8, FP8, FP16, BF16 and TF32 data types. Accelerated processing unit 1200 can include compute units 1234 that may be arranged in an array format, e.g., as a data-parallel-processor (DPP) array. Ultra-threaded dispatch processor 1242 can communicate with compute units 1234, and command processor 1244 can read commands that a host has written to memory-mapped registers in a system-memory address space (not shown) . Command processor 1244 can send hardware-generated interrupts to a host processor (e.g., a CCD) when a command is completed. Memory controller 1236 can also have direct access to all device memory and host-specified areas of system memory. To satisfy read and write requests, memory controller 1236 can perform functions of a direct-memory access (DMA) controller, including computing memory-address offsets based on a format of requested data in memory. For example, one or more of APIs described herein can, for example, get compiled into instructions that can be stored in instruction cache 1240 and then fetched by instruction fetch logic in processor 1240, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of processor 1200 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 1200, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.
[0261] An application can include 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 internal base-address and other configuration registers, specify a data domain on which accelerated processing unit 1200 can operate, invalidate and flush caches on accelerated processing unit 1200, and cause accelerated processing unit 1200 to begin execution of a program. Kernels can be referred to as programs executed by accelerated processing unit 1200. A kernel can be executed independently on every work item, or as groups of work-items that can be referred to as a wavefront, which can execute a kernel on all work-items in a group (e.g., 64) in one pass. Compute units 1234 can include a scalar arithmetic logic unit (ALU) , which can operates on one value per wavefront (common to all work items) , a vector ALU, which can operate on unique values per work-item, a local data share 1214, which can allow work-items within a workgroup to communicate and share data, a scalar memory (not shown) , which can transfer data between scalar general-purpose registers (SGPRs) and memory through a cache, and vector memory, which can transfer data between vector general-purpose registers (VGPRs) and memory, including sampling texture maps. Kernel control flow can be handled using scalar ALU instructions, which can includes if / else, branches and looping. Scalar ALU (SALU) and memory instructions can work on an entire wavefront and operate on one or more SGPRs. Vector memory and ALU instructions can operate on all work-items in a wavefront at one time.
[0262] In at least one embodiment, accelerated processing unit 1200 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0263] FIG. 13 illustrates a processor 1300, such as, but not limited to, a processor based on a Zen architecture (such as, e.g., Zen 1, 2, 3, 4, 5 or other) from AMD Corporation in Santa Clara, CA or another processor that shares at least some of the components described herein. Processor 1300 includes one or more CPU dies 1302 (1) -1302 (N) , where N is any integer greater than 1. CPU die 1302 can include any number of processor cores 1316 (e.g., to perform any of the operations described elsewhere herein) and any number of cache memories (e.g., to store instructions and other information to perform any of the operations described elsewhere herein) , in any combination. For example, L2 Cache units 1318 can be coupled to processor core (s) 1316, which can share and / or couple individually to L2 Cache units 1318. Processor cores 1316 can couple to L3 cache 1322 individually and / or share L3 Cache, which can be a lowest level cache (LLC) 1322 for access to data and other information used by processor cores 1316. One or more processor cores 1316 and one or more L2 Cache units 1318 can be included in a core complex (CCX) 1320 that can include (e.g., a 32 MB) shared cache (e.g., L3 cache 1322) . Core complex 1320 can be fabricated onto a die (CCD or CPU die) 1302. For example, up to 12 core complexes 1320 can be configured into a processor along with 8 CPU dies 1302 to provide up to 96 processor cores 1316 for processor 1300. A ‘Zen 4c’ core complex 1320, for example, can include up to eight cores 1316 and a shared 16 MB L3 cache 1322. Two of these core complexes 1320 can be combined onto a single CPU die 1302 for 16 cores per die and a total of 32 MB of L3 cache 1322 per die. Up to eight of CPU dies 1302 may be combined with an I / O unit 1304 to provide CPUs with up to 128 processor cores 1316. Up to four ‘Zen 4c’ dies described above can be combined to provide CPUs with up to 64 processor cores 1316.
[0264] Processor 1300 can include a variety of configurations for input / output operations that are described further herein. I / O unit 1304 can include one or more memory controllers 1306 that can manage memory usage (e.g., DDR5 memory) for processor 1300. I / O unit 1304 may include one or more SATA disk controllers for managing storage 1312 and one or more Compute Express Link (CXLTM) 1.1+ memory controllers 1314 that can provide CPU-to-device and CPU-to-memory connections and can be flexibly assigned to specific functions at server design time. I / O unit 1304 may include PCIe controller 1308 for connecting peripherals and other components connected to processor 1300. I / O unit 1304 may include USB ports 1310 for connecting to other components separate from processor 1300. CPU dies 1302 can support any number of connections, e.g., one or two connections, to I / O unit 1304. As shown, I / O unit 1304 can include components described further herein, and I / O unit 1304 can be a I / O die that houses several different components. Memory controller 1306, PCIe controller 1308, USB ports 1310, SATA controller 1312, and / or CXL controller 1314 can be integrated anywhere within processor 1300 either separately or in any groups or combinations thereof.
[0265] Processor 1300 can include Infinity Fabric 1324 interconnects (which can be similar to or based on PCIe architectures) that can provide connections among CPUs (e.g., CPU dies 1302 (1) -1302 (N) ) , graphics processor (s) 1326, inference engine (s) 1332, and other components in a multi-chip architecture, such as secure processor (s) 1328 and I / O unit 1304. One or more AMD Infinity FabricTM interconnects 1310 can connect to CPU dies 1302 (1) -1302 (N) and serve as a connection that is used between CPUs. One or more Infinity Fabric connections 1310 can connect each CPU die 1302 to I / O unit 1310.
[0266] In at least one embodiment, processor 1300 can include central processing units (CPUs) and other associated hardware and software described above and further herein. Processor 1300 can also include graphics processor (s) 1326. Graphics processor 1326 can be used for image generation and processing, as well as other computations and operations described further herein. Graphics processor 1326 can be based on RDNA 3 or 3.5 architecture from AMD in Santa Clara, CA. Graphics processor 1326 can include graphics compute dies (GCDs) and memory cache dies (MCDs) . GCDs can include any number of compute units (CUs) for graphics or other processing, such as operations performed by arithmetic logic units (ALUs) that are described further herein. Graphics processor 1326 can include L2 cache that can be used by compute units. MCDs (not shown) can include 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 1342 (1) - (N) , where N is an integer. Components within graphics processor 1326 can be connected using various approaches, such as using Infinity Fabric 1324 interconnects outside or within graphics processor 1326.
[0267] Inference engine 1332 can provide neural processing capabilities for processor 1300 for computational processes that are used for neural networks, deep learning, and other artificial intelligence-related operations described further herein. Processor 1300 can include secure processor (s) 1328 for managing security of processor 1300, display controller 1330 for controlling displays, a system management unit 1334 for managing and operating some or all of the components on processor 1300, multimedia engines 1336 for audio and video operations, fusion controller hub 1338 for managing USB, SATA and PCIe connections to processor 1300, and sensor fusion hub 1340 for managing sensors, such as accelerometers. Processor 1300 can also include memory 1342 (1) - (N) , where N is any integer. Memory can include different memory types, such as LPDDR5 and / or DDR5, or others described elsewhere herein.
[0268] For performing operations described further herein, processor 1300 can include an execution pipeline including a front-end that can include a cache (e.g., L1 cache) that stores instructions (not shown) . Flow of instructions can be modified by a branch predictor. Instructions can be decoded by a decoder, dispatched to a back-end for execution, and renamed. Instruction fetch and decode pipes, for example, can be dispatched to integer or floating point execution operations that can be scheduled by a scheduler and transferred to vector and / or general-purpose registers. Floating point multiplier and / or add operations can be processed, and arithmetic logic units (ALUs) can also be used to perform computations, such as arithmetic and logic operations. Outputs from computation units can be coupled to a load / store queue, which can be connected to cache, such as L1 cache and / or L2 cache.
[0269] With respect to processor 1300 and any of its components described above or elsewhere herein, one or more of APIs or equivalents described herein can, for example, get compiled into instructions or equivalents (e.g., AVX-512 instructions based on an SIMD model) , which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API(s) ) can be stored in any storage outside or inside of processor 1300 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 1300, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.
[0270] In at least one embodiment, processor 1300 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0271] FIG. 14 illustrates an example of a processing core 1400 that may implement Arm architecture (e.g., v9.0-A) or another processor that shares at least some of the components described herein. NeoverseTM V2 core 1400 can be implemented inside a DynamIQ Shared Unit (DSU) cluster via DSU-110 interconnect 1454 for connected one or more cores, e.g., for parallel processing. NeoverseTM V2 core may be implemented as a single core in a DSU cluster that is configured for Direct connect, with or without L3 cache, snoop filter, or Snoop Control Unit (SCU) logic (not shown) . NeoverseTM V2 core can include a CPU bridge 1452 that connects core 1400 to DSU-110 interconnect, which can also connect core 1400 to an external memory system and the rest of a system-on-a-chip. L1 instruction memory system 1402 can fetch instructions from an instruction cache 1404 and deliver instructions (e.g., one or more APIs described herein that may be compiled into instructions) to an instruction decode unit 1410, e.g., to perform some or all of operations described above or elsewhere herein. L1 instruction memory system 1402 may include L1 instruction cache 1404, e.g., with 64-byte cache lines, L1 instruction Translation Lookaside Buffer (TLB) 1406, e.g., with native support for 4KB, 16KB, 64KB, and 2MB page sizes, Macro-Operation Cache (MOP) 1408 (e.g., 1536-entry, 4-way skewed associative L0 MOP cache) , which can contain decoded and optimized instructions for higher performance. Instruction decode unit 1410 can decode AArch64 instructions into internal format. Register rename unit 1412 can perform register renaming to facilitate out-of-order execution and dispatches decoded instructions to various issue queues. Instruction issue unit 1414 can control when decoded instructions may be dispatched to execution pipelines, and it can include issue queues for storing instructions pending dispatch to execution pipelines. Integer execution pipeline 1416 can be included in an execution pipeline and include integer execute unit 1418 that can perform arithmetic and logical data processing operations. Vector execute unit 1420 can be included in an execution pipeline and can perform Advanced SIMD and floating-point operations (FPU) 1422, execute Scalable Vector Extension (SVE) and Scalable Vector Extension 2 (SVE2) instructions 1424, and can optionally execute cryptographic instructions (Crypto) 1426. Advanced SIMD can include media and signal processing architecture that adds instructions primarily for audio, video, 3D graphics, image, and speech processing. A floating-point architecture provides support for single-precision and double-precision floating-point operations. L1 data memory system 1430 can execute load and store instructions, as well as service memory coherency requests. L1 data memory system 1430 can include an L1 data cache 1432 and a fully associative L1 data TLB 1434 with native support for 4KB, 16KB and 64KB page sizes and 2MB and 512MB block sizes. Memory Management Unit (MMU) 1428 can provide fine-grained memory system control through a set of virtual-to-physical address mappings and memory attributes that can be held in translation tables, which can be saved into TLB 1434 when an address is translated. L2 memory system 1436 can include L2 cache 1438, and it can be connected to DSU-110 1454 through an asynchronous CPU bridge 1452. NeoverseTM V2 core 1400 can support a range of debug, test, and trace options including a trace unit 1442 and a trace buffer 1440, and an Embedded Logic Analyzer (ELA) 1448. NeoverseTM V2 core 1400 can implement Statistical Profiling Extension (SPE) 1444 to provide a statistical view of the performance characteristics of executed instructions that software writers can use to optimize their code for better performance. Performance Monitoring Unit (PMU) 1446 can provide performance monitors that can be configured to gather statistics on operation of each core and memory system. Information can be used for debug and code profiling. Generic Interrupt Controller (GIC) CPU interface 1450, 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 one CPU bridge 1452 between each NeoverseTM V2 core 1400 and DSU-110 1454. CPU bridge 1452 can control buffering and synchronization between core 1400 and DSU-1101454. CPU bridge 1452 can be asynchronous to allow different frequency, power, and area implementation points for each core 1400. CPU bridge 1452 can run synchronously without affecting other interfaces such as, but not limited to, debug and trace which can be asynchronous.
[0272] In at least one embodiment, core 1400 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0273] FIG. 15 illustrates one or more chips including one or more tensor processing units (TPUs) 1500, in accordance with at least one embodiment. TPUs 1500 in FIG. 15 can include application specific integrated circuits (ASICs) , e.g., to perform some or all of the operations described above or elsewhere herein, such as, but not limited to, accelerate machine learning workloads performing matrix operations. TPUs 1500 may be ASICs from Alphabet Corporation in Mountain View, CA. 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.
[0274] Chip 1500 can include any number of TPUs that can include tensor cores 1506. Tensor core 1506 can include one or more core sequencer 1508, vector processing unit (VPU) 1510, matrix multiply unit (MXU) 1512 (A) -1514 (N) , where N is any integer greater than 1, and a transpose permute unit 1516. Core Sequencer 1508 can fetch (e.g., VLIW (Very Long Instruction Word) ) instructions from core’s 1506 Instruction Memory (Imem) , execute scalar operations using a scalar data memory (Smem) and scalar registers (Sregs) (not shown) , and forward vector instructions to Vector Processing Unit (VPU) (1510. Instructions can, for example, launch eight operations: two scalar, two vector ALU, vector load and store, and a pair of slots that queue data to and from matrix multiply and transpose units. VPU 1510 can perform vector operations using a large on-chip vector memory (Vmem) , and vector registers (Vregs) . VPU 1510 can stream data to and from MXU through decoupling FIFOs. VPU 1510 can collect and distribute data to Vmem via data-level parallelism (2D matrix and vector functional units) and instruction-level parallelism (8 operations per instruction) . A large two-dimensional matrix multiply unit (MXU) 1512 (A) -1512 (N) can, e.g., use a systolic array to reduce area and energy plus large, software-controlled on-chip memories instead of caches. Transpose Reduction Permute Unit 1516 can do (e.g., 128x128) matrix transposes, reductions, and permutations of VPU 1510 lanes. High Bandwidth Memory 1504 can be used for applications on chip, and it can be coupled to host queue (s) 1502, e.g., over PCIe. One or more chips 1500 can be connected together for computing. For example, one or more chips 1500 can be connected as a torus, e.g., a 2D torus. Chip 1500 can also include any number (e.g., four) Inter-Core Interconnect (ICI) links 1518 that can enable direct connections between chips to form a supercomputer.
[0275] With respect to any processors in chip 1500 and any of its components described above or elsewhere herein, one or more of APIs or equivalents described herein can, for example, get compiled into instructions or equivalents, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of any processors in chip 1500 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of any processors in chip 1500, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.
[0276] In at least one embodiment, chip 1500 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0277] FIG. 16 illustrates a vector processor, in accordance with at least one embodiment. Vector processor 1600 may support a RISC-V standard. Vector processor 1600 can include one more cores 1610 (e.g., scalar units) with one or more Vector Processing Units (VPUs) 1642 (e.g., vector units) that can, e.g., perform some or all of the operations described above or elsewhere herein. Core 1610 may include Andes Custom Extension (ACE) 1616 that can be used for communication of customized instructions for processor 1600, for example, via ACP 1638. Core 1610 may include 1-cycle multiplier and 1-cycle instruction / data local memory (ILM / DLM) for increased parallelism by allowing simultaneous instruction fetches and data accesses. Memory management unit (MMU) 1624 may manage system memory and cache, and provide for branch execution, issuance of instruction pairs, L1 instruction / data caches and local memory storage. Core 1610 can include Physical memory protection and programmable physical memory attribute unit (PMP / PPMA) 1622. Core 1610 can include a digital signal processor (DSP) 1628, and a floating-point unit (FPU) 1626 as well as load-store unit (LSU) 1632 to interface with memory hierarchy (D$ 1634 and I$ 1630) . Core 1610 can include branch prediction unit 1618 and multiplier unit 1620.
[0278] Vector processing unit (VPU) 1642 can include one or more vector functional units (FUs) 1646 (A) -1646 (N) that can be chained together for parallel processing, independent memory paths for RISC-V vector (RVV) load / store via ACE-RVV 1648 and Andes Streaming port (ASP) 1644 load / store, and a vector load / store unit (VLSU) 1650.
[0279] Vector processor 1600 can include bus interfaces, such as, but not limited to, L2 cache memory port 1656 for cacheable access, a MMIO port 1654 for non-cacheable access, an input-output coherence Port (IOCP) 1658 for cacheless bus master, local memory access ports for ILM / DLM 1612, which can be coupled to SRAM 1606, and high-bandwidth vector memory (HVM) 1636 access, a shared peripheral port (SPP) 1652 for external peripherals. Other memory ports include LM slave port AXI 1602, HVM subordinate port AXI 1604, MEM (AXI) 1662, and AXI 1660. Trace I / F 1614 can capture, encode, and transmit off-chip via Inst. Trace I / F 1608, e.g., a record of executed processor instructions, which software tools can use to reconstruct the exact execution sequence of a program.
[0280] With respect to any processors in processor 1600 and any of its components described above or elsewhere herein, one or more of APIs or equivalents described herein can, for example, get compiled into instructions or equivalents, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of processor 1600 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 1600, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.
[0281] In at least one embodiment, vector processor 1600 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0282] FIG. 17A illustrates a diagram of an example many-core tiled processor microarchitecture. Many-core tiled processor in FIG. 17A can include a language processing processor. As illustrated in FIG. 17A, each “tile” of a processor architecture is a processing element tied together using a network-on-chip (NoC) that can be used, e.g., to perform some or all of the operations described above or elsewhere herein. For example, each tile may have an instruction dispatch 1704 and an integer (INT) 1706 and floating-point (FP) unit 1708 as well as load-store unit (LSU) 1712 to interface with memory hierarchy (data cache (D$) 1710 and instruction cache (I$) 1714) and network (NET) 1716 interface for communication with other tiles. Some tiles in processor 1700 may include memory controller 1702 for managing and controlling memory, as described further herein. Processor 1700 can have a functional slice architecture. Processor 1700 may be located on an application specific integrated circuit (ASIC) , and FIG. 17A may represent a layout of an ASIC. Processor 1700 can include a co-processor that is designed to execute instructions for a predictive model. A predictive model is any model that is configured to make a prediction from input data. A predictive model can use a classifier to make a classification prediction. A predictive model may be a machine learning model such as, but not limited to, a tensor flow model, and processor 1700 is a tensor streaming processor.
[0283] Processor 1700 can employ different microarchitectures, which disaggregates functional units shown in each tile in FIG. 17B. Instead, functional tiles 1724 of processor 1700 may be aggregated into a plurality of functional process units (hereafter referred to as “slices” ) 1704, each corresponding to a particular function type (e.g., FP / INT 1718, NET 1720, MEM 1722) . For example, as illustrated in FIG. 17B, each slice may correspond to a column of functional tiles extending in a north-south direction. In addition, processor 1700 also may include communication lanes to carry data between tiles of different slices, each running horizontally in an east-west direction. Each communication lane may be connected to each of slices 1704 of processor 1700.
[0284] Slices 1704 of processor 1700 may each correspond to a different function, and may include arithmetic logic slices (e.g., FP / INT1718) , lane switching slices (e.g., NET 1720) , and memory slices (e.g., MEM 1722) . Arithmetic logic units may execute one or more arithmetic and / or logic operations on data received via communication lanes to generate output data. Examples of arithmetic logic units may be matrix multiplication units and vector multiplication units. Memory slices include memory cells that store data. Memory slices can provide data to other slices through communication lanes. Memory slices can also receive data from other slices through communication lanes. Lane switching slices can configurably route data from one communication lane to any other communication lane. For example, data from a first lane can be provided to a second lane through a lane switching slice. In some embodiments, a lane switching slice can be implemented as a crossbar switch. Each slice 1704 also includes its own instruction queue (not shown) that stores instructions, and an instruction control unit (ICU) to control execution of instructions. Instructions in a given instruction queue may be executed only by tiles in its associated functional slice and may not be executed by other slice (s) of processor 1700.
[0285] By arranging tiles of processor 1700 into different functional slices 1704, on-chip instruction and control flow of processor 1700 can be decoupled from data flow. For example, one arrow in FIG. 17B illustrates flow of instructions within processor architecture, in accordance with some embodiments. Another arrow in FIG. 17B illustrates data flow within processor architecture, in accordance with at least one embodiment. As illustrated, instructions and control flow can flow in a first direction across tiles of processor 1700 (e.g., north-south, along a length of functional slices, as shown by the first arrow) , while data flows flow in a second direction across tiles of processor 1700 (e.g., east-west, across functional slices, as shown by the second arrow) that is perpendicular to the first direction.
[0286] Different functional slices of processor 1700 may correspond to MEM 1722 (memory) , VXM (vector execution module) , MXM (matrix execution module) , NIM (numerical interpretation module) , and SXM (switching and permutation module) . Each slice may include N tiles that may all be controlled by a same instruction control unit (ICU) (not shown) . Each slice may operate completely independently and can only be coordinated using barrier-like synchronization primitives or through a compiler by exploiting “tractable determinism. ” Each tile of processor 1700 can correspond to an execution unit organized as an ×M SIMD tile. For example, each tile of on-chip memory of processor 1700 may be organized to store an L-element vector atomically. As such, a MEM slice having N tiles may work together to store or process a large vector (e.g., having a total of N×M elements) .
[0287] Tiles in a slice may execute instructions in a “staggered” fashion where instructions may be issued tile-by-tile within a slice over a period of N cycles. Functional slices may be arranged physically on-chip to allow efficient data-flow for pipelined execution across hundreds of cycles for common patterns. Data flows can perform a single “u-turn” (change in direction) corresponding to a single matrix operation before being written back to memory, in some embodiments, a particular data flow may change direction multiple times (due to multiple matrix and vector operations) before resulting data is written back into memory.
[0288] When using processor 1700 (e.g., TSP) having a functional slice architecture, TSP compiler (not shown) generates an explicit plan for how processor 1700 can execute a program (e.g., a microprogram) . Compiler can specify when each operation will be executed, which functional slices will perform work, and which STREAM registers hold operands. Compiler can maintain a high-fidelity (cycle accurate) model of processor 1700 (e.g., TSP) hardware state so a microprogram can orchestrate data flow.
[0289] Processor 1700 (e.g., TSP) can use a Web-hosted compiler that takes as its input a model (e.g., a ML model such as, but not limited to, a TensorFlow model) and emits a proprietary instruction stream targeting processor 1700 (e.g., TSP) . Compiler is responsible for coordinating control and data flow of a program, and specifies any instruction-level parallelism by explicitly bundling instructions that can and should execute concurrently so that they may be dispatched together. Primary hardware structure includes an architecturally-visible streaming register file (STREAMs) , described in greater detail below, which serves as a conduit through which operands flow from MEM slices (e.g., SRAM) to functional slices and vice versa.
[0290] MEM 1722 of processor 1700 can serve as: (1) storage for model parameters, microprograms and data on which they operate, and (2) network-on-chip (NoC) for communicating data operands from MEM to functional slices and computed results back to MEM. In some embodiments, on-chip memory can consumes ≈75%of chip area of processor 1700. In some embodiments, due to bandwidth requirements of processor 1700, on-chip memory of MEM tiles may include SRAM, and not DRAM. On-chip memory capacity of processor 1700 can determine (i) number of ML models that can simultaneously reside on-chip, (ii) size of any given model, and (iii) partitioning of large models to fit into multi-chip systems. In some embodiments, MEM system of processor 1700 can provide a plurality of memory slices organized into two different hemispheres (referred to as “MEM WEST” and “MEM EAST” , respectively) .
[0291] Memory slices of each hemisphere may be mirrored, such that slices may be physically numbered {0, ... L} in an East hemisphere, and {L, ... 0} in a West hemisphere, such that memory slice 0 for each hemisphere corresponds to a slice closest to VXM slices between hemispheres, where each hemisphere comprises L slices. Direction of data transfer towards the center of a chip may be referred to as inwards, while data transfer toward the outer (Eastern or Western most) edge of a chip may be referred to as outwards. Although hemispheres of memory of processor 1700 may be referred to as east and west, it is understood that in other embodiments, other names may be used to refer to different hemispheres of memory.
[0292] In some embodiments, a streaming register file, referred to as STREAMS, transfers operands and results between SRAM of MEM slices and functional slices of processor 1700. In some embodiments, a plurality of MEM slices (e.g., between 2 and 10 adjacent MEM slices) may be physically organized as a set. Each set of slices may be located between a pair of STREAM register files, such that each slice is able to read or write to STREAM registers in either direction. By placing STREAM register files between sets of MEM slices, a number of cycles needed for data operands to be transmitted across a hemisphere is decreased (e.g., by a factor corresponding to a number of slices per set) . A number of slices per set may be configured based upon a distance over which data may be transmitted over a single clock cycle.
[0293] With respect to any processors in FIG. 17 and any components described above or elsewhere herein, one or more of APIs or equivalents described herein can, for example, get compiled into instructions or equivalents, which may be fetched by instruction fetch logic or equivalents, decoded by a processor decoder or equivalents, scheduled (e.g., in order or out of order) for execution by a scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of processor 1700 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 1700, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.
[0294] In at least one embodiment, processor 1700 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.SOFTWARE CONSTRUCTIONS
[0295] The following figures set forth, without limitation, examples of software constructs for implementing at least one embodiment.
[0296] FIG. 18 illustrates a software stack of a programming platform, in accordance with at least one embodiment. A programming platform can include a platform for leveraging hardware on a computing system to accelerate computational tasks. A programming platform may be accessible to software developers through libraries, compiler directives, and / or extensions to programming languages, in at least one embodiment. A programming platform may be CUDA, Radeon Open Compute Platform ( “ROCm” ) , OpenCL (OpenCLTM is developed by Khronos group) , SYCL, or Intel oneAPI.
[0297] A software stack 1800 of a programming platform can provide an execution environment for an application 1801. Application 1801 may include any computer software capable of being launched on software stack 1800. Application 1801 may 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.
[0298] Application 1801 and software stack 1800 run on hardware 1808. Hardware 1808 may include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of compute devices that support a programming platform. Software stack 1800 may be vendor specific and compatible with only devices from particular vendor (s) , such as CUDA, ROCm, OneAPI, OpenCL, or other implementations. Hardware 1808 can include a host connected to one more devices that can be accessed to perform computational tasks via application programming interface ( “API” ) calls. A device within hardware 1808 may include a GPU, FPGA, AI engine, or other compute device (but may also include a CPU) and its memory, as opposed to a host within hardware 1808 that may include a CPU (but may also include a compute device) and its memory, in at least one embodiment. With respect to any hardware 1808 described above or elsewhere herein, one or more of APIs described herein can, for example, get compiled into instructions, which may be fetched by instruction fetch logic, decoded by a processor decoder, scheduled (e.g., in order or out of order) for execution by a scheduler, executed by execution logic, reordered, and then retired by retirement logic. API (s) (and / or compiled instructions including API (s) ) can be stored in any storage outside or inside of hardware 1808 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of hardware 1808, including registers, DRAM, flash, SRAM, cache, or other memory. One or more of APIs described herein can receive a call. One or more of APIs described herein can communicate with a library or a portion of a library to perform a function described by the call. One or more of APIs described herein can receive a call and communicate with a library or portion of a library to perform a function described by the call.
[0299] Software stack 1800 of a programming platform can include a number of libraries 1803, a runtime 1805, an optional driver / interface 1807, and a device kernel driver 1808. Each of libraries 1803 may include data and programming code that can be used by computer programs and leveraged during software development. Libraries 1803 may include pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. Libraries 1803 can include functions that may be optimized for execution on one or more types of devices. Libraries 1803 may include functions for performing mathematical, deep learning, and / or other types of operations on devices. Libraries 1803 can be associated with corresponding APIs 1802, which may include one or more APIs, that expose functions implemented in libraries 1803. A processor (e.g. CPU, GPU) may perform, call, or otherwise use one or more APIs to prioritize kernels. For example, a first kernel (e.g., parent) can launch a second kernel (e.g., child kernel) , and said second kernel can be used by a processor to launch additional kernels (e.g., grandchildren kernels) independent of said first kernel. A processor may perform an API or calls an API from memory to be performed to support dynamic stream priority (e.g., updating priority while a stream is being used to perform operations) . For example, when a processor performs said API, it allows a programmer to copy stream priority from one stream to one or more other streams.
[0300] Software stack 1800 may include an API to support dynamic stream priority (e.g., updating priority while a stream is being used to perform operations) , which can allow a programmer to set priority of a stream at any time after creation. Software stack 1800 can include an API to support dynamic stream priority (e.g., updating priority while the stream is being used to perform operations) , which may allow a programmer to obtain current priority of a stream, where the priority is one of a plurality of attributes of a stream. Software stack 1800 can include an API to support dynamic stream priority (e.g., updating priority while the stream is being used to perform operations) , which may allow a programmer to obtain current priority of a stream as a single attribute. Software stack 1800 can include an API to support dynamic stream priority (e.g., updating priority while the stream is being used to perform operations) , which allows a programmer to launch a kernel to perform operations on a stream at a set priority, which may be different from the stream priority. Software stack 1800 may include an API to indicate whether an object (e.g., a thread synchronization object such as, but not limited to, a barrier) tracks whether all data movement operations for a set of threads operating on a GPU may be complete has a specified state after a specified period of time, where a specified state can be a state indicating that data has been moved and is ready for use, and is specified using an expected parity value as an input to the API.
[0301] Software stack 1800 can include one or more APIs to updated kernels. A processor can perform an API or call an API from memory to be performed to update to an existing API is to support context-free kernels, which may allow 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. Software stack 1800 may include one or more APIs to allow a programmer to obtain a kernel identifier and a graphics context as separate parameters from a kernel node, so that parameters to be obtained from kernels and from context-free kernels. Software stack 1800 can include one or more APIs to use parallel processor (s) , such as, but not limited to, one or more graphics processing units, to launch task graphs (e.g., task graphs) and to execute one or more task graphs (e.g., including one or more programs) .
[0302] Software stack 1800 may include one or more APIs to associate one or more instructions with one or more memory ordering operations, such as, but not limited to, a fence or membar operation. Instructions can be associated with one or more domains such that a memory ordering operation is executed in association to one or more particular domains without interfering with instructions of other domains. An API can indicate a thread has arrived (e.g., at a thread synchronization barrier) , or finished a stage of work in relation to asynchronous data movement operations on a GPU. Software stack 1800 may include one or more to allow programmers to manually indicate an expected transaction count when a thread has finished a stage of work, which can be used to update an object that tracks whether all data movement operations for a set of threads may be complete.
[0303] Application 1801 can be written as source code that is compiled into executable code, as discussed in greater detail below in conjunction with FIGs. 19 and 20. Executable code of application 1801 may run, at least in part, on an execution environment provided by software stack 1800. During execution of application 1801, code may be reached that needs to run on a device, as opposed to a host. In such a case, runtime 1805 may be called to load and launch requisite code on a device. Runtime 1805 may include any technically feasible runtime system that is able to support execution of application 1801.
[0304] Runtime 1805 can be implemented as one or more runtime libraries associated with corresponding APIs, which are shown as API (s) 1804. One or more of such runtime libraries may include functions for memory management, execution control, device management, error handling, and / or synchronization, among other things, . Memory management functions may include functions to allocate, deallocate, and copy device memory, as well as transfer data between host memory and device memory. Execution control functions may include functions to launch a function (sometimes referred to as a “kernel” when a function is a global function callable from a host) on a device and set attribute values in a buffer maintained by a runtime library for a given function to be executed on a device.
[0305] Runtime libraries and corresponding API (s) 1804 may be implemented in any technically feasible manner. One (or any number of) API may expose a low-level set of functions for fine-grained control of a device, while another (or any number of) API may expose a higher-level set of such functions. A high-level runtime API may be built on top of a low-level API. One or more of runtime APIs may be language-specific APIs that may be layered on top of a language-independent runtime API.
[0306] An optional driver or interface 1807 may be implemented, e.g., for CUDA and ROCm implementations, that are described further below. Optional driver / interface 1807 may be associated with optional driver or interface API (s) , such as, but not limited to, CUDA and / or ROCm API (s) .
[0307] One or more processors disclosed in “processing systems” can perform, access, or otherwise use software stack 1800. For example, system-on-a-chip 500, parallel processor 600, graphics multiprocessor 634, processor 700, processor 800, accelerator 900, neuromorphic processor 1005, supercomputer 1100, acceleration processing unit 1200, processor 1300, processor 1400, tensor processing unit 1500, processor 1600, and language processing unit 1700 can perform, use, call, or otherwise implement (e.g., through accessing a memory) one or more APIs included in software stack 1800.
[0308] Device kernel driver 1808 can be configured to facilitate communication with an underlying device. Device kernel driver 1808 may provide low-level functionalities upon which APIs, such as, but not limited to, API (s) 1804, and / or other software relies. Device kernel driver 1808 may be configured to compile intermediate representation ( “IR” ) code into binary code at runtime. For CUDA or other implementations such as, but not limited to, ROCm, OneAPI, or OpenCL, device kernel driver 1808 may compile Parallel Thread Execution ( “PTX” ) IR code that is not hardware specific into binary code for a specific target device at runtime (with caching of compiled binary code) , which is also sometimes referred to as “finalizing” code. Doing so may permit finalized code to run on a target device, which may not have existed when source code was originally compiled into PTX code. Alternatively, device source code may be compiled into binary code offline, without requiring device kernel driver 1808 to compile IR code at runtime.
[0309] Processors described elsewhere herein, such as, but not limited to, processors in FIGs. 5-17 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., software stack 1800 to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0310] In accordance with at least one embodiment, software stack 1800 of FIG. 18 can be performed in a CUDA implementation. A CUDA software stack 1800, on which an application 1801 may be launched, may include CUDA libraries 1803, a CUDA runtime 1805, a CUDA driver 1807, and a device kernel driver 1808. CUDA software stack 1800 can execute on hardware (e.g., graphics multiprocessor 634 that may include a GPU that supports CUDA and is developed by NVIDIA Corporation of Santa Clara, CA.
[0311] Application 1801, CUDA runtime 1805, and device kernel driver 1808 can perform functionalities that are described above and elsewhere herein. CUDA driver 1807 can include a library (libcuda. so) that may implement a CUDA driver API 1806. Similar to a CUDA runtime API 1804 implemented by a CUDA runtime library (cudart) , CUDA driver API 1806 may expose functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability, among other things. CUDA driver API 1806 can differ from CUDA runtime API 1804 in that CUDA runtime API 1804 simplifies device code management by providing implicit initialization, context (analogous to a process) management, and module (analogous to dynamically loaded libraries) management. In contrast to high-level CUDA runtime API 1804, CUDA driver API 1806 can be a low-level API providing more fine-grained control of a device, particularly with respect to contexts and module loading. CUDA driver API 1806 may expose functions for context management that may be not exposed by CUDA runtime API 1804. CUDA driver API 1806 may also be language-independent and support, e.g., OpenCL, in addition to CUDA runtime API 1804. Further, development libraries, including CUDA runtime 1805, may be considered as separate from driver components, including user-mode CUDA driver 1807 and kernel-mode device driver 1808 (also sometimes referred to as a “display” driver) .
[0312] CUDA libraries 1803 may include mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries, which parallel computing applications such as, but not limited to, application 1801 may utilize. CUDA libraries 1803 may include mathematical libraries such as, but not limited to, a cuBLAS library that is an implementation of Basic Linear Algebra Subprograms ( “BLAS” ) for performing linear algebra operations, a cuFFT library for computing fast Fourier transforms ( “FFTs” ) , and a cuRAND library for generating random numbers, among others. CUDA libraries 1803 may 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.
[0313] In at least one embodiment, processors described elsewhere herein, such as, but not limited to, processors in FIGs. 5-17 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., software stack 1800 to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0314] In accordance with at least one embodiment, software stack 1800 of FIG. 18 can be performed in a ROCm implementation. A ROCm software stack 1800, on which an application 1801 may be launched, includes a language runtime 1803, a system runtime 1805, a thunk 1807, and a ROCm kernel driver 1808. ROCm software stack 1800 executes on hardware 1809, which may include a GPU that supports ROCm and is developed by AMD Corporation of Santa Clara, CA.
[0315] Application 1801 may perform similar functionalities as discussed above in conjunction with FIG. 18. In addition, language runtime 1803 and system runtime 1805 may perform similar functionalities as runtime 1805 discussed above in conjunction with FIG. 18. Language runtime 1803 and system runtime 1805 may differ in that system runtime 1805 is a language-independent runtime that implements a ROCr system runtime API 1804 and makes use of a Heterogeneous System Architecture ( “HSA” ) Runtime API. HSA runtime API can include a thin, user-mode API that exposes interfaces to access and interact with an AMD GPU, including functions for memory management, execution control via architected dispatch of kernels, error handling, system and agent information, and runtime initialization and shutdown, among other things. In contrast to system runtime 1805, language runtime 1803 can be an implementation of a language-specific runtime API 1802 layered on top of ROCr system runtime API 1804. Language runtime API may include a Heterogeneous compute Interface for Portability ( “HIP” ) language runtime API, a Heterogeneous Compute Compiler ( “HCC” ) language runtime API, or an OpenCL API, among others. HIP language in particular is an extension of C++ programming language with functionally similar versions of CUDA mechanisms, and a HIP language runtime API may include functions that may be similar to those of CUDA runtime API discussed above in conjunction with FIG. 18, such as, but not limited to, functions for memory management, execution control, device management, error handling, and synchronization, among other things.
[0316] Thunk (ROCt) 1807 can be an interface 1806 that can be used to interact with underlying ROCm driver 1808. ROCm driver 1808 can be a ROCk driver, which is a combination of an AMDGPU driver and a HSA kernel driver (amdkfd) . AMDGPU driver can be a device kernel driver for GPUs developed by AMD that performs similar functionalities as device kernel driver 1809 discussed above in conjunction with FIG. 18. HSA kernel driver can be a driver permitting different types of processors to share system resources more effectively via hardware features.
[0317] Various libraries (not shown) may be included in ROCm software stack 1800 above language runtime 1803 and provide functionality similar to CUDA libraries 1803, discussed above in conjunction with FIG. 18. Various libraries may include mathematical, deep learning, and / or other libraries such as, but not limited to, a hipBLAS library that implements functions similar to those of CUDA cuBLAS, a rocFFT library for computing FFTs that is similar to CUDA cuFFT, among others.
[0318] Processors described elsewhere herein, such as, but not limited to, processors in FIGs. 5-17 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., software stack 1800 to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0319] In accordance with at least one embodiment, software stack 1800 of FIG. 18 can be performed in a OpenCL implementation. An OpenCL software stack 1800, on which an application 1801 may be launched, can include an OpenCL framework 1803, an OpenCL runtime 1805, and a driver 1808. OpenCL software stack 1800 may execute on hardware 1809 that is not vendor-specific. As OpenCL is supported by devices developed by different vendors, specific OpenCL drivers may be required to interoperate with hardware from such vendors.
[0320] Application 1801, OpenCL runtime 1805, device kernel driver 1808, and hardware 1809 may perform similar functionalities as other implementations of application 1801, runtime 1805, device kernel driver 1808, and hardware 1809, respectively, that are discussed above in conjunction with FIG. 18. Application 1801 can further include an OpenCL kernel (not shown) with code that is to be executed on a device.
[0321] OpenCL may define a “platform” that allows a host to control devices connected to a host. An OpenCL framework can provide a platform layer API and a runtime API, shown as platform API 1802 and runtime API 1804. Runtime API 1804 can use contexts to manage execution of kernels on devices. Each identified device may be associated with a respective context, which runtime API 1804 may use to manage command queues, program objects, and kernel objects, share memory objects, among other things, for that device. Platform API 1802 can expose functions that permit device contexts to be used to select and initialize devices, submit work to devices via command queues, and enable data transfer to and from devices, among other things. In addition, OpenCL framework can provide various built-in functions (not shown) , including math functions, relational functions, and image processing functions, among others.
[0322] A compiler (not shown) can also be included in OpenCL framework 1803. Source code may be compiled offline prior to executing an application or online during execution of an application. In contrast to CUDA and ROCm, OpenCL applications may be compiled online by a compiler that is representative of any number of compilers that may 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 may be compiled offline, prior to execution of such applications.
[0323] In at least one embodiment, processors described elsewhere herein, such as, but not limited to, processors in FIGs. 5-17 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., software stack 1800 to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0324] In accordance with at least one embodiment, software can be supported by a programming platform that is configured to support various programming models, middlewares and / or libraries, and frameworks that an application may rely upon. Application may be an AI / ML application implemented using, for example, 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, NVIDIA Collective Communications Library ( “NCCL” ) , and / or NVIDA Developer Data Loading Library ( “DALI” ) CUDA libraries to provide accelerated computing on underlying hardware.
[0325] Programming platform may be one of a CUDA, ROCm, or OpenCL platform described above in conjunction with FIG. 18. Programming platform can support multiple programming models, which may be abstractions of an underlying computing system permitting expressions of algorithms and data structures. Programming models may expose features of underlying hardware in order to improve performance. Programming models may include CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism ( “C++AMP” ) , Open Multi-Processing ( “OpenMP” ) , Open Accelerators ( “OpenACC” ) , and / or Vulkan Compute.
[0326] Libraries and / or middlewares may provide implementations of abstractions of programming models. Such libraries can include data and programming code that may be used by computer programs and leveraged during software development. Such middlewares can include software that provides services to applications beyond those available from programming platform. Libraries and / or middlewares may include cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. In addition, libraries and / or middlewares may include NCCL and ROCm Communication Collectives Library ( “RCCL” ) libraries providing communication routines for GPUs, a MIOpen library for deep learning acceleration, and / or an Eigen library for linear algebra, matrix and vector operations, geometrical transformations, numerical solvers, and related algorithms.
[0327] Application frameworks may depend on libraries and / or middlewares. Each of application frameworks can be a software framework used to implement a standard structure of application software. Returning to the AI / ML example discussed above, an AI / ML application may be implemented using a framework such as, but not limited to, Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet deep learning frameworks, for example.
[0328] In at least one embodiment, processors described elsewhere herein, such as, but not limited to, processors in FIGs. 5-17 can include one or more circuits to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software, e.g., programming platforms described herein, to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data or otherwise perform any of the operations described above or elsewhere herein.
[0329] FIG. 19 illustrates compiling code to execute on one of programming platforms of FIG. 18 described above, in accordance with at least one embodiment. A compiler 1901 is configured to receive source code 1900, compile source code 1900, and output an executable file 1910. Complier 1901 can be configured to convert source code 1900 into host executable code 1907 for execution on a host and device executable code 1908 for execution on a device. Source code 1900 may either be compiled offline prior to execution of an application, or online during execution of an application. Source code 1900 may include code in any programming language supported by compiler 1901, such as, but not limited to, C++, C, Fortran, etc. Source code 1900 may be included in a single-source file having a mixture of host code and device code, with locations of device code being indicated therein. A single-source file may be a . cu file that includes CUDA code or a . hip. cpp file that includes HIP code or a file in another format that includes both host code and device code. Alternatively, source code 1900 may include multiple source code files, rather than a single-source file, into which host code and device code may be separated. Compiler 1901 includes or has access to one or more libraries to recognize a sequence of API calls to perform a single fused API, where a single fused API is a combined API for two or more APIs. In at least one embodiment, compiler 1901 may be an NVIDIA CUDA compiler ( “NVCC” ) for compiling CUDA code in . cu files, or a HCC compiler for compiling HIP code in . hip. cpp files, or other compilers.
[0330] Compiler 1901 can be configured to compile source code 1900 into host executable code 1907 for execution on a host and device executable code 1908 for execution on a device. Compiler 1901 performs operations including parsing source code 1900 into an abstract system tree (AST) , performing optimizations, and generating executable code. When source code 1900 includes a single-source file, compiler 1901 may separate device code from host code in such a single-source file, compile device code and host code into device executable code 1908 and host executable code 1907, respectively, and link device executable code 1908 and host executable code 1907 together in a single file.
[0331] Compiler 1901 can include a compiler front end 1902, a host compiler 1905, a device compiler 1906, and a linker 1909. Compiler front end 1902 can be configured to separate device code 1904 from host code 1903 in source code 1900. Device code 1904 may be compiled by device compiler 1906 into device executable code 1908, which as described may include binary code or IR code, in at least one embodiment. Separately, host code 1903 may be compiled by host compiler 1905 into host executable code 1907. For NVCC other compilers, such as, but not limited to, those for oneAPI, ROCm, and OpenCL, host compiler 1905 may be a general purpose C / C++ compiler that outputs native object code, while device compiler 1906 may be a Low Level Virtual Machine ( “LLVM” ) -based compiler that forks a LLVM compiler infrastructure and outputs PTX code or binary code. For HCC, both host compiler 1905 and device compiler 1906 may be LLVM-based compilers that output target binary code.
[0332] Subsequent to compiling source code 1900 into host executable code 1907 and device executable code 1908, linker 1909 can link host and device executable code 1907 and 1908 together in executable file 1910. Native object code for a host and PTX or binary code for a device may be linked together in an Executable and Linkable Format ( “ELF” ) file, which is a container format used to store object code. Host executable code 1907 and device executable code 1908 may be in any suitable format, such as, but not limited to, binary code and / or IR code. In the case of CUDA, host executable code 1907 may include native object code and device executable code 1908 may include code in PTX intermediate representation, in at least one embodiment. In the case of ROCm, both host executable code 1907 and device executable code 1908 may include target binary code, in at least one embodiment. Other implementations, such as, but not limited to, oneAPI, OpenCL are contemplated and can be performed similarly to the CUDA and ROCm implementations above.
[0333] Source code 1900 may be translated prior to compiling source code. Source code is passed through a translation tool (not shown) , which translates source code 1900 into translated source code. A compiler 1901 can be used to compile translated source code into host executable code 1907 and device executable code 1908 in a process that is similar to compilation of source code 1900 by compiler 1901 into host executable code 1907 and device executable code 1908, as discussed above in conjunction with FIG. 19.
[0334] A translation performed by translation tool can be used to port source code 1900 for execution in a different environment than that in which it was originally intended to run. Translation tool may include a HIP translator that is used to “hipify” CUDA code intended for a CUDA platform into HIP code that can be compiled and executed on a ROCm platform. Translation of source code 1900 may include parsing source code 1900 and converting calls to API (s) provided by one programming model (e.g., CUDA) into corresponding calls to API(s) provided by another programming model (e.g., HIP) , as discussed in greater detail below in conjunction with FIG. 20. Returning to the example of hipifying CUDA code, calls to CUDA runtime API, CUDA driver API, and / or CUDA libraries may be converted to corresponding HIP API calls. Automated translations performed by translation tool 1901 may sometimes be incomplete, requiring additional, manual effort to fully port source code 1900.
[0335] One or more techniques described herein may utilize a variety of methods for converting one type of code to another type of code. For example, compiler 1901 or other compilers described herein can convert a high-level language (e.g., source code that is abstract to hardware) to a lower-level language (e.g., machine code or an intermediate representation) . Source code can be scanned, parsed, transformed into an abstract syntax tree semantically analyzed, then converted into an intermediate code, and then converted into machine code or assembly language. Compiler 1901 or other compilers described herein can include a transpiler, which can convert, for example, one type of source code to another type of source code or one type of machine code to another type of machine code. Source code can be parsed, and transformed into an abstract syntax tree, which can then be converted to an intermediate model that can be transformed into an abstract syntax tree of target language and code can be generated. Compiler 1901 or other compilers described herein can be used to enable interchangeability between different device architectures. For example, an application for one platform (e.g., a CUDA application) can be compiled into code for implementation on another platform (e.g., an AMD processor, Intel processor, or other processor) . Source code 1900 can include source code for one platform (e.g., CUDA) . Compiler 1901 can compile the source 1900 into an executable file 1910 that can be used by another platform (e.g., AMD or Intel) . Programming toolkits can allow applications for one platform (e.g., CUDA) to be compiled (e.g., natively) for another platform (e.g., AMD or Intel) . For example, a GPGPU programming toolkit can allow for CUDA applications to be natively compiled for AMD GPUs. Programs (e.g., CUDA programs) or its build system do not have to be modified or translated to another language before compiling to code for another platform. A compiler may accept the same command-line options and programming dialect (e.g., CUDA dialect) as another compiler (e.g., nvcc for CUDA) , serving as a drop-in replacement to impersonate an installation of a toolkit (e.g., NVIDIA CUDA Toolkit) , so existing build tools and scripts (e.g., like cmake) work 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. An example implementation includes SCALE from Spectral Compute in London, England. SCALE can allow programs written using CUDA language to be directly compiled to lower-level language (e.g., machine code) for AMD GPUs. SCALE can create one or more directories that can be used to impersonate NVIDIA CUDA Toolkit (from the point of view of a build system) by instructing a build system that a CUDA installation path is one provided by SCALE, rather than the one provided by NVIDIA. Additional implementations can include a Clang compiler that can provide a language front-end and tooling infrastructure for languages in the C language family (C, C++, Objective C / C++, OpenCL, CUDA, and RenderScript) . In at least one embodiment, compilers and / or transpilers described herein, such as, but not limited to compiler 1901, compiler 1905, and / or compiler 1906 can include one or more circuits to compile code (e.g., CUDA, HIP, OpenCL, OneAPI, or others) to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data and / or perform any of the operations described above or elsewhere herein. In at least one embodiment, compilers and / or transpilers described herein, such as, but not limited to compiler 1901, compiler 1905, and / or compiler 1906 can include one or more circuits to convert code (e.g., source code for CUDA) to one or more other types of code (e.g., machine code for CUDA and / or another platform, such as AMD or Intel processors) to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data and / or perform any of the operations described above or elsewhere herein.
[0336] FIG. 20 illustrates a system 2000 configured to compile and execute CUDA source code 2010 using different types of processing units, in accordance with at least one embodiment. System 2000 includes CUDA source code 2010, a CUDA compiler 2050, host executable code 2070 (1) , host executable code 2070 (2) , CUDA device executable code 2084, a CPU 2090, a CUDA-enabled GPU 2094, a GPU 2092, a CUDA to HIP translation tool 2020, HIP source code 2030, a HIP compiler driver 2040, an HCC 2060, and HCC device executable code 2082.
[0337] CUDA source code 2010 may be a collection of human-readable code 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, is executable in parallel on a device. A device may be a processor that is optimized for parallel instruction processing, such as, but not limited to, CUDA-enabled GPU 2090, GPU 2092, or another GPGPU, etc. Host code is source code that, after compilation, is executable on a host. A host is a processor that is optimized for sequential instruction processing, such as, but not limited to, CPU 2090.
[0338] CUDA source code 2010 can include any number (including zero) of global functions 2012, any number (including zero) of device functions 2014, any number (including zero) of host functions 2016, and any number (including zero) of host / device functions 2018. Global functions 2012, device functions 2014, host functions 2016, and host / device functions 2018 may be mixed in CUDA source code 2010. Each of global functions 2012 may be executable on a device and callable from a host. One or more of global functions 2012 may therefore act as entry points to a device. Each of global functions 2012 can be a kernel. In a technique known as dynamic parallelism, one or more of global functions 2012 can define a kernel that is executable on a device and callable from such a device. A kernel can be executed N (where N is any positive integer) times in parallel by N different threads on a device during execution.
[0339] Each of device functions 2014 can be executed on a device and callable from such a device only. Each of host functions 2016 can be executed on a host and callable from such a host only. Each of host / device functions 2016 may define both a host version of a function that is executable on a host and callable from such a host only and a device version of the function that is executable on a device and callable from such a device only.
[0340] CUDA source code 2010 may also include any number of calls to any number of functions that may be defined via a CUDA runtime API 2002. CUDA runtime API 2002 may include any number of functions that execute on a host to allocate and deallocate device memory, transfer data between host memory and device memory, manage systems with multiple devices, etc. CUDA source code 2010 may also include any number of calls to any number of functions that may be specified in any number of other CUDA APIs. A CUDA API may be any API that is designed for use by CUDA code. CUDA APIs can include CUDA runtime API 2002, a CUDA driver API, APIs for any number of CUDA libraries, etc, including any API (s) described elsewhere herein. Relative to CUDA runtime API 2002, a CUDA driver API can be a lower-level API but can provide finer-grained control of a device. Examples of CUDA libraries include cuBLAS, cuFFT, cuRAND, cuDNN, etc.
[0341] CUDA compiler 2050 may compile input CUDA code (e.g., CUDA source code 2010) to generate host executable code 2070 (1) and CUDA device executable code 2084. CUDA compiler 2050 may be, but is not limited to, NVCC. Host executable code 2070 (1) can be a compiled version of host code included in input source code that is executable on CPU 2090. CPU 2090 may be any processor that is optimized for sequential instruction processing.
[0342] CUDA device executable code 2084 may be a compiled version of device code included in input source code that is executable on CUDA-enabled GPU 2094. CUDA device executable code 2084 may include binary code. CUDA device executable code 2084 can include IR code, such as, but not limited to, PTX code, that is further compiled at runtime into binary code for a specific target device (e.g., CUDA-enabled GPU 2094) by a device driver. CUDA-enabled GPU 2094 may include any processor that is optimized for parallel instruction processing and that supports CUDA. CUDA-enabled GPU 2094 may be developed by NVIDIA Corporation of Santa Clara, CA.
[0343] CUDA to HIP translation tool 2020 can be configured to translate CUDA source code 2010 to functionally similar HIP source code 2030. HIP source code 2030 may include a collection of human-readable code in a HIP programming language. 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 includes functionally similar versions of CUDA mechanisms to define device code and distinguish between device code and host code. A HIP programming language may include a subset of functionality of a CUDA programming language. For example, a HIP programming language includes mechanism (s) to define global functions 2012, but such a HIP programming language may lack support for dynamic parallelism and therefore global functions 2012 defined in HIP code may be callable from a host only.
[0344] HIP source code 2030 may include any number (including zero) of global functions 2012, any number (including zero) of device functions 2014, any number (including zero) of host functions 2016, and any number (including zero) of host / device functions 2018. HIP source code 2030 may also include any number of calls to any number of functions that may be specified in a HIP runtime API 2032. HIP runtime API 2032 may include functionally similar versions of a subset of functions included in CUDA runtime API 2002. HIP source code 2030 may 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 may be any API that is designed for use by HIP code and / or ROCm. HIP APIs may include HIP runtime API 2032, a HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, etc.
[0345] CUDA to HIP translation tool 2020 can convert each kernel call in CUDA code from a CUDA syntax to a 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 include a call to a function specified in a CUDA API, and a HIP call can include a call to a function specified in a HIP API. CUDA to HIP translation tool 2020 may convert any number of calls to functions specified in CUDA runtime API 2002 to any number of calls to functions specified in HIP runtime API 2032.
[0346] CUDA to HIP translation tool 2020 can include a tool known as hipify-perl that executes a text-based translation process. CUDA to HIP translation tool 2020 can include a tool known as hipify-clang that, relative to hipify-perl, executes a more complex and more robust translation process that involves parsing CUDA code using clang (acompiler front-end) and then translating resulting symbols. Converting CUDA code to HIP code may include modifications (e.g., manual edits) in addition to those performed by CUDA to HIP translation tool 2020.
[0347] HIP compiler driver 2040 can include a front end that determines a target device 2046 and then configures a compiler that is compatible with target device 2046 to compile HIP source code 2030. Target device 2046 can include a processor that is optimized for parallel instruction processing. HIP compiler driver 2040 may determine target device 2046 in any technically feasible fashion.
[0348] If target device 2046 is compatible with CUDA (e.g., CUDA-enabled GPU 2094) , then HIP compiler driver 2040 can generate a HIP / NVCC compilation command 2042. HIP / NVCC compilation command 2042 can configure CUDA compiler 2050 to compile HIP source code 2030 using a HIP to CUDA translation header and a CUDA runtime library. In response to HIP / NVCC compilation command 2042, CUDA compiler 2050 may generate host executable code 2070 (1) and CUDA device executable code 2084.
[0349] If target device 2046 is not compatible with CUDA, then HIP compiler driver 2040 may generate a HIP / HCC compilation command 2044. HIP / HCC compilation command 2044 can configure HCC 2060 to compile HIP source code 2030 using an HCC header and a HIP / HCC runtime library. In response to HIP / HCC compilation command 2044, HCC 2060 may generate host executable code 2070 (2) and HCC device executable code 2082. HCC device executable code 2082 may be a compiled version of device code included in HIP source code 2030 that is executable on GPU 2092. GPU 2092 may be any processor that is optimized for parallel instruction processing, is not compatible with CUDA, and is compatible with HCC. GPU 2092 can be developed by AMD Corporation of Santa Clara, CA. GPU 2092 can include a non-CUDA-enabled GPU 2092.
[0350] For explanatory purposes only, three different flows that may be implemented in at least one embodiment to compile CUDA source code 2010 for execution on CPU 2090 and different devices are depicted in FIG. 20. A direct CUDA flow can compile CUDA source code 2010 for execution on CPU 2090 and CUDA-enabled GPU 2094 without translating CUDA source code 2010 to HIP source code 2030. An indirect CUDA flow can translate CUDA source code 2010 to HIP source code 2030 and then compiles HIP source code 2030 for execution on CPU 2090 and CUDA-enabled GPU 2094. A CUDA / HCC flow can translate CUDA source code 2010 to HIP source code 2030 and then can compile HIP source code 2030 for execution on CPU 2090 and GPU 2092.
[0351] A direct CUDA flow that may be implemented is depicted via dashed lines and a series of bubbles annotated A1-A3. As depicted with bubble annotated A1, CUDA compiler 2050 can receive CUDA source code 2010 and a CUDA compile command 2048 that can configure CUDA compiler 2050 to compile CUDA source code 2010. CUDA source code 2010 that can be used in a direct CUDA flow can be written in a CUDA programming language that is based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc. ) . In response to CUDA compile command 2048, CUDA compiler 2050 can generate host executable code 2070 (1) and CUDA device executable code 2084 (depicted with bubble annotated A2) . As depicted with bubble annotated A3, host executable code 2070 (1) and CUDA device executable code 2084 may be executed on, respectively, CPU 2090 and CUDA-enabled GPU 2094. CUDA device executable code 2084 can include binary code. CUDA device executable code 2084 can include PTX code and can be further compiled into binary code for a specific target device at runtime.
[0352] An indirect CUDA flow that may be implemented is depicted via dotted lines and a series of bubbles annotated B1-B6. As depicted with bubble annotated B1, CUDA to HIP translation tool 2020 can receive CUDA source code 2010. As depicted with bubble annotated B2, CUDA to HIP translation tool 2020 can translate CUDA source code 2010 to HIP source code 2030. As depicted with bubble annotated B3, HIP compiler driver 2040 can receive HIP source code 2030 and can determine that target device 2046 is CUDA-enabled.
[0353] As depicted with bubble annotated B4, HIP compiler driver 2040 can generate HIP / NVCC compilation command 2042 and can transmit both HIP / NVCC compilation command 2042 and HIP source code 2030 to CUDA compiler 2050. HIP / NVCC compilation command 2042 can configure CUDA compiler 2050 to compile HIP source code 2030 using a HIP to CUDA translation header and a CUDA runtime library. HIP to CUDA translation header can translate any number of mechanisms (e.g., functions) specified in any number of HIP APIs to any number of mechanisms specified in any number of CUDA APIs. CUDA compiler 2050 may use HIP to CUDA translation header in conjunction with a CUDA runtime library corresponding to CUDA runtime API 2002 to generate host executable code 2070 (1) and CUDA device executable code 2084. In response to HIP / NVCC compilation command 2042, CUDA compiler 2050 can generate host executable code 2070 (1) and CUDA device executable code 2084 (depicted with bubble annotated B5) . As depicted with bubble annotated B6, host executable code 2070 (1) and CUDA device executable code 2084 may be executed on, respectively, CPU 2090 and CUDA-enabled GPU 2094. CUDA device executable code 2084 can include binary code. CUDA device executable code 2084 can include PTX code and can be further compiled into binary code for a specific target device at runtime.
[0354] A CUDA / HCC flow that may be implemented is depicted via solid lines and a series of bubbles annotated C1-C6. As depicted with bubble annotated C1, CUDA to HIP translation tool 2020 can receive CUDA source code 2010. As depicted with bubble annotated C2, CUDA to HIP translation tool 2020 can translate CUDA source code 2010 to HIP source code 2030. As depicted with bubble annotated C3, HIP compiler driver 2040 can receive HIP source code 2030 and can determine that target device 2046 is not CUDA-enabled.
[0355] HIP compiler driver 2040 may generate HIP / HCC compilation command 2044 and may transmit both HIP / HCC compilation command 2044 and HIP source code 2030 to HCC 2060 (depicted with bubble annotated C4) . HIP / HCC compilation command 2044 can configure HCC 2060 to compile HIP source code 2030 using an HCC header and a HIP / HCC runtime library. HIP / HCC runtime library can correspond to HIP runtime API 2032. HCC header may include any number and type of interoperability mechanisms for HIP and HCC. In response to HIP / HCC compilation command 2044, HCC 2060 can generate host executable code 2070 (2) and HCC device executable code 2082 (depicted with bubble annotated C5) . As depicted with bubble annotated C6, host executable code 2070 (2) and HCC device executable code 2082 may be executed on, respectively, CPU 2090 and GPU 2092.
[0356] After CUDA source code 2010 is translated to HIP source code 2030, HIP compiler driver 2040 may subsequently be used to generate executable code for either CUDA-enabled GPU 2094 or GPU 2092 without re-executing CUDA to HIP translation tool 2020. CUDA to HIP translation tool 2020 can translate CUDA source code 2010 to HIP source code 2030 that is then stored in memory. HIP compiler driver 2040 can then configure HCC 2060 to generate host executable code 2070 (2) and HCC device executable code 2082 based on HIP source code 2030. In at least one embodiment, HIP compiler driver 2040 subsequently configures CUDA compiler 2050 to generate host executable code 2070 (1) and CUDA device executable code 2084 based on stored HIP source code 2030.
[0357] An example kernel may be translated by CUDA-to-HIP translation tool 2020 of FIG. 20, in accordance with at least one embodiment. CUDA source code 2010 partitions an overall problem that a given kernel is designed to solve into relatively coarse sub-problems that can independently be solved using thread blocks. Each thread block includes any number of threads. Each sub-problem can be partitioned into relatively fine pieces that can be solved cooperatively in parallel by threads within a thread block. Threads within a thread block can cooperate by sharing data through shared memory and by synchronizing execution to coordinate memory accesses.
[0358] CUDA source code 2010 can organize thread blocks associated with a given kernel into a one-dimensional, a two-dimensional, or a three-dimensional grid of thread blocks. Each thread block includes any number of threads, and a grid includes any number of thread blocks.
[0359] A kernel can be a function in device code that is defined using a ”__global__” declaration specifier. The dimension of a grid that executes a kernel for a given kernel call and associated streams may be specified using a CUDA kernel launch syntax. CUDA kernel launch syntax is specified as “KernelName<<<GridSize, BlockSize, SharedMemorySize, Stream>>> (KernelArguments) ; ” . An execution configuration syntax can include a “<<<... >>>” construct that is inserted between a kernel name ( “KernelName” ) and a parenthesized list of kernel arguments ( “KernelArguments” ) . CUDA kernel launch syntax can include a CUDA launch function syntax instead of an execution configuration syntax.
[0360] “GridSize” can be of a type dim3 and specify the dimension and size of a grid. Type dim3 may be a CUDA-defined structure that includes unsigned integers x, y, and z. If z is not specified, then z may default to one. If y is not specified, then y may default to one. The number of thread blocks in a grid can be equal to 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 may be equal to the product of BlockSize. x, BlockSize. y, and BlockSize. z. Each thread that executes a kernel may be given a unique thread ID that is accessible within the kernel through a built-in variable (e.g., ” threadIdx” ) .
[0361] With respect to CUDA kernel launch syntax, “SharedMemorySize” may be an optional argument that may specify a number of bytes in a shared memory that is dynamically allocated per thread block for a given kernel call in addition to statically allocated memory. With respect to CUDA kernel launch syntax, SharedMemorySize may default to zero. With respect to CUDA kernel launch syntax, “Stream” may be an optional argument that specifies an associated stream and defaults to zero to specify a default stream. A stream may be a sequence of commands (possibly issued by different host threads) that execute in order. Different streams may execute commands out of order with respect to one another or concurrently.
[0362] CUDA source code 2010 may include a kernel definition for an example kernel “MatAdd” and a main function. Main function may be host code that executes on a host and includes a kernel call that causes kernel MatAdd to execute on a device. Kernel MatAdd can add two matrices A and B of size NxN, where N is a positive integer, and store the result in a matrix C. Main function can define a threadsPerBlock variable as 16 by 16 and a numBlocks variable as N / 16 by N / 16. Main function can then specify kernel call “MatAdd<<<numBlocks, threadsPerBlock>>> (A, B, C) ; ” . As per CUDA kernel launch syntax, kernel MatAdd can be executed using a grid of thread blocks having a dimension N / 16 by N / 16, where each thread block has a dimension of 16 by 16. Each thread block can include 256 threads, a grid can be created with enough blocks to have one thread per matrix element, and each thread in such a grid may execute kernel MatAdd to perform one pair-wise addition.
[0363] While translating CUDA source code 2010 to HIP source code 2030, CUDA to HIP translation tool 2020 may translate each kernel call in CUDA source code 2010 from CUDA kernel launch syntax to a HIP kernel launch syntax and may convert any number of other CUDA calls in source code 2010 to any number of other functionally similar HIP calls. HIP kernel launch syntax can be specified as “hipLaunchKernelGGL (KernelName, GridSize, BlockSize, SharedMemorySize, Stream, KernelArguments) ; ” . Each of KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments can have the same meaning in HIP kernel launch syntax as in CUDA kernel launch syntax (described previously herein) . Arguments SharedMemorySize and Stream can be required in HIP kernel launch syntax and can be optional in CUDA kernel launch syntax.
[0364] A portion of HIP source code 2030 can be identical to a portion of CUDA source code 2010 depicted except for a kernel call that causes kernel MatAdd to execute on a device. Kernel MatAdd may be defined in HIP source code 2030 with the same ” __global__” declaration specifier with which kernel MatAdd is defined in CUDA source code 2010. A kernel call in HIP source code 2030 may be “hipLaunchKernelGGL (MatAdd, numBlocks, threadsPerBlock, 0, 0, A, B, C) ; ” , while a corresponding kernel call in CUDA source code 2010 is “MatAdd<<<numBlocks, threadsPerBlock>>> (A, B, C) ; ” .
[0365] Other implementations are contemplated and can be performed similarly to the CUDA and HIP implementations above, such as oneAPI, OpenCL, and other programming platforms. Code can be translated in any 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 or 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 to run on Intel or AMD processors using a translation tool, such as ZLUDA.
[0366] One or more techniques described herein can utilize a oneAPI programming model. A oneAPI programming model can refer to a programming model for interacting with various compute accelerator architectures. OneAPI may refer to an application programming interface (API) designed to interact with various compute accelerator architectures. A oneAPI programming model may utilize a DPC++ programming language. A DPC++ programming language may refer to a high-level language for data parallel programming productivity. A DPC++ programming language can be based at least in part on C and / or C++ programming languages. A oneAPI programming model can be a programming model such as, but not limited to, those developed by Intel Corporation of Santa Clara, CA.
[0367] OneAPI and / or oneAPI programming model can be utilized to interact with various accelerator, GPU, processor, and / or variations thereof, architectures. OneAPI may include a set of libraries that implement various functionalities. OneAPI may include at least a oneAPI DPC++ library, a oneAPI math kernel library, a oneAPI data analytics library, a oneAPI deep neural network library, a oneAPI collective communications library, a oneAPI threading building blocks library, a oneAPI video processing library, and / or variations thereof.
[0368] A oneAPI DPC++ library, also referred to as oneDPL, can be a library that implements algorithms and functions to accelerate DPC++ kernel programming. OneDPL may 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 API, and / or variations thereof. OneDPL can implement one or more classes and / or functions of a C++ standard library. OneDPL can implement one or more random number generator functions.
[0369] A oneAPI math kernel library, also referred to as oneMKL, can be a library that implements various optimized and parallelized routines for various mathematical functions and / or operations. OneMKL can implement one or more basic linear algebra subprograms (BLAS) and / or linear algebra package (LAPACK) dense linear algebra routines. OneMKL may implement one or more sparse BLAS linear algebra routines. OneMKL can implement one or more random number generators (RNGs) . OneMKL may implement one or more vector mathematics (VM) routines for mathematical operations on vectors. OneMKL may implement one or more Fast Fourier Transform (FFT) functions.
[0370] A oneAPI data analytics library, also referred to 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 analytics, in batch, online, and distributed processing modes of computation. OneDAL can implement various C++ and / or Java APIs and various connectors to one or more data sources. OneDAL may implement DPC++ API extensions to a traditional C++ interface and enables GPU usage for various algorithms.
[0371] A oneAPI deep neural network library, also referred to as oneDNN, can include a library that implements various deep learning functions. OneDNN may implement various neural network, machine learning, and deep learning functions, algorithms, and / or variations thereof.
[0372] A oneAPI collective communications library, also referred to as oneCCL, can include a library that implements various applications for deep learning and machine learning workloads. OneCCL can be built upon lower-level communication middleware, such as, but not limited to, message passing interface (MPI) and libfabrics. OneCCL can enable a set of deep learning specific optimizations, such as, but not limited to, prioritization, persistent operations, out of order executions, and / or variations thereof. OneCCL can implement various CPU and GPU functions.
[0373] A oneAPI threading building blocks library, also referred to as oneTBB, can include a library that implements various parallelized processes for various applications. OneTBB can be utilized for task-based, shared parallel programming on a host. OneTBB may implement generic parallel algorithms. OneTBB may implement concurrent containers. OneTBB may implement a scalable memory allocator. OneTBB may implement a work-stealing task scheduler. OneTBB may implement low-level synchronization primitives. OneTBB may be compiler-independent and usable on various processors, such as, but not limited to, GPUs, PPUs, CPUs, and / or variations thereof.
[0374] A oneAPI video processing library, also referred to as oneVPL, can include a library that is utilized for accelerating 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 discovery and selection in media centric and video analytics workloads. OneVPL can implement API primitives for zero-copy buffer sharing.
[0375] A oneAPI programming model may utilize a DPC++ programming language. A DPC++ programming language can include a programming language that can include functionally similar versions of CUDA mechanisms to define device code and distinguish between device code and host code. A DPC++ programming language may include a subset of functionality of a CUDA programming language. One or more CUDA programming model operations may be performed using a oneAPI programming model using a DPC++programming language.
[0376] Any application programming interface (API) described herein can be compiled into one or more instructions, operations, or any other signal by a compiler, interpreter, or other software tool. Compilation can include generating one or more machine-executable instructions, operations, or other signals from source code. An API compiled into one or more instructions, operations, or other signals, when performed, can cause one or more processors such as, but not limited to, processors described, e.g., in FIGs. 5-17, or any other logic circuit further described herein to perform one or more computing operations.
[0377] In at least one embodiment, translation tools described elsewhere herein, such as, but not limited to, can include one or more circuits to translate CUDA code to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data to HIP, oneAPI, OpenCL, or any other language used to perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to translate CUDA code to Compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data to HIP, oneAPI, OpenCL, or any other language used to perform any of the operations described above or elsewhere herein.AUTONOMOUS VEHICLE
[0378] FIG. 21 illustrates an example of an autonomous vehicle 2100, in accordance with at least one embodiment. Autonomous vehicle 2100 (alternatively referred to herein as “vehicle 2100” ) may be a passenger vehicle, such as, but not limited to, a car, a truck, a bus, and / or another type of vehicle that accommodates one or more passengers. In at least one embodiment, vehicle 2100 may be a semi-tractor-trailer truck used for hauling cargo. Vehicle 2100 may be an airplane, robotic vehicle, or other kind of vehicle.
[0379] Autonomous vehicles may be described in terms of automation levels, defined by National Highway Traffic Safety Administration ( “NHTSA” ) , a division of US Department of Transportation, and 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 on June 15, 2018, Standard No. J3016-201609, published on September 30, 2016, and previous and future versions of this standard) . In at least one embodiment, vehicle 2100 may be capable of functionality in accordance with one or more of Level 1 through Level 5 of autonomous driving levels. For example, in at least one embodiment, vehicle 2100 may be capable of conditional automation (Level 3) , high automation (Level 4) , and / or full automation (Level 5) , depending on embodiment.
[0380] Vehicle 2100 may 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 components of a vehicle. Vehicle 2100 may include a propulsion system 2150, such as, but not limited to, an internal combustion engine, hybrid electric power plant, an all-electric engine, and / or another propulsion system type. Propulsion system 2150 may be connected to a drive train of vehicle 2100, which may include a transmission, to enable propulsion of vehicle 2100. Propulsion system 2150 may be controlled in response to receiving signals from a throttle / accelerator (s) 2152.
[0381] A steering system 2154, which may include a steering wheel, is used to steer vehicle 2100 (e.g., along a desired path or route) when propulsion system 2150 is operating (e.g., when vehicle 2100 is in motion) . Steering system 2154 may receive signals from steering actuator (s) 2156. A steering wheel may be optional for full automation (Level 5) functionality. A brake sensor system 2146 may be used to operate vehicle brakes in response to receiving signals from brake actuator (s) 2148 and / or brake sensors.
[0382] Controller (s) 2136, which may include one or more system on chips ( “SoCs” ) and / or graphics processing unit (s) ( “GPU (s) ” ) , can provide signals (e.g., representative of commands) to one or more components and / or systems of vehicle 2100. For instance, controller (s) 2136 may send signals to operate vehicle brakes via brake actuator (s) 2148, to operate steering system 2154 via steering actuator (s) 2156, to operate propulsion system 2150 via throttle / accelerator (s) 2152. Controller (s) 2136 may include one or more onboard (e.g., integrated) computing devices that process sensor signals, and output operation commands (e.g., signals representing commands) to enable autonomous driving and / or to assist a human driver in driving vehicle 2100. Controller (s) 2136 may include a first controller for autonomous driving functions, a second controller for functional safety functions, a third controller for artificial intelligence functionality (e.g., computer vision) , a fourth controller for infotainment functionality, a fifth controller for redundancy in emergency conditions, and / or other controllers. A single controller may handle two or more of above functionalities, two or more controllers may handle a single functionality, and / or any combination thereof.
[0383] Controller (s) 2136 may provide signals for controlling one or more components and / or systems of vehicle 2100 in response to sensor data received from one or more sensors (e.g., sensor inputs) . Sensor data may be received from, for example, global navigation satellite systems ( “GNSS” ) sensor (s) 2158 (e.g., Global Positioning System sensor (s) ) , RADAR sensor (s) 2160, ultrasonic sensor (s) 2162, LIDAR sensor (s) 2164, inertial measurement unit ( “IMU” ) sensor (s) 2166 (e.g., accelerometer (s) , gyroscope (s) , a magnetic compass or magnetic compasses, magnetometer (s) , etc. ) , microphone (s) 2196, stereo camera (s) 2168, wide-view camera (s) 2170 (e.g., fisheye cameras) , infrared camera (s) 2172, surround camera (s) 2174 (e.g., 360 degree cameras) , long-range cameras 2198, mid-range camera (s) 2176, speed sensor (s) 2144 (e.g., for measuring speed of vehicle 2100) , vibration sensor (s) 2142, steering sensor (s) 2140, brake sensor (s) (e.g., as part of brake sensor system 2146) , and / or other sensor types.
[0384] One or more of controller (s) 2136 may receive inputs (e.g., represented by input data) from an instrument cluster 2132 of vehicle 2100 and provide outputs (e.g., represented by output data, display data, etc. ) via a human-machine interface ( “HMI” ) display 2134, an audible annunciator, a loudspeaker, and / or via other components of vehicle 2100. Outputs may include information such as, but not limited to, vehicle velocity, speed, time, map data (e.g., a High Definition map (not shown) , location data (e.g., vehicle’s 2100 location, such as, but not limited to, on a map) , direction, location of other vehicles (e.g., an occupancy grid) , information about objects and status of objects as perceived by controller (s) 2136, etc. For example, HMI display 2134 may display information about presence of one or more objects (e.g., a street sign, caution sign, traffic light changing, etc. ) , and / or information about driving maneuvers vehicle has made, is making, or will make (e.g., changing lanes now, taking exit 34B in two miles, etc. ) .
[0385] Each of components, features, and systems of vehicle 2100 in FIG. 21 may be connected via a bus 2102. Bus 2102 may include a CAN data interface (alternatively referred to herein as a “CAN bus” ) . A CAN may be a network inside vehicle 2100 used to aid in control of various features and functionality of vehicle 2100, such as, but not limited to, actuation of brakes, acceleration, braking, steering, windshield wipers, etc. Bus 2102 may be configured to have dozens or even hundreds of nodes, each with its own unique identifier (e.g., a CAN ID) . Bus 2102 may be read to find steering wheel angle, ground speed, engine revolutions per minute ( “RPMs” ) , button positions, and / or other vehicle status indicators. Bus 2102 may be a CAN bus that is ASIL B compliant.
[0386] In addition to, or alternatively from CAN, FlexRay and / or Ethernet protocols may be used. There may be any number of busses forming bus 2102, which may include zero or more CAN busses, zero or more FlexRay busses, zero or more Ethernet busses, and / or zero or more other types of busses using different protocols. Two or more busses may be used to perform different functions, and / or may be used for redundancy. For example, a first bus may be used for collision avoidance functionality and a second bus may be used for actuation control. Each bus of bus 2102 may communicate with any of components of vehicle 2100, and two or more busses of bus 2102 may communicate with corresponding components. Each of any number of system (s) on chip (s) ( “SoC (s) ” ) 2104 (such as, but not limited to, SoC 2104 (A) and SoC 2104 (B) ) , each of controller (s) 2136, and / or each computer within vehicle may have access to same input data (e.g., inputs from sensors of vehicle 2100) , and may be connected to a common bus, such CAN bus.
[0387] Any number of cameras can be positioned at any choice of camera locations and fields of view for autonomous vehicle 2100 of FIG. 21A, in accordance with at least one embodiment. Cameras and respective fields of view may be one example embodiment and are not intended to be limiting. For instance, additional and / or alternative cameras may be included and / or cameras may be located at different locations on vehicle 21...
Claims
1.A processor comprising:circuits to cause a compiler to compile source code comprising indications of a mapping between input data and tiles comprising portions of the input data, wherein the compiler:identifies one or more hardware features of a GPU, the one or more hardware features comprising one or more capabilities to perform blocks of threads in parallel; andgenerates code to perform operations on tiles of the input data using the one or more hardware features, including the capabilities to perform blocks of thread in parallel.