Tensor transposition

Abstract

Apparatuses, systems, and techniques to transpose a bit matrix by performing threads concurrently. In at least one embodiment, a processor comprising one or more circuits causes one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently.

Description

TENSOR TRANSPOSITIONTECHNICAL FIELD

[0001] At least one embodiment pertains to processing resources used to transfer data. At least one embodiment pertains to processors or computing systems used to transfer data to transpose a matrix.BACKGROUND

[0002] Multiple processing resources performing operations in parallel to transpose a matrix may cause inefficient computing. Techniques for performing operations in parallel by multiple processing resources can be improved.BRIEF DESCRIPTION OF DRAWINGS

[0003] FIG. 1 illustrates a matrix transpose system, in accordance with at least one embodiment;

[0004] FIG. 2 illustrates a system to transpose a matrix using software threads and an all-to-all communication pattern, in accordance with at least one embodiment;

[0005] FIG. 3 illustrates a system to transpose a matrix using software threads and partial results, in accordance with at least one embodiment;

[0006] FIG. 4 illustrates a process of transposing a matrix using two or more software threads to generate in parallel a different row or column of that transposed matrix, in accordance with at least one embodiment;

[0007] FIG. 5 illustrates test results of using a fast parallel bit matrix transpose algorithm, in accordance with at least one embodiment;

[0008] FIG. 6 illustrates a system used to transpose a matrix by performing software threads in parallel, in accordance with at least one embodiment;

[0009] FIG. 7 illustrates a system that includes a driver and / or runtime used to transpose a bit matrix by performing software threads in parallel, in accordance with at least one embodiment;

[0010] FIG. 8 illustrates an API call and response used to transpose a bit matrix by performing software threads in parallel, in accordance with one embodiment;

[0011] FIG. 9 illustrates a system including software and hardware used to transpose a bit matrix by performing software threads in parallel, in accordance with one embodiment;

[0012] FIG. 10 illustrates an example data center system, in accordance with at least one embodiment;

[0013] FIG. 11 illustrates an system-on-a-chip (SOC) , in accordance with at least one embodiment;

[0014] FIG. 12A illustrates a parallel processor, in accordance with at least one embodiment;

[0015] FIG. 12B illustrates a processing cluster, in accordance with at least one embodiment;

[0016] FIG. 12C illustrates a graphics multiprocessor, in accordance with at least one embodiment;

[0017] FIG. 13 illustrates an accelerator processor, in accordance with at least one embodiment;

[0018] FIG. 14A illustrate a central processing unit, in accordance with at least one embodiment;

[0019] FIG. 14B illustrates a core of central processing unit in FIG. 14A, in accordance with at least one embodiment;

[0020] FIG. 15 illustrates another accelerator processor, in accordance with at least one embodiment;

[0021] FIG. 16 illustrates a neuromorphic processor, in accordance with at least one embodiment;

[0022] FIG. 17 illustrates a supercomputer, in accordance with at least one embodiment;

[0023] FIG. 18 illustrates another accelerator processor, in accordance with at least one embodiment;

[0024] FIG. 19 illustrates another processor, in accordance with at least one embodiment;

[0025] FIG. 20 illustrates another accelerator processor, in accordance with at least one embodiment;

[0026] FIG. 21 illustrates a tensor processing unit, in accordance with at least one embodiment;

[0027] FIG. 22 illustrates a RISC-V-compatible processor, in accordance with at least one embodiment;

[0028] FIGS. 23A and 23B illustrate a language processing unit, in accordance with at least one embodiment;

[0029] FIG. 24 illustrates a software stack of a programming platform, in accordance with at least one embodiment;

[0030] FIG. 25 illustrates software that is supported by a programming platform, in accordance with at least one embodiment;

[0031] FIG. 26 illustrates compiling code to execute on programming platforms of FIG. 25, in accordance with at least one embodiment;

[0032] FIG. 27 illustrates an example of an autonomous vehicle and its system architecture, in accordance with at least one embodiment;

[0033] FIG. 28A illustrates inference and / or training logic, in accordance with at least one embodiment;

[0034] FIG. 28B illustrates inference and / or training logic, in accordance with at least one embodiment; and

[0035] FIG. 28C illustrates training and deployment of a neural network, in accordance with at least one embodiment.DETAILED DESCRIPTION

[0036] 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, and that any two or more aspects of any one or more embodiments described herein may be combined.

[0037] In at least one embodiment, a processor performs software threads to transpose a bit matrix, where each software thread concurrently performs operations to construct a different row or column of that transposed bit matrix. In at least one embodiment, a processor transposes a bit matrix using parallel software threads, where instructions of those software threads exchange portions of each row or column among each other’s registers until each software thread has accumulated a different transposed row or column of that bit matrix. In at least one embodiment, each software thread loads respective portions of every row or column into its registers then, in parallel, pairs of those software threads exchange portions of rows or columns so that, when those software threads are done exchanging portions of those rows or columns, each software thread has a corresponding row or column of a transposed matrix stored in its registers. In at least one embodiment, for example, a processor performs 32 software threads to transpose a bit matrix that represents an array of integers to produce 32 outputs, such as 32 unsigned integers, for all 32 software threads in 5 steps.

[0038] In at least one embodiment, a processor causes one or more matrix operands to be transposed based on two or more portions of those one or more matrix operands being loaded by two or more software threads concurrently, or as otherwise described herein. In at least one embodiment, a processor causes each of two or more software threads to generate, in parallel, a different row or column of one or more matrix operands in a transposed arrangement, or as otherwise described herein. In at least one embodiment, a processor causes one or more matrix operands to be transposed, where those one or more matrix operands comprise one or more arrays of one or more integers represented as bits, or as otherwise described herein. In at least one embodiment, a processor causes an identification, within two or more portions of one or more matrix operands to be transposed, of one or more bits to be operated upon by using one or more bitmasks, or as otherwise described herein. In at least one embodiment, a processor causes one or more matrix operands to be transposed by shifting one or more bits within one or more bit arrays of two or more portions of those one or more matrix operands, or as otherwise described herein. In at least one embodiment, a processor loads two or more portions of one or more matrix operands based on a size of a group of consecutive threads, or as otherwise described herein.

[0039] In at least one embodiment, a processor causes one or more matrix operands to be transposed as part of a data center used to perform one or more software applications, an algorithm used to assign a set of tasks to a set of computing systems, one or more autonomous vehicles, or some combination thereof. In at least one embodiment, a processor causes one or more matrix operands to be transposed to be used in one or more applications, such as image processing, data layout optimization, matrix multiplication, scientific simulations, cryptography, machine learning, or some combination thereof. In at least one embodiment, techniques described herein are applied to a network of autonomous vehicles to calculate inputs to control each autonomous vehicle more quickly. In at least one embodiment, techniques described herein are applied to machine learning processes to train a neural network more quickly. In at least one embodiment, a technical effect of embodiments described herein includes performing each thread of a warp to generate, or otherwise construct, a different column of a transposed matrix in parallel, instead of having each thread of that warp constructing an identical column of a transposed matrix in parallel before each thread of that warp repeats a process to construct another column of that transposed matrix in parallel. In at least one embodiment, a technical effect of embodiments described herein includes fewer operations, less time, and lower computational overhead required for a processor to perform a transpose of a matrix.

[0040] FIG. 1 illustrates a matrix transpose system 100 ( “system 100” ) , according to at least one embodiment. In at least one embodiment, system 100 is referred to as fast parallel transpose system 100. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 1 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 2-9. In at least one embodiment, system 100 includes at least a portion of, or is at least a portion of, system 200 of FIG. 2, system 300 of FIG. 3, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 600 of FIG. 6, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0041] In at least one embodiment, system 100 includes one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed based, at least in part, on two or more portions of those one or more matrix operands being loaded by two or more software threads concurrently, or to otherwise perform any operations described herein. In at least one embodiment, one or more processors perform one or more operations of system 100. In at least one embodiment, processor (s) 102 that performs one or more operations of system 100 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof. In at least one embodiment, system 100, or any other system described herein, is any computing system or combination of computing systems, such as those that make up one or more data centers or other facilities that house computing and networking devices.

[0042] In at least one embodiment, a logical processor refers to a virtualized processor core that an operating system can schedule tasks on. In at least one embodiment, a logical processor is a part of a processor’s architecture that allows for parallel processing. In at least one embodiment, a physical processor, such as a core, is an actual hardware component within a processor that performs computations. In at least one embodiment, a logical processor is a virtual representation of a physical core. In at least one embodiment, techniques such as Intel Hyper-ThreadingTM or AMD Simultaneous MultithreadingTM (SMT) splits each physical core of a processor into multiple logical processors. In at least one embodiment, this allows an operating system to treat each physical core as if that physical core were two or more separate cores, doubling a number of tasks that can be processed concurrently. In at least one embodiment, a logical processor can be created or otherwise implemented on any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0043] In at least one embodiment, processor (s) 102 receives or obtains input data of a tensor, such as input data depicted as a matrix in inputs and outputs 104. In at least one embodiment, embodiments described herein apply to tensors that include values other than bits, if an architecture of a processor so allows. In at least one embodiment, processing architectures that accommodate values other than bits are found in processors used in quantum computing, neuromorphic computing, optical computing, or some combination thereof. In at least one embodiment, techniques described herein apply to any data format, (e.g., INT8, FP8, FP16, FP32) where software instructions, such as those described further herein, are based on a number of bits used to represent a value in that data format. In at least one embodiment, techniques described herein apply to transposing tensors that include values of mixed precision, such as a tensor that includes values in both INT8 and FP32 data formats. In at least one embodiment, for example, to determine a value of a constant as described further herein at least in conjunction with FIG. 4, a dimensional block value is divided by 8 if integers of a bit matrix are represented in INT8.

[0044] In at least one embodiment, techniques described herein include a thread transferring one or more bits of data (e.g., 1 bit, 8 bits, 16 bits, 32 bits, 64 bits, 128 bits) using one or more instructions at a time to another thread. In at least one embodiment, techniques described herein include threads of a warp that perform one or more instructions in lockstep. In at least one embodiment, techniques described herein apply to transferring data between threads within a block of threads. In at least one embodiment, techniques described herein apply to sub-blocks of a block of threads. In at least one embodiment, techniques described herein apply to transferring data of blocks of threads to other blocks of threads.

[0045] In at least one embodiment, input data includes bits of a bit matrix, where each bit is referred to as an element of that bit matrix. In at least one embodiment, each value stored in a tensor, such as a matrix, is referred to as an element or an entry. In at least one embodiment, each row or column of a bit matrix represents an integer, where that matrix collectively represents an array of integers. In at least one embodiment, a bit matrix depicted in inputs and outputs 104 is a 32x32 matrix that includes 32-bit representations of integers. In at least one embodiment, each square of a bit matrix depicted in inputs and outputs 104 represents a register allocated specific software thread. In at least one embodiment, 32 software threads are assigned to each row or column of a matrix with index numbers ranging from 0 to 31. In at least one embodiment, each thread is allocated 32 registers, where each register stores one bit of data. In at least one embodiment, tensors of any dimensions are used to store input data, such as 128x128 matrices, three-dimensional (3D) tensors, or some combination thereof. In at least one embodiment, a register includes a register file, such as register file 1258 of FIG. 12.

[0046] In at least one embodiment, input data stored as a matrix is referred to as input matrix or input matrix operand. In at least one embodiment, input data stored as a matrix is stored in a host (e.g., CPU) memory. In at least one embodiment, processor (s) 102 includes one or more parallel processing units and receives or obtains input data and stores that input data in registers implemented as part of those one or more parallel processing units. In at least one embodiment, data of a matrix to be operated on by processor (s) 102 is referred to as a matrix operand. In at least one embodiment, a matrix operand refers to a matrix of data that is used as an input in a computational operation performed by processor (s) 102. In at least one embodiment, an operand is a data value, element, entry, or some combination thereof, on which an operation is performed. In at least one embodiment, a matrix operand is a matrix that is involved in operations such as addition, multiplication, or transposition.

[0047] In at least one embodiment, a parallel processing unit is referred to as a streaming multiprocessor (SM) . In at least one embodiment, SMs are referred to as streaming microprocessors, stream processors ( “SPs” ) , stream processing units ( “SPUs” ) , compute units ( “CUs” ) , execution units ( “EUs” ) , 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 and / or scheduler) .

[0048] In at least one embodiment, a software thread is referred to as a thread. In at least one embodiment, a thread is a sequence of computer instructions, or as otherwise described herein, including in conjunction with FIG. 11. In at least one embodiment, a thread is a software program or a portion thereof. In at least one embodiment, processor (s) 102 performs a thread by executing one or more instructions within a computing process to perform specific tasks or operations concurrently with other threads. In at least one embodiment, processor (s) 102 performs threads by executing identical threads concurrently, where identical threads refers to threads that each include one or more identical computer instructions but may use different values assigned to identical variables. In at least one embodiment, performing threads concurrently with other threads is referred to as performing threads in parallel or synchronously.

[0049] In at least one embodiment, a group of threads with identical instruction (s) to be performed in parallel is referred to as a thread warp or warp. In at least one embodiment, a group of thread warps is referred to as a thread block. In at least one embodiment, a thread warp includes 32 threads. In at least one embodiment, a thread represents a basic unit of execution. In at least one embodiment, a thread is referred to as a work-item. In at least one embodiment, a group of threads is referred to as a work-group or a thread group.

[0050] In at least one embodiment, processor (s) 102 operate on elements of a bit matrix illustrated in inputs and outputs 104 as further described herein at least in conjunction with FIGS. 1-9. In at least one embodiment, each thread of a warp is assigned to operate on a specific row of an input matrix depicted in an illustrated block labeled input matrix and outputs 104 of FIG. 1. In at least one embodiment, vertical arrays of inputs and outputs 104 represent rows of bits stored in registers allocated to a specific thread. In at least one embodiment, each thread of a warp is referred to as a lane. In at least one embodiment, each thread of a warp is assigned a thread ID, thread index, i, or a lane ID, where such a thread index, thread ID, or lane ID may be an integer. In at least one embodiment, a thread index, i, is not a unique thread identifier such as a thread ID. In at least one embodiment, as depicted in FIG. 1, each thread is assigned an index, such as i0, which represents a thread index of 0, which also corresponds to row 0 of a bit matrix. In at least one embodiment, as depicted in FIG. 1, each thread is assigned a thread position, such as t0, which represents a thread position of 0, which also corresponds to row 0 of a bit matrix.

[0051] In at least one embodiment, as indicated in FIG. 1, to generate an output, processor (s) 102 requires a warp of 32 threads to perform 5 different steps, or as otherwise described herein. In at least one embodiment, a step is referred to as a stage. In at least one embodiment, a step or stage includes one or more instructions, or as otherwise described herein. In at least one embodiment, as depicted in FIG. 1, processor (s) 102 outputs a different 32-bit unsigned integer to be stored in registers allocated to different threads of a warp by using operations described further herein at least in conjunction with FIGS. 1-9. In at least one embodiment, registers allocated to a thread are referred to as a thread’s corresponding registers, a thread’s registers, or registers corresponding to a thread. In at least one embodiment, notation such as ui32_0thread0 shown in FIG. 1 indicates that an unsigned 32-bit integer of column 0 is stored in registers corresponding to a thread with index 0. In at least one embodiment, outputs as depicted in FIG. 1 each represent a row of a transposed matrix.

[0052] In at least one embodiment, thread inputs and outputs 104 depicts a matrix configuration where input data from multiple threads is organized for processing. In at least one embodiment, this block illustrates a structured arrangement of bits, with each row representing a distinct thread's input data, labeled from i0 to i31, and each column corresponding to a specific bit position within a uint32 data type. In at least one embodiment, FIG. 1 depicts a final stage of a set of software instructions or algorithm and is labeled as Final Result, which signifies a completion of an all-to-all communication pattern, where each thread has successfully transposed its input data into output data, and as otherwise described herein. In at least one embodiment, a notation S_n = 32 indicates that a system has completed an all-to-all communication pattern using 32 threads, ensuring that each thread receives a complete set of transposed bits. In at least one embodiment, output data is organized into a series of uint32 values, each associated with a specific thread, labeled from ui32_0thread 0 to ui32_31thread 31, ensuring that each thread receives a comprehensive and correctly ordered set of transposed bits.

[0053] FIG. 2 illustrates system 200 that includes one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed based, at least in part, on an all-to-all communication method, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, system 200 is referred to as fast parallel transpose system 200. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 2 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1 and 3-9. In at least one embodiment, system 200 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 300 of FIG. 3, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 600 of FIG. 6, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0054] In at least one embodiment, one or more processors perform one or more operations of system 200. In at least one embodiment, processor (s) 202 that performs one or more operations of system 200 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0055] In at least one embodiment, processor (s) that perform one or more operations of system 200, perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 3, such as constructing a partial transpose matrix result based on previously constructed partial transpose matrix results. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 5, such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 6, such as operations to load elements of a bit matrix into registers implemented as part of a SM. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 7, such as API functions to perform a fast parallel bit matrix transpose. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 8, such as operations to perform a fast parallel transpose API call. In at least one embodiment, processor (s) that perform one or more operations of system 200 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0056] In at least one embodiment, processor (s) 202 perform operations to perform an all-to-all transfer of data between threads of a warp. In at least one embodiment, a transfer of data between threads occurs between two threads. In at least one embodiment, a transfer of data between threads is referred to as an exchange of data. In at least one embodiment, a technical effect of an all-to-all transfer of data to transpose a bit matrix is a distribution of bits of each thread of a warp to all threads in 5 steps for a 32-thread warp or in 2log (N) steps in general with N threads.

[0057] In at least one embodiment, an all-to-all transfer of bits across threads refers to each thread of a warp transferring one or more of its bits stored in that thread’s registers to registers of other threads in parallel. In at least one embodiment, a register is a referred to as a buffer. In at least one embodiment, an all-to-all transfer of data is referred to as an all-to-all communication method or pattern, all-to-all data exchange method or pattern, or some combination thereof. In at least one embodiment, an all-to-all communication method applies to subjects other than threads, such as computing nodes and computing devices. In at least one embodiment, an all-to-all communication method applies to data other than bits, such as messages, words, or some combination thereof. In at least one embodiment, an all-to-all communication method is one or more communication primitives that exchanges an increase in bandwidth or cables for a decrease of time or steps required to complete a task, such as a transposing a bit matrix.

[0058] In at least one embodiment, performance by processor (s) 202 of one or more operations to transpose a matrix by using an all-to-all communication pattern is referred to as a transformation. In at least one embodiment, a transformation operates as an all-to-all communication pattern among threads, where each thread sends a bit to each other thread within a warp. In at least one embodiment, each thread constructs a unsigned 32-bit integer, uint32, consisting of one bit from each thread, with bits ordered according to threads from which data originates, or as otherwise described herein.

[0059] In at least one embodiment, processor (s) 202 perform a process depicted as all-to-all transpose process 204, which describes a computing process where each thread within a warp sends and receives one bit to and from every other thread. In at least one embodiment, all-to-all transpose process 204, includes a structured communication pattern that facilitates efficient data exchange among threads, ensuring that each thread constructs a uint32 with bits sourced from all participating threads. In at least one embodiment, this process operates on a hardware level, utilizing processors to manage and execute bit transposition across threads.

[0060] In at least one embodiment, all-to-all transpose process 204 illustrates registers where input matrix elements are currently stored. In at least one embodiment, those matrix elements are represented by destination indicators, such as d0, d1, dk, and d31, which indicate specific threads and their corresponding registers those matrix elements are to be transferred to. In at least one embodiment, for example, a matrix element represented by d1 indicates that matrix element will be transferred to a register of a thread with index 1, or thread 1. In at least one embodiment, k is a variable that represents a thread index or its position within a warp.

[0061] In at least one embodiment, input matrix elements transferred to another thread’s registers are considered to be outputs and are represented by source indicators such as s0, s1, sk, and s31. In at least one embodiment, a source indicator indicates from which thread’s registers a matrix element was transferred. In at least one embodiment, for example, a matrix element represented by s0 indicates that element was transferred from a register corresponding to thread 0. In at least one embodiment, this transposition ensures that each thread, such as thread k, receives a complete set of bits from all other threads, which generates a complete transposed matrix stored in each thread’s registers. In at least one embodiment, a technical effect of an all-to-all communication pattern used to transpose a bit matrix is a reduced number of computer instructions required to transpose that bit matrix.

[0062] In at least one embodiment, bit arrays depicted on a left side of all-to-all transpose process 204 represent input data, specifically input message buffers of an all-to-all communication. In at least one embodiment, a resulting transpose on a right side holds one bit from each source thread, where thread k holds a kth column of bits from an input matrix of bits, forming an output buffer of an all-to-all communication. In at least one embodiment, labels in place of bits on a left side consist of a letter d and a number indicating a destination thread to which a corresponding bit is sent. In at least one embodiment, labels consisting of a letter s followed by a number on a right side indicate a source thread from which a bit originates.

[0063] In at least one embodiment, processor (s) 202 perform operations to generate outputs of threads as depicted in threads 206. In at least one embodiment, operations used to generate outputs of threads depicted in threads 206 are described in further detail herein at least in conjunction with FIG. 3. In at least one embodiment, a left thread subregion is displayed in a top row and labeled left, where a left and right thread subregion is described further herein at least in conjunction with FIG. 3. In at least one embodiment, outputs for threads in a right thread subregion are shown in bottom row and labeled right. In at least one embodiment, vertically-striped portions of threads 206 represent carry over bits that did not need to be shifted, as described further herein at least in conjunction with FIG. 3. In at least one embodiment, horizontally-striped portions of threads 206 represent bits that were received and consequently shifted, as described further herein at least in conjunction with FIG. 3. In at least one embodiment, for threads in left subregions and threads in right subregion, shifting bits happens only by received data. In at least one embodiment, data required to be selected is within a same range for both carry over data and received data, which is described further herein at least in conjunction with FIG. 3a. In at least one embodiment, for threads in a left subregion, selected bits will comprise a bit subarray on a left portion of a bit array, and as described further herein at least in conjunction with FIG. 3. In at least one embodiment, for threads in right subregion, selected bits will comprise a bit subarray on a right portion of a bit array, and as described further herein at least in conjunction with FIG. 3.

[0064] FIG. 3 illustrates system 300 that includes one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed based, at least in part, partial results of a transposed matrix, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, system 300 is referred to as fast parallel transpose system 300. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 3 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-2 and 4-9. In at least one embodiment, system 300 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 600 of FIG. 6, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0065] In at least one embodiment, one or more processors perform one or more operations of system 300. In at least one embodiment, one or more processors of system 300 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0066] In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations to generate tests results 500 described in conjunction with FIG. 5, such as operations 510 to test performance of bit matrix transpose techniques described herein. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 6, such as operations to load elements of a bit matrix into registers implemented as part of a SM. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 7, such as API functions to perform a fast parallel bit matrix transpose. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 8, such as operations to perform a fast parallel transpose API call. In at least one embodiment, processor (s) that perform one or more operations of system 300 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0067] In at least one embodiment, fast parallel transpose system 300 includes one or more processors to generate partial results of a transposed matrix. In at least one embodiment, a bit matrix transpose operation on a 32 × 32 bit matrix is performed in five steps, or 2log (N) steps for a more general N x N bit matrix, where N is a power of two. In at least one embodiment, steps are referred to as stages. In at least one embodiment, each stage n generates partial outputs of an all-to-all operation iteratively for each region each thread belongs to in order to be transferred to each future thread that will receive information in remaining stages. In at least one embodiment, a region refers to a group of consecutive threads. In at least one embodiment, regions are of size 2n, and a first stage is denoted by n = 1, thus containing two consecutive threads. In at least one embodiment, sizes of regions for five stages are 2, 4, 8, 16, 32, and numbers of regions are 16, 8, 4, 2, 1 respectively. In at least one embodiment, threads within a given region for iteration number n have ranks ranging from 0 to 2n-1 given by rank = tid %region_size.

[0068] In at least one embodiment, processor (s) of fast parallel transpose system 300 generates partial results of a transposed matrix by performing machine-level operations to transfer data between registers allocated to different threads of a warp. In at least one embodiment, data transfers between registers allocated to two different threads are referred to as exchanges of data or exchanges. In at least one embodiment, machine level-operations to exchange data between threads of a warp are referred to as warp shuffle operations, warp shuffles, warp-level primitives, warp intrinsics, intrinsics. In at least one embodiment, processor (s) of fast parallel transpose system 300 perform software instructions, such as warp shuffle commands, to cause those processor (s) or other processor (s) (e.g., GPUs) of fast parallel transpose system 300 to perform warp shuffle operations that exchange data between registers allocated to different threads of a warp.

[0069] In at least one embodiment, a warp shuffle are computing operations that allows threads within a warp to exchange data between threads. In at least one embodiment, exchanging data between threads refers to exchanging data between registers allocated to different threads. In at least one embodiment, a warp shuffle exchanges data between threads without using shared memory.

[0070] In at least one embodiment, shared memory is a high-speed memory space accessible by all threads within a block of threads of a SM. In at least one embodiment, shared memory is referred to other terms such as local memory or an L1 cache depending on an SM’s manufacturer or designer. In at least one embodiment, registers are high-speed memory spaces more quickly accessible by an SM than shared memory. In at least one embodiment, registers operate independently from shared memory on an SM, serving as storage locations allocated to threads to provide rapid access to frequently used variables, thereby enhancing computational efficiency.

[0071] In at least one embodiment, processor (s) perform warp shuffle operations by utilizing registers and instructions stored in a driver, firmware, instruction cache, or some combination thereof. In at least one embodiment, a driver, firmware, instruction cache, or some combination thereof, is installed on one or more processor (s) . In at least one embodiment, warp shuffle machine-level instructions are stored in firmware, an instruction cache, or some combination thereof, implemented on a parallel processing unit, such as a GPU, an SM of a GPU, or some combination thereof.

[0072] In at least one embodiment, each thread in a warp accesses a set of registers that store data values. In at least one embodiment, software interfaces with hardware through an API, such as NVIDIA CUDA and AMD ROCm , invoking warp shuffle commands, such as NVIDIA CUDA ’s __shfl_xor_sync () , to perform bitwise XOR operations on thread indices and determine threads that will be data exchange partners within a warp. In at least one embodiment, warp shuffle operations enable rapid data sharing without requiring shared memory access, enhancing computational efficiency. In at least one embodiment, hardware optimizes these operations by leveraging parallel processing capabilities inherent in CUDA architecture, ensuring minimal latency and maximal throughput during execution.

[0073] In at least one embodiment, a technical effect of embodiments disclosed here is an ability of a warp of threads to transform a bit matrix using warp shuffle operations and have each of those threads generate a different portion of a transformed bit matrix, instead of requiring two or more of those threads to each generate an identical portion of that transposed matrix. In at least one embodiment, a technical effect of embodiments disclosed herein is an improvement to computing efficiency by reducing a number of operations required to be performed by processor (s) of a computing system, such as fast parallel transpose system 300, when transposing a bit matrix using warp shuffle operations.

[0074] In at least one embodiment, a 32x32 bit matrix is transposed in five steps or stages. In at least one embodiment, with stage 1, a warp shuffle exchanges one uint32_t word or 4 bytes with its communicating thread 1 thread away. In at least one embodiment, a thread 1 thread away refers to a thread with a thread index or thread ID that is one away from a thread index or thread ID of a reference (e.g., calling) thread. In at least one embodiment, with stage 2, a warp shuffle exchanges one uint32_t word or 4 bytes with its communicating thread 2 threads away. In at least one embodiment, with stage 3, a warp shuffle exchanges one uint32_t word or 4 bytes with its communicating thread 4 threads away. In at least one embodiment, with stage 4, a warp shuffle exchanges one uint32_t word or 4 bytes with its communicating thread 8 threads away. In at least one embodiment, with stage 5, a warp shuffle exchanges one uint32_t word or 4 bytes with its communicating thread 16 threads away. In at least one embodiment, aspects of each stage can be described as follows:

[0075] In at least one embodiment, in a first stage, bits from a current region that are needed for threads of equal rank (or segindex) are extracted or filtered form received data (data_comp, comp for data from a complementary subregion) using a mask with a logical AND operation (data_comp &mask) , or as otherwise described herein;

[0076] In at least one embodiment, in a next stage, bits are shifted left or right depending on whether a current thread is in a left subregion or aright subregion, or as otherwise described herein;

[0077] In at least one embodiment, in a next stage, data that a thread itself holds from a previous stage data_seg (seg for segment, or a current subregion) is filtered using a mask (data_seg &mask) ;

[0078] In at least one embodiment, in a next stage, filtered data from data_seg, and filtered data and shifted data from data_comp are concatenated using a bitwise OR operator (e.g., (data_seg &mask) | data_comp_shft) to produce a result of a current region; and

[0079] In at least one embodiment, in a next stage, a result of a final stage will hold a transposed column corresponding to thread id ix %32.

[0080] In at least one embodiment, left in code corresponds to right << arrows in that code as bits are printed from right to left.

[0081] In at least one embodiment, a thread prepares bits for multiple columns instead of a single column. In at least one embodiment, techniques described herein attempt to fully utilize a bandwidth of a warp shuffle, meaning all bits that are sent are used meaningfully, rather than being left as zero or duplicated.

[0082] In at least one embodiment, bits prepared by a thread originate from a thread region it is a part of, specifically for threads receiving information from this thread. In at least one embodiment, during a first stage, all even threads (0, 2, 4, ...) and odd threads for threads (1, 3, 5, ...) participate.

[0083] In at least one embodiment, during a second stage, target threads include threads 0, 4, 8, ... for threads that are multiples of 4. In at least one embodiment, target threads include 1, 5, 9, ... for threads 1, 5, 9, ... In at least one embodiment, threads prepare partial columns for threads of a same rank.

[0084] In at least one embodiment, processor (s) of system 300 perform one or more operations to construct a transpose of a current thread region. In at least one embodiment, tid %region_size calculates a rank of a thread within its current read during a bit matrix transpose operation. In at least one embodiment, a current thread region is referred to as a current region. In at least one embodiment, a current thread region is a region of consecutive threads including a thread currently being performed by processor (s) of system 300. In at least one embodiment, a current thread region refers to a group of consecutive threads actively involved in a specific stage of bit matrix transpose operations, such as a warp shuffle operations, or as otherwise described herein. In at least one embodiment, a region is defined by a collection of threads, with a size determined by a stage number n, where each region contains 2n threads. In at least one embodiment, a first stage, denoted by n=1, includes two consecutive threads, and subsequent stages increase region sizes accordingly.

[0085] In at least one embodiment, aspects of threads that receive information or data from another thread are depicted in receiving threads 302. In at least one embodiment, for an i’ th region in an n’th stage, threads that are involved in that region are represented by i×sn, i×sn+1, ..., i×sn+sn-1. In at least one embodiment, i represents a thread index number. In at least one embodiment, sn represents a stride or step size used in parallel bitwise operations, such as warp shuffle operations. In at least one embodiment, i×sn, i×sn+1, …, i×sn+sn-1 is calculated based on an index number of a current thread and a stride, sn. In at least one embodiment, using a stride distributes work across multiple threads, where each thread processes a specific portion of data.

[0086] In at least one embodiment, a rank represents a position or order of a thread relative to other threads within a region of threads. In at least one embodiment, tid stands for thread identifier, and region_size denotes a size of a region in which a thread is operating. In at least one embodiment, a modulo operation %determines a remainder when tid is divided by region_size, effectively assigning a rank to each thread within its region. In at least one embodiment, this rank is used to organize and manage data communication between threads, as it helps identify which threads will interact and exchange information during a transpose process. In at least one embodiment, ranks are used to ensure that each thread has a unique position within its region.

[0087] In at least one embodiment, after each stage, each thread holds a uint32 containing resulting bits from its encompassing region for threads of region n rank equal to its own, or as otherwise described herein. In at least one embodiment, these threads receive information from this thread in remaining stages. In at least one embodiment, at an end of a final stage, each thread k holds all bits it needs-kth column of bits-for an entire 32-bit domain in order of bit-originating threads. In at least one embodiment, after a final stage, a transpose of each row of a matrix has been constructed.

[0088] In at least one embodiment, consider all four threads of rank 6 in stage n = 3: 6, 14, 22, 30, one for each region of size 8. In at least one embodiment, four threads of rank 6 are depicted in rank 6 threads 304. In at least one embodiment, when stage 3 ends, thread 6 holds bits of source arrays 0 -7. In at least one embodiment, thread 6’s first 8 bits are bits 6 of a column of inputs 0 -7. In at least one embodiment, thread 6’s second 8 bits are bits 14 of a column of inputs 0 -7. In at least one embodiment, thread 6’s third 8 bits are bits 22 of a column of inputs 0 -7. In at least one embodiment, thread 6’s fourth 8 bits consist of bit 30 of a column of inputs 0 -7.

[0089] In at least one embodiment, an output of a nth stage for thread k is a 32-bit word, and comprises partially constructed bit columns from data within region g = k / / 2n that thread k is part of.In at least one embodiment, an integer division is represented by a double slash,  / / . In at least one embodiment, with s = 2n region size for a nth stage, an output consists of m = 32 / snumber of partial sub bit arrays, where an m'th sub array holds results for final output thread k'= r +m'·s. In at least one embodiment, m'th consecutive s number of bits correspond to k'= r +m'·s-th column of bits of region g with r = k%sas generating a rank of thread k. In at least one embodiment, values of k' are depicted as final output threads 306. In at least one embodiment, a rank of a thread is referred to as a range of threads. In at least one embodiment, a rank of a thread is referred to as an index in a range of threads. In at least one embodiment, a rank or range of a thread refers to a value indicating a group of consecutive threads identified to exchange data with threads of a different rank or region of threads, or as otherwise described herein. In at least one embodiment, a range of threads can be referred to by a value indicating a group of consecutive threads in which pairs of threads exchange data with each other. In at least one embodiment, those pairs of threads are threads of equal rank in the preceding stage.

[0090] In at least one embodiment, computation of partial results is derived recursively. In at least one embodiment, aspects of constructing partial results from partial results of subregions of a previous stage are depicted in constructing partial results 308. In at least one embodiment, outputs for stage n + 1 are composed from results of previous stage n. In at least one embodiment, consider thread k in thread region g = gn+1 = k / / sn+1 = i such that a rank for stage n + 1 of thread k is r = k %sn+1 and its rank from a previous stage is %sn. In at least one embodiment, a thread region is divided into two subregions: left subregion corresponding to previous stage region and a right subregion corresponding to previous stage region both of which are half a size of stage n + 1 region g. In at least one embodiment, for g = i, a rank evaluates to if k is in a left subregion, and if k is in a right subregion, or as depicted as follows:

[0091] In at least one embodiment, inputs for current stage are outputs of previous stage for vertically-striped and gray-colored threads. In at least one embodiment, they are two threads within a current region that were of equal rank in previous stage. In at least one embodiment, output is subdivided into s bits for subsequent threads of equal rank, starting with corresponding thread in a first thread region. In at least one embodiment, threads labeled as k in left subregion and k in right subregion of communicating threads 310 are thread of region g. In at least one embodiment, threads of a left subregion with rank r are depicted with vertical stripes as shown in threads 312 and threads of a right subregion represented by (r + sn) %sn are represented in grey as shown in threads 313. In at least one embodiment, outputs are a concatenation of sub bit arrays for threads k′=r+m′s, as depicted with threads 314 shown with horizontal lines.

[0092] In at least one embodiment, let denote a rank of a current thread in a previous stage. In at least one embodiment, a rank of a previous stage for a communicating thread is equal to  as well. In at least one embodiment, an bit subarray is constructed for destination r+m'sn+1. In at least one embodiment, bits from a previous stage are required for a corresponding destination, which means obtaining an bitarray from a last stage for which:

[0093] In at least one embodiment, for threads in a left thread subregion,  then from equation (1) , a subarray index of previous stage is given by In at least one embodiment, a bit offset of a number of bits Sn is required from a previous stage and is equal to a bit offset of an m'th bit subarray that will be written to:

[0094] In at least one embodiment, a bit offset of a number of bits Sn is equal to a bit offset of an m'th bit subarray for both an input from a current thread with rank r and an input of a communicating thread that current thread is receiving data from.

[0095] In at least one embodiment, threads of threads 318 depicted in a top row and labeled as left: src includes a partial result subarray in a threads input that is required for an m′th partial result for this stage. In at least one embodiment, vertically-striped and horizontally-striped portions of threads 316 represent data to be used to construct a m′th bit subarray. In at least one embodiment, a vertically-striped portion is an input that is carried over within a thread and a horizontally-striped portions is a data portion that is received. In at least one embodiment, to construct a result for a stage including threads 318, bits are placed in order from lowest to highest originating thread. In at least one embodiment, to place bits in order from lowest to highest originating thread requires shifting inputs depicted in a bottom row of threads 318 and labeled as right: src. In at least one embodiment, a bottom row of threads 318 are shifted in a right-ward direction. In at least one embodiment, processor (s) of system 300 generate an output by concatenating those two subarray input portions depicted in threads 318. In at least one embodiment, concatenated subarray input portions are shown as destination outputs labeled as out_left.

[0096] In at least one embodiment, threads in a right thread subregion are depicted as threads 320 and has a rank of a previous stage In at least one embodiment, using equation (1) , a subarray index of a previous stage is given by In at least one embodiment, a bit offset of a number sn of bits from a previous stage is a second half of a current output's range, where

[0097] In at least one embodiment, bitoffsetdest is true for both input from a current thread with rank r and a communicating thread. In at least one embodiment, threads 320 depicts an incoming horizontally-striped portion of threads labeled as left : src and a carryover vertically-striped portion of threads labeled as right : src at bottom. In at least one embodiment, to arrange bits in increasing thread originating order, a thread labeled as left : src of threads 320 will be shifted left-ward. In at least one embodiment, processor (s) of system 300 concatenate two subarray input portions of threads 320 to construct an output subarray labeled as out_right. In at least one embodiment, composing partial results is repeated for all destinations and subarrays m' to generate outputs for threads as shown in threads 206 of FIG. 2.

[0098] In at least one embodiment, data exchanges between pairs of threads occur during one transaction or instruction. In at least one embodiment, if a first thread receives data from a second thread to which a first thread sends data, that is referred to as a data exchange. In at least one embodiment, a transaction or instruction refers to one warp shuffle instruction. In at least one embodiment, a region is subdivided further into sub regions of two consecutive and contiguous halves such that threads communicate with a thread with a corresponding index within another half.

[0099] In at least one embodiment, thread indexes, their location in a left or right sub region, and their location in a thread region (m’ ) during different stages are represented by comments inserted into code such as:

[0100] FIG. 4 illustrates process 400 performed by one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed in parallel, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 4 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-3 and 5-9. In at least one embodiment, a system that performs process 400 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, system 300 of FIG. 3, a system used to generate test results 500 of FIG. 5, system 600 of FIG. 6, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0101] In at least one embodiment, one or more processors perform one or more operations of process 400. In at least one embodiment, one or more processors that perform one or more operations of process 400 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0102] In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 3, such operations to construct partial transpose results. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations of process 500 described in conjunction with FIG. 5, such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 6, such as operations to load elements of a bit matrix into registers implemented as part of a SM. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 7, such as API functions to perform a fast parallel bit matrix transpose. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 8, such as operations to perform a fast parallel transpose API call. In at least one embodiment, processor (s) that perform one or more operations of process 400 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0103] In at least one embodiment, processor (s) begin process 400 by performing one or more operations to have each thread of a warp load an array of 32-bit unsigned integer (s) of a row of a bit matrix into corresponding registers with operation 402, or as otherwise described herein. In at least one embodiment, a pointer is uint_t *T is used to indicate to a thread in which row of a matrix an array of 32-bit unsigned integer (s) are stored, or as otherwise described herein. In at least one embodiment, processor (s) perform one or more operations to set up variables, constants, data structures, kernel functions, or some combination thereof, in addition to loading portions of an input bit matrix into registers allocated to threads. In at least one embodiment, performing one or more operations to set up variables, constants, data structures, kernel functions, or some combination thereof is referred to as performing or setting up utility functions. In at least one embodiment, software instructions performed by processor (s) to set up variables, constants, data structures, or some combination thereof, are represented in NVIDIA CUDA code as follows:

[0104] In at least one embodiment, any software instructions described herein are represented in code other than CUDA code, such as code used with AMD ROCm  (Radeon Open Compute) , Intel oneAPI , or Microsoft DirectCompute.

[0105] In at least one embodiment, inclusion of <cuda_runtime. h> facilitates CUDA runtime API usage. In at least one embodiment, inclusion of <cstdint> provides fixed-width integer types. In at least one embodiment, inclusion of <iomanip> enables formatted output manipulation. In at least one embodiment, inclusion of <iostream> allows standard input and output stream operations. In at least one embodiment, inclusion of <random> supports random number generation. In at least one embodiment, inclusion of <set> provides access to a set container. In at least one embodiment, inclusion of <vector> allows dynamic array management.

[0106] In at least one embodiment, #define BYSIZE_32 8 establishes a constant for a 32-bit size. In at least one embodiment, #define BYSIZE_64 8 sets a constant for a 64-bit size. In at least one embodiment, #define BYSIZE_96 6 defines a constant for a 96-bit size. In at least one embodiment, #define BYSIZE_128 4 specifies a constant for a 128-bit size.

[0107] In at least one embodiment, void check_cuda (cudaError_t result, char const *const func, const char *const file, int const line) defines a function for CUDA error checking. In at least one embodiment, if (result) checks if a CUDA error occurred. In at least one embodiment, std: : cerr << "CUDA error = " << static_cast<unsigned int> (result) << "at" << file << ": " << line << ″″′<< func << ″″′ << std: : endl; outputs error details. In at least one embodiment, std: : cerr << "CUDA error message =" << cudaGetErrorString (result) << std: : endl; displays an error message. In at least one embodiment, exit (99) ; terminates program execution with an error code.

[0108] In at least one embodiment, #define CUDA_CHECK (val) check_cuda ( (val) , #val, __FILE__, __LINE__) creates a macro for simplified CUDA error checking. In at least one embodiment, int get_bysize (const int bxsize) defines a function to determine bysize based on bxsize. In at least one embodiment, bxsize represents a compile-time constant that defines a block or matrix size in an x-dimension. In at least one embodiment, bxsize specifies how many elements or threads allocate along an x-axis of a grid or matrix. In at least one embodiment, bysize serves as a compile-time constant that defines a block or matrix size in a y-dimension. In at least one embodiment, bysize specifies how many elements or threads allocate along a y-axis.

[0109] In at least one embodiment, int bysize = -1; initializes bysize to -1. In at least one embodiment, switch (bxsize) initiates a switch statement for bxsize. In at least one embodiment, case 32: handles a case where bxsize is 32. In at least one embodiment, bysize =BYSIZE_32; assigns BYSIZE_32 to bysize. In at least one embodiment, break; exits a current case. In at least one embodiment, case 64: addresses a case where bxsize is 64. In at least one embodiment, bysize = BYSIZE_64; assigns BYSIZE_64 to bysize. In at least one embodiment, case 96: manages a case where bxsize is 96. In at least one embodiment, bysize =BYSIZE_96; assigns BYSIZE_96 to bysize. In at least one embodiment, case 128: covers a case where bxsize is 128. In at least one embodiment, bysize =BYSIZE_128; assigns BYSIZE_128 to bysize. In at least one embodiment, default: handles any other cases. In at least one embodiment, return bysize; returns a determined bysize.

[0110] In at least one embodiment, #define LAUNCH_TEST_KERNEL (kernel_name, N, ...) defines a macro for launching test kernels. In at least one embodiment, {begins a block of code. In at least one embodiment, dim3 block_size; declares a variable for block dimensions. In at least one embodiment, switch (N) initiates a switch statement for N. In at least one embodiment, case 32: handles a case where N is 32. In at least one embodiment, block_size =dim3 (32, BYSIZE_32, 1) ; sets block dimensions for 32. In at least oneembodiment, kernel_name<32, BYSIZE_32><<<1, block_size>>> (__VA_ARGS__) ; launches a kernel with specified parameters. In at least one embodiment, case 64: addresses a case where N is 64. In at least one embodiment, block_size = dim3 (64, BYSIZE_64, 1) ; sets block dimensions for 64. In at least one embodiment, kernel_name<64, BYSIZE_64><<<1, block_size>>> (__VA_ARGS__) ; launches a kernel with specified parameters.

[0111] In at least one embodiment, case 96: manages a case where N is 96. In at least one embodiment, block_size = dim3 (96, BYSIZE_96, 1) ; sets block dimensions for 96. In at least one embodiment, kernel_name<96, BYSIZE_96><<<1, block_size>>> (__VA_ARGS__) ; launches a kernel with specified parameters. In at least one embodiment, case 128: covers a case where N is 128. In at least one embodiment, block_size =dim3 (128, BYSIZE_128, 1) ; sets block dimensions for 128. In at least one embodiment, kernel_name<128, BYSIZE_128><<<1, block_size>>> (__VA_ARGS__) ; launches a kernel with specified parameters. In at least one embodiment, default: handles unsupported sizes. In at least one embodiment, std: : cerr << "Size" << N << "is not yet supported! " << std: : endl; outputs an error message. In at least one embodiment, } ends a block of code.

[0112] In at least one embodiment, template <typename dtype> begins a template definition for a data type dtype. In at least one embodiment, void allocate_gpu_memory (dtype **data, const int N) defines a function to allocate GPU memory. In at least one embodiment, CUDA_CHECK (cudaMalloc ( (void**) data, N *sizeof (dtype) ) ) ; allocates memory on GPU and checks for errors. In at least one embodiment, void delete_gpu_memory (dtype *d_data) defines a function to free GPU memory. In at least one embodiment, CUDA_CHECK (cudaFree (d_data) ) ; frees allocated GPU memory and checks for errors.

[0113] In at least one embodiment, void upload_array (dtype *d_data, const dtype *h_data, const int size) defines a function to upload data to GPU. In at least one embodiment, CUDA_CHECK (cudaMemcpy (d_data, h_data, sizeof (dtype) *size, cudaMemcpyHostToDevice) ) ; copies data from host to device and checks for errors. In at least one embodiment, void download_array (dtype *h_data, dtype *d_data, const int size) defines a function to download data from GPU. In at least one embodiment, CUDA_CHECK (cudaMemcpy (h_data, d_data, sizeof (dtype) *size, cudaMemcpyDeviceToHost) ) ; copies data from device to host and checks for errors.

[0114] In at least one embodiment, class GPUMem defines a class for GPU memory management. In at least one embodiment, public: specifies public access for class members. In at least one embodiment, GPUMem (const int bxsize) defines a constructor for GPUMem. In at least one embodiment, allocate (bxsize) ; calls an allocate function with bxsize. In at least one embodiment, ～GPUMem () ; declares a destructor for GPUMem. In at least one embodiment, uint32_t *d_M; declares a pointer for matrix data on GPU. In at least one embodiment, uint32_t *d_Mt_basic; declares a pointer for a basic transposed matrix on GPU. In at least one embodiment, uint32_t *d_Mt_opt; declares a pointer for an optimized transposed matrix on GPU. In at least one embodiment, float *d_times; declares a pointer for timing data on GPU.

[0115] In at least one embodiment, private: specifies private access for class members. In at least one embodiment, void allocate (const int bxsize) defines a function to allocate memory. In at least one embodiment, const int nints = bxsize  / 32; calculates a number of integers per row. In at least one embodiment, allocate_gpu_memory<uint32_t> (&d_M, bxsize *nints) ; allocates memory for matrix data.

[0116] In at least one embodiment, allocate_gpu_memory<uint32_t> (&d_Mt_basic, bxsize *nints) ; allocates memory for a basic transposed matrix. In at least one embodiment, allocate_gpu_memory<uint32_t> (&d_Mt_opt, bxsize *nints) ; allocates memory for an optimized transposed matrix. In at least one embodiment, allocate_gpu_memory<float> (&d_times, 2) ; allocates memory for timing data. In at least one embodiment, void cleanup () defines a function to free memory.

[0117] In at least one embodiment, delete_gpu_memory<uint32_t> (d_M) ; frees memory for matrix data. In at least one embodiment, delete_gpu_memory<uint32_t> (d_Mt_basic) ; frees memory for a basic transposed matrix. In at least one embodiment, delete_gpu_memory<uint32_t> (d_Mt_opt) ; frees memory for an optimized transposed matrix. In at least one embodiment, delete_gpu_memory<float> (d_times) ; frees memory for timing data.

[0118] In at least one embodiment, GPUMem: : ～GPUMem () defines a destructor for GPUMem. In at least one embodiment, cleanup () ; calls a cleanup function to free resources. In at least one embodiment, template <typename itype> begins a template definition for a data type itype. In at least one embodiment, bool compare_int_arrays (const std: : vector<itype> &a, const std: : vector<itype> &b, int size=-1) defines a function to compare integer arrays. In at least one embodiment, if (a. size () ! = b. size () ) return false; checks if array sizes differ. In at least one embodiment, bool equal = true; initializes an equality flag. In at least one embodiment, size =(size >= 0 ? size : a. size () ) ; determines a comparison size. In at least one embodiment, for (int i=0; i < size; ++i) {iterates over array elements. In at least one embodiment, equal = (equal && (a [i] == b [i] ) ) ; updates an equality flag. In at least one embodiment, return equal; returns a comparison result.

[0119] In at least one embodiment, void create_case (uint32_t *M, const int bxsize) defines a function to create a test case. In at least one embodiment, const int nints = bxsize  / 32; calculates a number of integers per row. In at least one embodiment, for (int y=0; y < bxsize; ++y) {iterates over rows. In at least one embodiment, for (int iint=0; iint < nints; ++iint) {iterates over integers. In at least one embodiment, M [iint *bxsize + y] = (uint32_t) 0; initializes a matrix element to zero. In at least one embodiment, std: : random_device rd; declares a random device. In at least one embodiment, std: : mt19937 gen (rd () ) ; initializes a random number generator.

[0120] In at least one embodiment, std: : uniform_int_distribution<> distbxsize (0, bxsize) ; defines a uniform distribution. In at least one embodiment, for (int y=0; y < bxsize; ++y) {iterates over rows. In at least one embodiment, int num_ones = distbxsize (gen) ; generates a random number of ones. In at least one embodiment, std: : set<int> pos; declares a set for positions. In at least one embodiment, for (int i=0; i < num_ones; ++i) {iterates over ones. In at least one embodiment, int index = std: : max (0, distbxsize (gen) -1) ; generates a random index. In at least one embodiment, while (pos. find (index) ! = pos. end () ) index = (index + 1) %bxsize; finds an unset bit. In at least one embodiment, pos. insert (index) ; inserts an index into a set. In at least one embodiment, int iint = index  / 32; calculates an integer index. In at least one embodiment, int ibit = index %32; calculates a bit index. In at least one embodiment, M [iint *bxsize + y] = (M[iint *bxsize + y] | (uint32_t) (1 << ibit) ) ; sets a bit in a matrix.

[0121] In at least one embodiment, void create_ref (uint32_t *Mt, uint32_t *M, const int bxsize) defines a function to create a reference matrix. In at least one embodiment, const int nints = bxsize  / 32; calculates a number of integers per row. In at least one embodiment, for (int y=0; y < bxsize; ++y) {iterates over rows. In at least one embodiment, for (int iint=0; iint < nints; ++iint) {iterates over integers. In at least one embodiment, Mt [iint *bxsize + y] = (uint32_t) 0;initializes a matrix element to zero. In at least one embodiment, for (int x=0; x < bxsize; ++x) {iterates over columns. In at least one embodiment, int iint_y = x  / 32; calculates an integer index for y. In at least one embodiment, int ibit_y = x %32; calculates a bit index for y.

[0122] In at least one embodiment, uint32_t Mt_x [nints] ; declares an array for a transposed row. In at least one embodiment, for (int iint=0; iint < nints; ++iint) iterates over integers. In at least one embodiment, Mt_x [iint] = (uint32_t) 0; initializes a transposed row element. In at least one embodiment, for (int y=0; y < bxsize; ++y) {iterates over rows. In at least one embodiment, int iint_x = y  / 32; calculates an integer index for x. In at least one embodiment, int ibit_x = y %32; calculates a bit index for x. In at least one embodiment, uint32_t bit = ( (M[iint_y *bxsize + y] & ( (uint32_t) 1 << ibit_y) ) > 0 ? (uint32_t) 1 : (uint32_t) 0) ; extracts a bit from a matrix. In at least one embodiment, Mt_x [iint_x] = (Mt_x [iint_x] | (bit <<ibit_x) ) ; sets a bit in a transposed row. In at least one embodiment, for (int iint = 0; iint < nints; ++iint) iterates over integers. In at least one embodiment, Mt [iint *bxsize + x] =Mt_x [iint] ; stores a transposed row in a matrix.

[0123] In at least one embodiment, processor (s) perform one or more operations to continue process 400 by performing each thread of a warp to identify whether a thread row index (iy) is less than a number of integers in a row with operation 404, or as otherwise described herein. In at least one embodiment, a constant based on a number of threads is referred to as nints as described further herein. In at least one embodiment, a constant based on a number of threads is calculated by dividing a number of threads along a dimension of a matrix by a number of bits used to represent an integer in a matrix to determine a number of integers represented in each row of a matrix, or as otherwise described herein .

[0124] In at least one embodiment, if processor (s) identify that a thread index iy is less than a constant with operation 404, processor (s) perform one or more operations to continue process 400 by performing each thread of a warp to identify another thread to exchange data with by using warp shuffle operations according to a stage number with operation 408, or as otherwise described herein. In at least one embodiment, warp shuffle operations allow pairs of threads to exchange data without using shared memory of an SM. In at least one embodiment, warp shuffle operations are machine-level instructions referred to as intrinsic operations. In at least one embodiment, if processor (s) identify that a thread index is equal to or greater than a constant with operation 404, process 400 ends.

[0125] In at least one embodiment, processor (s) perform one or more operations to continue process 400 by performing each thread of a warp to transfer data with an identified other thread with operation 410, or as otherwise described herein. In at least one embodiment, a thread transfers data to another thread identified with operation 408. In at least one embodiment, when a thread transfer data to another thread, that is referred to as an exchange of data, or as otherwise described herein.

[0126] In at least one embodiment, processor (s) perform one or more operations to continue process 400 by performing each thread to modify bits of transferred data by using bitwise mask, shift, and logic operations according to a stage number with operation 412, or as otherwise described herein. In at least one embodiment, a bitwise mask is used to isolate, identify, or otherwise select bits of an array to be operated on as part of a matrix transposition process, or as described further herein.

[0127] In at least one embodiment, processor (s) perform one or more operations to continue process 400 by performing each thread to perform an additional stage, if required, with operation 414, or as further described herein. In at least one embodiment, each stage of a matrix transposition process exchanges different portions of arrays assigned to different pairs of threads as described further herein. In at least one embodiment, stages are determined, at least in part, on a warp shuffle operation, such as __shfl_xor_sync () , or as described further herein.

[0128] In at least one embodiment, if processor (s) identify that more stages are required to be performed, processor (s) continue process 400 by performing operation 408. In at least one embodiment, if processor (s) identify that no more stages are required to be performed, processor (s) continue process 400 by storing in shared memory updated bits representing a transposed bit matrix with operation 416, and as described further herein. In at least one embodiment, processor (s) perform operations 404 through 416 described herein by performing software instructions represented as NVIDIA CUDA code as follows:

[0129] In at least one embodiment, code as shown above causes a thread of a warp to exchange selected data stored in registers allocated to that thread to be exchanged or overwrite data stored in registers allocated to another thread by using a combination of warp shuffle operations, logical operations (e.g., bitwise XOR, bitwise AND, bitwise OR) , or some combination thereof. In at least one embodiment, template <const int bxsize, const int bysize> initiates a template definition for constants bxsize and bysize, facilitating creation of specialized functions for different matrix dimensions. In at least one embodiment, __device__ __forceinline__void kernels<bxsize, bysize>: : bit_matrix_transpose_optimized defines a device function that performs an optimized bit matrix transposition using CUDA's warp shuffle operations.

[0130] In at least one embodiment, uint32_t *T serves as an input parameter, representing a matrix stored in registers where each thread holds a row consisting of integers. In at least one embodiment, data pointed to by T is being used as input and is stored in registers. In at least one embodiment, uint32_t is a data type that stores integer values from 0 to 232-1. In at least one embodiment, *T is a pointer that includes an indication of a contiguous block of memory where each thread accesses a specific segment (or "row" ) of this memorys. In at least one embodiment, uint32_t ds_Ttranspose [] [bxsize] acts as an output parameter, utilizing shared memory to store a transposed matrix, with each element representing an integer for a row of a transposed matrix.

[0131] In at least one embodiment, constexpr int nints = bxsize  / 32 calculates a number of integers per row, based on a block size. In at least one embodiment, int ix = threadIdx. x and int iy = threadIdx. y assign thread indices for x and y dimensions, respectively, to facilitate parallel processing.

[0132] In at least one embodiment, a conditional statement if (iy < nints) ensures that operations are performed only for valid integer segments within a row. In at least one embodiment, uint32_t transpose = (uint32_t) 0 initializes a variable to accumulate transposed data.

[0133] In at least one embodiment, a series of operations begins with calculating a segment index using int segindex = (ix %2) , which determines a position within a segment. In at least one embodiment, uint32_t data_segm = T [iy] retrieves a data segment from input matrix T. In at least one embodiment, uint32_t data_comp = __shfl_xor_sync (0xffffffff, data_segm, 1) performs a warp shuffle operation to exchange data between threads.

[0134] In at least one embodiment, a mask uint32_t mask = ( (uint32_t) 0b01010101010101010101010101010101 << segindex) is applied to isolate specific bits for rearrangement. In at least one embodiment, uint32_t data_comp_shft = ( (data_comp &mask) << (1-segindex) ) shifts and aligns data for merging. In at least one embodiment, data_comp_shft = (data_comp_shft >> segindex) adjusts data alignment further. In at least one embodiment, data_segm = ( (data_segm &mask) | data_comp_shft) combines original and shifted data to form a new segment.

[0135] In at least one embodiment, subsequent operations repeat this process for increasing segment sizes (4, 8, 16, 32) , each time recalculating segindex, is_right, and applying appropriate masks and shifts to refine data arrangement.

[0136] In at least one embodiment, ds_Ttranspose [ix / 32] [iy*32 + (ix %32) ] = transpose stores a computed transposed value in shared memory, ensuring efficient data access. In at least one embodiment, __syncthreads () synchronizes threads within a block, ensuring all threads complete their operations before proceeding, maintaining data consistency across a shared memory.

[0137] FIG. 5 illustrates a system that generates test results 500 includes one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed in parallel, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 5 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-4 and 6-9. In at least one embodiment, a system that generates test results 500 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, system 300 of FIG. 3, a system used to generate test results 500 of FIG. 5, system 600 of FIG. 6, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0138] In at least one embodiment, one or more processors perform one or more operations to generate test results 500. In at least one embodiment, one or more processors that perform one or more operations to generate test results 500 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0139] In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations described in conjunction with FIG. 3, such as construction partial transpose results. In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations of process 400 described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations. In at least one embodiment, processor (s) that perform one or more operations of to generate test results 500 perform one or more operations described in conjunction with FIG. 6, such as operations to load elements of a bit matrix into registers implemented as part of a SM. In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations described in conjunction with FIG. 7, such as API functions to perform a fast parallel bit matrix transpose. In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations described in conjunction with FIG. 8, such as operations to perform a fast parallel transpose API call. In at least one embodiment, processor (s) that perform one or more operations to generate test results 500 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0140] In at least one embodiment, processor (s) perform one or more operations to bit matrix transpose in 2log (N) steps for a network topology that requires 2log (N) N connections between threads or nodes. In at least one embodiment, a topology is represented by a crossbar on an SM, which connect threads within a warp such that each thread can send one message to one other thread in that warp.

[0141] In at least one embodiment, processor (s) perform one or more benchmarks to test performance of one or more embodiments described herein against other methods of transposing a bit matrix using a warp of threads. In at least one embodiment, table 1 shows a speedup using one or more embodiments described herein by comparing a time required to transpose a matrix of a certain size using a corresponding number of threads in a warp with a time required to transpose a matrix using other methods. In at least one embodiment, a time required to transpose a matrix using one or more embodiments described herein is listed under a column labeled optimized while a time required to transpose that same matrix using other methods is listed under a column labeled basic. In at least one embodiment, table 2 shows a speedup using one or more embodiments described herein when using a literature benchmark of a GPU Hungarian algorithm.

[0142] In at least one embodiment, benchmark tests comprise software instructions represented by NVIDIA CUDA code as follows:

[0143] In at least one embodiment, template <const int bxsize, const int bysize> initiates a template definition for constants bxsize and bysize, facilitating creation of specialized functions for different matrix dimensions. In at least one embodiment, __global__void bit_matrix_transpose_test_kernel defines a global function that tests bit matrix transposition, allowing execution on a GPU.

[0144] In at least one embodiment, uint32_t*d_Tt_basic, uint32_t*d_Tt_opt serve as output parameters for storing results of a basic and an optimized bit matrix transpose. In at least one embodiment, uint32_t*d_T acts as an input parameter representing a bit matrix. In at least one embodiment, float *d_time_transform serves as an output parameter for performance metrics.

[0145] In at least one embodiment, constexpr int nints = bxsize  / 32 calculates a number of integers per row, based on a block size. In at least one embodiment, int ix = threadIdx. x and int iy = threadIdx. y assign thread indices for x and y dimensions, respectively, to facilitate parallel processing.

[0146] In at least one embodiment, uint32_t T [nints] declares an array to hold a segment of a bit matrix for each thread. In at least one embodiment, __shared__uint32_t ds_Ttranspose [nints] [bxsize] allocates shared memory for storing a transposed matrix.

[0147] In at least one embodiment, a loop for (int iint=0; iint < nints; ++iint) initializes a bit matrix by assigning values from d_T to T [iint] , ensuring each thread holds a portion of a matrix. In at least one embodiment, __syncthreads () synchronizes threads within a block, ensuring all threads complete initialization before proceeding.

[0148] In at least one embodiment, uint64_t start and uint64_t stop declare variables for timing operations. In at least one embodiment, float clk = 0 initializes a variable to store clock cycles.

[0149] In at least one embodiment, __syncthreads () ensures synchronization before timing begins. In at least one embodiment, asm volatile ("mov. u64 %0, %%globaltimer; " : "=l" (start) ) captures a start time using a global timer. In at least one embodiment, kernels<bxsize, bysize>: : bit_matrix_transpose_basic (T, ds_Ttranspose) calls a basic transposition function, storing results in shared memory.

[0150] In at least one embodiment, __syncthreads () ensures synchronization after a basic transposition. In at least one embodiment, asm volatile ( "mov. u64 %0, %%globaltimer; " : "=l" (stop) ) captures a stop time. In at least one embodiment, clk = stop -start calculates elapsed clock cycles. In at least one embodiment, float time_in_microseconds = clk  /  (float) 1000.0 converts clock cycles to microseconds.

[0151] In at least one embodiment, a conditional statement if (ix == 0 &&iy == 0) ensures only one thread writes a timing result to d_time_transform [0] . In at least one embodiment, a conditional statement if (iy < nints) writes a transposed matrix to d_Tt_basic.

[0152] In at least one embodiment, __syncthreads () ensures synchronization before timing an optimized transposition. In at least one embodiment, asm volatile ( "mov. u64 %0, %%globaltimer; " : "=l" (start) ) captures a start time. In at least one embodiment, kernels<bxsize, bysize>: : bit_matrix_transpose_optimized (T, ds_Ttranspose) calls an optimized transposition function. In at least one embodiment, __syncthreads () ensures synchronization after an optimized transposition. In at least one embodiment, asm volatile ( "mov. u64 %0, %%globaltimer; " : "=l" (stop) ) captures a stop time. In at least one embodiment, clk = stop -start calculates elapsed clock cycles. In at least one embodiment, time_in_microseconds = clk  / prepares to convert clock cycles to microseconds, though a line appears incomplete.

[0153] FIG. 6 illustrates system 600 that includes one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed based, at least in part, partial results of a transposed matrix, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, system 600 is referred to as fast parallel transpose system 600. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 6 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-5 and 7-9. In at least one embodiment, system 600 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, system 300 of FIG. 3, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0154] In at least one embodiment, one or more processors perform one or more operations of system 600. In at least one embodiment, one or more processors of system 600 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0155] In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 3, such as operations to construct partial transpose results. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations of process 500 described in conjunction with FIG. 5, such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 7, such as API functions to perform a fast parallel bit matrix transpose. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 8, such as operations to perform a fast parallel transpose API call. In at least one embodiment, processor (s) that perform one or more operations of system 600 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0156] In at least one embodiment, system 600 includes CPU 602. In at least one embodiment, CPU 602 performs instructions to transpose a bit matrix using parallel threads. In at least one embodiment, CPU 602 utilizes fast parallel bit matrix transpose kernels module 604 to manage a distribution of matrix rows across multiple threads, with each thread handling a specific segment of a matrix.

[0157] In at least one embodiment, fast parallel bit matrix transpose kernels module 604 represents a software component that orchestrates a transposition of bit matrices by utilizing parallel processing techniques. In at least one embodiment, module 604 employs warp shuffle operations to enable communication between threads, allowing them to exchange data efficiently within a warp. In at least one embodiment, module 604 facilitates transformation of matrix rows into columns by coordinating an execution of threads and managing data dependencies.

[0158] In at least one embodiment, as used in any implementation described herein, unless otherwise clear from context or stated explicitly to contrary, terms such as “system, ” “device, ” “components, ” “agent, ” “manager, ” and “module, ” and nominalized verbs (e.g., coordinator, compiler, scheduler, manager, and / or other terms) each refer to any combination of software logic, firmware logic, hardware logic, and / or circuitry configured to provide functionality described herein. In at least one embodiment, any combination of software logic, firmware logic, hardware logic, and / or circuitry configured to provide functionality described herein is referred to as a component. In at least one embodiment, any component described herein is combined and / or communicatively connected with at least one other component, regardless of how such components are described to be combined and / or communicatively connected in other embodiments. In at least one embodiment, software may be embodied as a software package, code, and / or instruction set or instructions. In at least one embodiment, hardware includes, singly or in any combination, hardwired circuitry, programmable circuitry, state machine circuitry, fixed function circuitry, execution unit circuitry, and / or firmware that stores instructions executed by programmable circuitry. In at least one embodiment, modules may, collectively or individually, be 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, any one or more architectures of any circuits of one or more modules are represented as a register-transfer level (RTL) representation and / or another fabless representation that may be licensed and / or used in tape-out, a final phase in IC design before being used in manufacturing an IC.

[0159] In at least one embodiment, CPU memory 606 comprises a storage resource that stores bit matrices 608a-n, which are data structures to be transposed. In at least one embodiment, CPU memory 606 provides access to matrices 608a-n, enabling threads to receive elements of matrices 608a-n as part of a transposition process. In at least one embodiment, CPU memory is referred to as host memory.

[0160] In at least one embodiment, GPU 610 accelerates a transposition of bit matrices by utilizing its parallel processing capabilities. In at least one embodiment, GPU 610 contains multiple streaming multiprocessors, each capable of executing numerous threads concurrently. In at least one embodiment, an architecture of GPU 610 allows for efficient distribution of matrix transposition tasks across that GPU's resources.

[0161] In at least one embodiment, streaming multiprocessor (SM) 612a represents a processing unit within GPU 610 that executes threads responsible for transposing portions of a bit matrix. In at least one embodiment, SM 612a contains registers 614a and shared memory 616a, which store partial results and final results during a transposition process as described further herein at least in conjunction with FIGS. 1-9. In at least one embodiment, an architecture of SM 612a enables threads to perform warp shuffle operations, facilitating an exchange of data necessary to transpose matrix rows into columns.

[0162] In at least one embodiment, registers 614a comprise fast-access storage locations within streaming multiprocessor 612a that store data for individual threads during a transposition process. In at least one embodiment, registers 614a store portions of a bit matrix being transposed, allowing threads to perform bitwise operations and data manipulations efficiently In at least one embodiment, use of registers minimizes latency and enhances a speed of a transposition operation by providing immediate access to data.

[0163] In at least one embodiment, shared memory 616a comprises a memory resource within streaming multiprocessor 612a that stores transposed bit matrix 618a. In at least one embodiment, shared memory 616a allows threads to collaboratively store and access intermediate results during a transposition process. In at least one embodiment, shared memory 616a facilitates communication between threads, enabling them to synchronize their operations and ensure a correct assembly of a transposed matrix.

[0164] In at least one embodiment, transposed bit matrix 618a represents an output of a transposition process stored in shared memory 616a. In at least one embodiment, transposed bit matrix 618a is constructed by threads executing within streaming multiprocessor 612a, each contributing a portion of a final result. In at least one embodiment, matrix 618a represents a transformation of rows into columns, achieved through coordinated execution of parallel threads and a use of warp shuffle operations.

[0165] FIG. 7 illustrates system 700 that includes one or more processors comprising one or more circuits to cause one or more matrix operands to be transposed based, at least in part, partial results of a transposed matrix, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, system 700 is referred to as fast parallel transpose system 700. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 7 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-6 and 8-9. In at least one embodiment, system 700 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, system 300 of FIG. 3, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0166] In at least one embodiment, one or more processors perform one or more operations of system 700. In at least one embodiment, one or more processors of system 700 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0167] In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 3, such as operations to construct partial transpose results. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations of process 500 described in conjunction with FIG. 5, such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 6, such as storing elements of a bit matrix in registers of an SM. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 8, such as operations to perform a fast parallel transpose API call. In at least one embodiment, processor (s) that perform one or more operations of system 700 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0168] In at least one embodiment, system 700 includes a driver and / or runtime comprising one or more libraries to provide one or more application programming interfaces (APIs) to be performed by one or more processors comprising one or more circuits to perform fast parallel bit matrix transpose operations as described further herein. In at least one embodiment, a software program 702 is a software module. In at least one embodiment, a software program 702 comprises one or more software modules. In at least one embodiment, one or more APIs 710 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 710 are distributed or otherwise provided as a part of one or more runtimes 704, drivers 704, libraries 706, and / or any other grouping of software and / or executable code further described herein.

[0169] In at least one embodiment, one or more APIs 710 perform one or more computational operations in response to invocation by software programs 702. In at least one embodiment, a software program 702 is a collection of software code, commands, instructions, or other sequences of text to instruct a computing device to perform one or more computational operations and / or invoke one or more other sets of instructions, such as APIs 710 or function (s) 712, to be executed. In at least one embodiment, functionality provided by one or more APIs 710 include software functions, such as those usable to accelerate one or more portions of software programs 702 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.

[0170] In at least one embodiment, APIs 710 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 710 described herein are implemented as one or more circuits to perform one or more techniques described herein. In at least one embodiment, one or more software programs 702 comprise instructions that, if executed, cause one or more hardware devices and / or circuits to perform one or more techniques further described herein.

[0171] In at least one embodiment, software programs 702, such as user-implemented software programs, utilize one or more application programming interfaces (APIs) 710 to perform various computing 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 710 provide a set of callable function (s) 712, 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. For example, in an embodiment, one or more APIs 710 provide function (s) 712 to cause processor (s) to perform functions to identify a clock frequency at which one or more processors of a group of processors are to operate, or as otherwise described herein. In at least one embodiment, API (s) 710 provide one or more function (s) 712 that are one or more neural networks, such as a pre-trained LLM.

[0172] In at least one embodiment, one or more software programs 702 interact or otherwise communicate with one or more APIs 710 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 comprise 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 702 interact with one or more APIs 710 to facilitate parallel computing using a remote or local interface.

[0173] In at least one embodiment, an interface is software instructions that, if executed, provide access to one or more function (s) 712 provided by one or more APIs 710. In at least one embodiment, a software program 702 uses a local interface when a software developer compiles one or more software programs 702 in conjunction with one or more libraries 706 comprising or otherwise providing access to one or more APIs 710. In at least one embodiment, one or more software programs 702 are compiled statically in conjunction with pre-compiled libraries 706 or uncompiled source code comprising instructions to perform one or more APIs 710. In at least one embodiment, one or more software programs 702 are compiled dynamically and said one or more software programs utilize a linker to link to one or more pre-compiled libraries 706 comprising one or more APIs 710.

[0174] In at least one embodiment, a software program 702 uses a remote interface when a software developer executes a software program that utilizes or otherwise communicates with a library 706 comprising one or more APIs 710 over a network or other remote communication medium. In at least one embodiment, one or more libraries 706 comprising one or more APIs 710 are to be performed by a remote computing service, such as a computing resource services provider. In another embodiment, one or more libraries 806 comprising one or more APIs 710 are to be performed by any other computing host providing said one or more APIs 710 to one or more software programs 702.

[0175] In at least one embodiment, a processor performing or using one or more software programs 702 calls, uses, performs, or otherwise implements one or more APIs 710 to allocate and otherwise manage memory to be used by said software programs 702. In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 to allocate and otherwise manage memory to be used by one or more portions of said software programs 702 to be accelerated using one or more PPUs, such as GPUs or any other accelerator or processor further described herein. Those software programs 702 may be performed by one or more processors based, at least in part, on latency of interconnects coupled to one or more processors using function (s) 712 provided, in an embodiment, by one or more APIs 710.

[0176] In at least one embodiment, an API 710 is an API to facilitate parallel computing. In at least one embodiment, an API 710 is any other API further described herein. In at least one embodiment, an API 710 is provided by a driver and / or runtime 704. In at least one embodiment, an API 710 is provided by a CUDA user-mode driver. In at least one embodiment, an API 710 is provided by a CUDA runtime. In at least one embodiment, a driver 704 is data values and software instructions that, if executed, perform or otherwise facilitate operation of one or more function (s) 712 of an API 810 during load and execution of one or more portions of a software program 702. In at least one embodiment, a runtime 704 is data values and software instructions that, if executed, perform or otherwise facilitate operation of one or more function (s) 712 of an API 710 during execution of a software program 702. In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 implemented or otherwise provided by a driver and / or runtime 704 to perform combined arithmetic operations by said one or more software programs 702 during execution by one or more PPUs, such as GPUs.

[0177] In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 provided by a driver and / or runtime 704 to perform combined arithmetic operations of one or more PPUs, such as GPUs. In at least one embodiment, one or more APIs 710 provide combined arithmetic operations through a driver and / or runtime 704, as described above. In at least one embodiment, one or more software programs 702 utilize one or more APIs 710 provided by a driver and / or runtime 704 to allocate or otherwise reserve one or more blocks of memory 714 of one or more PPUs, such as GPUs. In at least one embodiment, one or more software programs 802 utilize one or more APIs 710 provided by a driver and / or runtime 704 to allocate or otherwise reserve blocks of memory. In at least one embodiment, one or more APIs 710 are to perform combined mathematical functions as described herein.

[0178] In at least one embodiment, to improve software programs 702 usability and / or optimization of one or more portions of said software programs 702 to be accelerated by one or more PPUs, such as GPUs, one or more APIs 810 provide one or more API function (s) 872 to perform a scheduling system usable or used by one or more computing devices as described herein. In at least one embodiment, a processor performs one or more software programs to combine two or more application programming interfaces (APIs) into a single API. In at least one embodiment, a processor uses an API to cause a scheduler to select a thread selection mechanism and / or otherwise perform operations described herein. In at least one embodiment, an API invokes a scheduler to cause a resource allocation. In at least one embodiment, a processor uses an exemplary API to schedule one or more instructions to be performed by one or more processors based, at least in part, on latency of one or more interconnects coupled to these one or more processors.

[0179] In at least one embodiment, memory 714 is system memory 1190 of SOC 1100. In at least one embodiment, memory 714 is processor memory. In at least one embodiment, memory 714 is any form of hardware that stores data and is referred to as storage or data storage. In at least one embodiment, memory 714 stores data used in various operations described herein, including elements of a bit matrix as described in conjunction at least with FIGS. 1-9.

[0180] In at least one embodiment, memory 714 is a computer readable storage medium and / or code stored on said computer readable storage medium in a form of a computer program including a plurality of computer readable instructions executable by one or more processors. In at least one embodiment, a computer readable storage medium is a non-transitory computer readable medium. In at least one embodiment, at least some computer readable instructions usable to perform operations described herein are not stored solely using transitory signals (e.g., a propagating transient electric or electromagnetic transmission) . In at least one embodiment, a non-transitory computer readable medium does not necessarily include non-transitory data storage circuitry (e.g., buffers, caches, and queues) within transceivers of transitory signals. In at least one embodiment, memory 714 is implemented as a non-transitory computer readable storage medium storing executable instructions that, if executed by one or more processors of a computer system, cause one or more processor (s) to cause one or more matrix operands to be transposed based, at least in part, on two or more portions of said one or more matrix operands being loaded by two or more software threads concurrently.

[0181] FIG. 8 illustrates system 800 that includes one or more API calls, that when performed by processors, cause one or more matrix operands to be transposed based, at least in part, partial results of a transposed matrix, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, system 800 is referred to as fast parallel transpose system 800. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 8 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-7 and 9.In at least one embodiment, system 800 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, system 300 of FIG. 3, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0182] In at least one embodiment, one or more processors perform one or more operations of system 800. In at least one embodiment, one or more processors of system 800 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0183] In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 3, such as operations to construct partial transpose results. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations of process 500 described in conjunction with FIG. 5, such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 6, such as storing elements of a bit matrix in registers of an SM. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 7, such as operations to perform to perform fast parallel bit matrix transpose API functions. In at least one embodiment, processor (s) that perform one or more operations of system 800 perform one or more operations described in conjunction with FIG. 9, such as operations used to perform a machine-level warp shuffle command.

[0184] In at least one embodiment, fast parallel transpose API call 802 is used (e.g., called by a user, application, or library) to receive one or more parameters that indicate a specific bit matrix, a thread block size, or some combination thereof, or as otherwise described herein. In at least one embodiment, fast parallel transpose API call 802 is an invocation of an API function of an API library used as part of a data center processor management system. In at least one embodiment, an API function is referred to as an API command. In at least one embodiment, a parameter received or otherwise obtained by an API is referred to as an input. In at least one embodiment, a parameter received or otherwise obtained by an API is referred to as an indication. In at least one embodiment, parameters received according to fast parallel transpose API call 802 are referred to as hints.

[0185] In at least one embodiment, fast parallel transpose API response 804 returns a portion of a transposed bit matrix stored in each thread’s corresponding registers. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to cause each thread of a warp to construct a row or column of a transposed bit matrix, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to cause each of two or more software threads to generate in parallel a different vector along a dimension of one or more transposed matrix operands, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to cause a transposition of a matrix operand comprising one or more arrays of one or more values represented as bits, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to use masks, such as bitmasks, to identify within two or more portions of a bit matrix one or more bits to be modified based, at least in part, on a conjunction operation, such as a logical AND operation, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to identify pairs of software threads used to exchange their respective data with each other based, at least in part, on a number that is a power of two representing a size of a group of consecutive threads within a warp of threads, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to transpose a bit matrix as part of a system operating one or more autonomous vehicles, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to shift one or more bits within one or more bit arrays of two or more portions of a matrix operand based, at least in part, on a size of a group of consecutive threads, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to identify within two or more portions of one or more matrix operands one or more bits to be operated upon, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to transpose one or more bit matrix operands representing one or more arrays of one or more integers, or as otherwise described herein. In at least one embodiment, processor (s) of system 800 perform fast parallel transpose API call 802 and / or fast parallel transpose API response 804 to cause each of two or more software threads to generate in parallel a different vector of one or more transposed matrix operands, or as otherwise described herein.

[0186] FIG. 9 illustrates system 900 that includes one or more API calls, that when performed by processors, cause one or more matrix operands to be transposed based, at least in part, partial results of a transposed matrix, or to otherwise perform any operations described herein, according to at least one embodiment. In at least one embodiment, system 900 is referred to as fast parallel transpose system 900. In at least one embodiment, one or more aspects of one or more embodiments described herein in conjunction with FIG. 9 are combined with one or more aspects of one or more embodiments described herein at least in conjunction with FIGS. 1-8. In at least one embodiment, system 900 includes at least a portion of, or is at least a portion of, system 100 of FIG. 1, system 200 of FIG. 2, system 300 of FIG. 3, a system used to perform process 400 of FIG. 4, a system used to generate test results 500 of FIG. 5, system 700 of FIG. 7, system 800 of FIG. 8, system 900 of FIG. 9, or some combination thereof.

[0187] In at least one embodiment, one or more processors perform one or more operations of system 900. In at least one embodiment, one or more processors of system 900 are any type of processor, portion of a processor, processor of a system, or combination of processors, described herein, including a logical processor, processor 1008 of FIG. 10, processor complex 1110 of FIG. 11, parallel processor 1200 of FIG. 12A, graphics multiprocessor 1234 of FIG. 12B, processor 1300 of FIG. 13, processor 1400 of FIG. 14A, core 1412 of FIG. 14B, accelerator 1500 of FIG. 15, processor 1655 of FIG. 16, processor 1732 of FIG. 17, accelerated processing unit 1800 of FIG. 18, processor 1900 of FIG. 19, core 2000 of FIG. 20, TPUs 2100 of FIG. 21, vector processor 2200 of FIG. 22, many-core tiled processor 2300 of FIG. 23A, hardware 2408 of FIG. 24, CPU 2690 of FIG. 26, streaming multiprocessors (SMs) of GPU (s) 2708 of FIG. 27, processor (s) 2710 of FIG. 27, a processor used in conjunction with logic 2815 illustrated in FIGS. 28A and 28B, a processor used in conjunction with training framework 2824 of FIG. 28C, or some combination thereof.

[0188] In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 1, such as operations to perform threads of a warp to transpose a bit matrix. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 2, such as transposing a matrix using threads and an all-to-all communication pattern. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 3, such as operations to construct partial transpose results. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 4, such as operation 410 to identify pairs of threads that exchange data stored in respective registers by using warp shuffle operations. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations of process 500 described in conjunction with FIG. 5, such as operations to test performance of a fast parallel matrix transpose system. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 6, such as storing elements of a bit matrix in registers of an SM. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 7, such as operations to perform to perform fast parallel bit matrix transpose API functions. In at least one embodiment, processor (s) that perform one or more operations of system 900 perform one or more operations described in conjunction with FIG. 8, such as operations used to perform a fast parallel transpose API call.

[0189] System 900 can include storage 902 and processor (s) 908. Storage 902 can include, for example, memory, cache, or other storage described further herein. Storage 902 can be separate from processor (s) 908, or storage 902 can be included in processor (s) 908 (e.g., in storage 912) . In at least one embodiment, software program 904 and / or software libraries (or instructions) 906 can be stored in memory, cache, or other storage and provided to processor (s) 908 to cause one or more circuits of processor (s) 908 to perform operations described herein. In at least one embodiment, software program 904 and / or software libraries (or instructions) 906 can be integrated into one or more circuits of processor (s) 908. Software program 904, which can be used to perform any one or more operations described herein, may be stored on storage 902.

[0190] In at least one embodiment, software program 904 can include one or more software modules. In at least one embodiment, software program 904 includes at least a portion of processor group sync API (s) module using workload variation 110 of FIG. 1. In at least one embodiment, software program 904 includes kernel functions to perform fast parallel bit matrix transform.

[0191] In at least one embodiment, software program 904 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 904 and / or software libraries (or instructions) 906 (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.

[0192] In at least one embodiment, system 900 can include one or more software libraries 906 that can, for example, provide one or more APIs and / or ISA instructions. In at least one embodiment, one or more software libraries 906 can be included in drivers and / or runtimes. In at least one embodiment, software libraries 906 (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) 908 to perform one or more computational operations, such as any one or more operations described herein. In at least one embodiment, software libraries 906 includes one or more machine-level warp shuffle operation commands. 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 906, 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 904.

[0193] Processor (s) 908 may include any number of processors and any suitableprocessing 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. 11-23. In at least one embodiment, processor (s) 908 can retrieve or fetch instructions (e.g., one or more APIs and / or ISA instructions) from storage 902 using, for example, instruction fetch 916 (e.g., for an Instruction Fetch stage) . Instructions can include instructions to identify a clock frequency at which a processor group is to operate by using measurements of activity levels, or as otherwise described herein. In at least one embodiment, processor (s) 908 can include storage 912 and instruction queue 910 to store and queue instructions fetched from storage 902. In at least one embodiment, fetched instructions can be decoded by decode 918 to determine what operation should be performed by processor (s) 908 (e.g., in an Instruction Decode stage) . In at least one embodiment, processor (s) 908 can fetch additional operands (data) that may be used for instructions, and operands can be stored, e.g., in registers or storage 912. In at least one embodiment, micro-operations 920 can perform operations on data stored in one or more registers or storage 912. For example, each step of instructions fetched by processor (s) 908 can be decomposed during execution so processor (s) 908 can execute instructions in steps through a series of micro-operations 920. In at least one embodiment, program counter (PC) 914 can hold an address for a next instruction and can be updated to point to a next instruction to be executed by processor (s) 908.

[0194] In at least one embodiment, processor (s) 908 can perform instructions (e.g., in an Execution stage) . For example, processor (s) 908 can perform an operation specified by one or more instructions, such as an arithmetic operation, a logical operation, or a data transfer. In at least one embodiment, compute unit (s) 922 can execute instructions to perform any one or more operations described herein. In at least one embodiment, compute unit (s) can include ALU (s) 924 (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) 926, which may be used for performing floating-point calculations. In at least one embodiment, other circuits 928 can be used to perform other operations, such as vector and / or scalar operations. In at least one embodiment, accelerator (s) 930 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 904 can utilize one or more APIs and / or ISA instructions to perform various computing operations with accelerator (s) 930, 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) 930 can include at least one or more groups of computing operations to be accelerated by execution at least in part by accelerator (s) 930, including to identify a clock frequency at which a processor group is to operate by using average activity levels of that processor group, or as otherwise described herein.

[0195] In at least one embodiment, system 900 can be used to perform one or more instructions that include functions or operations, such as those described in connection with FIGS. 1-8. In at least one embodiment, system 900 comprising one or more processors causes one or more circuits to cause one or more matrix operands tobe transposed based, at least in part, on two or more portions of those one or more matrix operands being loaded by two or more software threads concurrently, and / or to otherwise perform any one or more operations described herein, according to at least one embodiment. In at least one embodiment, system 900 comprising one or more processors causes one or more circuits to perform an API to cause one or more measurements of one or more activity levels of other processors to be stopped, and / or to otherwise perform any one or more operations described herein.

[0196] In at least one embodiment, system 900 is included in and / or otherwise includes systems illustrated or discussed in conjunction with FIGS. 1-8 to ca use oneor more matrix operands to be transposed based, at least in part, on two or more portions of said one or more matrix operands being loaded by two or more software threads concurrently, and / or to otherwise perform any of said operations described herein, according to at least one embodiment.

[0197] In at least one embodiment, system 900 includes one or more hardware illustrated in FIGS. 18-28C, such as to cause one or more matrix operands to be transposed based, at least in part, on two or more portions of said one or more matrix operands being loaded by two or more software threads concurrently, and / or to otherwise perform any of said operations described herein, according to at least one embodiment. DATA CENTER

[0198] FIG. 10 illustrates an example data center 1000, in accordance with at least one embodiment. Data center 1000 may include one or more rooms having racks 1002 and auxiliary equipment used to house one or more racks 1002 and one or more baseboards 1004. Rack 1002 can include one or more baseboards 1004. Rack 1002 can include a housing that receives and supports individual baseboards 1004. Operational aspects of rack 1002 may be regulated at a rack level, corresponding to a group of baseboards 1004, or at a baseboard level, corresponding to individual baseboards 1004, among other options. Rack 1002 or baseboards 1004 can have particularly selected maximum operating parameters, such as, but not limited to, power consumption, operating frequencies, and others. Data center 1000 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 1002. Baseboards 1004 within racks 1002 can get operational power from one or more power distribution units (PDUs; not shown) . PDUs may be arranged within racks 1002, for example between racks 1002 including baseboards 1004, or within racks 1002 that also house baseboards 1004.

[0199] Racks 1002 and baseboards 1004 can include sub-systems, modules, add-in cards, and other semiconductor components. Baseboards 1004 can include one or more computing units 1006 that can include one or more processors 1008, one or more memory 1010, and an interface controller 1012. Computing units 1006 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. 11-23. Computing units 1006 can include one or more memory storage devices 1010 (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 1006 may be a server having one or more of above-mentioned computing resources.

[0200] Computing units 1006 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 1014 may configure or otherwise control one or more computing units 1006 or groups of computing units. Resource orchestrator 1014 may include a software design infrastructure ( “SDI” ) management entity for data center 1000. Resource orchestrator 1014 may include hardware, software or some combination thereof.

[0201] Data center 1000 can include any one of or any combination of a framework layer 1020, a software layer 1030 and an application layer 1040. As shown in FIG. 10, framework layer 1020 includes a job scheduler 1022, a configuration manager 1024, a resource manager 1026 and a distributed file system 1028. Framework layer 1020 may include a framework to support software 1032 of software layer 1030 and / or one or more application (s) 1042 of application layer 1040. Software 1032 or application (s) 1042 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 1020 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 1028 for large-scale data processing (e.g., “big data” ) . Job scheduler 1022 may include a Spark driver to facilitate scheduling of workloads supported by various layers of data center 1000. Configuration manager 1024 may be capable of configuring different layers such as, but not limited to, software layer 1030 and framework layer 1020 including Spark and distributed file system 1028 for supporting large-scale data processing. Resource manager 1026 may be capable of managing clustered or grouped computing units 1006 mapped to or allocated for support of distributed file system 1028 and job scheduler 1022. Resource manager 1026 may coordinate with resource orchestrator 1014 to manage these mapped or allocated computing resources.

[0202] Software 1032 can be included in software layer 1030 and may include software used by at least portions of a computing unit 1006, one or more computing units 1006, groups of computing units 1006, and / or distributed file system 1028 of framework layer 1020. 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.

[0203] Application (s) 1042 can be included in application layer 1040 and may include one or more types of applications used by at least portions of a computing unit 1006, one or more computing units 1006, groups of computing units 1006, and / or distributed file system 1028 of framework layer 1020. 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.

[0204] Any of configuration manager 1024, resource manager 1026, and resource orchestrator 1014 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 1000 from making possibly bad configuration decisions and possibly avoiding underutilized and / or poor performing portions of a data center.

[0205] Data center 1000 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 1000. 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 1000 by using weight parameters calculated through one or more training techniques described herein.

[0206] Data center 1000 may use CPUs, application-specific integrated circuits (ASICs) , GPUs, FPGAs, or other hardware (e.g., embodiments in FIGs. 11-23) 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.

[0207] In at least one embodiment, processor 1008 can include one of the processors below and / or comprises one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. In at least one embodiment, processor 1008 is configured by software 1032 to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0208] PROCESSORS

[0209] The following figures set forth, without limitation, example processors and processing systems that can be used to cause each of the two or more software threads to generate in parallel a different row or column of the one or more transposed matrix operands, or otherwise perform some or all of processes, operations and / or and techniques described elsewhere herein.

[0210] Example processors and processing systems can be configured by software to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, 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 cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0211] FIGs. 28A and 28B illustrate logic 2815 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) .

[0212] FIG. 11 illustrates a processor which is a system-on-a-chip (SOC) 1100 (which may be referred to as system-on-chip, a superchip, or another name) , in accordance with at least one embodiment. SOC 1100 can include processor complex 1110 and processor complex 1140. SOC 1100 can include any number of processor complexes 1110 and / or processor complexes 1140 that may include any number of processors that are described herein, such as, but not limited to, those in FIGs. 11-23, in any combination. For example, processor 1110 may include a central processing unit (CPU) , and processor 1140 may include a graphics processor. Alternatively, processor 1110 may include a graphics processor, and processor 1140 may include a graphics processor. SOC 1100 may include any number of display controllers 1192, any number of multimedia engines 1194, any number of I / O Interfaces 1170, any number of memory controllers 1180, and any number of fabrics 1160 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 1100 can include a processor from Broadcom in Palo Alto, CA.

[0213] Processor complex 1110 can include a CPU, processor complex 1140 can include a GPU, and SOC 1100 can include a processing unit that integrates 1110 and 1140 onto a single chip. Some tasks may be assigned to processor complex 1110 and other tasks may be assigned to processor complex 1140. Processor complex 1110 can be configured to execute main control software associated with SOC 1100, such as, but not limited to, an operating system. Processor complex 1110 can be the master processor of SOC 1100, controlling and coordinating operations of other processors. Processor complex 1110 can issue commands that control the operation of processor complex 1140 to perform some or all of the operations described herein. Processor complex 1110 can be configured to execute host executable code derived from CUDA or other source code (e.g., HIP source code) , and processor complex 1140 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.

[0214] Processor complex 1110 can include cores 1120 (1) -1120 (4) and a cache (e.g., L3 cache) 1130 to store information to perform operations described herein. Processor complex 1110 may include any number of cores 1120 and any number and type of caches in any combination. Cores 1120 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 1120 can include a CPU core. Core 1120 (1) -1120 (4) can be referred to as a computing units or compute units. SOC 1100 can includes any number of processor complexes 1110, fabric 1160, I / O interfaces 1170, and memory controllers 1180.

[0215] Each core 1120 can include a fetch / decode unit 1122, an integer execution engine 1124, a floating point execution engine 1126, and an L2 cache 1128. Fetch / decode unit 1122 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 1124 and / or floating point execution engine 1126. Fetch / decode unit 1122 can concurrently dispatch one micro-instruction to integer execution engine 1124 and another micro-instruction to floating point execution engine 1126. Integer execution engine 1124 can execute integer and memory operations. Floating point engine 1126 can execute floating point and vector operations. Fetch-decode unit 1122 can dispatch micro-instructions to one or more execution engines that replaces both integer execution engine 1124 and floating point execution engine 1126.

[0216] Each core 1120 (i) , where i is an integer representing a particular instance of core 1120, may access L2 cache 1128 (i) included in core 1120 (i) . Each core 1120 included in core complex 1110 (j) , where j is an integer representing a particular instance of core complex 1110, can be connected to other cores 1120 included in core complex 1110 (j) via L3 cache 1130 (j) included in core complex 1110 (j) . Cores 1120 included in core complex 1110 (j) , where j is an integer representing a particular instance of core complex 1110, can access all of L3 cache 1130 (j) included in core complex 1110 (j) . L3 cache 1130 may include any number of slices.

[0217] Processor complex 1140 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 1140 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 1140 can be configured to execute operations unrelated to graphics, such as, but not limited to, neural network training and / or simulations. Processor complex 1140 can be configured to execute both operations related to graphics and operations unrelated to graphics.

[0218] Processor complex 1140 can include any number of compute units 1150 (1) -1150 (N) , where N is any integer greater than 1, and an L2 cache 1142. Compute units 1150 can share L2 cache 1142, which may store information to be used to perform some or all of the operations described herein. L2 cache 1142 can be partitioned. Processor complex 1140 can include any number of compute units 1150 and any number (including zero) and type of caches. Processor complex 1140 can include any amount of dedicated graphics hardware.

[0219] Each compute unit 1150 can include any number of SIMD units 1152 (1) -1152 (N) , where N is any integer greater than 1, and a shared memory 1154. Each SIMD unit 1152 can implement a SIMD architecture and can be configured to some or all of the operations described herein, in parallel. Each compute unit 1150 may execute any number of thread blocks, but each thread block can execute on a single compute unit 1150, 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 1152 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 1154. Each compute unit 1150 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) .

[0220] Fabric 1160 can be a system interconnect that facilitates data and control transmissions across processor complex 1110, processor complex 1140, I / O interfaces 1170, memory controllers 1180, display controller 1192, and multimedia engine 1194, e.g., to perform some or all of the operations described herein. SOC 1100 may include any amount and type of system interconnect in addition to or instead of fabric 1160 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 1100. I / O interfaces 1170 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 1170. Peripheral devices that can be coupled to I / O interfaces 1170 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.

[0221] Display controller 1192 may display images on one or more display device (s) , such as, but not limited to, a liquid crystal display ( “LCD” ) device. Multimedia engine 1194 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 1180 may facilitate data transfers between SOC 1100 and a unified system memory 1190. Processor complex 1110 and processor complex 1140 may share unified system memory 1190. Unified system memory 1190 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 1190 may include 3D stacked memory, including but not limited to high bandwidth memory (HBM) , HBM2e, or HDM3.

[0222] SOC 1100 may implement a memory subsystem that includes any amount and type of memory controllers 1180 and memory devices (e.g., shared memory 1154) that may be dedicated to one component or shared among multiple components in order to perform any of the operations described herein. SOC 1100 can implement a cache subsystem that includes one or more cache memories (e.g., L2 caches 1128, L3 cache 1130, and L2 cache 1142) that may each be private to or shared between any number of components (e.g., cores 1120, core complex 1110, SIMD units 1152, compute units 1150, and processor complex 1140) .

[0223] In at least one embodiment, SOC 1100 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0224] FIG. 12A illustrates a parallel processor 1200, in accordance with at least one embodiment. Parallel processor 1200 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. 11-23) to perform any of the operations described above or elsewhere herein.

[0225] Parallel processor 1200 can include a parallel processing unit 1202 to perform any of the operations described above or elsewhere herein. Parallel processing unit 1202 can include an I / O unit 1204 that enables communication with other devices, including other instances of parallel processing unit 1202. I / O unit 1204 may be directly connected to other devices. I / O unit 1204 may connect with other devices via use of a hub or switch interface, such as, but not limited to, a memory hub 1205. Connections between memory hub 1205 and I / O unit 1204 can form a communication link 1213. I / O unit 1204 may connect with a host interface 1206 and a memory crossbar 1216, where host interface 1206 receives commands directed to performing processing operations and memory crossbar 1216 receives commands directed to performing memory operations.

[0226] When host interface 1206 receives a command buffer via I / O unit 1204, host interface 1206 can direct work operations to perform those commands to a front end 1208. Front end 1208 can couple with a scheduler 1210 (which may be referred to as a sequencer) , which is configured to distribute commands or other work items to a processing cluster array 1212. Scheduler 1210 can ensure that processing cluster array 1212 is properly configured and in a valid state before tasks may be distributed to a cluster of processing cluster array 1212. Scheduler 1210 may be implemented via firmware logic executing on a microcontroller. Microcontroller-implemented scheduler 1210 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 1212. Host software can prove workloads for scheduling on processing cluster array 1212 via one of multiple graphics processing paths. Workloads can then be automatically distributed across processing array cluster 1212 by scheduler 1210 logic within a microcontroller including scheduler 1210.

[0227] Processing cluster array 1212 can perform any of the operations described above or elsewhere herein and can include up to “N” processing clusters (e.g., cluster 1214A, cluster 1214B, through cluster 1214N) , where “N” represents a positive integer (which may be a different integer “N” than used in other figures) . Each cluster 1214A-1214N of processing cluster array 1212 can execute a large number of concurrent threads. Scheduler 1210 can allocate work to clusters 1214A-1214N of processing cluster array 1212 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 1210, or can be assisted in part by compiler logic during compilation of program logic configured for execution by processing cluster array 1212. Different clusters 1214A-1214N of processing cluster array 1212 can be allocated for processing different types of programs or for performing different types of computations.

[0228] Processing cluster array 1212 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 1212 can be configured to perform general-purpose parallel compute operations. For example, processing cluster array 1212 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.

[0229] Processing cluster array 1212 can be configured to perform parallel graphics processing operations. Processing cluster array 1212 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 1212 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 1202 can transfer data from system memory via I / O unit 1204 for processing. During processing, transferred data can be stored to on-chip memory (e.g., parallel processor memory 1222) during processing, then written back to system memory.

[0230] When parallel processing unit 1202 is used to perform graphics processing, scheduler 1210 can be configured to divide a processing workload into approximately equal sized tasks, to better enable distribution of graphics processing operations to multiple clusters 1214A-1214N of processing cluster array 1212. Portions of processing cluster array 1212 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 1214A-1214N may be stored in buffers to allow intermediate data to be transmitted between clusters 1214A-1214N for further processing.

[0231] Processing cluster array 1212 can receive processing tasks to be executed via scheduler 1210, which receives commands defining processing tasks from front end 1208. 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 1210 may be configured to fetch indices corresponding to tasks or may receive indices from front end 1208. Front end 1208 can be configured to ensure processing cluster array 1212 is configured to a valid state before a workload specified by incoming command buffers (e.g., batch-buffers, push buffers, etc. ) is initiated.

[0232] Each of one or more instances of parallel processing unit 1202 can couple with a parallel processor memory 1222 to perform any of the operations described above or elsewhere herein. Parallel processor memory 1222 can be accessed via memory crossbar 1216, which can receive memory requests from processing cluster array 1212 as well as I / O unit 1204. Memory crossbar 1216 can access parallel processor memory 1222 via a memory interface 1218. Memory interface 1218 can include multiple partition units (e.g., partition unit 1220A, partition unit 1220B, through partition unit 1220N) that can each couple to a portion (e.g., memory unit) of parallel processor memory 1222. A number of partition units 1220A-1220N can be configured to be equal to a number of memory units, such that a first partition unit 1220A has a corresponding first memory unit 1224A, a second partition unit 1220B has a corresponding memory unit 1224B, and an N-th partition unit 1220N has a corresponding N-th memory unit 1224N. A number of partition units 1220A-1220N may not be equal to a number of memory units.

[0233] Memory units 1224A-1224N 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 1224A-1224N 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 1224A-1224N, allowing partition units 1220A-1220N to write portions of each render target in parallel to efficiently use available bandwidth of parallel processor memory 1222. A local instance of parallel processor memory 1222 may be excluded in favor of a unified memory design that utilizes system memory in conjunction with local cache memory.

[0234] Any one of clusters 1214A-1214N of processing cluster array 1212 can process data that will be written to any of memory units 1224A-1224N within parallel processor memory 1222. Memory crossbar 1216 can be configured to transfer an output of each cluster 1214A-1214N to any partition unit 1220A-1220N or to another cluster 1214A-1214N, which can perform additional processing operations on an output. Each cluster 1214A-1214N can communicate with memory interface 1218 through memory crossbar 1216 to read from or write to various external memory devices. Memory crossbar 1216 can have a connection to memory interface 1218 to communicate with I / O unit 1204, as well as a connection to a local instance of parallel processor memory 1222, enabling processing units within different processing clusters 1214A-1214N to communicate with system memory or other memory that is not local to parallel processing unit 1202. Memory crossbar 1216 can use virtual channels to separate traffic streams between clusters 1214A-1214N and partition units 1220A-1220N.

[0235] Multiple instances of parallel processing unit 1202 can be provided on a single add-in card, or multiple add-in cards can be interconnected. Different instances of parallel processing unit 1202 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 1202 can include higher precision floating point units relative to other instances. Systems incorporating one or more instances of parallel processing unit 1202 or parallel processor 1200 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.

[0236] FIG. 12A further includes a block diagram of a partition unit 1220, in accordance with at least one embodiment. Partition unit 1220 is an instance of one of partition units 1220A-1220N of FIG. 12A. Partition unit 1220 can include an L2 cache 1221, a frame buffer interface 1225, and a ROP 1226 (raster operations unit) . L2 cache 1221 can be a read / write cache that is configured to perform load and store operations received from memory crossbar 1216 and ROP 1226. Read misses and urgent write-back requests can be output by L2 cache 1221 to frame buffer interface 1225 for processing. Updates can also be sent to a frame buffer via frame buffer interface 1225 for processing. Frame buffer interface 1225 may interface with one of memory units in parallel processor memory, such as, but not limited to, memory units 1224A-1224N (shown as 1224) of FIG. 12A (e.g., within parallel processor memory 1222) .

[0237] ROP 1226 can be a processing unit that performs raster operations such as, but not limited to, stencil, z test, blending, etc. ROP 1226 can then output processed graphics data that is stored in graphics memory. ROP 1226 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 1226 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.

[0238] ROP 1226 can be included within each processing cluster (e.g., cluster 1214A-1214N of FIG. 12A) instead of within partition unit 1220. Read and write requests for pixel data may be transmitted over memory crossbar 1216 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 1200 of FIG. 12A.

[0239] In at least one embodiment, parallel processor 1200 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0240] FIG. 12B includes a block diagram of a processing cluster 1214 within a parallel processing unit, in accordance with at least one embodiment. A processing cluster can be an instance of one of processing clusters 1214A-1214N of FIG. 12A that can be used to perform any of the operations described above or elsewhere herein. Processing cluster 1214 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.

[0241] Operation of processing cluster 1214 can be controlled via a pipeline manager 1232 that distributes processing tasks to SIMT parallel processors. Pipeline manager 1232 can receive instructions from scheduler 1210 of FIG. 12A and manages execution of those instructions via a graphics multiprocessor 1234 and / or a texture unit 1236. Graphics multiprocessor 1234 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 1214. One or more instances of graphics multiprocessor 1234 can be included within a processing cluster 1214. Graphics multiprocessor 1234 can process data and a data crossbar 1240 can be used to distribute processed data to one of multiple possible destinations, including other shader units. Pipeline manager 1232 can facilitate distribution of processed data by specifying destinations for processed data to be distributed via data crossbar 1240.

[0242] Each graphics multiprocessor 1234 within processing cluster 1214 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.

[0243] Instructions transmitted to processing cluster 1214 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 1234. A thread group may include fewer threads than a number of processing engines within graphics multiprocessor 1234. 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 1234. When a thread group includes more threads than number of processing engines within graphics multiprocessor 1234, processing can be performed over consecutive clock cycles. Multiple thread groups can be executed concurrently on a graphics multiprocessor 1234.

[0244] Graphics multiprocessor 1234 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 1234 can forego an internal cache and use a cache memory (e.g., L1 cache 1248) within processing cluster 1214. Each graphics multiprocessor 1234 may also have access to L2 caches within partition units (e.g., partition units 1220A-1220N of FIG. 12A) that can be shared among all processing clusters 1214 and may be used to transfer data between threads. Graphics multiprocessor 1234 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 1202 may be used as global memory. Processing cluster 1214 can include multiple instances of graphics multiprocessor 1234 and can share common instructions and data, which may be stored in L1 cache 1248.

[0245] Each processing cluster 1214 may include an MMU 1245 (memory management unit) that can be configured to map virtual addresses into physical addresses. One or more instances of MMU 1245 may reside within memory interface 1218 of FIG. 12A. MMU 1245 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 1245 may include address translation lookaside buffers (TLB) or caches that may reside within graphics multiprocessor 1234 or L1 1248 cache or processing cluster 1214. 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.

[0246] A processing cluster 1214 may be configured such that each graphics multiprocessor 1234 is coupled to a texture unit 1236 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 1234 and can be fetched from an L2 cache, local parallel processor memory, or system memory, as needed. Each graphics multiprocessor 1234 can output processed tasks to data crossbar 1240 to provide processed task to another processing cluster 1214 for further processing or to store processed task in an L2 cache, local parallel processor memory, or system memory via memory crossbar 1216. A preROP 1242 (pre-raster operations unit) can be configured to receive data from graphics multiprocessor 1234, and direct data to ROP units, which may be located with partition units as described herein (e.g., partition units 1220A-1220N of FIG. 12A) . PreROP 1242 unit can perform optimizations for color blending, organizing pixel color data, and performing address translations.

[0247] In at least one embodiment, processing cluster 1214 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0248] FIG. 12C shows a graphics multiprocessor 1234, in accordance with at least one embodiment, e.g., to perform any of the operations described above or elsewhere herein. Graphics multiprocessor 1234 can couple with pipeline manager 1232 of processing cluster 1214. Graphics multiprocessor 1234 can include an execution pipeline including but not limited to an instruction cache 1252 (that, e.g., can store instructions, such as, not limited to compiled API instructions) , an instruction unit 1254, an address mapping unit 1256, a register file 1258, one or more general purpose graphics processing unit (GPGPU) cores 1262, and one or more load / store units 1266, where one or more load / store units 1266 can perform load / store operations to load / store instructions corresponding to performing an operation. GPGPU cores 1262 and load / store units 1266 can be coupled with cache memory 1272 and shared memory 1270 via a memory and cache interconnect 1268. GPGPU cores 1262 can be part of an SoC such as, but not limited to, part of integrated circuit 1100 in FIG. 11.

[0249] Instruction cache 1252 can receive a stream of instructions (e.g., to perform any of the operations described above or elsewhere herein) to execute from pipeline manager 1232. Instructions can be cached in instruction cache 1252 and dispatched for execution by an instruction unit 1254. Instruction unit 1254 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 1262. 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 1256 can be used to translate addresses in a unified address space into a distinct memory address that can be accessed by load / store units 1266.

[0250] Register file 1258 can provide a set of registers for functional units of graphics multiprocessor 1234. Register file 1258 may provide temporary storage for operands connected to data paths of functional units (e.g., GPGPU cores 1262, load / store units 1266) of graphics multiprocessor 1234. Register file 1258 may be divided between each of functional units such that each functional unit is allocated a dedicated portion of register file 1258. Register file 1258 can be divided between different warps (which may be referred to as wavefronts, subgroups, and / or waves or threads) being executed by graphics multiprocessor 1234.

[0251] GPGPU cores 1262 can each include floating point units (FPUs) and / or integer arithmetic logic units (ALUs) that can be used to execute instructions of graphics multiprocessor 1234. GPGPU cores 1262 can be similar in architecture or can differ in architecture. A first portion of GPGPU cores 1262 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 1234 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 1262 can also include fixed or special function logic.

[0252] GPGPU cores 1262 can include SIMD logic capable of performing a single instruction on multiple sets of data. GPGPU cores 1262 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.

[0253] Memory and cache interconnect 1268 can include an interconnect network that connects each functional unit of graphics multiprocessor 1234 to register file 1258 and to shared memory 1270. Memory and cache interconnect 1268 may be a crossbar interconnect that allows load / store unit 1266 to implement load and store operations between shared memory 1270 and register file 1258. register file 1258 can operate at a same frequency as GPGPU cores 1262, thus data transfer between GPGPU cores 1262 and register file 1258 can have very low latency. Shared memory 1270 can be used to enable communication between threads that execute on functional units within graphics multiprocessor 1234. Cache memory 1272 can be used as a data cache for example, to cache texture data communicated between functional units and texture unit 1236. Shared memory 1270 can also be used as a program managed cache. Threads executing on GPGPU cores 1262 can programmatically store data within shared memory in addition to automatically cached data that is stored within cache memory 1272.

[0254] 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.

[0255] In at least one embodiment, graphics multiprocessor 1234 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0256] FIG. 13 shows a processor 1300, in accordance with at least one embodiment. Processor 1300 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 1300 can include one or more Central Processing Unit (s) (CPU 1302) , one or more Graphics Processing Unit (s) (GPU 1306) , and / or one or more Neural Processing Unit (s) (NPU 1308) that can be, e.g., a dedicated AI accelerator that offloads artificial intelligence (AI) workloads from CPU 1302 and GPU 1306. Processor 1300 can use instructions that, if executed cause processor 1300 and / or any of its components to perform some or all of processes and techniques described elsewhere herein. Processor 1300 may include any number of memory and cache units 1310 to facilitate processing amongst different components of processor 1300. Memory and cache 1310 on processor 1300 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 1300 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 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. One or more of APIs described herein can include a call.

[0257] Processor 1300 can include compute engines as CPUs 1302 and can include any number of cores, such as, but not limited to, up to 16 cores / 22 threads. Cores in CPU 1302 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.

[0258] GPU 1306 can include any number of graphics engines, such as, but not limited to, Intel ArcTM graphics engines (Xe LPG) with 8 Xe cores (up to 128 Execution Units or EUs) . As shown in FIG. 13, GPU 1306 can include vector engines 1310 and matrix engines 1312, that, for example, can run FP, INT, and matrix operation tasks all at the same time or separately or in batches. GPU 1306 can include a load / store unit 1314, as well as other memory, such as, but not limited to, an instruction cache (I$) 1316 and L1 cache / subsystem local memory (SLM) 1318 that can, e.g., store instructions to perform any of the operations described above or elsewhere herein.

[0259] NPU 1304 can include one or more Intel AI Boost built-in neural processing unit (s) (NPUs) . NPU 1304 can be enumerated to a host processor as an integrated PCIe device. NPU 1304 can include one or more (e.g., two) Neural Compute Engine (NCE) tiles 1330. Each tile can be configured with any combination of, but not limited to, (e.g., 2000) Multiply Accumulate (MAC) Engines 1334, 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. 13. For general compute needs, Neural Compute Engines 1330 can include interference pipeline 1332, activation function (AF) 1336, data conversion 1338, load / store 1340, and Streaming Hybrid Architecture Vector Engines (SHAVE) 1328 for high performance parallel computing, which can include DMA (Direct Memory Access) engines 1324 to shuttle data between system memory DRAM (Dynamic Random Access Memory) 1326 and a software managed cache. Built-in device MMU (Memory Management Unit) 1322 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 1300 can also include a media unit (not shown) that is included on or separately from XCDs or other components of processor 1300 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.

[0260] A Intel Thread Director, which includes firmware that is built into processor 1300, 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. Intel Deep Learning Boost (Intel 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 1300 can be configured to execute an application program, such as, but not limited to, a CUDA program.

[0261] In at least one embodiment, processor 1300 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0262] Processor 1300 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 1304 as a Hexagon NPU, GPU 1306 as a Adreno GPU, CPU 1302 as a Kryo or Qualcomm Oryon CPU, as well as a Qualcomm Sensing Hub (not shown) and a memory subsystem 1310, in any combination. Hexagon NPU 1304 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 1306 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 1302 can perform AI workloads, and can handle contextualization for pervasive generative AI applications. CPU 1302 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 1300 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 1300 (e.g., in cache and / or memory) . Any number of CPU cores 1302 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 1302 can couple to memory subsystem 1310 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 1310 can include memory and cache on processor 1300, 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 1310 can be shared or used individually by any one or combinations of components (e.g., GPU 1306, NPU 1304, and CPU 1302) on processor 1300.

[0263] Qualcomm AI Engine 1300 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 Qualcomm 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 1306, OpenCL and DirectML may be supported. For CPU 1302, a LLVM compiler infrastructure optimizations enable accelerated and efficient AI inference. With respect to Qualcomm AI Engine 1300 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 1300 (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 1300, including registers, DRAM, flash, SRAM, cache, or other memory.

[0264] In at least one embodiment, processor 1300 or Qualcomm AI Engine 1300 can include one or more circuits to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands or otherwise perform any of the operations described above or elsewhere herein.

[0265] FIG. 14A illustrates a processor 1400, in accordance with at least one embodiment. Processor 1400 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 1400 can include one or more cores 1412 (1) -1412 (N) , where N is any integer greater than 1 that can perform the operations described elsewhere herein. Cores 1412 (1) -1412 (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 1412 (1) -1412 (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 1412 (1) -1412 (N) and can include a grid of converged mesh stops (CMS) that may be associated (e.g., 1: 1) with cores 1412 (1) -1412 (N) . Each core can be associated with one lower level cache (LLC) slice 1414 (1) -1414 (N) , or cores 1412 (1) -1412 (N) can share cache, e.g., lower level cache. LLCs 1414 (1) -1414 (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 Intel Ultra Path Interconnect (Intel UPI 1416) cache coherency functionality. UPI 1416 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.

[0266] Processor 1400 can also include System Agent 1410 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 1400 can include one or more integrated memory controller (s) (IMC) 1408. IMC 1408 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 1410 can include a display controller (not shown) to support display (s) . System Agent 1410 can also incorporate PCIe 1404 (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) 1406. System Agent 1410 can include an Image Processing Unit (IPU) (not shown) which incorporates an image signal processor (ISP) on-die. Fabric 1402 can provide scalability for connecting to other nodes (e.g., processors, such as processor 1400) , and can, for example, be used with Cornelis Networks, an element of Intel Scalable System Framework, that delivers the performance for high performance computing (HPC) workloads and the ability to scale to tens of thousands of nodes.

[0267] FIG. 14B illustrates components within core 1412, in accordance with at least one embodiment. Core 1412 can include front-end 1418, back-end or execution engine 1432, and memory subsystem 1442. Front-end 1418 can provide execution engine 1432 with operations (e.g., operations described elsewhere herein) by decoding instructions stored in memory. For example, front-end 1418 can include a micro-operations (μOps) cache path and / or a legacy path, along with branch prediction unit 1421 that can determine paths instructions. A legacy path for instructions may include fetching variable-length (e.g., x86) instructions from L1 instruction cache 1420 with instruction fetch and predecode 1422, queuing the instructions in instruction queue 1424, and decoding instructions using decoder 1426 into μOps that can be provided to allocation queue 1428. Alternatively, a μOPs cache path may include a cache containing already decoded μOps (μOps 1430) that can be sent to allocation queue 1428. Allocation queue 1428 can perform as an interface between front-end 1418 and execution engine 1432, and can provide instructions to execution engine 1432. 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 1418, execution engine 1432, and stored in memory subsystem 1442.

[0268] Execution engine 1432 can receive micro-operations into reorder buffer 1434, which can register allocation, rename, and retire μOPs. From reorder buffer, μOPs can be sent to scheduler 1436 that can be connected one or more different execution units 1438, which can be connected to address generation unit (AGU) 1440. Execution units 1438 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 1436 may manage queuing μOPs for one or more of execution units 1438 depending, e.g., on operations needed to be performed.

[0269] Memory subsystem 1442 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) 1444. Memory subsystem 1442 can also include shared or separate L1 data and instruction cache 1446, as well as L2 cache 1448 that can be used and shared by L1 data and instruction cache 1446. As described above for FIG. 14A, each core 1412 can be connected to a slice of a third level of cache (e.g., LLC 1414) that can be shared by all core 1412.

[0270] In at least one embodiment, processor 1400 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0271] FIG. 15 illustrates an AI accelerator 1500, in accordance with at least one embodiment. Processor 1500 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 1500 may use instructions that, if executed by AI accelerator 1500, cause AI accelerator 1500 to perform some or all of processes and techniques described elsewhere herein. For example, with respect to AI accelerator 1500 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 1500 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of AI accelerator 1500, including registers, DRAM, flash, SRAM, cache, or other memory. AI accelerator 1500 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.

[0272] In at least one embodiment, compute dies may include compute engines to perform AI computations. In at least one embodiment, AI accelerator 1500 compute dies may be split into any number of (e.g., four) clusters that may be referred to as a DCORE (Deep Learning Core)  1506 and contain any number of Matrix Multiplication Engines (MMEs) 1508, Tensor Processor  Cores (TPCs) 1510, memory management unit 1512, and L2 Cache 1514, in any combination. MME(s) 1508 can perform operations that use Matrix Multiplication, like fully connected layers, convolutions and batched-General Matrix Multiplications (GEMMs) . MMEs 1508 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) 1508 may be programmed with array dimensions, locations, data types, and various execution operands. MME (s) 1508 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) 1508 may push tensor C back to memory upon completion. TPC (s) 1510 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.

[0273] Memory on AI Accelerator 1500 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 1500 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 1508 and TPCs 1510. 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 1500 may include one or more Memory Management Units (MMUs) 1522 for managing memory, such as allowing AI accelerator 1500 memory subsystem to operate in a virtual space when accessing VRAM.

[0274] AI accelerator 1500 may include a communications port (e.g., a PCIe Gen5 X16 port) 1502 for communicating with a host and Scheduling and Synchronization Unit 1504. AI accelerator 1500 may include Media Unit 1516 that may include any number or combinations of Media Decoder Engines (DECs) 1520 and Rotator Engines (ROT) 1518. AI accelerator 1500 may include a network unit 1524 that may include any number or combinations of network ports 1526 and accompanied RDMA Engine (s) 1528, L2 Cache, and memory (e.g., HBM2e or HBM3) stacks. AI accelerator 1500 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.

[0275] AI accelerator 1500 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 1500 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 1500 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 1500 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.

[0276] RDMA 1528 over Converged Ethernet on AI accelerator 1500 may enable scaling from a single node (i.e., a single AI Accelerator 1500 to hundreds or thousands of nodes or AI Accelerators 1500) . NW Subsystem 1524 can include an Intel Gaudi 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 1526, a Layer2 MAC (not shown) , and RDMA Engines 1528. AI Accelerator 1500 can include Aggregation Engines for performing summing activities. All engines in processor 1500 can operate in parallel, e.g., MME(s) 1508, TPC (s) 1510 and NIC (s) 1526 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) .

[0277] AI Accelerator 1500 can be operated and controlled using software layer 1528 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 1528 may include implementations of algorithms, such as, but not limited to, Paged Attention, Flash Attention and more. Software layer 1528 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.

[0278] In at least one embodiment, AI accelerator 1500 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0279] 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. 16 is a simplified block diagram 1600 illustrating an example of at least a portion of such a neuromorphic computing device 1605, in accordance with at least one embodiment. Neuromorphic computing device 1605 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 1605 may be provided with a network 1610 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 1610 of spiking neural network cores may be provided in device 1605 and may each communicate via short packetized spike messages sent from core to core over network channels. Each core (e.g., 1615) 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 1605. With respect to neuromorphic computing device 1605 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 1605 (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 1605, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.

[0280] Continuing with the example of FIG. 16, neuromorphic computing device 1605 may additionally include processor 1620 and system memory 1625 to implement one or more components to manage and provide functionality of neuromorphic computing device 1605. For instance, system manager 1630 may be provided to manage global attributes and operations of neuromorphic computing device 1605 (e.g., attributes affecting network of cores 1610, multiple cores in network 1610, interconnections of neuromorphic computing device 1605 with other devices, manage access to global system memory 1625, among other potential examples) . In one example, system manager 1630 may manage the definition and provisioning of a specific routing tables to various routers in network 1610, orchestration of a network definition and attributes (e.g., weights, decay rates, etc. ) to be applied in network 1610, core synchronization and time multiplexing management, routing of inputs to appropriate cores, among other potential functions.

[0281] As another example, neuromorphic computing device 1605 may additionally include programming interface 1635 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 1610 of neuromorphic cores. A software-based programming tool may be provided with or separate from neuromorphic computing device 1605 through which a user may provide a definition for a particular neural network to be implemented using network 1610 of neuromorphic cores. Programming interface 1635 may take an input of a programmer to then generate corresponding routing tables and populate local memory of individual neuromorphic cores (e.g., 1615) with specified parameters to implement a corresponding, customized network of artificial neurons implemented by neuromorphic cores 1615.

[0282] In some cases, neuromorphic computing device 1605 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 1640 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 1640 may be utilized to accept input data from another device or external memory controller acting as a source of input data. External interface 1640 may be additionally or alternatively utilized to allow results or output of computations of a neural network implemented using neuromorphic computing device 1605 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.

[0283] As shown in FIG. 16, network 1610 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., 1615 a-d) . For instance, a number of neuromorphic cores (e.g., 1615 a-d) may be provided in a mesh, with each core being interconnected by a network including a number of routers (e.g., 1650) . In one implementation, each neuromorphic core (e.g., 1615 a-d) may be connected to a single one of routers (e.g., 1650) and routers may be connected to at least one other router (as shown at 1610 in FIG. 16) . As an example, in one particular implementation, four neuromorphic cores (e.g., 1615 a-d) may be connected to a single router (e.g., 1650) and each of routers 1650 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 1605. Moreover, as each neuromorphic core may be configured to implement multiple distinct neurons, router network of neuromorphic computing device 1605 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 1610 provided in neuromorphic computing device 1605.

[0284] FIG. 16 shows a block diagram illustrating internal components of one example implementation of neuromorphic core 1615. In one example, a single neuromorphic core may implement some number of neurons (e.g. 1024) that share architectural resources of neuromorphic core 1615 in a time-multiplexed manner. In one example, each neuromorphic core 1615 may include processor block 1655 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 1615 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 1615 may be tracked, parameters of each neuron implemented by core may 1615 be recorded, among other data and usage. Components, or architectural resources, of neuromorphic core 1615 may further include input interface 1665 to accept input spike messages generated by other neurons on other neuromorphic cores and output interface 1670 to send spike messages to other neuromorphic cores over mesh network 1610. In some instances, routing logic for neuromorphic core 1615 may be at least partially implemented using output interface 1670. Further, in some cases, core (e.g., 1615) 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 1615 may forego communication over routing fabric of neuromorphic computing device 1605 and may instead by managed locally at particular neuromorphic core 1615.

[0285] Each neuromorphic core may additionally include logic to implement, for each neuron 1675, artificial dendrite 1680 and artificial soma 1685 (referred to herein, simply, as “dendrite” and “soma” respectively) . Dendrite 1680 may be a hardware-implemented process that receives spikes from network 1610. Soma 1685 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 1680 may be defined for each connection receiving inputs from another source (e.g., another neuron) . In one implementation, dendrite process 1680 may receive and handle spike messages as they serially arrive in time-multiplexed fashion from network 1610. As spikes are received, neuron's activation (tracked using soma 1685 (and local memory 1660) ) may increase. When neuron's activation exceeds a threshold set for neuron 1675, neuron 1675 may generate a spike message that is propagated to a fixed set of fanout neurons via output interface 1670. 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.

[0286] As noted above, neuromorphic computing device 1605 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. 16 may advantageously supports all of these network models. As some or all cores of neuromorphic computing device 1605 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 1605 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 1610 to realize fully arbitrary connectivity graphs.

[0287] 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. 16, 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.

[0288] As an example, a neuromorphic processor may utilize time-multiplexed computation in both a spike communication network and neuron machinery of neuromorphic computing device 1605 to implement SNNs. Accordingly, physical circuitry of neuromorphic computing device 1605 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) .

[0289] 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.

[0290] In operation, neuromorphic computing device 1605 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.

[0291] Given this context, and as introduced above, a device (e.g., 1605) implementing a mesh 1610 of interconnected neuromorphic cores may be provided, with core 1615 implementing potentially multiple artificial neurons capable of being interconnected to implement an SNN. Each neuromorphic core (e.g., 1615) may provide two loosely coupled asynchronous processes: an input dendrite process (e.g., 1680) that receives spikes from network 1610 and applies them to an appropriate destination dendrite compartments at the appropriate future times, and output soma process (e.g., 1685) 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.

[0292] In at least one embodiment, neuromorphic computing device 1605 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0293] FIG. 17 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 1700 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 1700 can represent a data center. System 1700 may represent a server farm. System 1700 may represent a data cloud or a processing cloud. System 1700 can represent a supercomputer. System 17 may include tens, hundreds, or thousands of nodes. Nodes of system 1700 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. 11-23. With respect to any of processors in system 1700 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 1700 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.

[0294] One or more clients 1702 make requests over network 1704 to system 1700. Network 1704 represents one or more local networks, or wide area networks, or a combination. Clients 1702 can be human or machine clients, which generate requests for execution of operations by system 1700. System 1700 executes applications or data computation tasks requested by clients 1702.

[0295] System 1700 can include one or more racks, which represent structural and interconnect resources to house and interconnect multiple computation nodes. Rack 1710 can include multiple nodes 1730. Rack 1710 may host multiple blade components 1720 (0) to 1720 (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 1720 (0) to 1720 (N-1) can refer to computing resources on printed circuit boards (PCBs) , where a PCB houses hardware components for one or more nodes 1730. Blades 1720 (0) to 1720 (N-1) may or may not include a chassis or housing or other "box" other than that provided by rack 1710. Blades 1720 (0) to 1720 (N-1) may include housing with exposed connector to connect into rack 1710. System 1700 may or may not include rack 1710, and each blade (e.g., 1720 (0) ) can include a chassis or housing that can stack or otherwise reside in close proximity to other blades and allow interconnection of nodes 1730. System 1700 may include 10, 624 compute blades, which include 63, 744 Intel Max Series GPUs and 21, 248 Intel Xeon Max CPUs across 166 racks.

[0296] System 1700 can include fabric 1770, which represents one or more interconnectors for nodes 1730. Fabric 1770 can include multiple switches 1772 or routers or other hardware to route signals among nodes 1730. Additionally, fabric 1770 can couple system 1700 to network 1704 for access by clients 1702. In addition to routing equipment, fabric 1770 can be considered to include cables or ports or other hardware equipment to couples nodes 1730 together. Fabric 1770 can have one or more associated protocols to manage routing of signals through system 1700. A protocol or protocols is at least partly dependent on hardware equipment used in system 1700.

[0297] As illustrated, rack 1710 can include N blades (e.g., 1720 (0) to 1720 (N-1) ) . In addition to rack 1710, system 1700 can include rack 1750. As illustrated, rack 1750 may include M blades (e.g., 1760 (0) to 1760 (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 1700 over fabric 1770. Blades 1760 (0) to 1760 (M-1) can be the same or similar to blades 1720 (0) to 1720 (N-1) . Nodes 1730 can be any type of node as described herein, and may not be necessarily all the same type of node. System 1700 is not limited to being homogenous, nor is it limited to not being homogenous.

[0298] A node in blade 1720 (0) is illustrated in detail. However, other nodes in system 1700 can be the same or similar. At least some nodes 1730 may be computation nodes, with processor 1732 and memory 1740. 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 1730 can include storage server nodes with a server as processing resources 1732 and memory 1740. 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.

[0299] Node 1730 can include interface controller 1734, which can represent logic to control access by node 1730 to fabric 1770. Logic can include hardware resources to interconnect to physical interconnection hardware. Logic can include software or firmware logic to manage interconnection. Interface controller 1734 can include a host fabric interface, which can include a fabric interface in accordance with any embodiment described herein.

[0300] Node 1730 may include memory subsystem 1740. Memory 1740 can include memory computation resources (comp) 1742, which represent one or more capabilities by memory 1740 to perform memory computations. System 1700 enables remote memory operations, such as, but not limited to, the operations described elsewhere herein. Thus, nodes 1730 can request memory computations by remote nodes, where data for computation remains local to an executing node instead of being sent over fabric 1770 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.

[0301] Processor 1732 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 1740 can be or include memory devices and a memory controller.

[0302] 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.

[0303] 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.

[0304] In at least one embodiment, system 1700 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0305] FIG. 18 illustrates accelerated processing unit 1800, in accordance with at least one embodiment. Accelerated processing unit 1800 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 1800 can include one or more accelerator complex dies (XCDs) 1804 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 1800 can include one or more complex compute dies (CCDs) 1806 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 1802) , 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 1808 and, in some embodiments, high-bandwidth memory (e.g., HMB3) . Accelerated processing unit 1800 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.

[0306] Accelerated processing unit 1800 can include one or more input / output (I / O) interfaces. For example, XCDs 1804 and CCDs 1806 can be together on one or more input-output dies (IODs) 1810 that can include one or more I / O interfaces. IODs 1810 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 1870. I / O interfaces from IODs 1810 can also be used for connected one or more accelerated processing units 1800, e.g., in a server architecture.

[0307] Accelerated processing unit 1800 can include one or more memory units 1802 for storing instructions and other information used to perform operations described elsewhere herein. Memory units 1802 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 1800 (e.g., memory units 1802) can include system memory that can be used, for example, for commands, instructions and constants, and inputs and outputs. Memory units 1802 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 1802 can be linked to one or more IODs 1810. In at least on embodiment, L1 cache 1820 starts a memory hierarchy that includes shared L2 cache 1828, e.g., within XCDs. AMD Infinity CacheTM, which is a last level cache (LLC) located on an active I / O die (IOD) . CCDs 1806 and XCDs 1804 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.

[0308] As shown in FIG. 18, an XCD 1804 can include a shared set of global resources 1830, which can include hardware scheduler 1832 and Asynchronous Compute Engines (ACE) 1824 that send tasks (e.g., compute shader workgroups) to Compute Units (CUs or cores) 1834. ACEs 1824 (e.g., four) can be each associated with CUs 1834 (e.g., 40 CUs) , and some of CUs 1834 can be disabled for yield management. CUs 1834 can have dedicated cache or share cache (e.g., L2 cache) 1828 that may be used to coalesce all memory traffic for a die. CUs 1834 can include threaded and parallel processor cores including instruction fetching and scheduling with Scheduler (S) 1812, matrix core unit (MCU) 1816 and shader core (SC) 1818 (e.g., execution units for scalar, vector and matrix data types) , as well as load / store pipelines with an L1 cache 1820 and Local Data Share (LDS) 1814. 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 1840 (e.g., for storing and providing instructions for performing operations described elsewhere herein) and a constant cache 1838 can be connected to one or more CUs and can be shared between two CUs. Matrix cores 1816 can process a variety of data types, such as, but not limited to, INT8, FP8, FP16, BF16 and TF32 data types. Accelerated processing unit 1800 can include compute units 1834 that may be arranged in an array format, e.g., as a data-parallel-processor (DPP) array. Ultra-threaded dispatch processor 1842 can communicate with compute units 1834, and command processor 1844 can read commands that a host has written to memory-mapped registers in a system-memory address space (not shown) . Command processor 1844 can send hardware-generated interrupts to a host processor (e.g., a CCD) when a command is completed. Memory controller 1836 can also have direct access to all device memory and host-specified areas of system memory. To satisfy read and write requests, memory controller 1836 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 1840 and then fetched by instruction fetch logic in processor 1840, 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 1800 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 1800, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.

[0309] 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 1800 can operate, invalidate and flush caches on accelerated processing unit 1800, and cause accelerated processing unit 1800 to begin execution of a program. Kernels can be referred to as programs executed by accelerated processing unit 1800. 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 1834 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 1814, 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.

[0310] In at least one embodiment, accelerated processing unit 1800 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein.

[0311] One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0312] FIG. 19 illustrates a processor 1900, 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 1900 includes one or more CPU dies 1902 (1) -1902 (N) , where N is any integer greater than 1. CPU die 1902 can include any number of processor cores 1916 (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 1918 can be coupled to processor core (s) 1916, which can share and / or couple individually to L2 Cache units 1918. Processor cores 1916 can couple to L3 cache 1922 individually and / or share L3 Cache, which can be a lowest level cache (LLC) 1922 for access to data and other information used by processor cores 1916. One or more processor cores 1916 and one or more L2 Cache units 1918 can be included in a core complex (CCX) 1920 that can include (e.g., a 32 MB) shared cache (e.g., L3 cache 1922) . Core complex 1920 can be fabricated onto a die (CCD or CPU die) 1902. For example, up to 12 core complexes 1920 can be configured into a processor along with 8 CPU dies 1902 to provide up to 96 processor cores 1916 for processor 1900. A ‘Zen 4c’ core complex 1920, for example, can include up to eight cores 1916 and a shared 16 MB L3 cache 1922. Two of these core complexes 1920 can be combined onto a single CPU die 1902 for 16 cores per die and a total of 32 MB of L3 cache 1922 per die. Up to eight of CPU dies 1902 may be combined with an I / O unit 1904 to provide CPUs with up to 128 processor cores 1916. Up to four ‘Zen 4c’ dies described above can be combined to provide CPUs with up to 64 processor cores 1916.

[0313] Processor 1900 can include a variety of configurations for input / output operations that are described further herein. I / O unit 1904 can include one or more memory controllers 1906 that can manage memory usage (e.g., DDR5 memory) for processor 1900. I / O unit 1904 may include one or more SATA disk controllers for managing storage 1912 and one or more Compute Express Link (CXLTM) 1.1+ memory controllers 1914 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 1904 may include PCIe controller 1908 for connecting peripherals and other components connected to processor 1900. I / O unit 1904 may include USB ports 1910 for connecting to other components separate from processor 1900. CPU dies 1902 can support any number of connections, e.g., one or two connections, to I / O unit 1904. As shown, I / O unit 1904 can include components described further herein, and I / O unit 1904 can be a I / O die that houses several different components. Memory controller 1906, PCIe controller 1908, USB ports 1910, SATA controller 1912, and / or CXL controller 1914 can be integrated anywhere within processor 1900 either separately or in any groups or combinations thereof.

[0314] Processor 1900 can include Infinity Fabric 1924 interconnects (which can be similar to or based on PCIe architectures) that can provide connections among CPUs (e.g., CPU dies 1902 (1) -1902 (N) ) , graphics processor (s) 1926, inference engine (s) 1932, and other components in a multi-chip architecture, such as secure processor (s) 1928 and I / O unit 1904. One or more AMD Infinity FabricTM interconnects 1910 can connect to CPU dies 1902 (1) -1902 (N) and serve as a connection that is used between CPUs. One or more Infinity Fabric connections 1910 can connect each CPU die 1902 to I / O unit 1910.

[0315] In at least one embodiment, processor 1900 can include central processing units (CPUs) and other associated hardware and software described above and further herein. Processor 1900 can also include graphics processor (s) 1926. Graphics processor 1926 can be used for image generation and processing, as well as other computations and operations described further herein. Graphics processor 1926 can be based on RDNA 3 or 3.5 architecture from AMD in Santa Clara, CA. Graphics processor 1926 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 1926 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 1942 (1) - (N) , where N is an integer. Components within graphics processor 1926 can be connected using various approaches, such as using Infinity Fabric 1924 interconnects outside or within graphics processor 1926.

[0316] Inference engine 1932 can provide neural processing capabilities for processor 1900 for computational processes that are used for neural networks, deep learning, and other artificial intelligence-related operations described further herein. Processor 1900 can include secure processor (s) 1928 for managing security of processor 1900, display controller 1930 for controlling displays, a system management unit 1934 for managing and operating some or all of the components on processor 1900, multimedia engines 1936 for audio and video operations, fusion controller hub 1938 for managing USB, SATA and PCIe connections to processor 1900, and sensor fusion hub 1940 for managing sensors, such as accelerometers. Processor 1900 can also include memory 1942 (1) - (N) , where N is any integer. Memory can include different memory types, such as LPDDR5 and / or DDR5, or others described elsewhere herein.

[0317] For performing operations described further herein, processor 1900 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.

[0318] With respect to processor 1900 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 1900 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 1900, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.

[0319] In at least one embodiment, processor 1900 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0320] FIG. 20 illustrates an example of a processing core 2000 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 2000 can be implemented inside a DynamIQ Shared Unit (DSU) cluster via DSU-110 interconnect 2054 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 2052 that connects core 2000 to DSU-110 interconnect, which can also connect core 2000 to an external memory system and the rest of a system-on-a-chip. L1 instruction memory system 2002 can fetch instructions from an instruction cache 2004 and deliver instructions (e.g., one or more APIs described herein that may be compiled into instructions) to an instruction decode unit 2010, e.g., to perform some or all of operations described above or elsewhere herein. L1 instruction memory system 2002 may include L1 instruction cache 2004, e.g., with 64-byte cache lines, L1 instruction Translation Lookaside Buffer (TLB) 2006, e.g., with native support for 4KB, 16KB, 64KB, and 2MB page sizes, Macro-Operation Cache (MOP) 2008 (e.g., 1536-entry, 4-way skewed associative L0 MOP cache) , which can contain decoded and optimized instructions for higher performance. Instruction decode unit 2010 can decode AArch64 instructions into internal format. Register rename unit 2012 can perform register renaming to facilitate out-of-order execution and dispatches decoded instructions to various issue queues. Instruction issue unit 2014 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 2016 can be included in an execution pipeline and include integer execute unit 2018 that can perform arithmetic and logical data processing operations. Vector execute unit 2020 can be included in an execution pipeline and can perform Advanced SIMD and floating-point operations (FPU) 2022, execute Scalable Vector Extension (SVE) and Scalable Vector Extension 2 (SVE2) instructions 2024, and can optionally execute cryptographic instructions (Crypto) 2026. 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 2030 can execute load and store instructions, as well as service memory coherency requests. L1 data memory system 2030 can include an L1 data cache 2032 and a fully associative L1 data TLB 2034 with native support for 4KB, 16KB and 64KB page sizes and 2MB and 512MB block sizes. Memory Management Unit (MMU) 2028 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 2034 when an address is translated. L2 memory system 2036 can include L2 cache 2038, and it can be connected to DSU-110 2054 through an asynchronous CPU bridge 2052. NeoverseTM V2 core 2000 can support a range of debug, test, and trace options including a trace unit 2042 and a trace buffer 2040, and an Embedded Logic Analyzer (ELA) 2048. NeoverseTM V2 core 2000 can implement Statistical Profiling Extension (SPE) 2044 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) 2046 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 2050, 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 2052 between each NeoverseTM V2 core 2000 and DSU-110 2054. CPU bridge 2052 can control buffering and synchronization between core 2000 and DSU-1102054. CPU bridge 2052 can be asynchronous to allow different frequency, power, and area implementation points for each core 2000. CPU bridge 2052 can run synchronously without affecting other interfaces such as, but not limited to, debug and trace which can be asynchronous.

[0321] In at least one embodiment, core 2000 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0322] FIG. 21 illustrates one or more chips including one or more tensor processing units (TPUs) 2100, in accordance with at least one embodiment. TPUs 2100 in FIG. 21 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 2100 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.

[0323] Chip 2100 can include any number of TPUs that can include tensor cores 2106. Tensor core 2106 can include one or more core sequencer 2108, vector processing unit (VPU) 2110, matrix multiply unit (MXU) 2112 (A) -2114 (N) , where N is any integer greater than 1, and a transpose permute unit 2116. Core Sequencer 2108 can fetch (e.g., VLIW (Very Long Instruction Word) ) instructions from core’s 2106 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) (2110. 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 2110 can perform vector operations using a large on-chip vector memory (Vmem) , and vector registers (Vregs) . VPU 2110 can stream data to and from MXU through decoupling FIFOs. VPU 2110 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) 2112 (A) -2112 (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 2116 can do (e.g., 128x128) matrix transposes, reductions, and permutations of VPU 2110 lanes. High Bandwidth Memory 2104 can be used for applications on chip, and it can be coupled to host queue (s) 2102, e.g., over PCIe. One or more chips 2100 can be connected together for computing. For example, one or more chips 2100 can be connected as a torus, e.g., a 2D torus. Chip 2100 can also include any number (e.g., four) Inter-Core Interconnect (ICI) links 2118 that can enable direct connections between chips to form a supercomputer.

[0324] With respect to any processors in chip 2100 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 2100 (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 2100, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.

[0325] In at least one embodiment, chip 2100 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0326] FIG. 22 illustrates a vector processor, in accordance with at least one embodiment. Vector processor 2200 may support a RISC-V standard. Vector processor 2200 can include one more cores 2210 (e.g., scalar units) with one or more Vector Processing Units (VPUs) 2242 (e.g., vector units) that can, e.g., perform some or all of the operations described above or elsewhere herein. Core 2210 may include Andes Custom Extension (ACE) 2216 that can be used for communication of customized instructions for processor 2200, for example, via ACP 2238. Core 2210 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) 2224 may manage system memory and cache, and provide for branch execution, issuance of instruction pairs, L1 instruction / data caches and local memory storage. Core 2210 can include Physical memory protection and programmable physical memory attribute unit (PMP / PPMA) 2222. Core 2210 can include a digital signal processor (DSP) 2228, and a floating-point unit (FPU) 2226 as well as load-store unit (LSU) 2232 to interface with memory hierarchy (D$ 2234 and I$ 2230) . Core 2210 can include branch prediction unit 2218 and multiplier unit 2220.

[0327] Vector processing unit (VPU) 2242 can include one or more vector functional units (FUs) 2246 (A) -2246 (N) that can be chained together for parallel processing, independent memory paths for RISC-V vector (RVV) load / store via ACE-RVV 2248 and Andes Streaming port (ASP) 2244 load / store, and a vector load / store unit (VLSU) 2250.

[0328] Vector processor 2200 can include bus interfaces, such as, but not limited to, L2 cache memory port 2256 for cacheable access, a MMIO port 2254 for non-cacheable access, an input-output coherence Port (IOCP) 2258 for cacheless bus master, local memory access ports for ILM / DLM 2212, which can be coupled to SRAM 2206, and high-bandwidth vector memory (HVM) 2236 access, a shared peripheral port (SPP) 2252 for external peripherals. Other memory ports include LM slave port AXI 2202, HVM subordinate port AXI 2204, MEM (AXI) 2262, and AXI 2260. Trace I / F 2214 can capture, encode, and transmit off-chip via Inst. Trace I / F 2208, e.g., a record of executed processor instructions, which software tools can use to reconstruct the exact execution sequence of a program.

[0329] With respect to any processors in processor 2200 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 2200 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 2200, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.

[0330] In at least one embodiment, vector processor 2200 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0331] FIG. 23A illustrates a diagram of an example many-core tiled processor microarchitecture. Many-core tiled processor in FIG. 23A can include a language processing processor. As illustrated in FIG. 23A, 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 2304 and an integer (INT) 2306 and floating-point (FP) unit 2308 as well as load-store unit (LSU) 2312 to interface with memory hierarchy (data cache (D$) 2310 and instruction cache (I$) 2314) and network (NET) 2316 interface for communication with other tiles. Some tiles in processor 2300 may include memory controller 2302 for managing and controlling memory, as described further herein. Processor 2300 can have a functional slice architecture. Processor 2300 may be located on an application specific integrated circuit (ASIC) , and FIG. 23A may represent a layout of an ASIC. Processor 2300 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 2300 is a tensor streaming processor.

[0332] Processor 2300 can employ different microarchitectures, which disaggregates functional units shown in each tile in FIG. 23B. Instead, functional tiles 2324 of processor 2300 may be aggregated into a plurality of functional process units (hereafter referred to as “slices” ) 2304, each corresponding to a particular function type (e.g., FP / INT 2318, NET 2320, MEM 2322) . For example, as illustrated in FIG. 23B, each slice may correspond to a column of functional tiles extending in a north-south direction. In addition, processor 2300 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 2304 of processor 2300.

[0333] Slices 2304 of processor 2300 may each correspond to a different function, and may include arithmetic logic slices (e.g., FP / INT2318) , lane switching slices (e.g., NET 2320) , and memory slices (e.g., MEM 2322) . 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 2304 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 2300.

[0334] By arranging tiles of processor 2300 into different functional slices 2304, on-chip instruction and control flow of processor 2300 can be decoupled from data flow. For example, one arrow in FIG. 23B illustrates flow of instructions within processor architecture, in accordance with some embodiments. Another arrow in FIG. 23B 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 2300 (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 2300 (e.g., east-west, across functional slices, as shown by the second arrow) that is perpendicular to the first direction.

[0335] Different functional slices of processor 2300 may correspond to MEM 2322 (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 2300 can correspond to an execution unit organized as an ×M SIMD tile. For example, each tile of on-chip memory of processor 2300 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) .

[0336] 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.

[0337] When using processor 2300 (e.g., TSP) having a functional slice architecture, TSP compiler (not shown) generates an explicit plan for how processor 2300 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 2300 (e.g., TSP) hardware state so a microprogram can orchestrate data flow.

[0338] Processor 2300 (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 2300 (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.

[0339] MEM 2322 of processor 2300 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 2300. In some embodiments, due to bandwidth requirements of processor 2300, on-chip memory of MEM tiles may include SRAM, and not DRAM. On-chip memory capacity of processor 2300 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 2300 can provide a plurality of memory slices organized into two different hemispheres (referred to as “MEM WEST” and “MEM EAST” , respectively) .

[0340] 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 2300 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.

[0341] 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 2300. 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.

[0342] With respect to any processors in FIG. 23 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 2300 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of processor 2300, including registers, DRAM, flash, SRAM, cache, or other memory equivalents.

[0343] In at least one embodiment, processor 2300 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, or otherwise perform any of the operations described above or elsewhere herein. One or more circuits can be configured by software to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0344] Software Constructions

[0345] The following figures set forth, without limitation, examples of software constructs for implementing at least one embodiment.

[0346] FIG. 24 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.

[0347] A software stack 2400 of a programming platform can provide an execution environment for an application 2401. Application 2401 may include any computer software capable of being launched on software stack 2400. Application 2401 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.

[0348] Application 2401 and software stack 2400 run on hardware 2408. Hardware 2408 may include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of compute devices that support a programming platform. Software stack 2400 may be vendor specific and compatible with only devices from particular vendor (s) , such as CUDA, ROCm, OneAPI, OpenCL, or other implementations. Hardware 2408 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 2408 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 2408 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 2408 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 2408 (e.g., in cache and / or memory) . A result of API (s) can then be stored in storage within or outside of hardware 2408, 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.

[0349] Software stack 2400 of a programming platform can include a number of libraries 2403, a runtime 2405, an optional driver / interface 2407, and a device kernel driver 2408. Each of libraries 2403 may include data and programming code that can be used by computer programs and leveraged during software development. Libraries 2403 may include pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. Libraries 2403 can include functions that may be optimized for execution on one or more types of devices. Libraries 2403 may include functions for performing mathematical, deep learning, and / or other types of operations on devices. Libraries 2403 can be associated with corresponding APIs 2402, which may include one or more APIs, that expose functions implemented in libraries 2403. 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.

[0350] Software stack 2400 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 2400 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 2400 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 2400 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 2400 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.

[0351] Software stack 2400 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 2400 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 2400 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) .

[0352] Software stack 2400 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 2400 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.

[0353] Application 2401 can be written as source code that is compiled into executable code, as discussed in greater detail below in conjunction with FIGs. 25 and 26. Executable code of application 2401 may run, at least in part, on an execution environment provided by software stack 2400. During execution of application 2401, code may be reached that needs to run on a device, as opposed to a host. In such a case, runtime 2405 may be called to load and launch requisite code on a device. Runtime 2405 may include any technically feasible runtime system that is able to support execution of application 2401.

[0354] Runtime 2405 can be implemented as one or more runtime libraries associated with corresponding APIs, which are shown as API (s) 2404. 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.

[0355] Runtime libraries and corresponding API (s) 2404 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.

[0356] An optional driver or interface 2407 may be implemented, e.g., for CUDA and ROCm implementations, that are described further below. Optional driver / interface 2407 may be associated with optional driver or interface API (s) , such as, but not limited to, CUDA and / or ROCm API (s) .

[0357] One or more processors disclosed in “processing systems” can perform, access, or otherwise use software stack 2400. For example, system-on-a-chip 1100, parallel processor 1200, graphics multiprocessor 1234, processor 1300, processor 1400, accelerator 1500, neuromorphic processor 1605, supercomputer 1700, acceleration processing unit 1800, processor 1900, processor 2000, tensor processing unit 2100, processor 2200, and language processing unit 2300 can perform, use, call, or otherwise implement (e.g., through accessing a memory) one or more APIs included in software stack 2400.

[0358] Device kernel driver 2408 can be configured to facilitate communication with an underlying device. Device kernel driver 2408 may provide low-level functionalities upon which APIs, such as, but not limited to, API (s) 2404, and / or other software relies. Device kernel driver 2408 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 2408 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 2408 to compile IR code at runtime.

[0359] In at least one embodiment, processors in FIGs. 11-23 can include one or more circuits to cause one or more matrix operands to be transposed by loading two or more portions of those one or more matrix operands by two or more software threads concurrently, 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 2400, to cause each of the two or more software threads to generate in parallel a different row or column of one or more transposed matrix operands, or otherwise perform any of the operations described above or elsewhere herein.

[0360] In accordance with at least one embodiment, software stack 2400 of FIG. 24 can be performed in a CUDA implementation. A CUDA software stack 2400, on which an application 2401 may be launched, may include CUDA libraries 2403, a CUDA runtime 2405, a CUDA driver 2407, and a device kernel driver 2408. CUDA software stack 2400 can execute on hardware (e.g., graphics multiprocessor 1234 that may include a GPU that supports CUDA and is developed by NVIDIA Corporation of Santa Clara, CA.

[0361] Application 2401, CUDA runtime 2405, and device kernel driver 2408 can perform functionalities that are described above and elsewhere herein. CUDA driver 2407 can includ...

Claims

A processor comprising:one or more circuits to cause one or more matrix operands to be transposed based, at least in part, on two or more portions of the one or more matrix operands being loaded by two or more software threads concurrently.The processor of claim 1, wherein the one or more circuits are to cause each of the two or more software threads to generate, in parallel, a different row or column of the one or more transposed matrix operands.The processor of claim 1, wherein the one or more matrix operands comprise one or more arrays of one or more integers.The processor of claim 1, wherein the one or more circuits are to identify, within the two or more portions, one or more bits to be operated upon based, at least in part, on one or more bitmasks.The processor of claim 1, wherein the one or more circuits are to cause one or more bits within one or more bit arrays of the two or more portions to be shifted.The processor of claim 1, wherein the one or more circuits are to cause the one or more matrix operands to be transposed as part of a data center used to perform one or more software applications.The processor of claim 1, wherein the one or more circuits are to load the two or more portions based, at least in part, on a size of a group of consecutive threads.A system, comprising:one or more processors to cause one or more matrix operands to be transposed based, at least in part, on two or more portions of the one or more matrix operands being loaded by two or more software threads concurrently.The system of claim 8, wherein the one or more processors are to cause each of the two or more software threads to generate, in parallel, a different vector of the one or more transposed matrix operands.The system of claim 8, wherein the one or more matrix operands comprise one or more bit matrix operands representing one or more arrays of one or more integers.The system of claim 8, wherein the one or more processors are to use one or more masks to identify within the two or more portions one or more bits to be operated upon.The system of claim 8, wherein the one or more processors are to shift one or more bits within one or more bit arrays of the two or more portions based, at least in part, on a rank of consecutive threads.The system of claim 8, wherein the one or more processors are to cause the one or more matrix operands to be transposed as part of an algorithm used to assign a set of tasks to a set of computing systems.The system of claim 8, wherein the one or more processors are to load the two or more portions based, at least in part, on pairs of software threads identified to exchange their respective data with each other.A method, comprising:causing one or more matrix operands to be transposed based, at least in part, on two or more portions of the one or more matrix operands being loaded by two or more software threads concurrently.The method of claim 15, further comprising causing each of the two or more software threads to generate, in parallel, a different vector along a dimension of the one or more transposed matrix operands.The method of claim 15, wherein the one or more matrix operands comprise one or more arrays of one or more values represented as bits.The method of claim 15, further comprising using one or more masks to identify within the two or more portions one or more bits to be modified based, at least in part, on a conjunction operation.The method of claim 15, further comprising identifying pairs of software threads used to exchange their respective data with each other based, at least in part, on a number that is a power of two representing a size of a range of consecutive threads within a warp of threads.The method of claim 15, wherein the one or more matrix operands are to be transposed as part of a system operating one or more autonomous vehicles.

Images