Matrix multiplication with block scaling

Abstract

The present disclosure relates to matrix multiplication with block scaling, specifically disclosing apparatuses, systems, and techniques for performing matrix multiply-accumulate operations in which specified portions of an input matrix are block scaled by scaling factors. In at least one embodiment, an input matrix is multiplied with a scaled matrix that is generated by multiplying elements of a second input matrix by one or more corresponding scaling factors stored in a tensor memory.

Description

[0001] Claiming priority This application claims the benefit of Indian Provisional Application No. 202511032739 entitled “MATRIX MULTIPLICATION INSTRUCTIONS”, filed on 2 April 2025, the entire contents of which are incorporated herein by reference. Technical Field

[0002] At least one embodiment relates to processing resources for performing matrix multiplication on scaled operands. Background Technology

[0003] Matrix multiplication is used in neural networks to compute numerical quantities for backpropagation, such as the output of a network layer, the gradient of the weights, and the error signal. As an example, the weights of a layer can be represented as a weight matrix, while a batch of multiple inputs can be represented as an input matrix. The output of the layer is computed by multiplying the weight matrix by the input matrix. The value of each matrix element is stored in memory in floating-point decimal data format, for example, 32 bits per value. Larger data sizes (e.g., 32 bits) can provide higher computational precision but consume more computational resources compared to smaller data sizes. Smaller data sizes, such as 16-bit floating-point format, can be used to accelerate computation on devices with resource constraints (e.g., memory bandwidth bottlenecks). Smaller data sizes, such as 8-bit format, can be used to significantly reduce memory usage and improve inference speed, but this smaller data format also significantly reduces the precision of the results. It can improve the amount of time, memory, or computational resources used in computer programs to perform matrix-multiplication-accumulation operations. Attached Figure Description

[0004] Figure 1 The illustration shows a processor according to at least one embodiment for executing block-scaled matrix multiply accumulate (MMA) instructions, the inputs and outputs of which are stored in memory; Figure 2 The illustration shows an interaction in a processor for executing block scaling MMA instructions according to at least one embodiment, the inputs and outputs of which are stored in memory; Figure 3 The illustration shows a block scaling MMA operation according to at least one embodiment, wherein the input and output matrices are stored in tensor memory or shared memory; Figure 4A The diagram illustrates an instruction descriptor for a block scaling MMA instruction according to at least one embodiment; Figure 4BThe illustration shows a block-scaled data layout in a tensor memory for various scaling vector sizes according to at least one embodiment; Figure 5 The illustration shows scaling factors selected from tensor memory for various scaling factor selector values ​​according to at least one embodiment; Figure 6 The illustration shows a processor for performing a block scaling MMA call according to at least one embodiment, wherein the inputs and outputs of the block scaling MMA call are stored in registers; Figure 7 The illustration shows the interaction in a processor for executing block scaling MMA instructions according to at least one embodiment, wherein the input and output matrices of the block scaling MMA instructions are stored in registers; Figure 8 The illustration shows a block scaling MMA operation according to at least one embodiment, with its inputs and outputs stored in registers; Figure 9 The illustration shows a block scaling MMA operation using a scaling matrix selected from scaling data, according to at least one embodiment; Figure 10 The illustration shows a block scaling MMA operation using a scaling matrix with a scaling vector size of 1X, according to at least one embodiment; Figure 11 The illustration shows a block scaling MMA operation using a scaling matrix with a scaling vector size of 2X, according to at least one embodiment; Figure 12 The illustration shows a block scaling MMA operation using a scaling matrix with a scaling vector size of 4X, according to at least one embodiment; Figure 13 The illustration depicts an example process of performing block scaling MMA operations using data stored in tensor memory and / or shared memory, according to at least one embodiment; Figure 14 The illustration depicts an example process of performing block scaling MMA operations using a matrix stored in a register, according to at least one embodiment; Figure 15 The illustration shows an example process for selecting scaling matrix elements from scaling data according to at least one embodiment; Figure 16 The illustration shows scaling factors selected from a register for various scaling factor selector values ​​according to at least one embodiment; Figure 17A The illustration shows a fragment layout of an input matrix according to at least one embodiment and associated thread IDs of a thread bundle, the thread bundle including threads that contribute scaling factors from the scaling data; Figure 17BThe diagram illustrates a pair of threads contributing scaling factors from scaling data A to form a scaling matrix A according to at least one embodiment; Figure 18 The diagram illustrates a thread that contributes scaling factors from scaling data B to form a scaling matrix B according to at least one embodiment; Figure 19 An example of a system including a processor and modules according to at least one embodiment is illustrated; Figure 20 An example data center system according to at least one embodiment is shown; Figure 21 A system-on-a-chip (SOC) according to at least one embodiment is shown. Figure 22A A parallel processor according to at least one embodiment is shown; Figure 22B A processing cluster according to at least one embodiment is shown; Figure 22C A graphics multiprocessor according to at least one embodiment is shown; Figure 23 An accelerator processor according to at least one embodiment is shown; Figure 24A A central processing unit according to at least one embodiment is shown; Figure 24B The illustration shows an embodiment according to at least one of the embodiments. Figure 24A The core of the central processing unit; Figure 25 Another accelerator processor according to at least one embodiment is shown; Figure 26 A neuromorphic processor according to at least one embodiment is shown; Figure 27 A supercomputer according to at least one embodiment is shown; Figure 28 Another accelerator processor according to at least one embodiment is shown; Figure 29 Another processor according to at least one embodiment is shown; Figure 30 Another accelerator processor according to at least one embodiment is shown; Figure 31 A tensor processing unit according to at least one embodiment is shown; Figure 32 A RISC-V compatible processor according to at least one embodiment is shown; Figure 33A and Figure 33B A language processing unit according to at least one embodiment is shown; Figure 34 A software stack of a programming platform according to at least one embodiment is shown; Figure 35 Software supported by a programming platform according to at least one embodiment is shown; Figure 36 A method for using at least one embodiment is shown. Figure 35 Compiled code executed on the programming platform; Figure 37 An example of an autonomous vehicle and its system architecture according to at least one embodiment is shown; Figure 38A The inference and / or training logic according to at least one embodiment is illustrated; Figure 38B The inference and / or training logic according to at least one embodiment is shown; and Figure 38C Training and deployment of a neural network according to at least one embodiment are illustrated. Detailed Implementation

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

[0006] Neural networks use numerical matrices to represent information such as network inputs, weights, and network outputs. Matrix multiplication is widely used in neural networks. For example, network outputs are computed by performing matrix multiplication on the matrices representing the network inputs and weights. As another example, neural network training involves using matrix multiplication to compute updated weights from previous layers. Each matrix contains a set of numbers, called matrix elements, arranged in rows and columns. Each matrix element can be a floating-point value represented using a specific number of bits. Each matrix element can be represented using 32 bits for high precision. Smaller matrix elements can be used to improve performance by reducing the amount of computation and memory required for each matrix element. For example, 8-bit or 4-bit reduced-size formats can be used for each matrix element to significantly reduce the amount of memory and computational resources used by the neural network. However, the numerical accuracy of this reduced-size format is lower than that of the larger format. Therefore, using a reduced-size format significantly reduces the precision of the matrix multiplication results, and consequently, the accuracy of the neural network output.

[0007] The systems, methods, and apparatuses described herein execute matrix multiplication-accumulation (MMA) processor instructions to multiply the elements of two input matrices by a scaling factor, and then multiply the resulting scaled input matrices using matrix multiplication to produce an output matrix. The output matrix from the matrix multiplication operation can be scaled up to a larger input size using a scaling factor. Matrix multiplication where the multiplied matrices have already been multiplied by a scaling factor is called “block scaling,” which reduces the size of the input matrix elements, allowing matrix multiplication to be performed using smaller matrix elements. The input matrix can be a fragment of a large matrix, so different fragments can be multiplied by different scaling factors. For large matrices (e.g., weight matrices), using the same scaling factor may not necessarily result in sufficiently accurate outputs from matrix multiplication, because different parts of the matrix may have element values ​​within different ranges that cannot be represented by a single scaling factor. Using different scaling factors for different parts of the input matrix (referred to herein as “blocks”) can improve accuracy for element sizes of 8 bits or smaller. The scaling factor can be determined using appropriate techniques, for example, based on the expected range of matrix element values ​​determined in previous inference and / or training operations of the neural network. As an example, block scaling reduces the size of input matrix elements from a larger format (e.g., 32-bit floating-point values) to a reduced-size format (e.g., 8-bit, 6-bit, or 4-bit floating-point formats) so that they can be provided as input for matrix multiplication operations. The specific reduced-size format used for matrix elements is specified by the operand of the instruction. The input matrix can reside in memory (e.g., tensor memory or shared memory) or registers.

[0008] The scaling factor to be used for different parts or under different conditions of the input matrix can be selected by the block scaling MMA instruction. The block scaling MMA instruction has input operands that specify scaling factor selector values ​​that select a specific region of the scaling factor data (“scaling data”), from which the scaling factor is retrieved for each call to one of these instructions. More specifically, the block scaling MMA instruction has a scaling vector size qualifier that specifies that each input matrix will be partitioned into a specified number of regions (referred to herein as “chunks”), and the elements of each chunk will be multiplied by the scaling factor from the corresponding scaling vector. Thus, different chunks of the input matrix can be multiplied by the elements of different vectors that represent different regions of the scaling data. The scaling vector can be a row vector or a column vector, depending on whether the scaling vector is to be multiplied by the rows or columns of the input matrix. For example, one scaling vector can be multiplied by each column of the first half of the input matrix A, and another scaling vector can be multiplied by each column of the second half of the input matrix A. In at least one embodiment, the multiplication of the vector of the input matrix with the scaling vector is performed by multiplying the individual elements to produce the elements of the resulting scaled vector. Two or more scaling vectors may be grouped together to form a scaling matrix. In at least one embodiment, the scaling vector size qualifier specifies the number of different scaling vectors to be multiplied with the input matrix, rather than the size of an individual scaling vector. In at least one embodiment, the scaling vector size specifies the number of scaling vectors in the scaling matrix. For example, if there are two scaling vectors, the scaling vector size qualifier is 2X. In at least one embodiment, the scaling vector size qualifier specifies the size of the scaling vector, such as the number of elements in the scaling vector. The scaling vector is selected from predetermined scaling data. A scaling factor selector value provided as an operand of the instruction can specify which portion of the scaling data to select to form the scaling vector. For example, the scaling data may be an 8-bit value in tensor memory or a register. In at least one embodiment, the block scaling MMA instruction also has an input operand for specifying a thread selector value for selecting which threads (e.g., a thread bundle) of threads performing the MMA operation will be used to retrieve the scaling factor from tensor memory or a register.

[0009] Figure 1The illustration shows a processor 102 according to at least one embodiment for executing a block-scaled matrix multiplication-accumulation (MMA) instruction 112, the inputs and outputs of which are stored in memory. The processor 102 can execute the MMA instruction 112 to multiply the elements of two input matrices by a scaling factor, and then multiply the resulting scaled input matrices using matrix multiplication to produce an output matrix. The inputs and / or outputs of the MMA instruction 112 are stored in tensor memory (TMEM) 130 and shared memory (SMEM) 150. The processor 102 may be a graphics processing unit (GPU), a central processing unit (CPU), or other device suitable for executing the MMA instruction 112. The processor 102 includes one or more streaming multiprocessors 104 (“SM”). Each SM 104 includes a TMEM 130 and an SMEM 150 for storing data, and one or more tensor cores 180 for performing computational operations such as MMA operations. Each streaming multiprocessor 104 also includes registers (not shown) for storing data.

[0010] In at least one embodiment, tensor core 180 includes scaling factor selection logic 170A, which selects a scaling factor from scaling data 134 identified by scaling data operand 120 based on scaling data selector operand 122. Alternatively or additionally, scaling factor selection logic 170B may also be included in SM 104 outside of tensor core 180. In at least one embodiment, tensor core 180 includes block scaling MMA computation logic 182, which performs computation on input matrix A 132 (in TMEM) or input matrix A 152A (in SMEM), input matrix B 152B (in SMEM), and D 138 (in TMEM). The input matrices are identified by matrix A operand 114A or matrix A operand 116 (SMEM), matrix B operand 114B (SMEM), and matrix D operand 118. In at least one embodiment, block-scaled MMA computation logic 182 multiplies input matrices 132 or 152A and 152B by corresponding selected scaling factors from scaling data 134A, 134B to obtain scaled input matrices, and performs MMA operations on the scaled input matrices to compute an MMA result, which is stored in an input / output (e.g., accumulator) matrix D 138 identified by matrix D operand 118. Thus, matrix D provides the input to the MMA operation and receives its output.

[0011] In at least one embodiment, processor 102 includes a cooperative thread array (CTA) 110, which may be a group of GPU threads that work together to execute blocks of code and share resources such as memory and synchronization primitives. In at least one embodiment, the CTA 110 executed by processor 102 initiates MMA instructions 112 via instruction scheduler 126 and pipeline manager 128 so that the instructions can be executed on or otherwise implemented on the tensor core 180 of streaming multiprocessor 104. A group of threads (e.g., 32 threads) is referred to as a warp, which may also be called a subgroup or wavefront, wherein each thread in the warp may belong to a single thread block and be configured to process different datasets based on a single set of instructions. CTA 110 may include multiple warps.

[0012] When processor 102 executes MMA instruction 112, the MMA instruction 112 causes processor 102 to perform an MMA operation that calculates the product of input matrix A 132 (in TMEM) or A 152A (in SMEM) and input matrix B 152B (in SMEM). The MMA operation can be a "block-scaled" MMA operation, which scales one or more elements of at least one input matrix by multiplying each element of that matrix by a scaling factor, thereby generating one or more scaled input values. These scaled input values ​​are used to calculate the product of matrices A and B, instead of the corresponding unscaled elements of matrices A and B being received as input. For example, the MMA operation can calculate a scaled matrix A by multiplying the elements of matrix A 132 by the corresponding scaling factor. The MMA operation can also calculate a scaled matrix B by multiplying the elements of matrix B 152B (in SMEM) by the corresponding scaling factor. The MMA operation can further calculate the product of scaled matrices A and B.

[0013] MMA instruction 112 can be configured via input parameters to multiply specific matrix elements by a selected scaling factor. For example, there can be more matrix elements than the scaling factor, in which case the same “block” scaling factor (e.g., a set of scaling factors arranged by rows and columns) will be multiplied by two or more corresponding blocks of the input matrix (e.g., fragments of the input matrix). Multiplying by the scaling factor can be performed by multiplying each scaling factor (e.g., a value) in the block by each matrix element (e.g., a value) in the block. The result of multiplying the scaling factor by the matrix elements can be used as input for matrix multiplication, or it can replace matrix elements in the matrix, depending on the implementation details. Specific implementations can leverage optimizations, such as by merging the multiplication of the scaling factor and matrix elements with one or more subsequent operations used in matrix multiplication, in which case the one or more values ​​described herein are not necessarily stored in memory or registers. For example, for illustrative purposes, the product of the scaling factor and the input matrix elements is described as a value in the examples herein, but implementations do not necessarily store this product as a value in memory or registers.

[0014] In at least one embodiment, the MMA operation is a mathematical operation that, if executed by the processor 102, calculates a value. D : D = AB + D in A , B and D It is a matrix, submatrix, submatrix fragment or tile, or set of values. A and B This is the input matrix, where D on the right is an optional input matrix (if specified, it is the accumulator matrix), and on the left... D This is the output matrix, and the result of the MMA operation is stored in this output matrix. AB It is a matrix A With matrix B The product of. A The size is M × K ( M Multiply K List), B The size is K × N ( K Multiply N (list), and D The size is M × N ( M Multiply N(Column). The shape of an MMA operation with these dimensions is... M × N × K In this article, it is also referred to as shape. mMnNkK ,in M , N and K These are integer values ​​that indicate the dimensions of the input matrices A and B, and the dimension of the result matrix D, respectively.

[0015] The block scaling MMA operation can be a mathematical operation that, if executed by processor 102, calculates a value. D : in and It is a scaling matrix, * indicates matrix A One or more elements of another matrix (e.g.) The element-wise multiplication of one or more elements of ), and It is by scaling the matrix A (calculated as) ) and the scaled matrix B (which is calculated as The matrix is ​​obtained by multiplication. In element-wise multiplication, the matrix... A Each element is related to The corresponding elements of the matrices are multiplied to produce a scaled element. This scaled element can be used in subsequent matrix multiplication operations, replacing elements of matrix A. The scaled element corresponds to an element of matrix A. In at least one embodiment, the scaled element replaces elements of matrix A in memory and is used as input to a matrix multiplication operation, replacing elements of matrix A. In at least one embodiment, the scaled element does not replace elements of matrix A in memory, but is used as input to a matrix multiplication operation, replacing elements of matrix A.

[0016] The MMA instruction 112 has a set of operands and a set of qualifiers. The operands provide parameters that can be determined at runtime for the MMA instruction 112, and the qualifiers provide configuration parameters for the MMA instruction 112. The operands of the MMA instruction 112 include matrix A operand 114A, alternative matrix A operand 116, matrix B operand 114B, and matrix D operand 118. The operands also include scaling data A operand 120A, scaling data B operand 120B, scaling data A selector operand 122A, and scaling data B selector operand 122B.

[0017] about A Which elements are related to The determination of which elements are multiplied depends on the configuration input provided to the MMA instruction 112, such as the scaling vector size qualifier 142 as described above. In at least one embodiment, the scaling vector size qualifier 142 at least partially determines... The size of the matrix and / or The size of the matrix, as discussed in this article... Figure 2 Further description. For example, The number of rows ( M )and A The number of rows is the same, but The number of columns can be less than A In other words, The number of columns is relatively small, such as 1, 2, or 4. For the purposes of explanation, this article will... The columns are called vectors, where Each vector has M The column vector of the row. For example, if the scaling vector size qualifier 142 is specified as 1X for MMA instruction 112, then Having one column, in this case, The elements of the column and A Multiply the corresponding elements of each column. As another example, if the scaling vector size qualifier 142 is specified as 2X, then It has two columns. In this case, The elements of the first column and A The first half (e.g., the first half) K Multiply the corresponding elements of each column in ( / 2 columns), and The elements in the second column and A The latter half (e.g., the latter part) K The corresponding elements of each column in the matrix (column 2) are multiplied together. For illustrative purposes, these matrix columns will also be referred to as column vectors in this paper. Furthermore, this paper will... A The two halves are called A "Combine".

[0018] One or more operands in MMA operands 114-122 identify the data to be operated on by MMA instruction 112. Each of operands 114-122 can be a shared memory descriptor for identifying an address in shared memory, a tensor memory address including rows and columns for identifying an address in tensor memory, a register (specified as one or more register names, such as {%Ra0}), and / or an intermediate value.

[0019] In at least one embodiment, the operand identifying the data includes a reference to the data, for example, in the form of a memory address. This memory address could be a tensor memory address 130 or a shared memory address 150. In at least one embodiment, the operand identifying data in the tensor memory is a tensor memory address that specifies the channel index and column index in the tensor memory (e.g., see [reference to...]). Figure 4B (Example tensor memory layout described). In at least one embodiment, the operand identifying matrix data in shared memory is a matrix descriptor that describes the properties of the matrix in shared memory, including the matrix's location in shared memory and the current CTA 110. The shared memory descriptor is a 64-bit value stored in a register that includes the matrix's starting address and other information related to the matrix's representation (e.g., base offset, leading dimension stride mode, and swapping mode related to the rearrangement of the matrix elements in shared memory).

[0020] Instruction qualifiers can be specified for MMA instruction 112 to modify its behavior. For example, instruction qualifiers can define the shape of the input matrix and other parameters for block-scaled MMA computation. Qualifiers for MMA instruction 112 include MMA type qualifier 124, which, for example, specifies the number of vectors of scaling factors in the scaling factor matrix. In at least one embodiment, the qualifiers for MMA instruction 112 (e.g., MMA type qualifier 124) are the input to the instruction.

[0021] As an example, instruction scheduler 126 extracts shared memory descriptors from MMA instruction 112, stores these shared memory descriptors in SMEM 150, and then loads these shared memory descriptors from SMEM 150 into registers. As another example, before MMA instruction 112 is executed on or otherwise implemented on Tensor Core 180, processor 102 instructs instruction scheduler 126 to load operands from SMEM 150 or TMEM 130 (depending on whether the operands are located in SMEM 150 or TMEM 130) into registers. Registers are fast storage locations that can temporarily hold operands during the execution of MMA instruction 112. After the operands are loaded into registers, instruction scheduler 126 schedules MMA instruction 112 for execution on Tensor Core 180 of streaming multiprocessor 104 based on the operands and qualifiers of MMA instruction 112.

[0022] Matrix A operand 114A specifies the location of matrix A in TMEM 130. Alternatively, matrix A operand 116 specifies the location of matrix A in SMEM 150. Matrix B operand 114B specifies the location of matrix B in shared memory 150. Matrix D operand 118 specifies the location of matrix D in tensor memory 130.

[0023] Scaling data A operand 120A and scaling data B operand 120B specify the positions of scaling data A and scaling data B in TMEM 130, respectively. Scaling data A contains one or more scaling matrices A (e.g., ...). The value of ), and the scaled data B contains one or more scaling matrices B (e.g. The scaling data A selector operand 122A indicates which scaling vectors to select from scaling data A, for example, to form scaling matrix A, and the scaling data B selector operand 122B indicates which scaling vectors to select from scaling data B, for example, to form scaling matrix B. Selector operand 122 also indicates which threads will contribute (e.g., selected from the scaling data) the scaling vectors, as discussed in this article. Figure 6 and Figure 7 The number of scaling vectors selected from the scaling data is specified by the scaling vector size qualifier 142, for example, it can be 1X, 2X, or 4X, as described above. Figure 2 As stated above.

[0024] The matrix element data types of the elements of matrices A and B, as well as the scaling factor data types (also known as scaling factors) of the elements of matrices A and B, are specified by the MMA type qualifier 124 of the MMA instruction 112. The MMA type qualifier 124 can be mxf4, mxf8f6f4, or mf4nvf4. Table 1 lists the matrix element data types and scaling factor data types available for each valid combination of the MMA type qualifier 124 and the scaling vector size qualifier 142. The type `mxf8f6f4` indicates 8-bit, 6-bit, or 4-bit matrix element data types, as well as a scaling factor data type (`ue8m0`) with an 8-bit exponent and no mantissa. Specific matrix element data types are defined by [the relevant specification]. Figure 4AThe description describes the operand selection for the element types of matrix A and matrix B. An 8-bit type can be e4m3 (4-bit exponent, 3-bit mantissa, and 1-bit sign) or e5m2 (5-bit exponent, 2-bit mantissa, and 1-bit sign). A 6-bit type can be e3m2 or e2m3. A 4-bit type is e2m1 (2-bit exponent, 1-bit mantissa, and 1-bit sign). In at least one embodiment, for mxf8f6f4, the 8-bit, 6-bit, and 4-bit elements of matrices A and B are stored in an 8-bit container, which uses 8 bits of storage for each matrix element, even if the data of the matrix element is less than 8 bits. The type mxf8f6f4 can be used with a scaling vector size of 1X. For example, the container is a memory region with an associated size that can be larger than the number of bits used to store the values ​​stored in the container. In at least one embodiment, the container is the smallest unit of storage that can be used to represent the elements of the input matrix for MMA operations.

[0025] The type `mxf4` indicates a 4-bit matrix element data type (e2m1) and a scaling factor data type (ue8m0) with an 8-bit exponent and no mantissa. `mxf4` can be used with a scaling vector size of 2X. For `mxf4`, the 4-bit elements of matrices A and B are stored in a 4-bit container, which uses 4 bits of data for each matrix element.

[0026] The type `mxf4nvf4` indicates a 4-bit element type, and an 8-bit exponent scaling factor data type (ue8m0) to be used with a 2X scaling vector size, or a 4-bit exponent, 3-bit mantissa, and 1-bit sign scaling factor data type (ue4m3) to be used with a 4X scaling vector size. For type `mxf4nvf4`, the 4-bit elements of matrices A and B are stored in a 4-bit container, which uses 4 bits of data for each matrix element.

[0027] Matrix A 132, matrix D 138, and scaling data 134 can be stored in the tensor memory 130 of the streaming multiprocessor 104. Scaling data 134 includes scaling data A 134A and scaling data B 134B, where scaling data A 134A contains a scaling factor (“SF_A”) to be multiplied by matrix A, and scaling data B 134B contains a scaling factor (“SF_B”) to be multiplied by matrix B. Matrix A 132 can be identified by matrix A operand 114A, and matrix D 138 can be identified by matrix D operand 118 of MMA instruction 112. Scaling data A 134A and scaling data B 134B can be identified by scaling data A selector operand 122A and scaling data B selector operand 122B, respectively. Matrix A 152A (in SMEM) can be used as a replacement for matrix A 132 and can be stored in the shared memory 150 of the streaming multiprocessor 104, where matrix B 152B can also be stored.

[0028] In at least one embodiment, combined Figure 1 At least one feature described is achieved by utilizing at least one feature of any system and / or process depicted in and / or described in conjunction with other figures. In at least one embodiment, in conjunction with Figure 1 At least one feature described is used to implement at least a portion of any system and / or process depicted in or in conjunction with the other figures.

[0029] Combination Figure 1 At least one feature described is achieved by utilizing at least one feature of any system and / or process depicted in and / or described in conjunction with other figures. Figure 1 At least one feature described is used to implement at least a portion of any system and / or process depicted in and / or described in conjunction with other figures.

[0030] Figure 2 The illustration depicts an interaction in a processor 202 according to at least one embodiment for executing block-scale MMA instructions having block-scale MMA inputs 212 for matrices stored in tensor memory 230 and / or shared memory 250. The processor 202 includes tensor memory (TMEM) 230 and shared memory (SMEM) 250. Furthermore, the processor 202 also includes block-scale MMA inputs 212. Block-scale MMA inputs 212 include operands referencing matrices stored in TMEM 230 and / or SMEM 250. Block-scale MMA inputs 212 also include MMA type qualifiers 224 (e.g., combined with...). Figure 1 The matrix A operand 214, matrix A operand 216A (SMEM), matrix B operand 216B (SMEM), and matrix D operand 218 are described above. Further details regarding... Figure 1 Describes matrix operands. Block-scaled MMA input 2^12 also includes M-dimensional 2^26 operands (e.g., matrix...). A and D The number of rows), N-dimensional 228 operands (e.g., matrix) B and D The processor 102 also includes scaling factor selection logic 270 and a tensor core 280, which are similar to those referenced herein. (The text also mentions column count, scaling vector size 242 qualifier, scaling data operands 220, and scaling factor selector operands 222.) Figure 1 The scaling factor selection logic 170 and tensor kernel 180 are mentioned above.

[0031] In at least one embodiment, the scaling vector size 242 specifies the number of scaling vectors in the scaling matrix 244. Hereinafter, the term "vector" is used to refer to a row or column of a matrix. For example, a scaling vector can be a column vector in scaling matrix A 244A, or a row vector in scaling matrix B 244B. Therefore, the scaling vector size 242 specifies the number of columns in scaling matrix A 244A or the number of rows in scaling matrix B 244B. A scaling vector size 242 of 1X indicates that there is one scaling vector in each scaling matrix. Therefore, if the scaling vector size qualifier 242 is 1X, then scaling matrix A has one column and scaling matrix B has one row. A scaling vector size 242 of 2X indicates that there are two scaling vectors in scaling matrix 244. Therefore, if the scaling vector size qualifier 242 is 2X, then scaling matrix A has two columns and scaling matrix B has two rows. A scaling vector size 242 of 4X indicates that there are four scaling vectors (e.g., scaling matrix A has four columns and scaling matrix B has four rows). Although a specific scaling vector size is described in the examples in this article, other scaling vector sizes, such as 3X or 6X, may be used in other examples.

[0032] The scaling factor to be used in a block scaling operation performed in response to a call to the block scaling MMA instruction 112 is stored in scaling data 234. More specifically, the scaling factor to be multiplied by elements of input matrix A 232 (or input matrix A 252A in TMEM) is stored in scaling data A 234A, and the scaling factor to be multiplied by elements of input matrix B 252B is stored in scaling data B 234B. In at least one embodiment, processor 200 includes scaling factor selection logic 270 located outside tensor core 280, such as Figure 2As shown. Alternatively or additionally, the tensor kernel 280 also includes scaling factor selection logic 270. Scaling factor selection logic 270 retrieves a specific scaling factor to be used in the block scaling operation based on the block scaling MMA input 212. Different scaling factors can be used for different rows and / or columns of the input matrices A and B by storing different scaling factors in different rows of scaling data A 234A and / or scaling data B 234B. Furthermore, different portions of scaling data 234 can be used for different parts of the input matrix by providing different values ​​for scaling data operands 220 pointing to different locations in TMEM 230. Using different scaling factors for different portions of the input matrix (referred to herein as blocks) can improve accuracy for input matrix element sizes of 8 bits or less. The scaling factor is determined using appropriate techniques, such as from the expected range of matrix element values ​​determined from previous inference and / or training operations of the neural network. For example, the expected range of values ​​for the input matrix elements (e.g., magnitude range) can be determined, and a scaling factor can be calculated to convert the input matrix elements from a larger element size (e.g., 32 bits) capable of representing values ​​within a larger range to a smaller element size (e.g., 4 bits, 6 bits, or 8 bits) capable of representing elements within the expected range. The scaling factor to be used in the block-scale MMA operation caused by executing MMA instruction 112 is stored in scaling data 234, the storage location of which is specified by scaling data operand 220. Specific portions of scaling data 234 to be used in the block-scale MMA operation are specified by operands of MMA instruction 112, such as scaling factor value selector operand 222.

[0033] In response to the execution of MMA instruction 112, scaling factor selection logic 270 selects scaling matrix A 244A from scaling data A 234A according to the scaling vector size qualifier 242 and scaling factor value selector operand 222A. Alternatively or additionally, scaling factor selection logic 270 selects scaling matrix B 244B from scaling data B 234B according to the scaling vector size qualifier 242 and scaling factor value selector operand 222B.

[0034] Scaling factor selection logic 270 provides the selected scaling factor as input to block scaling MMA logic 282. Block scaling MMA logic 282 multiplies the scaling factor by elements of the input matrix A to generate scaled elements. Block scaling MMA logic 282 provides the scaled elements (e.g., scaled elements of input matrices A and B) to the MMA operation. Block scaling MMA logic 282 performs the MMA operation to generate an accumulator matrix D 238. To perform block scaling, elements of the scaling vector are multiplied by elements of the vector of the input matrix to produce corresponding result elements. The result elements correspond to elements of the input matrix and can be used as input for matrix multiplication in the MMA computation. In at least one embodiment, in block scaling, vector multiplication is element-wise multiplication, which produces a vector with the same number of elements, where each element is the product of two corresponding elements of the input vector. In at least one embodiment, in a particular embodiment, this element-wise multiplication does not involve calculating the sum of element-wise products.

[0035] In shape M × N × K In block scaling MMA operations, the input matrix A The size is M × K Input matrix B The size is K × N Furthermore, the scaling matrix A 244A includes 1, 2, 4, or other quantities. X The scaling vectors, and each scaling vector contains M Each element of the scaling vector in scaling matrix A244A is multiplied by the element in the corresponding row and column of the input matrix A244A. Similarly, scaling matrix B244B includes 1, 2, 4, or other numbers. X The scaling vectors, and each scaling vector contains N Each element of the scaling vector in the scaling matrix B244B will be multiplied by the element in the corresponding row and column of the input matrix B244B. As an example, for a matrix named... Scaling vector, input column vector A and named The resulting scaled vector, the scaled input vector Each element [i] is calculated as: [i] = ,in i The value range is 0 to M Scaled vector This will be used as input for matrix multiplication operations. Similarly, as an example, for a matrix named... Scaling vector, input column vector B and named The resulting scaled vector, the scaled input vector Each element [i] is calculated as: [i] = ,in i The value range is 0 to N Scaled vector It will be used as input for matrix multiplication operations. Figure 9 The illustration shows an example of block scaling MMA, where the scaling vector is multiplied by a vector of the input matrix.

[0036] In at least one embodiment, scaling factors stored in scaling data 234 are arranged by rows and columns, where each row contains one or more scaling factors that will be multiplied by a thread of processor 202 (executed by tensor core 280) with corresponding rows of input matrices A and / or B. In at least one embodiment, different columns of scaling data 234 may contain different scaling factors that will be multiplied by different columns of input matrix A and / or different rows of input matrix B in response to MMA operations invoked by MMA instruction 112. In at least one embodiment, different columns of scaling data 234 may contain scaling factors that will be multiplied by one or more columns of input matrix A and / or one or more rows of input matrix B in response to different invocations of MMA operations (e.g., by different calls to MMA instruction 112).

[0037] In at least one embodiment, the scaling data B 234B is arranged with the same row and column orientation as the scaling data B 234B. Due to the orientation of matrices A and B used for matrix multiplication, the scaling factors stored in the rows of the scaling data B 234B are transformed into column format so as to be included in the scaling matrix B 244B. Each column of the scaling data B may be transformed (e.g., transposed from column to row) and included in or otherwise used to generate the scaling matrix B 244B. Figure 9 The example shown illustrates block scaling, which demonstrates the scaling data and scaling matrix selected from the scaling data.

[0038] If the scaling vector size 242 is 1X, then scaling factor selection logic 270 selects one column of scaling factors from scaling data A 234A and one column of scaling factors from scaling data B 234B. The column to be selected from scaling data A 234A is specified by selector operand 222A, and the column to be selected from scaling data B 234B is specified by selector operand 222B. The column selected from scaling data A 234A forms one column of scaling matrix A 244A (in the 1X case, scaling matrix A 244A has only one column). Furthermore, the column selected from scaling data B 234B forms one row of matrix B 244B (in the 1X case, matrix B 244B has only one row). For example, scaling factor selection logic 270 transposes the column selected from scaling data B 234B to form the corresponding row of scaling matrix B 244B. Scaling factor selection logic 270 transposes the columns, for example, by converting the columns to rows, which may involve swapping the index of each element in the column. Scaling factor selection logic 270 includes the resulting rows in the scaling matrix B244B. (Regarding...) Figure 4B and Figure 5 The use of selector operand 222 to select the vector of scaled data 234 is further described. Figure 10 An example of block scaling MMA is shown, illustrating the multiplication of a scaling vector with an input matrix of size 1X.

[0039] If the scaling vector size 242 is 2X, then scaling factor selection logic 270 selects two columns of scaling factors from scaling data A 234A and two columns of scaling factors from scaling data B 234B. The two columns to be selected from scaling data A 234A are specified by selector operand 222A, and the two columns to be selected from scaling data B 234B are specified by selector operand 222B. The two columns selected from scaling data A 234A form two columns of scaling matrix A 244A (in the case of 2X, scaling matrix A 244A has two columns). Furthermore, the two columns selected from scaling data B 234B form rows of scaling matrix B 244B (in the case of 2X, scaling matrix B 244B has two rows). For example, scaling factor selection logic 270 transposes each of the two selected columns in scaling data B 234B to form the corresponding row. Scaling factor selection logic 270 includes the resulting two rows in scaling matrix B 244B. Figure 11 An example of block scaling MMA is shown, illustrating the multiplication of a scaling vector with an input matrix of size 2X.

[0040] If the scaling vector size 242 is 4X, then scaling factor selection logic 270 selects four columns of scaling factors from scaling data A 234A and four columns of scaling factors from scaling data B 234B. The four columns selected from scaling data A 234A are specified by selector operand 222A, and the four columns selected from scaling data B 234B are specified by selector operand 222B. The four columns selected from scaling data A 234A form the four columns of scaling matrix A 244A. Furthermore, the four columns selected from scaling data B 234B form the four rows of scaling matrix B 244B. For example, scaling factor selection logic 270 transposes each of the four selected columns in scaling data B 234B to form the corresponding row. Scaling factor selection logic 270 includes the resulting four rows in scaling matrix B 244B. Figure 12 An example of block scaling MMA is shown, illustrating the multiplication of a scaling vector with an input matrix of size 4X.

[0041] In at least one embodiment, if the scaling vector size 242 is a specified number X, then scaling factor selection logic 270 selects X columns of scaling factors from scaling data A 234A and X columns of scaling factors from scaling data B 234B. The X columns selected from scaling data A 234A are specified by selector operand 222A, and the X columns selected from scaling data B 234B are specified by selector operand 222B. The X columns selected from scaling data A 234A constitute X columns of scaling matrix A 244A. Furthermore, the X columns selected from scaling data B 234B constitute X rows of scaling matrix B 244B. For example, scaling factor selection logic 270 transposes each column of the selected X columns in scaling data B 234B to form the corresponding row of scaling matrix B 244B. Scaling factor selection logic 270 includes the resulting X rows in scaling matrix B 244B.

[0042] Block scaling MMA logic 282 performs scaling multiplication on input matrix A and / or input matrix B 252B to generate scaled input matrix A and / or scaled input matrix B. To perform scaling multiplication on input matrix A, block scaling MMA logic 282 multiplies the elements of scaled matrix A 244A with the corresponding elements of input matrix A. Input matrix A can be input matrix 232 or input matrix A 252A, depending on whether matrix A operand 214 (in TMEM) or matrix A operand 216A (in SMEM) is specified in the call to MMA instruction 112.

[0043] With a scaling vector size of 1X, the block scaling MMA logic 282 performs scaling multiplication on the input matrix A by multiplying the elements of a single column of the scaling matrix A 244A with the corresponding elements of each column of the input matrix A. The single column of the scaling matrix A 244A is selected by the scaling factor selection logic 270, as described herein. Furthermore, with a scaling vector size of 1X, the block scaling MMA logic 282 performs scaling multiplication on the input matrix B 252B by multiplying the elements of a single row of the scaling matrix B 244B with the corresponding elements of each row of the input matrix B. This single row is selected by the scaling factor selection logic 270, as described herein. In at least one embodiment, the result of multiplying the elements of the input matrices A and / or B with the elements of the scaling matrices A and / or B (rather than the corresponding elements of the matrices) is used as the input to the MMA operation.

[0044] With a scaling vector size of 2X, the block scaling MMA logic 282 performs scaling multiplication on the input matrix A 232 (or the input matrix A 252A in SMEM) by partitioning the input matrix A into a first block and a second block (each block containing K / 2 columns). Furthermore, the scaling matrix A 244A comprises 2 columns. The block scaling MMA logic 282 multiplies the elements of the first column of the scaling matrix A 244A with the corresponding elements of each column in the first block of the input matrix A 232 (or the input matrix A 252A in SMEM). Additionally, the block scaling MMA logic 282 multiplies the elements of the second column of the scaling matrix A 244A with the corresponding elements of each column in the second block of the input matrix A 232 (or the input matrix A 252A in SMEM).

[0045] Furthermore, in the 2X case, the block scaling MMA logic 282 performs scaling multiplication on the input matrix B 252B by partitioning the input matrix B into a first block and a second block (each block comprising K / 2 rows). Additionally, the scaling matrix B 244B comprises 2 rows. The block scaling MMA logic 282 multiplies the elements of the first row of the scaling matrix B 244B with the corresponding elements of each row in the first block of the input matrix B 252B. Furthermore, the block scaling MMA logic 282 multiplies the elements of the second row of the scaling matrix B 244B with the corresponding elements of each row in the second block of the input matrix B 252B.

[0046] With a scaling vector size of 4X, the block scaling MMA logic 282 performs scaling multiplication on the input matrix A 232 (or input matrix A 252A in SMEM) by partitioning each input matrix A 232 (or input matrix A 252A in SMEM) and / or input matrix B into four blocks, each block comprising K / 4 input vectors. These input vectors are columns in input matrix A or rows in input matrix B. Furthermore, each scaling matrix 244 comprises four scaling vectors. These scaling vectors are columns in scaling matrix A 244A or rows in scaling matrix B 244B. Therefore, four vectors with scaling factors are selected from each scaling matrix 244 (instead of the two vectors in the 2X case). For scaling matrix A 244A, these vectors are columns of scaling matrix A 244A. For scaling matrix B 244B, the vectors are rows of scaling matrix B 244B. Each input matrix A or B is divided into 4 blocks, each block having K / 4 vectors (instead of the two blocks in the 2X case). Each block corresponds to one of the four vectors of scaling matrix 244. For each of the four vectors of scaling matrix 244, the block scaling MMA logic 282 multiplies the element of that vector with each of the K / 4 vectors in the block of input matrix A or B (instead of multiplying the element of each of the two vectors of scaling matrix 244 with each of the K / 2 vectors in the block, as in the 2X case).

[0047] In at least one embodiment, if the scaling vector size 242 is a specified number X, then X vectors (columns or rows) with scaling factors are selected from each scaling matrix 244. Each input matrix A or B is partitioned into X blocks, each block having K / N vectors (columns of input matrix A or rows of input matrix B). Each block corresponds to one vector among the X vectors of the scaling matrix 244. For each vector among the X vectors of the scaling matrix 244, the block scaling MMA logic 282 multiplies the element of that vector with the element of each of the K / N vectors in the block of input matrix A or B. The result of the multiplication is a scaled A value 236 (or a scaled A value 256A in SMEM) and a scaled B value 256B for input matrix A. The scaled A values ​​236 (or 256A) and 256B are used as inputs to the MMA operation performed by the block scaling MMA logic 282. Block scaling MMA logic 282 generates an accumulator matrix D 238 due to the block scaling MMA operation performed on the scaled value.

[0048] In at least one embodiment, if the scaling vector size 242 is a specified value X (e.g., scaling vector size = XX), then X vectors (columns or rows) with scaling factors are selected from each scaling matrix A 244A or B 244B. Each input matrix A or B is partitioned into X blocks, each block comprising 1 / X vectors of input matrix A or B. These vectors are columns of input matrix A or rows of input matrix B. Each block of input matrix A comprises 1 / X columns of input matrix A. More specifically, since input matrix A has M rows and K columns, each block of input matrix A has M rows and K / X columns. Each block of input matrix B comprises 1 / X rows of input matrix B. More specifically, since input matrix B has K rows and N columns, each block of input matrix B has K / X rows and N columns.

[0049] When the scaling vector size is XX, for the input matrix A 252A, the block scaling MMA logic 282 performs scaling multiplication using the scaling matrix A 244A by multiplying the elements of each column of the scaling matrix A 244A with the corresponding elements of each column in the K / X columns of the block of the input matrix A 252A. As an example, the block of the input matrix A comprises the K / X columns of the input matrix A. There are X blocks, and a specific block of the input matrix A corresponds to a specific column of the scaling matrix A 244A. For each column in a specific block of the input matrix A, the block scaling MMA logic 282 multiplies each element in that column of the input matrix A with each corresponding element in the specific column of the scaling matrix A 244A. The result of the multiplication is used as the input to the MMA operation.

[0050] When the scaling vector size = XX, for the input matrix B 252B, the block scaling MMA logic 282 performs scaling multiplication using the scaling matrix B 244B by multiplying the elements of each row of the scaling matrix B 244B with the corresponding elements of each row in the K / X rows of the block of the input matrix B. As an example, the block of the input matrix B comprises K / X rows of the input matrix B. There are X blocks, and a specific block of the input matrix B corresponds to a specific row of the scaling matrix B 244B. For each row in a specific block of the input matrix A, the block scaling MMA logic 282 multiplies each element of that row of the input matrix A with each corresponding element of the specific row of the scaling matrix A 244A. The result of the multiplication is used as the input to the MMA operation. Although specific numbers of rows and columns, specific scaling factors (e.g., 1, 2, and 4), and a specific number of scaling factors are described in the examples here, other numbers of rows and columns, as well as other scaling factors, can be used in other examples; simply change these numbers and scaling factors according to the description herein.

[0051] As described above, the block scaling MMA logic 282 generates the scaled A value 236 (or the scaled A value 256A in SMEM 250) in TMEM 230 by multiplying the elements of the scaling matrix A 244A by the elements of the input matrix A 232 (or input matrix A 252A). Similarly, the block scaling MMA logic 282 generates the scaled B value 256B in SMEM by multiplying the elements of the scaling matrix B 244B by the elements of the input matrix B 252B. The block scaling MMA logic 282 uses a tensor kernel 280 to generate the resulting accumulator matrix D 238 of dimension M×N by performing the following operation: performing the operation of the tensor kernel 280 of dimension M×N. M × K The matrix multiplication operation, which multiplies a scaled A-value 236 matrix (or a scaled A-value 256A matrix) by a scaled B-value 256B matrix, adds the element values ​​of the accumulator matrix D 238, which is received as input, to the product of A and B, and stores the result in the accumulator matrix D 238.

[0052] Combination Figure 2 At least one feature described is implemented using at least one feature of any system and / or process depicted in or in conjunction with other figures. Figure 2 At least one feature described is used to implement at least a portion of any system and / or process depicted in or in conjunction with the other figures.

[0053] Figure 3 The illustration shows a block-scaled MMA operation 300 in which input and output matrices, according to at least one embodiment, are stored in tensor memory or shared memory. Operation 300 represents one or more computational operations that enable a processor (e.g., a GPU) to perform or otherwise implement MMA (matrix multiplication-accumulation) instructions by adjusting the scaling of the elements of the input matrix using block scaling.

[0054] Operation 300 includes block scaling MMA instruction fetch 302 and MMA instruction response 378. In at least one embodiment, fetch 302 is a fetch of instructions to cause one or more processors to perform one or more block scaling MMA computation operations (“MMA operations”) and generate response 378. In at least one embodiment, fetch 302 is a call to instructions (“instruction call”) to cause one or more processors 202 to perform one or more block scaling MMA operations and generate response 378.

[0055] In at least one embodiment, call 302 is an API call to cause one or more processors 802 to perform one or more computational operations and generate a response 378. In at least one embodiment, call 302 is a function call to be executed by one or more software programs (e.g., CUDA kernels) on a parallel processing unit (e.g., a GPU). Call 302 specifies parameters and configurations for performing block-scaled MMA operations on MMA operands 310, thereby achieving performance optimization through block scaling techniques (e.g., multiplying MMA input matrix elements by scaling factors), which may differ for different rows, columns, and / or segments of the input matrix.

[0056] Call 302 includes MMA qualifiers 304, which specify parameters used to configure certain aspects of the MMA operation, which typically remain unchanged during program execution. MMA qualifiers 304 include a CTA group qualifier 306, an MMA type qualifier 324, a block scaling qualifier 308, and a scaling vector size qualifier 342. The CTA group qualifier 306 indicates which CTA group will perform the MMA operation. CTA group 306 can be cta_group::1, indicating that CTA group 1 will perform the MMA operation, or it can be cta_group::2, indicating that CTA group 2 will perform the MMA operation.

[0057] The MMA type qualifier 324 specifies the size of the matrix elements in input matrices A and B. MMA type qualifier 324 can be, for example, MXF4, MXF8F6F4, or MF4NVF4, etc., as will be referred to below. Figure 1 Further description is provided in Table 1. The block scaling qualifier 308 (if present) indicates that the MMA operation is used to perform block scaling on the input matrix. If the block scaling qualifier 308 is not present, block scaling is invoked for the block scaling MMA 302, and no block scaling is performed. The scaling vector size qualifier 342 specifies the scaling vector size, which indicates the number of scaling vectors (e.g., 1, 2, 4, or other size X) of the scaling matrix 244 to be used for block scaling. For example, this document will discuss... Figure 2 Used to describe the scaling vector size qualifier 342.

[0058] Recall 302 receives MMA operand 310 as input. MMA operand 310 may include input matrix A operand 314, which specifies input matrix A 132 in tensor memory (TMEM) 130. MMA operand 310 may also include input matrix A operand 316A, which specifies matrix A152A in shared memory 150. Therefore, matrix A can reside in either tensor memory 130 or shared memory 150, so either input matrix A operand 314 or input matrix A operand 316A will be specified in call 302. The term "matrix A operand" as used herein may refer to either input matrix A operand 314 or input matrix A operand 316A, whichever is specified in the specific block-scale MMA call 302.

[0059] MMA operand 310 also includes input matrix B operand 316B, which specifies an auxiliary input matrix B in shared memory 150 that will be used in the MMA operation along with matrix A. Accumulator matrix D operand 318 specifies accumulator matrix D 138 in TMEM 130. Accumulator matrix D 138 accumulates the results of matrix multiplication operations across multiple computation cycles (e.g., multiple block-scale MMA calls 302). Accumulator matrix D 138 improves performance by allowing intermediate results to be cached and reused across multiple computation cycles, thereby minimizing the need to fetch intermediate results from main memory or write them back to main memory.

[0060] Scaling data A operand 320A refers to scaling matrix A 244A in TMEM 230, and scaling data B operand 320B refers to scaling matrix B 244B, also in TMEM 230. Instruction descriptor 376 refers to a set of additional input operands for block scaling MMA call 302, such as input matrix element type, scaling matrix element type, input matrix dimension, and scaling factor selector. Examples of operands specified in instruction descriptor 376 are as follows... Figure 4A As shown.

[0061] Input matrix A operand 314, accumulator matrix D operand 318, scaled data A operand 320A, and scaled data B operand 320B can be tensor memory addresses, which can be used to specify operand data in TMEM 230. Input matrix operands 316A and 316B can be shared memory descriptors, which can be used to specify operand data in SMEM 250. The enable input of matrix D operand 374 refers to a configuration parameter indicating whether the result of the MMA operation is added to the accumulator matrix D 238, which is received as an input to the MMA operation.

[0062] The block scaling MMA response 378 indicates the output of the MMA instruction execution, which includes the accumulator matrix D result 338. If the enable input of the D matrix operand 374 is true (or other specified value), the MMA operation adds the matrix result calculated by the MMA operation to the accumulator matrix D 238. If the enable input of the D matrix operand 374 is false (or other specified value), the MMA operation overwrites the accumulator matrix D 238 with the matrix result calculated by the MMA operation.

[0063] In at least one embodiment, if call 302 is used to call an instruction, the output is the data that the instruction will output. In at least one embodiment, if call 302 is used to call an API, the output is the data that the API will output. In at least one embodiment, the block scaling MMA response 378 includes a status indicator (not shown) indicating the success or failure of the MMA operation, such as whether the MMA operation completed successfully or encountered an error. The accumulator matrix D result 338 represents the final output data matrix produced by the MMA operation.

[0064] In at least one embodiment, the scaling vector size qualifier 342 is 1X, 2X, 4X, or other suitable numbers. The number of scaling vectors in each scaling matrix, as well as other aspects of block scaling, are determined by the scaling vector size qualifier 342, for example, as per [reference to...]. Figure 2 As described.

[0065] If the scaling vector size qualifier 342 is 2X, then each of the two scaling vectors will be multiplied by a portion of the input matrix. More specifically, these portions of the input matrix are referred to herein as “chunks” of the input matrix. The input matrix is ​​divided into several logical chunks, the number of which is specified by the scaling vector size qualifier 342. For example, if the scaling vector size qualifier 342 is 2X, then the input matrix is ​​arranged into two chunks, and each chunk comprises half a vector (e.g., rows or columns) of the input matrix. One of the two scaling vectors will be multiplied by each vector in the first chunk of the input matrix, and the other scaling vector will be multiplied by each vector in the second chunk. Figure 11 An example of block scaling is shown, illustrating the multiplication of a scaling vector with an input matrix of size 2X.

[0066] If the scaling vector size qualifier 342 is 4X, then each of the four scaling vectors will be multiplied by a portion of the input matrix. More specifically, the input matrix is ​​arranged into four blocks, and each block includes a vector (e.g., a row or column) of a quarter of the input matrix. Each of the four scaling vectors will be multiplied by a vector of the corresponding block in the four blocks of the input matrix. Figure 12 An example of block scaling is shown, illustrating the multiplication of a scaling vector with an input matrix of size 4X.

[0067] The example syntax for calling MMA 302 for block scaling is as follows, where input matrices A and B are located in shared memory, the accumulator matrix is ​​located in tensor memory, and the scaling data values ​​are also located in tensor memory: tcgen05.mma.cta_group.kind.block_scale{.scale_vec_size} [d-tmem], a-desc, b-desc, idesc, [scale-A-tmem], [scale-B-tmem], enable-input-d; The example syntax for calling MMA 302 for block scaling is as follows, where input matrix A is located in tensor memory, input matrix B is located in shared memory, the accumulator matrix is ​​located in tensor memory, and the scaling data values ​​are also located in tensor memory: tcgen05.mma.cta_group.kind.block_scale{.scale_vec_size} [d-tmem], [a-tmem], b-desc, idesc, [scale-A-tmem], [scale-B-tmem], enable-input-d; In at least one embodiment, the type (kind) is kind::mxf8f6f4, kind::mxf4, or kind::mxf4nvf4, .cta_group is cta_group::1 or .cta_group::2, and .scale_vec_size is .scale_vec::1X, .scale_vec::2X, or .scale_vec::4X. In at least one embodiment, the example tcgen05.mma instruction is an asynchronous instruction that initiates... M × N × K Matrix multiplication and accumulation operations are used to calculate... ,in A The matrix is M × K , B The matrix is K × N , D The matrix is M × N Because the tcgen05.mma instruction is asynchronous, the calling thread can continue to perform other operations while executing the MMA operation.

[0068] The 32-bit register operand idesc is the instruction descriptor, as shown in the example below. Figure 4A The instruction descriptor 476 describes this. This instruction descriptor specifies the shape, specific type, sparsity, and other details of the input matrix, output matrix, and matrix multiplication and accumulation operations. The qualifier `kind` specifies the general type of the element type of the multiplicand matrix. The specific type of the elements of the input and output matrices for each MMA type is specified in instruction descriptor MMA 476.

[0069] Example block scaling MMA call 302 using operands in tensor memory and shared memory is as follows: tcgen05.mma.cta_group::1.kind::mxf8f6f4 [taddr2],[taddr1], bdesc, idesc,[tmem_scaleA], [tmem_scaleB], p; In the example call, CTA group qualifier 306 is "1", MMA type qualifier 324 is .kind::mxf8f6f4, accumulator matrix D operand 318 is "[taddr2]", input matrix A operand 314 is "[taddr1]", input matrix B operand 316B is "bdesc", instruction descriptor 376 is "idesc", scale data A operand 320A is "[tmem_scaleA]", scale data B operand 320B is "[tmem_scaleB]", and the enable input for D matrix operand 374 is "p".

[0070] Combination Figure 3 At least one characteristic described is achieved using at least one characteristic of any system and / or process depicted in other figures and / or in combination with other figures. Figure 3 At least one feature described is used to implement at least a portion of any system and / or process depicted in other figures and / or in combination with other figures.

[0071] Figure 4AA block scaling MMA instruction descriptor 476 is illustrated according to at least one embodiment. The instruction descriptor 476 includes matrix A element type 412, matrix B element type 414, scaling matrix element type 416, M-dimensionality of matrix A 426, N-dimensionality of matrix B 428, and scaling factor selector operand 432. Matrix A element type 412 specifies the element type of the input matrix A 314 or 316A. As shown in Table 1 herein, if the scaling vector size is 1X, then matrix A element type 412 can be E4M3, E5M2, E3M2, E2M3, or E2M1. If the scaling vector size is 2X, then matrix A element type 412 is E2M1. Matrix B element type 414 specifies the element type of the input matrix B 316B. Matrix B element type 414 follows the same rules as matrix A element type 412.

[0072] The scaling matrix element type 416 specifies the element type of scaling matrix 244. If the scaling vector size is 1X or 2X, then scaling matrix element type 416 is UE8M0. If the scaling vector size is 4X, then scaling matrix element type 416 is UE4M3. The M dimension of matrix A 426 is the number of rows in the input matrix A 132 or 152A. The N dimension of matrix B 428 is the number of columns in the input matrix B152B.

[0073] Scaling factor selector operand 432 specifies scaling factor selectors that specify which scaling factors to select from the scaling data 134. Scaling factor selector operand 432 includes scaling factor selector A 424A and scaling factor selector B 424B. Scaling factor selector A 424A includes an "SFA_ID" value, which is used to select the scaling factor to be multiplied by the elements of matrix A. Similarly, scaling factor selector B 424B specifies an "SFB_ID" value, which is used to select the scaling factor to be multiplied by the elements of matrix B. This document will discuss... Figure 4B and Figure 5 The scaling factor selectors 424A and 424B are described in further detail.

[0074] Figure 4B The illustration shows a block-scaled data layout for various scaling vector sizes in a tensor memory 230 according to at least one embodiment. Tensor memory columns 464 are columns in tensor memory 230, where each channel 446 (e.g., row) specifies up to four scaling factors. These four scaling factors are selected from corresponding sub-columns 444 of tensor memory column 464. SFA_ID and SFB_ID selector values ​​(from scaling factor selectors A 424A and B 424B) specify which sub-columns 444 of the scaled data 434 are to be selected as scaling factors for matrices A and B, respectively. The scaled data 434 may correspond to scaled data 234A or scaled data 234B.

[0075] Tensor memory column 464 is one of the column sequences 464A, 464B to 464N of tensor memory channels 446 in tensor memory 130. Tensor memory channel 446 is one of the channel sequences 446A, 446B and 446C in tensor memory 130. Each column 464 includes sub-columns 444. A scaling factor 448 (“SF”) is stored in each sub-column 444. More specifically, column 464 stores four sub-columns 444, and each sub-column 444 stores one byte of data representing the scaling factor 448. Thus, as Figure 4B As shown, column 464 can store up to four scaling factors 448A-448D. The four sub-columns 444 are named byte-3, byte-2, byte-1, and byte-0, respectively. The four scaling factors stored in bytes-3, byte-2, byte-1, and byte-0 of column 464 are SF_3 448D, SF_2 448C, SF_1 448B, and SF_0 448A, respectively.

[0076] The scaling vector size qualifier 342 specifies the scaling vector size, indicating how many of the four sub-columns 444 are used to provide the scaling factor. The specific sub-column 444 used to provide the scaling factor is specified by the scaling factor selector operands 424A and 424B. Scaling factor selector A 424A selects a scaling factor 448 from channel 446 of the scaling data A 234A to be multiplied with one or more elements of matrix A. Scaling factor selector B 424B selects a scaling factor from channel 446 of the scaling data B 234B to be multiplied with matrix B. An example of scaling factor selection is shown below. Figure 5 As shown.

[0077] In case 460A, where the scaling vector size is 1X, a scaling factor 448 is selected from tensor memory column 464 of tensor memory channel 446A. The selected scaling factor 448 is multiplied by each vector of the input matrix. If the scaling factor selector 424 is 0, scaling factor #0 448A is selected. Similarly, scaling factor selector 424 being 1, 2, or 3 will result in the selection of scaling factor #1 448B, scaling factor #2 448C, or scaling factor #3 448C. Figure 10 An example of block scaling MMA is shown, illustrating the multiplication of a scaling vector with an input matrix of size 1X.

[0078] In case 460B, where the scaling vector size is 2X, two scaling factors 448 are selected from the tensor memory column 464 of the tensor memory channel 446B. Each of these two selected scaling factors 448 is multiplied by each vector in the corresponding block of the input matrix. If the scaling factor selector 424 is 0, scaling factor #1 448B and scaling factor #0 448A are selected to multiply by each vector in each of the two corresponding blocks of the input matrix. If the scaling factor selector 424 is 1, scaling factor #3 448D and scaling factor #2 448C are selected to multiply by each vector in each of the two corresponding blocks of the input matrix.

[0079] In the case 460C where the scaling vector size is 4X, four scaling factors 448 are selected from the tensor memory column 464 of the tensor memory channel 446C. Each of these four selected scaling factors 448 is multiplied by each vector in the corresponding block of the input matrix (selected from four blocks). In the 4X case, the effective value of the scaling factor selector 424 is 0, which indicates that these four scaling factors 448 will be selected.

[0080] As described above, the scaling factor can be an 8-bit (1-byte) value. For scaling vector sizes 1X and 2X, it is represented as type UE8M0 (8-bit exponent, no mantissa); for scaling vector size 4X, it is represented as type UE4M3 (4-bit exponent, 3-bit mantissa, 1 sign bit). Scaling factor 3 (SF_3) 448D is the scaling factor value in the fourth byte of tensor memory column 464. Scaling factor 2 (SF_2) 448C is the scaling factor value in the third byte of tensor memory column 464. Scaling factor 1 (SF_1) 448B is the scaling factor value in the second byte of tensor memory column 464. Scaling factor 0 (SF_0) 448A is the scaling factor value in the first byte of tensor memory column 464.

[0081] Combination Figures 4A-4B At least one feature described is achieved using at least one feature of any system and / or process depicted in other figures and / or in combination with other figures. Figures 4A-4B At least one feature described is used to implement at least a portion of any system and / or process depicted in other figures and / or in combination with other figures.

[0082] Figure 5 The illustration shows scaling factors selected from tensor memory 230 according to at least one embodiment, for various scaling factor selector values ​​424. The scaling factor selector values ​​424 include scaling factor selector A 424A and scaling factor selector B 424B, which are as follows... Figure 3 and Figures 4A-4BThe block scaling MMA call 302 uses operands. Scaling factor selector 424 specifies which scaling factors SF0, SF1, SF2, and SF3 will be selected from tensor memory column 564 by scaling factor selection logic 270. Tensor memory column 564 can reside in scaling data A 234A or scaling data B 234B. For example, different tensor memory columns 564 can exist in scaling data A 234A and scaling data B 234B.

[0083] In at least one embodiment, tensor memory column 564 is a column of tensor memory 230. Tensor memory column 564 includes sub-columns 544A-544D. Each sub-column 544 may be a 1-byte tensor memory, in which case each tensor memory column 564 is 32 bits of data. Each scaling factor SF0, SF1, SF2, and SF3 is stored in the corresponding sub-columns 544A, 544B, 544C, and 544D of tensor memory column 564.

[0084] Scaling factor selection logic 270 selects one or more scaling factors (“scaling factor A”) for block scaling of input matrix A 232 (in TMEM 230) or input matrix A 252A (in SMEM 250). Scaling factor selection logic 270 can cause scaling factor A to be included in the scaling matrix A 244A (e.g., by referencing or copying scaling factor A). As an example, scaling factor A may include one, two, or four scaling factors SF0, SF1, SF2, and / or SF3. Scaling factor selector A 424A (referred to herein as “SFA_ID”) is used by scaling factor selection logic 270 to select one, two, or four scaling factors SF0, SF1, SF2, and / or SF3 from tensor memory column 564 for multiplication with input matrix A 232 or 252A.

[0085] Scaling factor selection logic 270 selects one or more scaling factors (“scaling factor B”) for block scaling of the input matrix B 252B. Scaling factor selection logic 270 can cause scaling factor B to be included in the scaling matrix B 244B (e.g., by referencing or copying scaling factor B). As an example, scaling factor B may include one, two, or four scaling factors SF0, SF1, SF2, and / or SF3. A scaling factor selector B 424B, referred to herein as “SFB_ID”, is used by scaling factor selection logic 270 to select one, two, or four scaling factors SF0, SF1, SF2, and / or SF3 from tensor memory column 564 for multiplication with the input matrix B.

[0086] In at least one embodiment, tensor memory column 564 includes a number of rows corresponding to the number of rows in input matrix A. Scaling factor selection logic 270 selects a scaling factor from each row of tensor memory column 564 to multiply with the elements of the corresponding row of input matrix A. Scaling factor selection logic 270 may cause each selected scaling factor value (e.g., SF0, SF1, SF2, or SF3) to be included in a column of scaling matrix A 244A. In at least one embodiment, tensor memory column 564 includes a number of rows corresponding to the number of columns in input matrix B. Scaling factor selection logic 270 selects a scaling factor from each row to multiply with the elements of the corresponding column of input matrix B. Scaling factor selection logic 270 may cause each selected scaling factor value to be included in a row of scaling matrix B 244B.

[0087] As shown in case 560A where the scaling vector size is 1X, if the scaling vector size specified by scaling vector size qualifier 342 is 1X, a scaling factor is selected from tensor memory column 564 in the scaling data A 234A used for input matrix A. Similarly, a scaling factor is selected from tensor memory column 564 in the scaling data B 234B used for input matrix B. More specifically, as shown in example tensor memory column 564A, in the case of 1X, if the scaling factor selector SFA_ID is 0, the scaling factor selection logic 270 selects the scaling factor SF0 from sub-column 544A to be multiplied with the elements of matrix A. Furthermore, as shown in example tensor memory columns 564B, 564C, and 564D, in the case of 1X, scaling factor selector SFA_ID values ​​of 1, 2, and 3 cause the scaling factor selection logic 270 to select scaling factors SF1, SF2, and SF3 from sub-columns 544B, 544C, and 544D, respectively, to be multiplied with the elements of matrix A. Similarly, the scaling factor selector SFB_ID value of 0, 1, 2 or 3 causes scaling factor selection logic 270 to select scaling factors SF0, SF1, SF2 or SF3 respectively for multiplication with the elements of matrix B.

[0088] As shown in case 560B where the scaling vector size is 2X, if the scaling vector size is 2X, two scaling factors are selected from tensor memory column 564 in the scaling data A 234A used for input matrix A. Similarly, two scaling factors are selected from tensor memory column 564 in the scaling data B 234B used for input matrix B. More specifically, as shown in example tensor memory column 564E, in the 2X case, if the scaling factor selector SFA_ID is 0, the scaling factor selection logic 270 selects two scaling factors SF1 and SF0 from subcolumns 544B and 544A to be multiplied with the elements of matrix A. Furthermore, as shown in example tensor memory column 564F, in the 2X case, the scaling factor selector SFA_ID value is 2 such that the scaling factor selection logic 270 selects scaling factors SF3 and SF2 from subcolumns 544D and 544C respectively to be multiplied with the elements of matrix B. Similarly, a scaling factor selector SFB_ID value of 0 causes scaling factor selection logic 270 to select scaling factors SF1 and SF0 for multiplication with the elements of matrix B, and a scaling factor selector SFB_ID value of 2 causes scaling factor selection logic 270 to select scaling factors SF3 and SF2 for multiplication with the elements of matrix B.

[0089] As shown in case 560C where the scaling vector size is 4X, four scaling factors are selected from tensor memory column 564 in the scaling data A 234A used for matrix A. Similarly, four scaling factors are selected from tensor memory column 564 in the scaling data B 234B used for input matrix B. More specifically, as shown in example tensor memory column 564G, in the case of 4X, if the scaling factor selector SFA_ID is 0, then scaling factor selection logic 270 selects four scaling factors SF3, SF2, SF1, and SF0 from subcolumns 544D, 544C, 544B, and 544A to be multiplied with the elements of matrix A. Similarly, as shown in example tensor memory column 564G, in the case of 4X, the scaling factor selector SFB_ID value of 0 causes scaling factor selection logic 270 to select scaling factors SF3, SF2, SF1, and SF0 from subcolumns 544D, 544C, 544B, and 544A, respectively, for multiplication with the elements of matrix B. While specific numbers of rows and columns, specific scaling factors (e.g., 1, 2, and 4), and a specific number of scaling factors are described in this example, other numbers of rows and columns, as well as other scaling factors, can be used in other examples by changing the numbers and scaling factors according to the description herein.

[0090] Figure 6 The illustration shows a processor 602 according to at least one embodiment, which is used to perform block scaling MMA calls, with its inputs and outputs stored in register 660. Processor 602 and... Figure 1 The processor 102 is similar, but executes an MMA instruction 612 with input operands that identify data stored in register 660, rather than data in tensor memory TMEM 130 and / or SMEM 150. More specifically, the MMA instruction 612 has input matrix operands 614A, 614B, 616, and 618 that refer to input matrices 632A, 632B, 640, and 638 located in register 660, and scaling data operands 620A and 620B that refer to scaling data 634A and 634B located in register 660.

[0091] The MMA instruction 612 can be executed by processor 602 to multiply the elements of two input matrices by a scaling factor, and then multiply the resulting scaled input matrices using matrix multiplication to generate an output matrix. The inputs and / or outputs of the MMA instruction 612 are stored in register 660. Processor 602 can be a graphics processing unit (GPU), a central processing unit (CPU), or other device suitable for executing the MMA instruction 612. Processor 602 includes one or more SMs 604. Each SM 604 includes register 660 for storing data and one or more tensor cores 680 for performing computational operations (such as MMA operations). Each streaming multiprocessor 604 also includes TMEM and SMEM (not shown) for storing data.

[0092] In at least one embodiment, tensor core 680 includes scaling factor selection logic 670A, which selects a scaling factor from scaling data identified by scaling data operand 620 based on scaling data selector operand 622. Alternatively or additionally, scaling factor selection logic 670B may be included in SM 604 outside of tensor core 680. In at least one embodiment, tensor core 680 includes block scaling MMA computation logic 682, which performs computation on input matrices A632A, B632B, C640, and D638 stored in register 660. The input matrices are identified by matrix A operand 614A, matrix B operand 614B, matrix C operand 616, and matrix D operand 618. In at least one embodiment, block scaling MMA computation logic 682 multiplies each of input matrices A632A and B632B by a scaling factor selected from scaling data A634A and B634B to obtain a scaled input matrix. Block-scale MMA computation logic 182 performs MMA operations on scaled input matrices A and B to compute the MMA result (e.g., accumulator) matrix D 638.

[0093] In at least one embodiment, processor 602 includes a cooperative thread array (CTA) (not shown), which may be a group of GPU threads that work together to execute blocks of code and share resources such as memory and synchronization primitives. In at least one embodiment, the CTA executed by processor 602 initiates MMA instruction 612 via instruction scheduler 126 and pipeline manager 128 so that the instruction can be executed on or otherwise implemented on the tensor core 680 of streaming multiprocessor 604.

[0094] When processor 602 executes MMA instruction 612, MMA instruction 612 causes processor 602 to perform an MMA operation that calculates the product of matrix A (register) 632A and matrix B (register) 632B. These MMA operations can be "block-scaled" MMA operations, which perform scaling on one or more elements of at least one input matrix by multiplying each of one or more elements by a scaling factor to generate one or more scaled input values. The scaled input values ​​are used to calculate the product of matrices A and B, instead of using the corresponding unscaled elements of matrices A and B received as input. For example, an MMA operation can calculate a scaled matrix A by multiplying the elements of matrix A 632A by the corresponding scaling factor. An MMA operation can also calculate a scaled matrix B by multiplying the elements of matrix B 632B by the corresponding scaling factor. An MMA operation can further calculate the product of scaled matrices A and B.

[0095] The MMA instruction 612 can be configured via input parameters to multiply specific matrix elements by a selected scaling factor. For example, there may be more matrix elements than the scaling factor, in which case the same “block” scaling factor (e.g., a set of scaling factors arranged by rows and columns) is multiplied by two or more corresponding blocks (e.g., fragments of the input matrix) of the input matrix. Multiplying by the scaling factor can be performed by multiplying each scaling factor (e.g., a numerical value) in the block by each matrix element (e.g., a numerical value) in the block of the matrix. The result of multiplying the scaling factor by the matrix element can be used as input to matrix multiplication, or it can replace matrix elements in the matrix, depending on the implementation details. Certain implementations may employ optimizations, such as by merging the multiplication of the scaling factor by the matrix element with one or more subsequent operations used in the matrix multiplication, in which case one or more of the numerical values ​​described herein are not necessarily stored in memory or registers. For example, for illustrative purposes, the product of the scaling factor and the input matrix element is described as a numerical value in the examples herein, but in an implementation, this product is not necessarily stored as a numerical value in memory or registers.

[0096] In at least one embodiment, the MMA operation is a mathematical operation that, if executed by the processor 602, calculates a value.D : D=AB+C in A , B , C and D It is a matrix, submatrix, submatrix fragment or tile, or a set of values. A and B This is the input matrix, and C on the right is an optional input matrix. If specified, it is used in conjunction with the matrix... A Multiply by matrix B Add the results together, the left side D This is the output matrix, where the results of MMA operations are stored. For example, matrix multiplication. AB The calculation method is as follows: Figure 1 As stated above.

[0097] The block scaling MMA operation can be a mathematical operation that, if executed by the processor 602, calculates a value. D : in and It is a scaling matrix, * indicates a matrix A One or more elements with another matrix Element-wise multiplication of one or more elements, and It is by scaling the matrix A (calculated as) ) and the scaled matrix B (calculated as) The calculation is obtained by multiplying ( ). For example, the block scaling MMA operation can be calculated about ( ). Figure 1 Further description.

[0098] The MMA instruction 612 has a set of operands and a set of qualifiers. The operands provide parameters that can be determined at runtime for the MMA instruction 612, and the qualifiers provide configuration parameters for the MMA instruction 612. The instruction scheduler 126 schedules the MMA instruction 612 for execution on the tensor core 680 of the streaming multiprocessor 604 based on the operands and qualifiers of the MMA instruction 612.

[0099] The operands of MMA instruction 612 include matrix A operand 614A, matrix B operand 614B, matrix C operand 616, and output matrix operation matrix D operand 618. Operands also include scaling data A operand 620A, scaling data B operand 620B, scaling data A selector operand 622A, and scaling data B selector operand 622B. Which elements of A are involved with... The determination of which elements are multiplied depends on the configuration input provided to the MMA instruction 612, such as the scaling vector size qualifier 642, which determines... The size, as described in this article regarding Figure 2 Further described. One or more of the MMA operands 614A-622B identify the data to which the MMA instruction 612 will perform operations. Each of the operands 614A-622B can be a register (specified as one or more register names, such as {%Ra0}), and / or an intermediate value.

[0100] Instruction qualifiers can be specified for MMA instruction 612 to modify its behavior. MMA instruction qualifiers include MMA shape qualifier 606, MMA type qualifier 624, and scaling vector size qualifier 642. MMA shape qualifier 606 specifies the M, N, and K dimensions for the MMA calculation; for example, it can be m16n8k32 (M=16, N=8, K=32), m16n8k64 (M=16, N=8, K=64), or other suitable shapes.

[0101] The MMA type qualifier 624 specifies the size of the matrix elements in input matrices A and B. The MMA type qualifier 624 can be, for example, MXF4, MXF8F6F4, or MF4NVF4, etc., and is referenced in this document. Figure 1 Further description is provided in Table 1. The scaling vector size qualifier 642 specifies the scaling vector size, which indicates how many scaling vectors (e.g., 1, 2, or 4) of the scaling matrix 244 will be used for block scaling.

[0102] The input matrix A operand 614A specifies one or more registers where matrix A 632A is stored. The input matrix B operand 614B specifies one or more registers where matrix B 632B is stored. The matrix C operand 616 specifies one or more registers where matrix C 640 is stored. The matrix D operand 618 specifies one or more registers, and the MMA operation stores matrix D (register) 638 in these registers.

[0103] Both scaling data operand 620A and scaling data operand 620B specify that scaling data A 634A and scaling data B 634B are stored in one or more registers. Scaling data A 634A contains one or more scaling matrices A (e.g., ...). The value of scaling data B 634B contains one or more scaling matrices B (e.g., The scaling data A selector operand 622A indicates which columns to select from scaling data A 634A, for example, to form the columns of scaling matrix A, and the scaling data B selector operand 622B indicates which columns to select from scaling data B 634B, for example, to be converted into the rows of scaling matrix B. The use of scaling data selector operand 622 to indicate which columns to select from scaling data 634A is related to... Figure 7 and Figure 16 Further description. The selector operand 622 also indicates which threads will contribute (e.g., select from the scaled data) the scaling vector, as described in this article regarding... Figure 6 and Figure 7 The number of scaling vectors selected from the scaling data is specified by the scaling vector size qualifier, which can be 1X, 2X, or 4X, for example, as described regarding Figure 2 As stated above.

[0104] In at least one embodiment, combined Figure 6 At least one feature described is achieved by utilizing at least one feature of any system and / or process depicted and / or described in conjunction with other figures. In at least one embodiment, in conjunction with Figure 6 At least one feature described is used to implement at least a portion of any system and / or process depicted in other figures and / or in combination with other figures.

[0105] Figure 7 The illustration depicts an interaction in a processor 702 according to at least one embodiment for executing a block scaling MMA instruction 612, the input and output matrices of which are stored in register 760. In at least one embodiment, the block scaling MMA input 712 includes qualifiers for the MMA instruction 612, as well as operands that also include the MMA instruction 612. These operands specify the input and output matrices stored in register 760. The qualifiers include an MMA shape qualifier 706, an MMA type qualifier 724, and a scaling vector size qualifier 742, which are related to... Figure 6 The descriptions of MMA shape qualifier 606, MMA type qualifier 624, and scaling vector size qualifier 642 are similar. These operands include matrix A operand (register) 714A, matrix B operand (register) 714B, matrix C operand (register) 716, and matrix D operand (register) 718. (About...) Figure 6 Further description of matrix operands. The block-scaled MMA input 712 also includes M-dimensional 726 operands (e.g., matrix...). A , C and D The number of rows), N-dimensional 728 operands (e.g., matrix) B , C and DThe number of columns), scaling data operands 720 and scaling factor selector operands 722.

[0106] Scaling factor selection logic 770 and block scaling MMA logic 782 and Figure 2 The scaling factor selection logic 270 and block scaling MMA logic 282 are similar, but they access the input matrices 732 and 740 in register 760 and store the accumulator matrix D 738 in register 760 (instead of...). Figure 2 (That is, accessing the matrix in memory and storing the matrix in memory). Block scaling MMA logic 782 can be executed by one or more tensor cores 780. Furthermore, scaling factor selection logic 770 can select a specific thread to contribute scaling factor values ​​from scaling data A 734A and scaling data B 734B. Scaling factor selection logic 770 is used in conjunction with... Figure 2 The scaling factor selection logic 270 in the code describes an operation similar to the one described above, used to select the scaling matrix 744 from the scaling data 734. Furthermore, the block scaling MMA logic 782 multiplies the scaling matrix 744 with the input matrix 743 and provides the resulting scaled input matrix as input to the MMA operation, for use with respect to... Figure 2 The block scaling MMA logic 282 describes an operation similar to generating an accumulator matrix D 738.

[0107] See Figure 7 The scaling data operand 720A specifies that scaling data A 734A is stored in one or more of these registers. Additionally, the scaling data operand 720B specifies that scaling data B 734B is stored in one or more of these registers. The scaling factor selector operand 722A specifies a byte-id-a selector value, which the scaling factor selection logic 770 uses to select a column (also called a byte) from the registers where scaling data A 734A is stored. The scaling factor selection logic 770 includes (e.g., stores) the selected column of scaling data into the scaling matrix A 744A. This byte-id-a selector value is similar to that of... Figure 4A The scaling factor selector A424A specifies the SFA_ID value.

[0108] The scaling factor selector operand 722B specifies the byte-id-b selector value, which the scaling factor selection logic 770 uses to select a column (also called a byte) from the register where the scaling data B 734B is stored. The scaling factor selection logic 770 converts the selected column of the scaling data B 734B into a row and includes (e.g., stores) that row in the scaling matrix A744A. The byte-id-b selector value is similar to that of... Figure 4AThe scaling factor selector B 424B specifies the SFB_ID value. For example, the byte-id-a and byte-id-b selectors are related to... Figure 16 It was described further.

[0109] The scaling factor selection logic 770 and / or block scaling MMA logic 782 use threads of processor 702 to contribute (e.g., select and / or retrieve) columns of scaling data 734 to be included in scaling matrix 744. In at least one embodiment, two threads contribute scaling factors from scaling data A (register) 734A, and one thread contributes scaling factors from scaling data B (register) 734B. The thread contributing scaling factors from scaling data A 734A is selected as specified by the thread-id-a portion of selector operand 722A. The thread contributing scaling factors from scaling data B 734B is selected as specified by the thread-id-b portion of selector operand 722B. For example, the thread-id-a and thread-id-b selectors are related to... Figure 17A , Figure 17B and Figure 18 It was described further.

[0110] The block scaling MMA logic 782 multiplies selected elements of scaling matrix A 744A with corresponding elements of input matrix A 732 to generate a scaled A value (register) 736A. Additionally or alternatively, the block scaling MMA logic 782 multiplies elements of scaling matrix B 744B with corresponding elements of input matrix B 752B to generate a scaled B value (register) 736B, for example, as per [reference to...]. Figure 2 The block scaling MMA logic 282 is described in the text.

[0111] Combination Figure 7 At least one feature described is achieved using at least one feature of any system and / or process depicted and / or described in conjunction with other figures. Figure 7 At least one feature described is used to implement at least a portion of any system and / or process depicted in other figures and / or in combination with other figures.

[0112] Figure 8 The illustration shows a block-scaling MMA call 802 according to at least one embodiment, wherein the input matrix and output matrix are stored in registers. Operation 800 represents one or more computational operations that cause a processor (e.g., a GPU) to execute or otherwise implement MMA (matrix multiplication-accumulation) instructions using block scaling to adjust the scaling of the input matrix elements stored in the registers.

[0113] Operation 800 includes a block scaling MMA call 802 and a block scaling MMA response 878. In at least one embodiment, call 802 is a call to instructions to cause one or more processors to perform one or more block scaling MMA computation operations (“MMA operations”) and generate a response 878. In at least one embodiment, call 802 is a call to instructions (“instruction call”) to cause one or more processors 702 to perform one or more block scaling MMA operations and generate a response 878.

[0114] In at least one embodiment, call 802 is an API call to cause one or more processors 702 to perform one or more computational operations and generate a response 878. In at least one embodiment, call 802 is a function call to be executed by one or more software programs (e.g., CUDA kernels) on a parallel processing unit (e.g., a GPU). Call 802 defines parameters and configurations for performing block-scaled MMA operations on MMA operands 810, thereby achieving performance optimization through block scaling techniques (e.g., multiplying MMA input matrix elements by scaling factors), which may differ for different rows, columns, and / or segments of the input matrix.

[0115] Call 802 includes MMA qualifiers 804, which specify parameters for configuring aspects of the MMA operation that typically remain unchanged during program execution. MMA qualifiers 804 include an MMA type qualifier 824, a block scaling qualifier 808, and a scaling vector size qualifier 842, which are respectively associated with... Figure 3 The MMA type qualifier 324, block scaling qualifier 308, and scaling vector size qualifier 342 are similar. MMA qualifier 804 also includes the element types of scaling matrix element type qualifier 816, MMA shape qualifier 830, and input and output matrix qualifiers 832. Scaling matrix element type qualifier 816 is used in conjunction with this document. Figure 4A The scaling matrix element type described is similar to 416. The MMA shape qualifier 830 is similar to the one described in this article. Figure 6 The description of the MMA shape qualifier 606 is similar. The data matrix element type qualifier 832 includes qualifiers that specify the element types of matrices A 732A, B 732B, C 740, and D 738.

[0116] Call 802 receives MMA operand 810 as input. MMA operand 810 includes input matrix A operand 814A, input matrix B operand 814B, input matrix C operand 836, and matrix D operand 818. Input matrix A operand 814A specifies one or more registers 760 containing input matrix A 732A. Input matrix B operand 814B specifies one or more registers 760 containing input matrix B 732B, which will be used together with matrix A in the MMA operation. Input matrix C operand 836 specifies one or more registers 760 containing input matrix C 740 to be added to the result of multiplying scaled matrices A and B. Accumulator matrix D operand 818 specifies one or more registers 760 in which the accumulator matrix D 738 generated by the MMA operation will be stored. The matrix D operand 818 can be used in conjunction with the input matrix C operand 836 to accumulate the results of matrix multiplication operations across multiple computation cycles (e.g., multiple block-scaling MMA fetches 802). Accumulating MMA results in this way can improve performance because it allows intermediate results to be cached and reused across multiple computation cycles, thereby minimizing the need to fetch intermediate results from main memory or write them back to main memory.

[0117] MMA operand 810 also includes scaling data A operand (register) 820A, scaling data B operand (register) 820B, and scaling factor selector operand 822. Scaling data A operand 820A specifies one or more registers 760 containing scaling matrix A 744A. Scaling data B operand 820B specifies one or more registers 760 containing scaling matrix B 744B. Scaling factor selector operand 822 includes scaling factor A byte and thread selector 834A and scaling factor B byte and thread selector 834B. Scaling factor A byte and thread selector 834A specifies byte-id-a and / or thread-id-a selectors, and scaling factor B byte and thread selector 834B specifies byte-id-b and / or thread-id-b selectors, for example, as described herein. Figure 7 As stated above.

[0118] The block-scaling MMA response 878 represents the output of the MMA instruction execution, which includes the matrix D result 838 stored in one or more registers 760. In at least one embodiment, if call 802 is to call an instruction, the output is the data that the instruction will output. In at least one embodiment, if call 802 is to call an API, the output is the data that the API will output. In at least one embodiment, the block-scaling MMA response 878 includes a status indicator (not shown) indicating the success or failure of the MMA operation, such as whether the operation completed successfully or encountered an error. The accumulator matrix D result 838 represents the final output data matrix produced by the MMA operation.

[0119] The following is an example syntax for calling 802 using the MMA for block scaling, where matrices A, B, C, D and the scaling data values ​​are in registers: mma.sync.aligned.shape.row.col.kind.block_scale {.scale_vec_size}.dtype.atype.btype.ctype.stype d, a, b, c, scale-a-data, {byte-id-a, thread-id-a}, scale-b-data, {byte-id-b, thread-id-b}; In this example call syntax, the "shape" qualifier specifies the shape of the MMA operation, in the format mMnNkK, where M, N, and K are the dimensions of the matrix. The MMA instruction 612 with the block_scale qualifier initiates a matrix multiplication operation using block scaling. This operation calculates... ,in A The matrix is M × K , B The matrix is K × N , C and D All matrices are M × N In at least one embodiment, the "sync" qualifier specifies that the MMS instruction is synchronous, for example, the calling thread waits until the MMA operation is complete.

[0120] In the example syntax, the "row" qualifier indicates that matrix A is in row-major order, the "col" qualifier indicates that matrix B is in column-major order, the "kind" qualifier corresponds to the MMA type qualifier 624, the "scale_vec_size" qualifier corresponds to the scaling vector size qualifier 342, the MMA shape is m16n8k64, the type is mxf4, and the scaling vector size is 2X.

[0121] For example, if kind is specified as mxf4, then shape is m16n8k64, scale_vec_size is 2X, dtype is f32, atype and btype are e2m1, ctype is f32, and stype is ue8m0. In another example, if kind is specified as mxf4nvf, then shape is m16n8k64, scale_vec_size is 2X or 4X, dtype is f32, atype and btype are both e2m1, ctype is f32, and stype is ue8m0 (if scale_vec_size is 2X) or ue4m3 (if scale_vec_size is 4X). In another example, if kind is specified as mxf8f6f4, then shape is m16n8k32, scale_vec_size is 1X, dtype is f32, atype and btype are one of e4m3, e5m2, e3m2, e2m3 or e2m1, ctype is f32, and stype is ue8m0.

[0122] In at least one embodiment, the `block_scale` qualifier specifies that matrices A and B are scaled using matrices `scale_A` and `scale_B`, respectively, before performing matrix multiplication and accumulation operations. The data type corresponding to each element in matrices `scale_A` and `scale_B` is specified using `stype`. The `scale_vec_size` qualifier specifies the number of columns in matrix `scale_A` and the number of rows in matrix `scale_B`.

[0123] In at least one embodiment, Table 1 describes valid combinations of kind, stype, and scale_vec_size. For an MMA using kind::mxf4, scale_vec_size defaults to 2X when the qualifier scale_vec_size is not specified. Conversely, when kind is specified as kind::mxf8f6f4, the qualifier scale_vec_size defaults to 1X. However, for kind::mxf4nvf4, a valid scale_vec_size will be specified. A thread bundle executing the mma.sync instruction with shape=m8n8k4 computes four matrix multiplications and accumulations. For other shapes, mma.sync computes a single matrix multiplication and accumulation per thread bundle.

[0124] In at least one embodiment, operands a and b represent two multiplicand matrices A and B, while c and d represent the accumulator matrix and the destination matrix, distributed across threads in the thread bundle. When the block_scale qualifier is specified, operands scale-a-data and scale-b-data represent scaling matrix metadata (e.g., scaling data) containing matrices scale_A and scale_B, respectively. Triples {byte-id-a, thread-id-a} and {byte-id-b, thread-id-b} represent selectors for choosing matrices scale_A and scale_B from their corresponding metadata arguments scale-a-data and scale-b-data, respectively. Operands scale-a-data and scale-b-data are of type b32. Operands byte-id-a, thread-id-a, byte-id-b, and thread-id-b are unsigned 16-bit integer values. A register in each thread holds a fragment of the matrix. The qualifiers dtype, atype, btype, and ctype indicate the data type of the elements in matrices D, A, B, and C, respectively. The qualifier stype indicates the data type of the elements in matrices scale_A and scale_B.

[0125] In at least one embodiment, a specific shape may have the meaning of the storage format of matrix elements. When the type is kind::mxf8f6f4 or kind::f8f6f4, each 4-bit and 6-bit floating-point type element is packed into a 6-bit container. Matrix elements of type e2m1 are located in the middle 4 bits of this 6-bit container and padded with the high 2 bits and low 2 bits of the container. When the type of the matrix element is e3m2 or e2m3, the matrix element is located in the low 6 bits of the 6-bit container and padded with the high 2 bits of the container. Conversely, when using mma with kind::mxf4 or kind::mxf4nvf4, explicit padding is not required even if the type of the matrix element is e2m1. In at least one embodiment, when using mma with kind::mxf4 or kind::mxf4nvf4, the elements of input matrix 952 (e.g., Figure 9 shown A[0, c] or B[r, 0] The elements have a reduced size, such as 4 bits, and these elements are represented using the amount of data specified by the reduced size. For example, when using an mma with kind::mxf4 or kind::mxf4nvf4, the elements of the input matrix 952 are stored in data units referred to herein as “containers”, which have a reduced size, such as 4 bits or other bit lengths less than 8 bits, such as 6 bits.

[0126] Example block scaling MMA call 612 using operands in registers (there are three lines of register declarations before the mma instruction call) is as follows: .reg .b32 %Ra <4> , %Rb <4> ; .reg .f32 %Rc <4> , %Rd <4> ; .reg .b32 scaleAData, scaleBData; mma.sync.aligned.m16n8k64.row.col.kind::mxf4.block_scale.f32.e2m1.e2m1.f32.ue8m0 {%Rd0, %Rd1, %Rd2, %Rd3}, {%Ra0, %Ra1, %Ra2, %Ra3}, {%Rb0, %Rb1}, {%Rc0, %Rc1, %Rc2, %Rc3}, scaleAData, {2, 1}, scaleBData, {2, 3}; In this example block scaling MMA call 613, the "shape" qualifier is m16n8k64, the "row" qualifier specifies that matrix A is in row-major order, the "col" qualifier specifies that matrix B is in column-major order, the "kind::mxf4" qualifier specifies that the scaling vector size is 4X, the "block_scale" qualifier specifies that block scaling should be performed using scale-a-data and scale-b-data operands, the "f32" qualifier specifies that the element type of matrix D is f32, the e2m1 qualifier specifies that the element type of matrices A and B is e2m1, the f32 qualifier specifies that the element type of matrix C is f32, and the ue8m0 qualifier specifies that the element type of scaling matrices A and B is ue8m0.

[0127] The operands {%Rd0,%Rd1,%Rd2,%Rd3} specify that the resulting matrix D will be stored in registers Rd0, Rd1, Rd2, and Rd3, each register (e.g., Rd0) containing 32 bits. Matrix D has 16 × 8 elements, and each element occupies 32 bits (because the element type of matrix D is f32), therefore the number of bits representing matrix D is... In a thread bundle comprising 32 threads, each thread processes... Bit data. Since each matrix element occupies 32 bits, the operand size is 32 bits. To fit the 128 bits processed by each thread into 32-bit operands, four 32-bit operands are needed. Therefore, matrix D is represented by four 32-bit registers: %Rd0, %Rd1, %Rd2, and %Rd3. Furthermore, the operand {%Ra0, %Ra1, %Ra2, %Ra3} specifies that input matrix A will be retrieved from registers Ra0, Ra1, Ra2, and Ra3. The operand {%Rb0, %Rb1} specifies that input matrix B will be retrieved from registers Rb0 and Rb1. The operand {%Rc0, %Rc1, %Rc2, %Rc3} specifies that input matrix C will be retrieved from registers Rc0, Rc1, Rc2, and Rc3.

[0128] The operand `scaleAData` specifies a 32-bit (b32) register from which scaled data A will be retrieved. The operand `{2,1}` specifies byte-id-a as 2 and thread-id-a as 1. The operand `scaleBData` specifies a 32-bit (b32) register from which scaled data B will be retrieved. The second operand `{2,3}` specifies byte-id-b as 2 and thread-id-b as 3.

[0129] Combination Figure 8 At least one feature described is achieved using at least one feature of any system and / or process depicted and / or described in conjunction with other figures. Figure 8At least one feature described is used to implement at least a portion of any system and / or process depicted in other figures and / or in combination with other figures.

[0130] Figure 9 The illustration depicts a block scaling MMA operation performed according to at least one embodiment using scaling matrices 944A and 944B selected from scaling data 934. In at least one embodiment, the block scaling MMA operation includes: selecting scaling matrices 944A and 944B from scaling data 934; performing block scaling on input matrices 952A and 952B using scaling matrices 944A and 944B to compute scaled element values ​​corresponding to unscaled element values ​​of input matrices 952A and 952B; and performing matrix multiplication using the scaled element values ​​instead of one or more unscaled element values ​​of input matrices 952A and 952B. The matrix multiplication computes an accumulator matrix D 938. The matrix multiplication operation may include addition operations, and / or may be followed by addition operations to accumulate the results of previous matrix multiplication operations, for example, if the accumulator matrix D 938 generated by the previous matrix multiplication operation is provided as an operand to MMA instruction 112.

[0131] Scaling factor selection logic 970 selects scaling matrices 944A, 944B from one or more portions of scaling data 934 specified by scaling factor selector operand 922. In at least one embodiment, scaling factor selection logic 970 selects scaling matrix A 944A from scaling data A 934A and scaling matrix B 944B from scaling data B 934B. Block scaling MMA logic 984 uses element-wise multiplication to multiply the input element values ​​of each input matrix 952 with the element values ​​of the corresponding scaling matrix 944 to compute scaled element values. For example, this document relates to... Figure 1 Describes element-wise multiplication. The block-scaled MMA logic 984 performs matrix multiplication by using scaled element values ​​as elements of the input matrix, instead of using the corresponding unscaled input element values, to compute the accumulator matrix D 938.

[0132] Scaling factor selection logic 970 selects one or more scaling vectors 982 from the tensor memory column 964 of the scaling data 934, representing the scaling matrix 944. The number of scaling vectors 982 selected from the tensor memory column 964 is specified by the scaling vector size 942. In at least one embodiment, scaling factor selection logic 970 selects one or more scaling factors SF0 from a sub-column of the tensor memory column 964, as described herein. Figure 4B (where the scaling factor is selected from sub-column 444) and Figure 5(Where the scaling factor is selected from sub-column 544) as described. Each scaling factor SF0 can have a different value. For example, three scaling factors SF0, SF0, and SF0 (where 0 indicates column #0) with different values ​​are shown in three rows of a sub-column of tensor memory column A 964A, but any suitable number of scaling factors SF0 can be selected from this sub-column. Scaling factors SF1 (column #1), SF2 (column #2), and SF3 (column #3) are also shown in other sub-columns of tensor memory column A 964A. The number of sub-columns from which scaling factors are selected is specified by scaling vector size 942. For example, if scaling vector size 942 is 1X, then one sub-column of scaling factors is selected. The sub-columns from which scaling factors SF0, SF1, SF2, and / or SF3 are selected are specified by scaling factor selector operand 922. For example, if scaling vector size 942 is 1X and scaling factor A selector operand 922A is SF2 (e.g., a value of 2), then as per the description... Figure 5 The second column (from left to right) containing the scaling factors of SF1 is selected from tensor memory column A964A. Scaling factor selection logic 970 includes each selected column of scaling factors in scaling matrix 944 as a scaling vector 982, which is either a column in scaling matrix A944A or a row in scaling matrix B944B. For example, if the scaling vector size 942 is 1X, the number of rows M in input matrix A952A is 16, and the scaling factor A selector operand 922A is SF2, then a vector of M=16 scaling factors SF1 from the second sub-column of tensor memory column A964A is included in scaling matrix A944A as scaling vector A982A. That is, if M=16 rows, then 16 scaling factors SF1 with different values ​​can be selected from the second sub-column of tensor memory column A964A and included in scaling matrix A944A.

[0133] Scaling factor selection logic 970 identifies one or more scaling vectors A 982A (which are column vectors) to be included in the scaling vector A 982A from one or more tensor memory columns A 964A. The number of scaling factors is specified by the scaling vector size 942. The row of scaling factors selected from tensor memory columns A 964A is specified by the scaling factor A selector operand 922A. In at least one embodiment, the scaling matrix A 944A references the identified tensor memory sub-columns in tensor memory column A 964A, thus eliminating the need to make a copy of the identified sub-columns. In another embodiment, scaling factor selection logic 970 copies the identified sub-columns to columns of the scaling matrix A 944A, which may be stored in a different memory location than the scaling data 934.

[0134] Similarly, scaling factor selection logic 970 identifies one or more scaling vectors B 982B (which are row vectors) to be included in the scaling vector B 982B from one or more tensor memory columns B 964B. The number of scaling factors is specified by the scaling vector size 942. The column of the scaling factor selected from the tensor memory column B 964B is specified by the scaling factor B selector operand 922B. In at least one embodiment, the scaling matrix B 944B references the identified tensor memory sub-columns in the tensor memory column B 964B, so there is no need to make a copy of the identified sub-columns. In another embodiment, scaling factor selection logic 970 copies the identified sub-columns to the rows of the scaling matrix B 944B (e.g., by transposing the sub-columns to form rows), and these rows can be stored in a different memory location than the scaling data 934.

[0135] To perform the scaling operation, the block scaling MMA logic 984 multiplies each scaling vector A 982A with one or more input vectors A 980A of the input matrix A 952A, as described in this paper. Figure 2 As described above. Similarly, the block scaling MMA logic 984 multiplies each scaling vector B 982B with one or more input vectors B 980B of the input matrix B 952B, as described herein. Figure 2 As stated above.

[0136] For example, if the scaling vector size 942 is 1X, then the scaling matrix A 944A contains a single scaling vector A 982A (one column), which is multiplied by each input vector A 980A (another column) to produce a scaled vector. For example, if k=64, then the scaling vector A 982A is multiplied by each of the 64 input vectors A 980A to produce 64 scaled input vectors. These 64 scaled input vectors are provided as inputs to the scaling matrix A to the matrix multiplication operation to be performed by the block scaling MMA logic 984. Similarly, when the scaling vector size 942 is 1X, the scaling matrix B 944B contains a single scaling vector B 982B (one row), which is multiplied by each input vector A 980B (another row) to produce a scaled vector. For example, if k=64, then the scaling vector B 982B is multiplied by each of the 64 input vectors A 980A to produce 64 scaled input vectors. These 64 scaled input vectors are provided as inputs to the scaled matrix A to the block scaling MMA logic 984. In at least one embodiment, the block scaling MMA logic 984 combines (e.g., merges) the scaling factor multiplication operation with the addition, multiplication, and / or other operations involved in matrix multiplication. Figure 12 An example of scaling operation with a scaling vector size of 942 = 1X is shown.

[0137] As another example, if the scaling vector size 942 is 4X, then the scaling matrix A944A contains four scaling vectors A982A (four columns). These four scaling vectors are multiplied by the four corresponding input vectors A980A (columns) of the input matrix A952A to produce four scaled vectors. These four corresponding input vectors A980A are located in four corresponding blocks of the input matrix A952A. This column multiplication is repeated for each additional column in each block of the input matrix A952A, as described in this paper regarding... Figure 12 As described. For example, if the input matrix A 952A has 64 columns, each of the four scaling vectors A982A of the scaling matrix A 944A is multiplied by 64 / 4 = 16 corresponding columns in a block of four blocks of the input matrix A 952A. The resulting 64 scaled input vectors A 980A are provided as inputs to the scaling matrix A to the block scaling MMA logic 984.

[0138] Similarly, when the scaling vector size 942 is 4X, the scaling matrix B 944B contains four scaling vectors B 982A (four rows), which are multiplied by the four corresponding input vectors B 980B (rows) of the input matrix B 952B to produce four scaled vectors. These four corresponding input vectors B 980B are located in four corresponding blocks of the input matrix B 952B. This row multiplication is repeated for each additional row in each block of the input matrix B 952B, as described in this paper regarding... Figure 12 As described. For example, if the input matrix B 952B has 64 rows, each of the four scaling vectors B 982B of the scaling matrix B 944B is multiplied by 64 / 4 = 16 rows of the corresponding block (a block of four blocks) of the input matrix B 952B. The resulting 64 scaled input vectors B 980B are provided as inputs to the scaling matrix B to the block scaling MMA logic 984. In at least one embodiment, the block scaling MMA logic 984 combines (e.g., fuses) the scaling factor multiplication operation with addition, multiplication, and / or other operations involved in matrix multiplication.

[0139] The block-scale MMA logic 984 performs matrix multiplication accumulation by multiplying the scaled element values ​​of the scaled input vector A980A corresponding to the input matrix A952A with the scaled element values ​​of the scaled input vector B980B corresponding to the input vector B980B. If the input accumulator matrix is ​​provided as input (e.g., as an operand of MMA instruction 112), the block-scale MMA logic 984 adds the result of the matrix multiplication to the input accumulator matrix. The block-scale MMA logic 984 outputs the resulting product matrix as the accumulator matrix D938. If the input accumulator matrix is ​​provided, the block-scale MMA logic 984 outputs the sum of the resulting product matrix and the accumulator matrix as the accumulator matrix D938.

[0140] Combination Figure 9 At least one feature described is achieved using at least one feature of any system and / or process depicted and / or described in conjunction with other figures. Figure 9 At least one feature described is used to implement at least a portion of any system and / or process depicted in other diagrams and / or in combination with other diagrams. For example, scaling factor selection logic 970 may correspond to scaling factor selection logic 170, and block scaling MMA logic 984 may correspond to Figure 1 Block scaling MMA calculation logic 182.

[0141] Figure 10 The illustration depicts a block scaling MMA operation using scaling matrix 1044 according to at least one embodiment, wherein the scaling vector size 1042 is 1X. In at least one embodiment, the block scaling MMA operation includes: selecting scaling matrices 1044A and 1044B from scaling data 1034; performing block scaling on input matrices 1052A and 1052B using scaling matrices 1044A and 1044B to compute scaled element values ​​corresponding to the unscaled element values ​​of the input matrices 1052A and 1052B; and performing matrix multiplication using the scaled element values ​​instead of one or more unscaled element values ​​of the input matrices 1052A and 1052B. This matrix multiplication computes an accumulator matrix D 1038. Since in Figure 10 In the example, the scaling vector size is 1042, which is 1X. Therefore, the scaling matrix A1044A is a single column vector, and the scaling matrix B1044B is a single row vector.

[0142] The scaling factor selection logic 1070 selects scaling matrix A 1044A from one or more columns of scaling data 1034 according to the scaling factor A selector operand 1022A, and selects scaling matrix B 1044B from one or more columns of scaling data 1034 according to the scaling factor B selector operand 1022B. The scaling factor selection logic 1070 converts the columns of scaling data 1034 into rows to be included in scaling matrix B 1044B, for example, by performing a transpose operation to transpose the columns of scaling data 1034 into a row vector of scaling matrix B 1044B.

[0143] To perform the scaling operation, the block scaling MMA logic 1084 uses element-wise multiplication to multiply a single column vector in the scaling matrix A1044A with each input vector A1080A of the input matrix A1052A to compute a scaled A value 1056A. The scaled A value 1056A is used as input to the MMA operation, instead of the element values ​​of the input matrix A1052A. In at least one embodiment, the scaled A value 1056A may be a matrix “Scaled” with the same dimensions as the input matrix A1052A, where each matrix element Scaled[row,column] is computed by multiplying A[row,column] by Scale[row,0], where “A” is the input matrix A1052A and “Scale” is the scaling matrix A1044A. The scaled matrix is ​​used as the input matrix for the MMA operation performed by the block scaling MMA logic 1084. In at least one embodiment, the scaled A value 1056A is not stored in the matrix, but is calculated by the block-scaled MMA logic 1084 for each element of the input matrix A 1052A and used as input for the MMA operation. For example, this document regarding Figure 1 Further describe element-wise multiplication of vectors.

[0144] As an example, scaling matrix A 1044A contains a single scaling vector A 1082A (one column), which is multiplied by each input vector A 1080A (another column) to produce a scaled vector. For example, if k=64, the scaling vector A 1082A is multiplied by each of the 64 input vectors A 1080A to produce 64 scaled input vectors, which are shown as scaled A values ​​1056A. The scaled A value 1056A is provided as input to the scaling matrix A to the matrix multiplication operation to be performed by the block scaling MMA logic 1084. Similarly, scaling matrix B 1044B contains a single scaling vector B 1082B (one row), which is multiplied by each input vector B 1080B (one row) to produce a scaled vector. For example, if k=64, the scaling vector B 1082B is multiplied by each of the 64 input vectors B 1080B to generate 64 scaled input vectors, which are represented as scaled B values ​​1056B. The scaled B values ​​1056B are provided as inputs to the scaled matrix B to the MMA operation to be performed by the block scaling MMA logic 1084. In at least one embodiment, the block scaling MMA logic 1084 combines (e.g., fuses) the scaling factor multiplication operation with addition, multiplication, and / or other operations involved in matrix multiplication.

[0145] The block scaling MMA logic 1084 uses the scaled A value 1056A as an element of the scaled input matrix A and the scaled B value 1056B as an element of the scaled matrix B to perform matrix multiplication to compute the accumulator matrix D1038.

[0146] Figure 11 The illustration depicts a block scaling MMA operation using scaling matrix 1144 according to at least one embodiment, wherein the scaling vector size 1142 is 2X. In at least one embodiment, the block scaling MMA operation includes: selecting scaling matrices 1144A and 1144B from scaling data 1134; performing block scaling on input matrices 1152A and 1152B using scaling matrices 1144A and 1144B to compute scaled element values ​​corresponding to the unscaled element values ​​of the input matrices 1152A and 1152B; and performing matrix multiplication using the scaled element values ​​instead of one or more unscaled element values ​​of the input matrices 1152A and 1152B. This matrix multiplication computes an accumulator matrix D 1138. Since in Figure 11 In the example, the scaling vector size is 2X, so the scaling matrix A1144A consists of two column vectors, and the scaling matrix B1144B consists of two row vectors.

[0147] Scaling factor selection logic 1170 selects scaling matrix A 1144A from one or more columns of scaling data 1134 according to the scaling factor selector operand 1122, and also selects scaling matrix B 1144B from one or more columns of scaling data 1134 according to the scaling factor selector operand 1122. Scaling factor selection logic 1170 converts the columns of scaling data 1134 into rows to be included in scaling matrix B 1144B, for example, by performing a transpose operation to transpose the columns of scaling data 1134 into a row vector of scaling matrix B 1144B.

[0148] Block scaling MMA logic 1184 performs block scaling on the input matrix A 1152A to compute a scaled A value 1156A. To perform scaling on the input matrix A 1152A, block scaling MMA logic 1184 multiplies the elements of each column vector 1182 of the scaling matrix A 1144A with the elements of each input vector A 1180A (e.g., column) in the block 1110 of the input matrix A 1152A. In at least one embodiment, block scaling MMA logic 1184 stores the result of the multiplication in the scaled A value 1156A. Because in Figure 11 In the example, the scaling vector size 1142 is 2X, so the input matrix A1152A is divided into 2 blocks, each with K / 2 columns. For illustrative purposes, the example input matrix A1152A is shown as having K=8 columns, and each block has K / 2=4 columns. In at least one embodiment, K is 64, and each block has 64 / 2=32 columns. Furthermore, each scaling vector A1182A in the scaling matrix A1144A has M elements, where M is the number of rows in the input matrix A1152A.

[0149] The block scaling MMA logic 1184 performs block scaling on the input matrix B 1152B to compute the scaled B value 1156B. To perform scaling on the input matrix B 1152B, the block scaling MMA logic 1184 performs element-wise multiplication of each row vector of the scaling matrix B 1144B with each input vector 1180B in a block of the input matrix B 1152B. The block scaling MMA logic 1184 stores the result of the multiplication in the scaled B value 1156B. Since the scaling vector size is 2X, the input matrix B 1152B is divided into 2 blocks, each with K / 2 rows. For illustrative purposes, an example input matrix B 1152B is shown as having K=8 rows, with each block having K / 4=2 rows. In one example, K=64, and each block has 64 / 2=32 rows. Furthermore, each scaling vector B1182B in the scaling matrix B1144B has N elements, where N is the number of columns in the input matrix B1152B.

[0150] The block-scaled MMA logic 1184 calculates the scaled A value 1156A by multiplying the columns of the input matrix A 1152A with the columns of the scaling matrix A 1144A, as follows. For the first block of the input matrix A 1152A (e.g., columns 0-3), the scaling vector A 1182A (which is the first column of the scaling matrix A 1144A) is multiplied by each of the four columns in the first block (columns 0-3) of the input matrix A 1152A to calculate a portion of the scaled A value 1156A corresponding to the first four columns of the input matrix A 1152A. In at least one embodiment, the scaled A value 1156A is a matrix, in which case the scaled A value 1156A may be, for example, the first four columns of the scaled A value 1156A. In at least one embodiment, the scaled A value 1156A are individual scaled values ​​used in the MMA operation without being stored in a matrix. Furthermore, for the second block of the input matrix A 1152A (e.g., columns 4-7), another scaling vector A 1182A (which is the second column of the scaling matrix A 1144A) is multiplied by each of the four columns in the second block (columns 4-7) of the input matrix A 1152A to compute a portion of the scaled A value 1156A corresponding to the last four columns of the input matrix A 1152A.

[0151] The block-scaled MMA logic 1184 computes the scaled B value 1156B by multiplying the rows of the input matrix B 1152B with the rows of the scaling matrix B 1144B, as follows. For the first block of the input matrix B 1152B (e.g., rows 0-3), the scaling vector B 1182B (which is the first row of the scaling matrix B 1144B) is multiplied by each of the four rows in the first block (rows 0-3) of the input matrix B 1152B to compute a portion of the scaled B value 1156B corresponding to the first four rows of the input matrix B 1152B. In at least one embodiment, the scaled B value 1156B is a matrix, in which case the scaled B value 1156B may be, for example, the first four rows of the scaled B value 1156B. In at least one embodiment, the scaled B value 1156B are individual scaled values ​​used in the MMA operation without being stored in a matrix. Furthermore, for the second block of the input matrix A 1152B (e.g., column 4-7), another scaling vector B 1182B (which is the second row of the scaling matrix B 1144B) is multiplied by each of the four rows in the second block of the input matrix B 1152B (column 4-7) to compute a portion of the scaled B value 1156B corresponding to the last four rows of the input matrix B 1152B.

[0152] The block scaling MMA logic 1184 uses the scaled A value 1156A as an element of the scaled input matrix A and the scaled B value 1156B as an element of the scaled matrix B to perform matrix multiplication to compute the accumulator matrix D1138.

[0153] Figure 12 The illustration shows a block scaling MMA operation 1200 using a scaling matrix 1244 according to at least one embodiment, wherein the scaling vector size is 4X. In at least one embodiment, the block scaling MMA operation includes: selecting scaling matrices 1244A and 1244B from scaling data 1234; performing block scaling on input matrices 1252A and 1252B using scaling matrices 1244A and 1244B to compute scaled element values ​​corresponding to the unscaled element values ​​of the input matrices 1252A and 1252B; and performing matrix multiplication using the scaled element values ​​instead of one or more unscaled element values ​​of the input matrices 1252A and 1252B. The matrix multiplication computes an accumulator matrix D 1238. Since in Figure 12 In the example, the scaling vector size is 4X, so the scaling matrix A 1244A consists of four column vectors, and the scaling matrix B 1244B consists of four row vectors.

[0154] Scaling factor selection logic 1270 selects scaling matrix A 1244A from one or more columns of scaling data 1234 according to the scaling factor selector operand 1222, and also selects scaling matrix B 1244B from one or more columns of scaling data 1234 according to the scaling factor selector operand 1222. Scaling factor selection logic 1270 converts the columns of scaling data 1234 into rows for inclusion in scaling matrix B 1244B, for example, by performing a transpose operation to transpose the columns of scaling data 1234 into a row vector of scaling matrix B 1244B.

[0155] The block scaling MMA logic 1284 performs block scaling on the input matrix A 1252A to compute the scaled A value 1256A. To perform scaling on the input matrix A 1252A, the block scaling MMA logic 1284 multiplies the element of each column vector of the scaled matrix A 1244A with the element of each input vector (e.g., column) in block 1210 of the input matrix A 1252A. The block scaling MMA logic 1284 stores the result of the multiplication in the scaled A value 1256A. Because in Figure 12In the example, the scaling vector size is 4X, so the input matrix A 1252A is divided into 4 blocks 1210, each block having K / 4 columns. For illustrative purposes, the example input matrix A 1252A is shown as having K=8 columns, with each block 1210 having K / 4=2 columns. In at least one embodiment, K is 64, and each block 1210 has 64 / 4=16 columns. Furthermore, each scaling vector A in the scaling matrix A 1244A has M elements, where M is the number of rows in the input matrix A 1252A.

[0156] The block scaling MMA logic 1284 performs block scaling on the input matrix B 1252B to compute the scaled B value 1256B. To perform scaling on the input matrix B 1252B, the block scaling MMA logic 1284 performs element-wise multiplication of each row vector of the scaling matrix B 1244B with each input vector in block 1212 of the input matrix B 1252B. The block scaling MMA logic 1284 stores the result of the multiplication in the scaled B value 1256B. Since the scaling vector size is 4X, the input matrix B 1252B is divided into 4 blocks 1212, each block having K / 4 rows. For illustrative purposes, an example input matrix B 1252B is shown with K=8 rows, and each block 1212 has K / 4=2 rows. In one example, K=64, and each block 1212 has 64 / 4=16 rows. Furthermore, each scaling vector B in the scaling matrix B1244B has N elements, where N is the number of columns in the input matrix B1252B.

[0157] The block scaling MMA logic 1284 calculates the scaled value A1256A by multiplying the columns of the input matrix A1252A by the columns of the scaling matrix A1244A, as shown below. Since the scaling vector size is 4X, the input matrix A1252A has four blocks 1210, and the scaling matrix A1244A has four column vectors. In the example shown, each block 1210 of the input matrix A1252A has K / 4 = 2 column vectors. Each block 1210 corresponds to one column vector in the scaling matrix A1244A, and the column vector of block 1210 is multiplied by the corresponding column vector of the scaling matrix A1244A, as shown by a pair of arrows pointing to the "x" (multiplication) sign from the two vectors in each block 1210 of the input matrix A1252A to the corresponding column vector of the scaling matrix A1244A.

[0158] For the first block 1210 (e.g., columns 0-1) of the input matrix A 1252A, the scaling vector A (which is the first column of the scaling matrix A1244A) is multiplied by each of the two columns in the first block 1210 (columns 0-1) of the input matrix A 1252A to compute a portion of the scaled A value 1256A corresponding to the first two columns of the input matrix A 1252A. In at least one embodiment, the scaled A value 1256A is a matrix, in which case the scaled A value 1256A may be, for example, the first two columns of the scaled A value 1256A. In at least one embodiment, the scaled A value 1256A are individual scaled values ​​used in the MMA operation without needing to be stored in the matrix. Furthermore, for the second block 1210 (e.g., columns 2-3) of the input matrix A 1252A, another scaling vector A (which is the second column of the scaling matrix A 1244A) is multiplied by each of the two columns in the second block 1210 (columns 2-3) of the input matrix A 1252A to compute a portion of the scaled A value 1256A corresponding to the last four columns of the input matrix A 1252A. Similarly, for the third and fourth blocks 1210 of the input matrix A 1252A, the third and fourth columns of the scaling matrix A 1244A are multiplied by the third and fourth blocks 1210 of the input matrix A 1252A.

[0159] The block scaling MMA logic 1284 calculates the scaled B value 1256B by multiplying the rows of the input matrix B 1252B with the rows of the scaling matrix B 1244B, as shown below. Since the scaling vector size is 4X, the input matrix B 1252B has four blocks 1212, and the scaling matrix B 1244B has four row vectors. In the example shown, each block 1212 of the input matrix B 1252B has K / 4 = 2 row vectors. Each block 1212 corresponds to one of the row vectors of the scaling matrix B 1244B, and the row vector of block 1212 is multiplied by the corresponding row vector of the scaling matrix B 1244B, as indicated by a pair of arrows pointing from two vectors in each block 1212 of the input matrix B 1252B to the "x" (multiplication) sign of the corresponding column vector from the scaling matrix B 1244B.

[0160] For the first block 1212 (e.g., rows 0-1) of the input matrix B 1252B, the scaling vector B (which is the first row of the scaling matrix B 1244B) is multiplied by each of the two rows in the first block 1212 (e.g., rows 0-1) of the input matrix B 1252B to compute a portion of the scaled B value 1256B corresponding to the first two rows of the input matrix B 1252B. In at least one embodiment, the scaled B value 1256B is a matrix, in which case the scaled B value 1256B may be, for example, the first two rows of the scaled B value 1256B. In at least one embodiment, the scaled B value 1256B are individual scaled values ​​used in the MMA operation without needing to be stored in the matrix. Furthermore, for the second block 1212 (e.g., column 2-3) of the input matrix A 1252B, another scaling vector B (which is the second row of the scaling matrix B 1244B) is multiplied by each of the two rows in the second block 1212 (column 2-3) of the input matrix B 1252B to compute a portion of the scaled B value 1256B corresponding to the last two rows of the input matrix B 1252B. Similarly, for the third and fourth blocks 1212 of the input matrix B 1252B, the third and fourth rows of the scaling matrix B 1244B are multiplied by the third and fourth blocks 1212 of the input matrix B 1252B.

[0161] The block scaling MMA logic 1284 uses the scaled A value 1256A as an element of the scaled input matrix A and the scaled B value 1256B as an element of the scaled matrix B to perform matrix multiplication to compute the accumulator matrix D1238.

[0162] Figure 13An example process 1300 according to at least one embodiment is illustrated, which performs block scaling MMA operations using data stored in tensor memory and / or shared memory. In at least one embodiment, part or all of process 1300 (or any other process described herein, or variations and / or combinations thereof) is executed under the control of one or more computer systems configured with computer-executable instructions and is implemented as code (e.g., computer-executable instructions, one or more computer programs, or one or more application programs) executed jointly on one or more processors by hardware, software, or a combination thereof. In at least one embodiment, the code is stored in the form of a computer program on a computer-readable storage medium comprising a plurality of computer-readable instructions executable by one or more processors. In at least one embodiment, the computer-readable storage medium is a non-transitory computer-readable medium. At least some of the computer-readable instructions available for executing process 1300 are not stored using only transient signals (e.g., propagating transient electrical or electromagnetic signals). Non-transitory computer-readable media do not necessarily include non-transitory data storage circuitry (e.g., buffers, caches, and queues) within a transient signal transceiver.

[0163] In at least one embodiment, process 1300 is executed by one or more systems, such as the systems described in this disclosure. In at least one embodiment, process 1300 is performed by, for example, systems combined with... Figure 1-12 The system execution is described above. In at least one embodiment, process 1300 is executed by one or more programming models (e.g., CUDA, HIP, oneAPI, and / or variations thereof). One or more processes of process 1300 are executed in any suitable order, including sequential, parallel, and / or variations thereof, and are executed by any suitable system (e.g., the system described herein) and using any suitable processing unit, such as CPU, GPU, PPU, GPGPU, and / or variations thereof. Although process 1300 is described as a series of steps or operations, process 1300 may include modified or reordered steps or operations, or omit certain steps or operations, unless otherwise explicitly stated or logically required, such as when the output of one step or operation is used as the input of another step or operation. In at least one embodiment, process 1300 is executed by one or more systems, such as in combination with... Figure 19-38C The system described. One or more steps of process 1300 may be performed as part of the execution of one or more APIs (such as the API described herein).

[0164] The system executing at least a portion of process 1300 includes code for receiving at least 1302 instructions for matrix multiplication and accumulation (MMA) operations, the MMA operations having operands indicating a first input matrix, a second input matrix, and one or more scaling data storage locations storing scaling data. The instructions may be MMA instruction 112, which, when invoked in block scaling MMA call 302, has MMA operands 310 specifying matrices stored in tensor memory 130 and / or shared memory 150. These matrices include input matrix A 132 (e.g., in tensor memory 130) or input matrix A 152A (e.g., in shared memory 150). These matrices also include input matrix B 152B (e.g., in shared memory 150).

[0165] MMA instruction 112 is associated with or includes MMA type qualifier 124. If MMA type qualifier 124 is mxf4 or mxf4nvf4, it indicates that the matrix element type of the input matrix is ​​a 4-bit quantity. That is, the mxf4 or mxf4nvf4 qualifier indicates that each element of matrix A and / or matrix B is stored in 4 bits of a storage device (e.g., in memory or a register). Furthermore, the mxf4 and mxf4nvf4 qualifiers specify that the 4-bit elements of matrices A and B will be stored in a 4-bit container that uses 4 bits of data for each matrix element, without the need for additional data. In at least one embodiment, using a 4-bit container is more efficient than using a larger container (e.g., an 8-bit container). In at least one embodiment, mxf8f6f4 specifies that the matrix element type of matrix A and / or matrix B is 8-bit, 6-bit, or 4-bit, and these elements will be stored in an 8-bit container. In at least one embodiment, using an 8-bit container is less efficient than using a 4-bit container because less memory is used to store each matrix element, and the memory bandwidth and / or communication bandwidth (e.g., the amount of data accessed, sent, received, or otherwise processed per unit time) used for operations performed on matrix elements is less than that used for operations performed on matrix elements when the container size is 4 bits, compared to a larger container size (e.g., 8 bits).

[0166] In at least one embodiment, the first input matrix corresponds to input matrix A 132 or 152A, the second input matrix corresponds to input matrix B 152B, and the scaled data storage location corresponds to... Figure 1The scaling data is A 134A and / or B 134B. In at least one embodiment, the first input matrix corresponds to input matrix A 632A stored in register 660, and the second input matrix corresponds to input matrix B (register) 632B stored in register 660. In at least one embodiment, the scaling data storage location corresponds to scaling data operand 720A and / or scaling data operand 720B. In at least one embodiment, the scaling data storage location corresponds to scaling data A 634A and / or scaling data B 634B stored in register 660.

[0167] The system executing process 1300 includes at least a portion of code for selecting a scaling matrix 1304 from one or more scaling data storage locations based on a scaling vector size input and a scaling data selector operand. For example, the scaling matrix could be... Figure 7 The scaling matrix A is 744A and / or the scaling matrix B is 744B. The scaling vector size input can be a scaling vector size qualifier 742, and the scaling data selector operand can be a scaling factor selector operand 722A and / or a scaling factor selector operand 722B. In at least one embodiment, the system uses the scaling vector size qualifier 742, the scaling data operand 720, and / or the scaling factor selector operand 722 to select the scaling matrix. For example, the system can use the scaling matrix as described herein. Figure 9 The scaling matrix is ​​selected as described above. The scaling matrix may correspond to scaling matrix A 944A, which includes at least one scaling vector A 982A. The scaling matrix may alternatively or additionally correspond to scaling matrix B 944B, which includes at least one scaling vector B 982B.

[0168] The system performing at least a portion of process 1300 includes code for identifying, at least 1306, one or more scaling vectors indicated by a scaling vector size input (e.g., scaling vector size qualifier 742) within a scaling matrix. The scaling matrix includes one or more scaling vectors, which can be row vectors or column vectors, depending on whether the scaling matrix is ​​used to scale input matrix A or input matrix B. For scaling input matrix A, the scaling matrix contains column vectors of elements multiplied with the elements of input matrix A. For scaling input matrix B, the scaling matrix contains row vectors of elements multiplied with the elements of input matrix B. The number of vectors selected from the scaling matrix is ​​specified by the scaling vector size input. In at least one embodiment, the system identifies, for example, those described herein with respect to… Figure 1 and Figure 9 The scaling vector mentioned above.

[0169] The system executing at least a portion of process 1300 includes code for generating at least 1308 multiple scaled values ​​corresponding to elements of the first input matrix. Each scaled value is the product of a corresponding matrix element in the first input matrix and a scaling element in one or more identified scaling vectors. Furthermore, the position of the scaling element in the one or more scaling vectors is determined based on the position of the corresponding matrix element in the first input matrix. Scaled values ​​can be generated by multiplying matrix elements by their corresponding scaling elements in the scaling vectors.

[0170] In at least one embodiment, the system generates a scaled A value 236 in tensor memory, or a scaled A value 256A in shared memory, and / or a scaled B value 256B in shared memory. For example, this document also relates to... Figure 10 The scaled A value 1056A and scaled B value 1056B in the figure describe the scaled values. In at least one embodiment, the scaled matrices A and B are generated using element-wise multiplication, wherein vectors of input matrices A and B are multiplied by the corresponding scaling factors.

[0171] The system executing at least a portion of process 1300 includes code for performing at least 1310 MMA operations, which generate a result matrix by performing one or more matrix multiplication operations based on a plurality of scaled values ​​and one or more elements of a second input matrix. In at least one embodiment, the system performs matrix multiplication operations. For example, block-scaled MMA logic 282 multiplies a scaled matrix A having the same dimensions as the first input matrix, the scaled matrix A containing scaled values ​​in its elements located at positions (e.g., row, column positions) in the matrix corresponding to the positions of the elements of the first input matrix A, these elements being multiplied by a scaling factor to compute the scaled values. In at least one embodiment, the matrix multiplication operation for multiplying an input of dimension M×K by a scaling matrix A and an input of dimension K×N by a scaling matrix B includes: generating an output matrix D of dimension M×N, and for each element D[i,j] of the output matrix D, calculating the product of each element of row i of the scaling matrix A with the corresponding element of column j of the scaling matrix B, then calculating the sum of these products, and storing the sum of products in D[i,j]. The sum of the products of each row and each column is also called the dot product of the row vector and the column vector. In at least one embodiment, the input matrix A... In another example, block scaling MMA logic 282 replaces the elements of the first input matrix with scaled values ​​calculated by multiplying the elements of the first input matrix by a scaling value. In at least one embodiment, block scaling MMA logic 282 uses matrix multiplication to multiply the elements of a matrix having a scaled A value 236 with the elements of a second input matrix to generate an accumulator matrix D 238.

[0172] Figure 14 An example process 1400 according to at least one embodiment is illustrated, which performs a block scaling MMA operation using data stored in registers. In at least one embodiment, part or all of process 1400 (or any other process described herein, or variations and / or combinations thereof) is executed under the control of one or more computer systems configured with computer-executable instructions and is implemented as code (e.g., computer-executable instructions, one or more computer programs, or one or more application programs) executed jointly on one or more processors by hardware, software, or a combination thereof. In at least one embodiment, the code is stored in the form of a computer program on a computer-readable storage medium comprising multiple computer-readable instructions executable by one or more processors. In at least one embodiment, the computer-readable storage medium is a non-transitory computer-readable medium. At least some of the computer-readable instructions available for executing process 1400 are not stored using only transient signals (e.g., propagating transient electrical or electromagnetic signals). Non-transitory computer-readable media do not necessarily include non-transitory data storage circuitry (e.g., buffers, caches, and queues) within a transient signal transceiver.

[0173] In at least one embodiment, process 1400 is performed by one or more systems, such as the systems described in this disclosure. In at least one embodiment, process 1400 is performed by, for example, systems combined with Figure 1-12 The system execution is as shown. In at least one embodiment, process 1400 is executed by one or more programming models (e.g., CUDA, HIP, oneAPI, and / or variations thereof). One or more processes in process 1400 are executed in any suitable order, including sequential, parallel, and / or variations thereof, and are executed by any suitable system (e.g., the system described herein) and using any suitable processing unit, such as CPU, GPU, PPU, GPGPU, and / or variations thereof. Although process 1400 is described as a series of steps or operations, process 1400 may include modified or reordered steps or operations, or omit certain steps or operations, unless otherwise explicitly stated or logically required, such as when the output of one step or operation is used as the input of another step or operation. In at least one embodiment, process 1400 is executed by one or more systems, such as in combination Figure 19-38C The system described. One or more steps of process 1400 may be performed as part of the execution of one or more APIs (such as the API described herein).

[0174] The system executing at least a portion of process 1400 includes code for receiving instructions from 1402 for matrix multiplication and accumulation (MMA) operations, the MMA operations having inputs indicating a first input matrix, first scaling data, a second input matrix, second scaling data, and a scaling vector size. The instructions may be MMA instruction 612, which, when invoked in block scaling MMA call 802, has MMA operands 810 specifying matrices stored in register 660. These matrices include input matrix A 632A and input matrix B 632B. MMA operand 810 also specifies first scaling data (e.g., scaling data A 634A) and second scaling data (e.g., scaling data B 634B). MMA instruction 612 is also associated with or includes a scaling vector size, such as scaling vector size qualifier 642.

[0175] The system executing at least a portion of process 1400 includes code for identifying at least one or more blocks of the first input matrix 1404, wherein each block includes one or more first input vectors of the first input matrix, and the number of blocks is at least partially based on the scaling vector size. The system is as described herein with respect to, for example... Figure 2 and Figure 11 The aforementioned identification block, Figure 2 and Figure 11 The illustration shows an example block 1110 of input matrix A 1152A and a block 1112 of input matrix B 1152B. In at least one embodiment, the MMA instruction 612 has an MMA type qualifier 824 that specifies the number of blocks in input matrices A and B.

[0176] The system executing at least a portion of process 1400 includes code for selecting at least one thread 1406 from a set of threads associated with a portion of the first input matrix, based on instruction-based thread selection input. The thread selection input may, for example, be as described herein... Figure 8 The scaling factor selector operand 822 includes one or more scaling factor selector operands. Scaling factor selector operand 822 includes a scaling factor A byte and a thread selector 834A, the latter including a byte selector (“byte-id-a”) and a thread selector (“thread-id-a”). Similarly, scaling factor selector operand 822 includes a scaling factor B byte and a thread selector 834B, which includes a byte selector (“byte-id-b”) and a thread selector (“thread-id-b”).

[0177] The system selects two threads from the four-thread thread quad 1752 associated with the first input matrix (e.g., input matrix A). The two threads for the first input matrix are selected based on the thread-id-a selector value, for example, as described in this paper regarding... Figure 17A and Figure 17B The system also selects a thread from the four-thread quadruple 1752 associated with the second input matrix (e.g., input matrix B). For example, if the thread-id-a selector is 0, the system uses threads T0 and T1 from the thread quadruple 1752 to contribute the scaling factor for scaling input matrix A; if the thread-id-a selector is 1, threads T2 and T3 are used to contribute the scaling factor. Furthermore, the system uses one of the threads from the thread quadruple 1752 to contribute the scaling factor for scaling input matrix B. The system selects thread T0, T1, T2, or T3 from the quadruple 1752 based on the thread-id-b selector value, as described herein. Figure 18 As described.

[0178] The system executing at least a portion of process 1400 includes code for retrieving one or more elements of one or more scaling vectors of a scaling matrix 1408 from first scaling data using at least one identified thread. The number of scaling vectors in the scaling matrix is ​​based on the scaling vector size, and one or more scaling vectors are selected from the scaling matrix based on instruction byte selector operands.

[0179] For example, if the input matrix A is to be scaled using a scaling matrix, the system uses two threads to contribute scaling factors for scaling the input matrix A. One thread contributes the scaling factors from the scaling data A 734A to the first half of the scaling matrix A 1780, and the other thread contributes the scaling factors from the scaling data A 734A to the second half of the scaling matrix A 1782, as described in this paper. Figure 17B As described above, the scaling factor contributed to scaling matrix A is selected from one or more sub-columns of scaling data column 1664 of scaling data A 734A according to the byte-id-a selector, as described in this article regarding... Figure 4B , Figure 5 and Figure 16 For example, a thread specified by thread-id-a can be used to contribute 1, 2, 4, or other number of subcolumns of the scaled data A 734A as column vectors to (e.g., load and / or include in) the scaled matrix A 744A, where the number of subcolumns is specified by the scaled vector size qualifier 742.

[0180] Furthermore, for example, if the input matrix B is to be scaled by a scaling matrix, the system uses a thread to contribute the scaling factor for scaling the input matrix B. Figure 18 As shown, the thread contributes the scaling factor from scaling data B 734B to scaling matrix B 1844B. The scaling factor contributed to scaling matrix B is selected from one or more sub-columns of scaling data column 1664 of scaling data B 734B according to the byte-id-b selector. For example, a thread specified by thread-id-b can contribute 1, 2, 4, or other numbers of sub-columns of scaling data B 734B as row vectors to (e.g., load, transpose, and / or include) scaling matrix B 744B, where the number of sub-columns is specified by scaling vector size qualifier 742.

[0181] The system performing at least a portion of process 1400 includes code for generating at least one or more scaled input vectors, each scaled input vector comprising one or more scaled values. The system generates the scaled values ​​by computing one or more products of one or more elements of the corresponding input vector of the first input matrix and one or more corresponding elements of the scaling vector of the scaling matrix. The input vectors reside in a set of blocks of the first input matrix, and the scaling vectors correspond to that set of blocks of the first input matrix. For example, if the scaling vector size is 2X, the first input matrix corresponds to an input matrix A1152A having two sets of blocks 1110 (each set having K / 2 columns). Furthermore, the input vectors may correspond to input vector A1180A, the sets of blocks of the first input matrix may correspond to a set of blocks 1100 containing the first four columns of the input matrix A1152A, and the scaling vector may be a scaling vector A1182A, where scaling vector A is the first column of the scaling matrix A1144A. The scaling vector A1182A is multiplied by each of the four columns in that set of blocks.

[0182] The system executing at least a portion of process 1400 includes code for performing at least 1412 MMA operations to generate a plurality of elements of a result matrix. Each of the plurality of elements is generated using matrix multiplication techniques, wherein the system computes the dot product of a scaled input vector corresponding to the first input matrix and a second scaled input vector corresponding to the second input matrix.

[0183] Figure 15An example process 1500 according to at least one embodiment is illustrated for selecting scaling matrix elements from scaling data. In at least one embodiment, part or all of process 1500 (or any other process described herein, or variations and / or combinations thereof) is executed under the control of one or more computer systems configured with computer-executable instructions and is implemented as code (e.g., computer-executable instructions, one or more computer programs, or one or more application programs) executed jointly on one or more processors by hardware, software, or a combination thereof. In at least one embodiment, the code is stored in the form of a computer program on a computer-readable storage medium comprising multiple computer-readable instructions executable by one or more processors. In at least one embodiment, the computer-readable storage medium is a non-transitory computer-readable medium. At least some of the computer-readable instructions available for executing process 1500 are not stored using only transient signals (e.g., propagated transient electrical or electromagnetic signals). Non-transitory computer-readable media do not necessarily include non-transitory data storage circuitry (e.g., buffers, caches, and queues) within a transient signal transceiver.

[0184] In at least one embodiment, process 1500 is executed by one or more systems, such as the systems described in this disclosure. In at least one embodiment, process 1500 is performed by, for example, systems combined with Figure 1-12 The system execution is as shown. In at least one embodiment, process 1500 is executed by one or more programming models (e.g., CUDA, HIP, oneAPI, and / or variations thereof). One or more processes in process 1500 are executed in any suitable order, including sequential, parallel, and / or variations thereof, and are executed by any suitable system (e.g., the system described herein) and using any suitable processing unit, such as CPU, GPU, PPU, GPGPU, and / or variations thereof. Although process 1500 is described as a series of steps or operations, process 1500 may include modified or reordered steps or operations, or omit certain steps or operations, unless otherwise explicitly stated or logically required, such as when the output of one step or operation is used as the input of another step or operation. In at least one embodiment, process 1500 is executed by one or more systems, such as in combination Figure 19-38C The system described. One or more steps of process 1500 may be performed as part of the execution of one or more APIs (such as the API described herein).

[0185] The system executing at least a portion of process 1500 includes code for identifying at least one or more rows in tensor memory 1502, these rows beginning with the row indicated by the scaled data operand of the instruction, wherein the number of selected rows is determined based on the number of rows in the input matrix. For example, these rows may correspond to lane 446 of tensor memory 430, as described herein. Figure 4B As described, the system can identify M rows of tensor memory 430 corresponding to M channels of tensor memory, where the M channels begin with channel 446A indicated by a scaling data operand (e.g., scaling data A 234A). For example, scaling data A 234A may be a shared memory descriptor that includes a channel number and a column number.

[0186] The system executing at least a portion of process 1500 includes at least code for selecting one or more sub-columns of one or more columns in tensor memory, starting from the column indicated by the scaling data operand, wherein the one or more sub-columns are selected using the scaling vector size operand and the scaling data selection operand of the instruction. For example, the system selects one or more sub-columns 444 of column 464A in tensor memory 430, wherein the position (e.g., column number) of column 464A in tensor memory 432 is specified by the scaling data operand (e.g., scaling data A 234A).

[0187] The system executing at least a portion of process 1500 includes at least one or more matrix elements, including code 1506 in a scaling matrix, for retrieving from an identified sub-column of an identified row in a tensor memory. For example, as Figure 4B and Figure 5 As shown, if the scaling vector size is 2X and SFA_ID=0, the system selects two scaling factors from two sub-columns 544A and 544B in channel 446A of tensor memory column 464A, for example, SF_0 448A and SF_1 448B. The system includes the two scaling factors SF_0 and SF_1 in one row of scaling matrix A 244A as two matrix elements in two corresponding columns. See also Figure 11 The two scaling factors can be stored in the two elements of the top row of the scaling matrix A 1144A.

[0188] Figure 16 The illustration shows scaling factors selected from register 660 for various scaling factor selector operands 822 according to at least one embodiment. Scaling factor selector operand 822 is as follows... Figure 8The block scaling MMA call 802 has an operand 822. The scaling factor selector operand 822. The scaling factor selector 834 specifies which scaling factors SF0, SF1, SF2, and SF3 will be selected from the scaling data column 1664 by the scaling factor selection logic 770. In at least one embodiment, the scaling data column 1664 may reside in scaling data A 734A or scaling data B (register) 734B. For example, scaling data A 734A and scaling data B 734B may have different scaling data columns 1664.

[0189] In at least one embodiment, the scaling data column 1664 is stored in one or more registers 660. The scaling data column 1664 includes four sub-columns 1644A-1644D. Each sub-column 1644 may be 1 byte of register storage space, in which case each scaling data column 1664 is 32 bits of data. Each scaling factor SF0, SF1, SF2, and SF3 is stored in the corresponding sub-columns 1644A, 1644B, 1644C, and 1644D of the scaling data column 1664.

[0190] Scaling factor selection logic 770 selects one or more scaling factors (“scaling factor A”) for block scaling of the input matrix A 732A. Scaling factor selection logic 770 can cause scaling factor A to be included in the scaled matrix A 744A (e.g., by referencing or copying scaling factor A). For example, scaling factor A can include one, two, or four scaling factors SF0, SF1, SF2, and / or SF3.

[0191] The scaling factor selector operand 822 includes a scaling factor A byte and a thread selector 834A, the latter including a “byte-id-a” selector. The byte-id-a selector is used by the scaling factor selection logic 770 to select one, two, or four scaling factors SF0, SF1, SF2, and / or SF3 from the scaling data column 1664. The scaling factor A byte and thread selector 834A may also include a thread-id-a selector, which is used by the scaling factor selection logic 770 to select one or more threads that will be used to contribute to the scaling factor indicated by the byte-id-a selector, for example, as described herein. Figure 17A and Figure 17B As stated above.

[0192] Scaling factor selection logic 770 selects one or more scaling factors (“scaling factor B”) for block scaling of the input matrix B 732B. Scaling factor selection logic 770 can cause scaling factor B to be included in the scaled matrix B 744B (e.g., by referencing or copying scaling factor B). As an example, scaling factor B may include one, two, or four scaling factors SF0, SF1, SF2, and / or SF3. Scaling factor selector operand 822 includes a scaling factor B byte and thread selector 834B, which includes a “byte-id-b” selector that will be used by scaling factor selection logic 770 to select one, two, or four scaling factors SF0, SF1, SF2, and / or SF3 from the scaled data column 1664 for multiplication with the input matrix B. The scaling factor B byte and thread selector 834B may also include a thread-id-b selector, which is used by the scaling factor selection logic 770 to select one or more threads that will contribute the scaling factor indicated by the byte-id-b selector, for example, as described herein. Figure 17A and Figure 18 As stated above.

[0193] In at least one embodiment, the number of rows in the scaling data column 1664 corresponds to the number of rows in the input matrix A 732A. Scaling factor selection logic 770 selects a scaling factor from each row of the scaling data column 1664 to multiply with the element of the corresponding row of the input matrix A. Scaling factor selection logic 770 can cause each selected scaling factor value (e.g., SF0, SF1, SF2, or SF3) to be included in a column of the scaling matrix A 744A. In at least one embodiment, the number of rows in the scaling data column 1664 corresponds to the number of columns in the input matrix B 732B. Scaling factor selection logic 770 selects a scaling factor from each row of the scaling data column 1664 to multiply with the element of the corresponding column of the input matrix B. Scaling factor selection logic 770 can cause each selected scaling factor value to be included in a row of the scaling matrix B 744B.

[0194] As shown in case 1660A (scaling vector size = 1X), if the scaling vector size specified by scaling vector size qualifier 842 is 1X, a scaling factor is selected from scaling data column 1664 of scaling data A 734A of input matrix A. Similarly, a scaling factor is selected from scaling data column 1664 of scaling data B 734B of input matrix B. More specifically, as shown in example scaling data column 1664A, in the case of 1X, if the scaling factor selector byte-id-a is 0, the scaling factor selection logic 770 selects a scaling factor SF0 from sub-column 1644A to be multiplied by the elements of matrix A. Furthermore, as shown in the example scaled data columns 1664B, 1664C, and 1664D, in the case of 1X, the scaling factor selector byte-id-a values ​​of 1, 2, and 3 cause scaling factor selection logic 770 to select scaling factors SF1, SF2, and SF3 from subcolumns 1644B, 1644C, and 1644D, respectively, for multiplication with the elements of matrix A. Similarly, the scaling factor selector byte-id-b values ​​of 0, 1, 2, or 3 cause scaling factor selection logic 770 to select scaling factors SF0, SF1, SF2, or SF3, respectively, for multiplication with the elements of matrix B.

[0195] As shown in example scaling vector size 1660B (2X), if the scaling vector size is 2X, two scaling factors are selected from scaling data column 1664 of the scaling data A 734A of the input matrix A. Similarly, two scaling factors are selected from scaling data column 1664 of the scaling data B 734B of the input matrix B. More specifically, as shown in example scaling data column 1664E, in the 2X case, if the scaling factor selector byte-id-a is 0, the scaling factor selection logic 770 selects two scaling factors SF1 and SF0 from subcolumns 1644B and 1644A to be multiplied with the elements of matrix A. Furthermore, as shown in example scaling data column 1664F, in the 2X case, the scaling factor selector byte-id-a value is 2 such that the scaling factor selection logic 770 selects scaling factors SF3 and SF2 from subcolumns 1644D and 1644C respectively to be multiplied with the elements of matrix B. Similarly, a scaling factor selector byte-id-b value of 0 causes scaling factor selection logic 770 to select scaling factors SF1 and SF0 for multiplication with the elements of matrix B, and a scaling factor selector byte-id-b value of 2 causes scaling factor selection logic 270 to select scaling factors SF3 and SF2 for multiplication with the elements of matrix B.

[0196] As shown in case 1660C where the scaling vector size is 4X, if the scaling vector size is 4X, then four scaling factors are selected from scaling data column 1664 of scaling data A 734A of matrix A. Similarly, four scaling factors are selected from scaling data column 1664 of scaling data B734B of input matrix B. More specifically, as shown in example scaling data column 1664G, in the case of 4X, if the scaling factor selector byte-id-a is 0, then scaling factor selection logic 770 selects four scaling factors SF3, SF2, SF1, and SF0 from subcolumns 1644D, 1644C, 1644B, and 1644A to be multiplied with the elements of input matrix A. Similarly, as shown in the scaled data column 1664G in the 4X case, a scaling factor selector byte-id-b value of 0 causes scaling factor selection logic 770 to select scaling factors SF3, SF2, SF1, and SF0 from subcolumns 1644D, 1644C, 1644B, and 1644A, respectively, for multiplication with the elements of the input matrix B. Although specific numbers of rows and columns, specific scaling factors (e.g., 1, 2, and 4), and a specific number of scaling factors are described in this example, other numbers of rows and columns, as well as other scaling factors, can be used in other examples by changing the numbers and scaling factors according to the description herein.

[0197] Figure 17A The illustration shows a fragment layout of an input matrix 732A according to at least one embodiment, and the associated thread ID of a thread bundle including threads contributing scaling factors from scaling data 734A. Matrix A fragment 1750 includes, for example, 32 columns and 8 rows of the input matrix A 732A. The input matrix A 732A will be used for MMA operations, having an example shape of m16n8k64, where M is 16, n is 8, and k is 64. Since the input matrix A 732A has 16 rows and 64 columns, matrix A fragment 1750, with 8 rows (dimension M) and 32 columns (dimension K), represents one-quarter of the input matrix A 732A. Each column of matrix A fragment 1750 contains 32 bits of data and can store eight 4-bit elements of the input matrix A 732A.

[0198] Matrix A segment 1750 is generated by 4 × 8 = 32 threads (e.g., a thread bundle) using matrix A segment 1750. Figure 17A The mapping of threads shown is used for processing. Each row of segment 1750 of matrix A is processed by 4 threads, which form a thread quadruple 1752. For example, two threads in thread quadruple 1752 retrieve the scaling factor from the scaling data A 734A, as described herein.

[0199] The columns of matrix A segment 1750 are divided into four groups of eight columns each, with each column used for MMA calculations (e.g., matrix multiplication calculations performed after block scaling). Each group comprises eight elements 1770 from a row of matrix A segment 1750. Therefore, each row of matrix A segment 1750 has four groups of eight elements each. For example, the eight elements 1770 are eight adjacent elements in the eight corresponding columns of a group in matrix A segment 1750. The first group of eight elements 1770A comprises eight matrix elements in columns 0-7 of matrix A segment 1750. The second group of eight elements 1770B comprises eight matrix elements in columns 8-15. The third group of eight elements 1770C comprises eight matrix elements in columns 16-23. The fourth group of eight elements 1770D comprises eight matrix elements in columns 24-31.

[0200] exist Figure 17A In the example, the MMA calculation is performed by 32 threads (e.g., a single thread bundle). Each element is a 4-bit value (e.g., E2M1), so each group of eight elements 1770 is 32 bits in size. Each group of eight elements 1770 is processed by one of the threads performing the MMA calculation. The four groups per row are processed by four corresponding threads, referred to as thread quadruples 1752. For example, thread quadruples 1752A, which processes the first row of matrix A fragment 1750 in the MMA calculation, includes threads T0, T1, T2, and T3. To perform matrix multiplication and accumulation, thread T0 processes columns 0-7, thread T1 processes columns 8-15, thread T2 processes columns 16-23, and thread T3 processes columns 24-31. Thread quadruples 1752B processes the second row, and so on, processing rows three through eight. The eight rows are processed by thread quadruples 1752H. Since each row of matrix A segment 1750 is processed by 4 threads, and matrix A segment 1750 has eight rows, 32 threads (e.g., a thread bundle) are used to perform MMA operations on the 8x32 matrix A segment 1750 in parallel. Furthermore, since matrix A segment 1750 is one-quarter of the input matrix A 732A in Example 16x64, the complete input matrix A 732A can be processed in parallel using four thread bundles.

[0201] For each thread, the thread ID and associated fragment IDs a0,…,a7 are shown in the form Tx:{a0,…,a7}, referred to herein as thread ID and fragment ID 1754, where fragment IDs a0,…,a7 indicate the eight columns that the thread is to process. For example, thread ID and fragment ID 1754A includes the thread ID T28 of a thread (which has ID=28) used to perform MMA calculations on the first set of eight columns in row 7, and a sequence of eight fragment IDs corresponding to the first set of eight columns. As another example, thread ID and fragment ID 1754B includes the thread ID T29 of a thread (which uses the thread to perform MMA calculations on the first set of eight columns in row 7), and a sequence of eight fragment IDs corresponding to the second set of eight columns.

[0202] The fragment ID sequence {a0,…,a7} comprises eight fragment IDs a0…a7. Each fragment ID identifies a column of the matrix (e.g., a column of the input matrix A 732A). The combination of thread ID and fragment ID corresponds to an element of matrix A. For example, thread ID T0 and fragment ID a7 correspond to the element in row #0 and column #7 of matrix A. As another example, in an MMA operation of size M=16, N=8, K=64 (MMA shape is m16n8k64), the input matrix A 732A is 16x64, the input matrix B 732B is 64x8, and the resulting matrix D 138 obtained by multiplying A and B is 16x8.

[0203] In at least one embodiment, block scaling is performed before matrix multiplication and accumulation operations. Thread quad 1752A includes four threads T0-T3, where two threads contribute (e.g., retrieve) scaling factors from scaling data A 734A to be multiplied by the elements of row 0 of matrix segment A 1750. For example, two selected threads in thread quad 1752A may retrieve scaling factors from scaling data A 734A and include the retrieved scaling factors in scaling matrix A 744A to be multiplied by the elements of row 0 of the column block of matrix A. Furthermore, thread quad 1752B includes four threads T4-T7, where two threads are selected to contribute scaling factors from scaling data A 734A to be multiplied by the elements of row 1 of matrix segment A 1750. Furthermore, the thread quadruple 1752H includes four threads T28-T31, two of which are selected to contribute the scaling factor of the scaled data A 734A for multiplication with the element of row 7 of the matrix A fragment 1750. In at least one embodiment, the scaling factor of the scaled data B 734B is contributed by a single thread because the number of scaling factor multiplications performed for matrix B (e.g., it is 64x8) is only half that for matrix A (e.g., it is 16x64, in the m16n8k64 example shape).

[0204] The thread pairs used to contribute the scaling factor A 734A are selected using a thread-id-a selector, which is part of the scaling factor A byte and the thread selector 834A, as referenced. Figure 17B As described, the thread used to contribute the scaling factor B 734B is selected using a thread-id-b selector, which is part of the scaling factor B byte and the thread selector 834B, as described in the reference. Figure 18 As described.

[0205] Figure 17B The illustration depicts a pair of threads contributing scaling factors from scaling data A 1734A to form a scaling matrix A 1744A according to at least one embodiment. In at least one embodiment, the scaling data A 1734A and the scaling matrix A 1744A respectively correspond to... Figure 7 The scaling data 734A and scaling matrix A 744A are included. In at least one embodiment, one or more portions of the scaling matrix A 744A, which contains scaling factors, are contributed by two threads in the thread quadruple 1752 (e.g., loaded from scaling data A 734A). To perform the block scaling operation, scaling factors SF_A and / or scaling factors SF_B are retrieved from scaling data A 734A and / or scaling data B 734B, respectively. For example, each of the scaling factors SF_A and SF_B can be an 8-bit value (e.g., UE8M0 or UE4M3).

[0206] In the block scaling operation, the elements of matrix A are multiplied by the corresponding scaling factor SF_A selected from scaling matrix A 1744A to compute the scaled elements of matrix A. Similarly, the elements of matrix B are multiplied by the corresponding scaling factor SF_B selected from scaling matrix B 744B to compute the scaled elements of matrix B. The scaled elements can be stored in any suitable register 760 and / or memory location (e.g., tensor memory 230 or shared memory 250). The scaled elements correspond to the elements about... Figure 11 The scaled A value 1156A and / or scaled B value 1156B are described. The scaled elements are used as input for matrix multiplication and accumulation operations. Matrix multiplication and accumulation operations are performed by one or more threads processing matrix A segment 1750. Specific columns or rows of scaled data 734A, 734B to be retrieved from register 760, and specific threads contributing scaling factors SF_A and SF_B from scaled data B 734B and / or performing other operations to contribute scaling factors, are specified by the scaling factor A byte and thread selector 834A and the scaling factor B byte and thread selector 834B operands, respectively.

[0207] exist Figure 17BIn the example shown, two threads in thread quadruple 1752 are used to retrieve the scaling factor SF_A for matrix A, and one thread in thread quadruple 1752 is used to retrieve the scaling factor SF_B for matrix B. In one example, the scaling factor SF_A is arranged in a scaling matrix A 1744A with 16 rows and 4 columns (but 1, 2, 4, or other numbers of columns can also be used for block scaling, depending on whether the scaling vector size qualifier 842 is 1X, 2X, 4X, or some other value). For example, for each row (e.g., 16 rows) of scaling matrix A 244A, the scaling factor SF_A in columns 1, 2, or 4 of scaling matrix A 744A will be multiplied by the corresponding element of the corresponding row of a matrix such as input matrix A 732A. The specific column in each row of scaling matrix A 244A used to provide the scaling factor SF_A is selected by the scaling factor A byte and the byte-id-a portion of the thread selector 834A operand.

[0208] like Figure 18 As shown, the scaling factor SF_B is arranged in a scaling matrix B 1844B with 4 rows and 8 columns (which corresponds to scaling matrix B 744B). In at least one embodiment, block scaling is performed using rows 1, 2, or 4 of the 4 rows, depending on whether the scaling vector size qualifier 342 is 1X, 2X, or 4X. For example, for each column (e.g., 8 columns) of scaling matrix B 244B, the scaling factor SF_B in rows 1, 2, or 4 of scaling matrix B 244B is multiplied by the corresponding element in the corresponding row of the input matrix B 252B. The specific column in each row of scaling matrix A 244A used to provide the scaling factor SF_B is selected by the byte-id-b portion of the scaling factor selector B 424B operand.

[0209] See you again Figure 17BThe first half of scaling matrix A1780 (which includes rows 0-7 of scaling matrix A744A) is contributed from scaling data A734A by thread 0 (e.g., T0) in thread quad 1752 (if the thread-id-a portion of scaling factor selector A424A is 0); or by thread 2 (e.g., T2) in thread quad 1752 (if the thread-id-a portion is 1). In other words, if scaling factor selector A424A specifies thread-id-a as 0, then the next two threads T0 and T1 in thread quad 1742 contribute scaling factor SF_A to both halves of scaling matrix A1744A. If scaling factor selector A424A specifies thread-id-a as 1, then the previous two threads T2 and T3 in thread quad 1742 contribute scaling factor SF_A to both halves of scaling matrix A1744A. As another example, if thread-id-a is 0, then thread T0 contributes (e.g., loads) rows 0-7 of the scaling matrix 1744A (which has four 8-bit columns) from the scaling data A 234A, and thread T1 contributes rows 8-15 from the scaling data A 234A. As yet another example, if thread-id-a is 1, then thread T2 contributes rows 0-7 of the scaling matrix 1744A from the scaling data A 234A, and thread T3 contributes rows 8-15 from the scaling data A 234A.

[0210] The latter half of scaling matrix A 1782 (which comprises rows 8-15 of scaling matrix A 244A) is contributed by thread 1 (e.g., T1) of thread quad 1752 (if the thread-id-a portion of scaling factor selector A 424A is 0); or by thread 3 (e.g., T3) of thread quad 1752 (if the thread-id-a portion of scaling factor selector A 424A is 1). Although scaling matrix A 1744A is described as having 1, 2, or 4 columns in the examples described herein, in other examples, scaling matrix A 1744A can have any suitable number of columns.

[0211] Figure 18 The illustration depicts a thread that contributes a scaling factor from scaling data B 1834B to form a scaling matrix B 1844B according to at least one embodiment. The scaling data B 1834B includes scaling data to be used for scaling the input matrix B 252B. In at least one embodiment, the scaling data B 1834B and the scaling matrix B 1844B respectively correspond to... Figure 7The scaling data B 734 and scaling matrix B 744B are specified in the data set. Scaling matrix B 1844B is contributed from scaling data B 1834B by one thread in thread quad 1752 (e.g., selected, loaded, or otherwise accessed). In at least one embodiment, scaling matrix B 644B is otherwise accessed from scaling data B 634B by one or more threads in thread quad 1752. Scaling matrix B 1844B contains a scaling factor SF_B. If the scaling factor B byte and the thread-id-b of the thread selector 834B are 0, scaling matrix B 1844B is contributed from thread quad 1752 by thread 0 (e.g., T0); if thread-id-b is 1, it is contributed from thread quad 1752 by thread T1; if thread-id-b is 2, it is contributed from thread quad 1752 by thread T2; and if thread-id-b is 3, it is contributed from thread quad 1752 by thread T3.

[0212] Figure 19 An example of a system 1900 according to at least one embodiment is illustrated. The system 1900 may include software and hardware for executing instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any other operations described herein. The system 1900 may include a storage device 1902 and a processor 1908. The storage device 1902 may include, for example, memory, cache, or other storage devices further described below. The storage device 1902 may be separate from the processor 1908, or the storage device 1902 may be included within the processor 1908 (e.g., in storage device 1912). In at least one embodiment, software program 1904 and / or software instructions 1906 may be stored in memory, cache, or other storage devices and provided to the processor 1908 to cause one or more circuits of the processor 1908 to perform the operations described herein. In at least one embodiment, software program 1904 and / or software library (or instructions) 1906 may be integrated into one or more circuits of one or more processors 1908. Software program 1904, which can be used to perform any of the operations described herein, may be stored on storage device 1902.

[0213] In at least one embodiment, unless the context explicitly specifies otherwise, in any implementation described herein, "module" means any combination of software logic, firmware logic, hardware logic, and / or circuitry configured to provide the functionality described herein. In at least one embodiment, "software" is embodied in software packages, code, and / or instruction sets or instructions, while "hardware" as used in any implementation described herein includes (alone or in any combination) such as hardwired circuitry systems, programmable circuitry systems, state machine circuitry systems, fixed-function circuitry systems, execution unit circuitry systems, and / or firmware storing instructions executed by the programmable circuitry system. In at least one embodiment, a module may be embodied collectively or individually as a circuitry system forming part of a larger system, such as an integrated circuit (IC), a system-on-a-chip (SoC), etc. In at least one embodiment, a module is combined with any suitable processing unit and / or combination of processing units (e.g., one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof, including those further described herein) to perform one or more processes.

[0214] In at least one embodiment, software program 1904 may include a collection of software code, commands, instructions, or other text sequences for instructing a computing device to perform one or more computational operations and / or invoke one or more other instruction sets, such as APIs or API functions or Instruction Set Architecture (ISA) level instructions, for execution or otherwise implementation. Instructions (e.g., hardware instructions) or microcode may include ISA-level instructions, which may include native ISA instructions or non-native ISA instructions. Software program 1904 and / or software instructions 1906 may be distributed across multiple processors that communicate via buses, networks, writes to shared memory, and / or any suitable communication process (such as the communication processes described herein).

[0215] Processor 1908 may include any number of processors and any suitable processing units and / or combinations of processing units, such as, but not limited to, a central processing unit (“CPU”), a graphics processing unit (“GPU”), 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, those described herein. Figure 21-32The processor 1908 is a processor in the memory. In at least one embodiment, the processor 1908 may retrieve or fetch instructions (e.g., one or more API and / or ISA instructions) from the memory device 1902 using, for example, instruction fetch 1916 (e.g., for the instruction fetch phase). Instructions may include instructions for performing instructions to perform one or more matrix multiplication-accumulation (MMA) operations to multiply a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors. In at least one embodiment, the processor 1908 may include a memory device 1912 and an instruction queue 1910 for storing and queuing instructions fetched from the memory device 1902. In at least one embodiment, the fetched instructions may be decoded by decode 1918 to determine what operation the processor 1908 should perform (e.g., in the instruction decode phase). In at least one embodiment, the processor 1908 may fetch additional operands (data) that can be used for the instructions, and the operands may be stored, for example, in a register or memory device 1912. In at least one embodiment, micro-operation 1920 can perform operations on data stored in one or more registers or storage devices 1912. For example, each step of an instruction fetched by processor 1908 can be broken down during execution so that processor 1908 can execute the instruction step-by-step through a series of micro-operations 1920. In at least one embodiment, program counter (PC) 1914 can hold the address of the next instruction and can be updated to point to the next instruction to be executed by processor 1908.

[0216] In at least one embodiment, processor 1908 may execute instructions (e.g., during the execution phase). For example, processor 1908 may perform operations specified by the instructions, such as arithmetic operations, logical operations, or data transfers. In at least one embodiment, computing unit 1922 may execute instructions to perform any of the operations described herein. In at least one embodiment, computing unit may include an arithmetic logic unit (ALU) 1924, which can be used to perform arithmetic and logical operations. In at least one embodiment, computing unit may include a floating-point unit (FPU) 1926, which can be used to perform floating-point calculations. In at least one embodiment, other circuitry 1928 may be used to perform other operations, such as vector operations and / or scalar operations. In at least one embodiment, accelerator 1930 may include one or more matrix multiplication accelerators, one or more parallel processing units (PPUs) (e.g., GPUs), or any other accelerators or processors further described herein. In at least one embodiment, software program 1904 may utilize one or more API and / or ISA instructions to perform various computational operations, such as matrix multiplication, arithmetic operations, or any other computational operations further described herein, with accelerator 1930. In at least one embodiment, one or more computational operations using accelerator 1930 may include at least one or more sets of computational operations that will be accelerated by execution at least in part by accelerator 1930, including execution instructions to perform one or more matrix multiplication-accumulation (MMA) operations to multiply a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors.

[0217] In at least one embodiment, system 1900 may be used to execute one or more instructions, including instructions such as Figure 1-18 The functions or operations described herein. In at least one embodiment, the system 1900, including one or more processors, causes one or more circuits to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, thereby multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, and / or otherwise performing the operations described herein. In at least one embodiment, the system 1900 is included Figure 1-18 The system shown, and / or otherwise includes Figure 1-18The system shown is used to cause one or more circuits to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, thereby multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, and / or otherwise performing the operations described herein. In at least one embodiment, system 1900 includes Figure 20-33B One or more pieces of hardware, for example, are used to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, multiply a first input matrix by a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, and / or otherwise perform the operations described herein.

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

[0219] Rack 2002 and substrate 2004 may include subsystems, modules, add-in cards, and other semiconductor components. Substrate 2004 may include one or more computing units 2006, each computing unit 2006 including one or more processors 2008, one or more memories 2010, and an interface controller 2012. Computing units 2006 may include any number of processors, such as, but not limited to, central processing units (“CPU”), 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, those described herein. Figure 21-33B The processor in the computing unit 2006 may include one or more memory storage devices 2010 (e.g., dynamic read-only memory, solid-state storage devices, or disk drives), as well as network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power supply modules, and cooling modules, etc. One or more computing units 2006 may be a server having one or more of the aforementioned computing resources.

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

[0221] Data center 2000 may include any one or any combination of the framework layer 2020, software layer 2030, and application layer 2040. For example... Figure 20As shown, the framework layer 2020 includes a job scheduler 2022, a configuration manager 2024, a resource manager 2026, and a distributed file system 2028. The framework layer 2020 may include a framework for supporting software 2032 of the software layer 2030 and / or one or more applications 2042 of the application layer 2040. Software 2032 or application 2042 may respectively include web-based service software or applications, such as, but not limited to, software or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. The framework layer 2020 may be a type of free and open-source software web application framework, such as, but not limited to, Apache Spark™ (hereinafter “Spark”), which can leverage the distributed file system 2028 for large-scale data processing (e.g., “big data”). The job scheduler 2022 may include Spark drivers, which facilitate the scheduling of workloads supported by the various layers of the data center 2000. Configuration Manager 2024 can configure different layers, such as, but not limited to, the software layer 2030 and the framework layer 2020 (which includes Spark and the distributed file system 2028 for supporting large-scale data processing). Resource Manager 2026 can manage clustered or grouped compute units 2006 mapped to or allocated to support the distributed file system 2028 and the job scheduler 2022. Resource Manager 2026 can coordinate with Resource Coordinator 2014 to manage these mapped or allocated compute resources.

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

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

[0224] Any of Configuration Manager 2024, Resource Manager 2026, and Resource Coordinator 2014 can implement any number and type of self-modification actions based on any amount and type of data obtained in any technically feasible manner. Self-modification actions can reduce the burden on Data Center 2000 data center operators to make potentially incorrect configuration decisions and may avoid underutilized and / or poorly performing parts of the data center.

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

[0226] Data Center 2000 can use CPUs, application-specific integrated circuits (ASICs), GPUs, FPGAs, or other hardware (e.g., Figure 21-33B The embodiments described herein can be used to perform some or all of the processes and techniques described elsewhere, such as, but not limited to, training and / or inference using the resources described above. Furthermore, one or more of the software and / or hardware resources described above can be configured as a service to allow a user to train or perform information inference, such as, but not limited to, image recognition, speech recognition, or other artificial intelligence services.

[0227] In at least one embodiment, processor 2008 may include one of the processors described below and / or include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. In at least one embodiment, processor 2008 is configured by software 2032 to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. Data center 2000 may use logic, CPU, application-specific integrated circuit (ASIC), GPU, FPGA, or other hardware (e.g., Figure 21-33B (In the embodiments) perform any of the operations described above or elsewhere in this document.

[0228] processor The following figures illustrate, but are not limited to, example processors and processing systems that can be used to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, thereby multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing some or all of the processes, operations, and / or techniques described elsewhere herein. The example processors and processing systems can be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, thereby multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. Processors and processing systems may include logic, a central processing unit (CPU), an application-specific integrated circuit (ASIC), a graphics processing unit (GPU), a field-programmable gate array (FPGA), an XPU (i.e., any computing architecture best suited to the needs of the application), or other hardware (e.g., Figure 21-33B The embodiments described herein are used to perform any of the operations described above, below, or elsewhere herein. The processor and / or processing system described herein may include one or more circuits that can be configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein. As used herein, one or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein. Figure 38A and Figure 38BThe illustration depicts logic 3815 according to at least one embodiment, 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. For example, logic can refer to any combination of software logic, hardware logic, and / or firmware logic that provides the functionality and / or operations described herein, wherein the logic can be collectively or individually embodied as part of a circuit system forming a larger system, such as an integrated circuit (IC), application-specific integrated circuit (ASIC), field-programmable gate array (FPGA), system-on-a-chip (SoC), or one or more processors (e.g., CPU, GPU).

[0229] Figure 21 A processor according to at least one embodiment is illustrated, which is a system-on-a-chip (SOC) 2100 (which may be referred to as a system-on-a-chip, superchip, or other names). SOC 2100 may include processor complex 2110 and processor complex 2140. SOC 2100 may include any number of processor complexes 2110 and / or processor complexes 2140, which may include any number of processors described herein in any combination, such as, but not limited to, processors described herein in any combination. Figure 21-33B The processor in the system. For example, processor 2110 may include a central processing unit (CPU), and processor 2140 may include a graphics processor. Alternatively, processor 2110 may include a graphics processor, and processor 2140 may include a graphics processor. The SOC 2100 may include any number of display controllers 2192, any number of multimedia engines 2194, any number of I / O interfaces 2170, any number of memory controllers 2180, and any number of fabrics 2160 in any combination. For ease of explanation, this document uses reference numbers identifying objects and bracket numbers identifying instances (if needed) to denote multiple instances of similar objects. The SOC 2100 may include a processor from Broadcom Inc., Palo Alto, California.

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

[0231] Processor complex 2110 may include cores 2120(1)-2120(4) and cache (e.g., L3 cache) 2130 for storing information for performing the operations described herein. Processor complex 2110 may include any number of cores 2120 in any combination and any number and type of cache. Cores 2120 may be configured to execute instructions of a particular instruction set architecture (“ISA”) to perform some or all of the operations described herein. Each core 2120 may include a CPU core. Cores 2120(1)-2120(4) may be referred to as compute units or arithmetic units. SOC 2100 may include any number of processor complexes 2110, architecture 2160, I / O interface 2170, and memory controller 2180.

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

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

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

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

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

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

[0238] Display controller 2192 can display images on one or more display devices, such as, but not limited to, liquid crystal displays (“LCD” devices. Multimedia engine 2194 can include any number and type of circuitry related to multimedia, such as, but not limited to, video decoders, video encoders, image signal processors, etc. Memory controller 2180 can facilitate data transfer between SOC 2100 and unified system memory 2190. Processor complex 2110 and processor complex 2140 can share unified system memory 2190. Unified system memory 2190 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 2190 can include 3D stacked memory, including but not limited to high bandwidth memory (HBM), HBM2e, or HDM3.

[0239] The SOC 2100 can implement a memory subsystem including any number and type of memory controllers 2180 and memory devices (e.g., shared memory 2154), which may be dedicated to a single component or shared among multiple components to perform any of the operations described herein. The SOC 2100 can implement a cache subsystem including one or more cache memories (e.g., L2 cache 2128, L3 cache 2130, and L2 cache 2142), each cache memory may be dedicated to any number of components (e.g., core 2120, core complex 2110, SIMD unit 2152, compute unit 2150, and processor complex 2140), or may be shared among any number of components (e.g., core 2120, core complex 2110, SIMD unit 2152, compute unit 2150, and processor complex 2140).

[0240] In at least one embodiment, the SOC 2100 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

[0241] Figure 22A A parallel processor 2200 according to at least one embodiment is illustrated. The parallel processor 2200 may be implemented using one or more circuits and may be referred to as a programmable processor (e.g., CPU and / or GPU), logic, application-specific integrated circuit (ASIC), field-programmable gate array (FPGA), or other hardware (e.g., Figure 21-33B (The embodiments in the document) are used to perform any of the operations described above or elsewhere herein.

[0242] Parallel processor 2200 may include parallel processing unit 2202 for performing any of the operations described above or elsewhere herein. Parallel processing unit 2202 may include I / O unit 2204 that enables communication with other devices, including other instances of parallel processing unit 2202. I / O unit 2204 may be directly connected to other devices. I / O unit 2204 may be connected to other devices via the use of a hub or switch interface, such as, but not limited to, memory hub 2205. The connection between memory hub 2205 and I / O unit 2204 may form a communication link 2213. I / O unit 2204 may be connected to host interface 2206 and memory crossbar switch 2216, wherein host interface 2206 receives commands directed to perform processing operations, and memory crossbar switch 2216 receives commands directed to perform memory operations.

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

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

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

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

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

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

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

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

[0251] Any cluster 2214A-2214N in the processing cluster array 2212 can process data to be written to any memory cell 2224A-2224N within the parallel processor memory 2222. The memory crossbar switch 2216 can be configured to transfer the output of each cluster 2214A-2214N to any partition cell 2220A-2220N, or to another cluster 2214A-2214N on which additional processing operations can be performed. Each cluster 2214A-2214N can communicate with the memory interface 2218 via the memory crossbar switch 2216 to read from or write to various external memory devices. The memory crossbar switch 2216 can be connected to the memory interface 2218 to communicate with the I / O unit 2204, or to a local instance of the parallel processor memory 2222, enabling processing units within different processing clusters 2214A-2214N to communicate with system memory or other memory local to the non-parallel processing unit 2202. The memory crossbar switch 2216 can use virtual channels to separate traffic flows between clusters 2214A-2214N and partition units 2220A-2220N.

[0252] Multiple instances of the parallel processing unit 2202 can be mounted on a single add-in card, or multiple add-in cards can be interconnected. Even if different instances of the parallel processing unit 2202 have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences, these different instances can be configured to interoperate. For example, some instances of the parallel processing unit 2202 may include higher-precision floating-point units relative to other instances. Systems including one or more instances of the parallel processing unit 2202 or the parallel processor 2200 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.

[0253] Figure 22A It also includes a block diagram of a partitioning unit 2220 according to at least one embodiment. The partitioning unit 2220 is... Figure 22AAn example of one of the partition units 2220A-2220N in the parallel processor memory. Partition unit 2220 may include L2 cache 2221, frame buffer interface 2225, and ROP 2226 (raster operation unit). L2 cache 2221 may be a read / write cache configured to perform load and store operations received from memory crossbar switch 2216 and ROP 2226. Read misses and urgent write-back requests may be output from L2 cache 2221 to frame buffer interface 2225 for processing. Updates may also be sent to the frame buffer via frame buffer interface 2225 for processing. Frame buffer interface 2225 may interface with one of the memory cells in the parallel processor memory, such as, but not limited to, Figure 22A The memory cells 2224A-2224N (shown as 2224) are located in the parallel processor memory 2222.

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

[0255] ROP 2226 can be included in each processing cluster (e.g., Figure 22A The data is stored within clusters 2214A-2214N, not within partition units 2220. Read and write requests for pixel data (not pixel fragment data) can be transferred via memory crossbar switch 2216. Processed graphics data can be displayed on a monitor and routed for further processing by the processor, or routed to... Figure 22A One of the processing entities within the parallel processor 2200 in the system is further processed.

[0256] In at least one embodiment, the parallel processor 2200 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

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

[0258] The operation of cluster 2214 can be controlled via pipeline manager 2232, which distributes processing tasks to SIMT parallel processors. Pipeline manager 2232 can... Figure 22A The scheduler 2210 receives instructions and manages the execution of these instructions via the graphics multiprocessor 2234 and / or texture unit 2236. The graphics multiprocessor 2234 may be an example instance of a SIMT parallel processor. However, the processing cluster 2214 may include various types of SIMT parallel processors with different architectures. The processing cluster 2214 may include one or more instances of the graphics multiprocessor 2234. The graphics multiprocessor 2234 can process data and can use the data cross switch 2240 to distribute the processed data to one of several possible destinations, including other shader units. The pipeline manager 2232 can facilitate the distribution of processed data by specifying the destination of the processed data to be distributed via the data cross switch 2240.

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

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

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

[0262] Each processing cluster 2214 may include an MMU 2245 (Memory Management Unit), which can be configured to map virtual addresses to physical addresses. One or more instances of the MMU 2245 may reside in... Figure 22A The MMU 2245 is located within the memory interface 2218. It may include a set of page table entries (PTEs) for mapping virtual addresses to physical addresses on tiles, and optional cache line indexes. The MMU 2245 may include address translation lookup buffers (TLBs) or caches that may reside within the graphics multiprocessor 2234 or L1 2248 cache or processing cluster 2214. Physical addresses can be processed to distribute surface data access locally, allowing for efficient request interleaving between partition units. The cache line indexes can be used to determine whether a request for a cache line is a hit or a miss.

[0263] Processing cluster 2214 can be configured such that each graphics multiprocessor 2234 is coupled to a texture unit 2236 to perform texture mapping operations, such as determining texture sample locations, reading texture data, and filtering texture data. Texture data can be read from an internal texture L1 cache (not shown) or an L1 cache within the graphics multiprocessor 2234, and can be retrieved as needed from an L2 cache, local parallel processor memory, or system memory. Each graphics multiprocessor 2234 can output processed tasks to a data crossbar switch 2240 to provide the processed tasks to another processing cluster 2214 for further processing, or store the processed tasks in an L2 cache, local parallel processor memory, or system memory via a memory crossbar switch 2216. Pre-ROP 2242 (pre-raster operation unit) can be configured to receive data from the graphics multiprocessor 2234 and direct the data to ROP units, which can be associated with partitioning units (e.g., ...) described herein. Figure 22A The PreROP 2242 unit is located together with partition units 2220A-2220N. The PreROP 2242 unit can perform color blending optimization, organize pixel color data, and perform address translation.

[0264] In at least one embodiment, the processing cluster 2214 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

[0265] Figure 22C A graphics multiprocessor 2234 according to at least one embodiment is illustrated, for example, to perform any of the operations described above or elsewhere herein. The graphics multiprocessor 2234 may be coupled to a pipeline manager 2232 of a processing cluster 2214. The graphics multiprocessor 2234 may include an execution pipeline including, but not limited to, an instruction cache 2252 (e.g., which may store instructions, such as, but not limited to, compiled API instructions), instruction units 2254, address mapping units 2256, register files 2258, one or more general-purpose graphics processing unit (GPGPU) cores 2262, and one or more load / store units 2266, one or more of which may perform load / store operations to load / store instructions corresponding to the execution operations. The GPGPU cores 2262 and load / store units 2266 may be coupled to cache memory 2272 and shared memory 2270 via a memory and cache interconnect 2268. The GPGPU cores 2262 may be part of a SoC, such as, but not limited to, [other components]. Figure 21 It is part of the integrated circuit 2100.

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

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

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

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

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

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

[0272] In at least one embodiment, the graphics multiprocessor 2234 may include one or more circuitry configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuitry may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

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

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

[0275] The GPU 2306 can include any number of graphics engines, such as, but not limited to, the Intel® Arc™ Graphics Engine (Xe LPG) with 8 Xe cores (up to 128 execution units or EUs). Figure 23 As shown, GPU 2306 may include vector engine 2310 and matrix engine 2312, which, for example, can run FP, INT, and matrix operation tasks simultaneously, individually, or in batches. GPU 2306 may include load / store unit 2314, as well as other memories, such as, but not limited to, instruction cache (I$) 2316 and L1 cache / subsystem local memory (SLM) 2318, which may, for example, store instructions for performing any of the operations described above or elsewhere herein.

[0276] The NPU 2304 may include one or more Intel® AI Boost built-in Neural Processing Units (NPUs). The NPU 2304 can be enumerated as an integrated PCIe device to the host processor. The NPU 2304 may include one or more (e.g., two) Neural Computation Engine (NCE) tiles 2330. Each tile may be configured with any combination of, but not limited to, the following: (e.g., 2000) Multiply-Accumulate (MAC) engines 2334, a post-processing engine (not shown), an AI DSP processor (not shown), and memory per tile (2MB dedicated SRAM), such as... Figure 23 As shown. For general computing needs, the neural computing engine 2330 may include a disturbance pipeline 2332, an activation function (AF) 2336, a data transformation 2338, a load / store 2340, and a streaming hybrid architecture vector engine (SHAVE) 2328 for high-performance parallel computing, which may include a DMA (Direct Memory Access) engine 2324 for transporting data between system memory DRAM (Dynamic Random Access Memory) 2326 and a software-managed cache. The built-in device MMU (Memory Management Unit) 2322, plus the IOMMU (Input-Output Memory Management Unit) (not shown), can support multiple concurrent hardware contexts and provide secure isolation between execution contexts according to the MCDM (Microsoft Computing Driver Model) architecture. The processor 2300 may also include a media unit (not shown), which may be included on or separate from the XCD or other components of the processor 2300 to enable video playback and video processing of compressed or uncompressed data, such as hardware-accelerated decoding support for HEVC, AV1, VP9, ​​and AVC, and hardware-accelerated encoding support for HEVC, VP9, ​​and AVC.

[0277] The Intel® Thread Director (which includes firmware built into the processor 2300) prioritizes and manages workload distribution, sending tasks to optimized cores. For example, the Thread Director can tie P-cores, E-cores, and / or LP-E-cores (as described above) together with task scheduling capabilities and the ability to send less demanding tasks to E-cores or LP-E-cores. Intel® Deep Learning Acceleration (Intel® DLBoost) (not shown) provides 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 can be optimized using the OpenVINO™ toolkit and oneAPI to accelerate INT8 inference. For example, a software stack as described elsewhere in this document can be used to implement AI inference using the OpenVINO™ toolkit. The processor 2300 can be configured to execute applications, such as, but not limited to, CUDA programs.

[0278] In at least one embodiment, processor 2300 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

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

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

[0281] In at least one embodiment, the processor 2300 or the Qualcomm AI engine 2300 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

[0282] Figure 24AA processor 2400 according to at least one embodiment is illustrated. The processor 2400 may include a Scalable family processor from Intel Corporation, Santa Clara, California, or other processors that share at least some of the components described herein. The processor 2400 may include one or more cores 2412(1)-2412(N) capable of performing the operations described elsewhere herein, where N is any integer greater than 1. Cores 2412(1)-2412(N) may be interconnected using ring and / or mesh interconnects. Utilizing a mesh interconnect architecture, arrays of vertical and horizontal communication paths may allow traversal from one core to another 2412(1)-2412(N) via the shortest path (jumping to the correct row along the vertical path and to the correct column along the horizontal path). For the mesh interconnect, a die may accommodate cores 2412(1)-2412(N) and may include a Converging Mesh Stop Point (CMS) grid that may be associated with cores 2412(1)-2412(N) (e.g., 1:1). Each core may be associated with a low-level cache (LLC) slice 2414(1)-2414(N), or cores 2412(1)-2412(N) may share a cache, such as a low-level cache. LLC 2414(1)-2414(N) may be inclusive or non-inclusive (having blocks not present in the high-level cache) by merging blocks from a higher-level cache (e.g., L2 cache). Each core and LLC slice may include a caching and home agent (CHA) (not shown) that maintains cache coherency by providing resource scalability via mesh interconnect, thereby enabling cache coherency functionality of Intel® HyperPath Interconnect (Intel® UPI 2416). UPI 2416 provides a coherent interconnect for scalable systems and allows multiple processors to share a single shared address space via links, such as, but not limited to, two or three UPI links per processor.

[0283] The processor 2400 may also include a system agent 2410, which may house and / or perform various functions, such as, but not limited to, memory management, display functions, and / or input / output (I / O) functions. For example, the processor 2400 may include one or more integrated memory controllers (IMCs) 2408. The IMC 2408 may control and manage memory, such as, but not limited to, different memory types, such as DDR RAM, such as DDR4, or other memory described elsewhere herein. The system agent 2410 may include a display controller (not shown) for supporting one or more displays. The system agent 2410 may also integrate a PCIe 2404 (e.g., up to 20 PCIe lanes), which may, for example, be connected to an external dedicated graphics connector via a DMI bus (e.g., Intel's DMI 3.0 bus) 2406. The system agent 2410 may include an image processing unit (IPU) (not shown) that integrates an on-die image signal processor (ISP). Structure 2402 provides scalability to connect to other nodes (e.g., processors, such as processor 2400) and can be used, for example, with Cornelis Networks (an element of the Intel® Scalable System Framework), which provides performance for high-performance computing (HPC) workloads and the ability to scale to tens of thousands of nodes.

[0284] Figure 24B Components within a core 2412 according to at least one embodiment are illustrated. The core 2412 may include a front-end 2418, a back-end or execution engine 2432, and a memory subsystem 2442. The front-end 2418 may provide operations (e.g., operations described elsewhere herein) to the execution engine 2432 by decoding instructions stored in memory. For example, the front-end 2418 may include micro-operation (µOps) cache paths and / or traditional paths, and a branch prediction unit 2421 capable of determining path instructions. A traditional path for instructions may include fetching variable-length (e.g., x86) instructions from an L1 instruction cache 2420 and instruction fetch and pre-decode 2422, queuing these instructions into an instruction queue 2424, and decoding the instructions into µOps that can be provided to an allocation queue 2428 using a decoder 2426. Alternatively, the µOps cache path may include a cache that includes decoded µOps (µOps 2430) that can be sent to the allocation queue 2428. The allocation queue 2428 can act as an interface between the front end 2418 and the execution engine 2432, and can provide instructions to the execution engine 2432. For example, one or more APIs described herein can be compiled into instructions that can be stored, processed, and executed by the front end 2418 and the execution engine 2432, and stored in the memory subsystem 2442.

[0285] Execution engine 2432 can receive micro-operations into reordering buffer 2434, which can register allocations, renames, and decommissions of µOPs. µOPs can be sent from the reordering buffer to scheduler 2436, which can be connected to one or more different execution units 2438, which can be connected to address generation units (AGUs) 2440. Execution units 2438 can perform operations such as basic arithmetic logic unit (ALU) operations, multiplication, division, and / or more complex operations, such as, but not limited to, various vector operations. Scheduler 2436 can manage the queuing of µOPs for one or more execution units 2438 based on, for example, the operations that need to be performed.

[0286] The memory subsystem 2442 can handle load and store requests as well as sorting operations. For example, µOPs may be associated with memory accesses (such as load and store), and these µOPs can be sent through dedicated scheduler ports that can perform these memory operations. For example, store and load operations can be sent to load and store buffers 2444. The memory subsystem 2442 may also include shared or separate L1 data and instruction caches 2446, and an L2 cache 2448 that can be used and shared by the L1 data and instruction caches 2446. As described above for... Figure 24A Each core 2412 can be connected to a slice of a level 3 cache (e.g., LLC2414), which can be shared by all cores 2412.

[0287] In at least one embodiment, processor 2400 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

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

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

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

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

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

[0293] The AI ​​accelerator 2500's RDMA 2528 based on converged Ethernet enables scaling from a single node (i.e., from a single AI accelerator 2500 to hundreds or thousands of nodes or AI accelerators 2500). The network subsystem 2524 may include the Intel® Gaudi® Communications Library (IGCL), a master controller coordinating data movement, and a programmable scheduling mechanism that enables smooth engine activation while maintaining task dependencies. The accelerator network subsystem may include a Gigabit Ethernet NIC port 2526, a Layer 2 MAC (not shown), and the RDMA engine 2528. The AI ​​accelerator 2500 may include an aggregation engine for performing summation activities. All engines in the processor 2500 can run in parallel; for example, the MME 2508, TPC 2510, and NIC 2526 can all operate simultaneously. Dependencies may exist between operations running on different engines; for example, the output of one engine may be used as the input of another engine, and / or the MME, TPC, and NIC may be scheduled to run in parallel. Once one engine has finished its operation, another engine can be scheduled to begin working on the next operation (executed immediately once its input is ready).

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

[0295] In at least one embodiment, the AI ​​accelerator 2500 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above or elsewhere herein.

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

[0297] continue Figure 26For example, neuromorphic computing device 2605 may also include processor 2620 and system memory 2625 for implementing one or more components to manage and provide the functionality of neuromorphic computing device 2605. For instance, a system manager 2630 may be provided to manage the global attributes and operations of neuromorphic computing device 2605 (e.g., attributes affecting core network 2610, multiple cores in network 2610, interconnection of neuromorphic computing device 2605 with other devices, managing access to global system memory 2625, and other potential examples). In one example, system manager 2630 may manage the definition and configuration of specific routing tables for individual routers in network 2610, orchestration of network definitions and attributes to be applied to network 2610 (e.g., weights, attenuation rates, etc.), core synchronization and time multiplexing management, routing input to appropriate cores, and other potential functions.

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

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

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

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

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

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

[0304] In improved implementations of systems capable of supporting SNNs, for example, but not limited to... Figure 26 The example illustrates a very large-scale integrated circuit (VLSI) hardware device that can provide high-speed, reliable circuitry to implement SNNs (Spiritual Neural Networks) to model the information processing algorithms used by the brain, but in a more programmable way. For instance, while a biological brain can only perform a specific set of defined behaviors (a consequence of years of development), a neuromorphic processor device can provide the ability to rapidly reprogram all neural parameters. Therefore, a single neuromorphic processor can be used to implement a wider range of behaviors than a single slice of biological brain tissue. This distinction can be achieved by employing neuromorphic processors with neuromorphic designs that are distinctly different from those found in natural neural circuitry.

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

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

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

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

[0309] In at least one embodiment, the neuromorphic computing device 2605 may include one or more circuits configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or otherwise performing any of the operations described above or elsewhere herein. The one or more circuits may be software-configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations, such as multiplying a first input matrix with a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors, or performing any of the operations described above o...

Claims

1. One or more processors, comprising: A circuit system for executing instructions to perform one or more matrix multiplication-accumulation (MMA) operations to multiply a first input matrix by a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors stored in a tensor memory.

2. The processor of claim 1 or more, wherein the input of the instruction indicates that the elements of the second input matrix are stored in 4 bits of tensor memory.

3. The processor of claim 2 or more, wherein the elements of the second input matrix are stored in a container represented by a 4-bit tensor memory, wherein the container is a storage unit representing the elements of the second input matrix.

4. One or more processors according to claim 1, wherein the one or more corresponding scaling factors are retrieved from scaling data stored in tensor memory based at least in part on scaling data selection operands of the instructions, wherein the scaling data is indicated by scaling data operands of the instructions.

5. The processor of claim 1, wherein the one or more corresponding scaling factors are arranged as a plurality of vectors indicated by the input of the instruction, and one or more scaling factors of each vector are multiplied by one or more elements of each row or each column of a portion of the second input matrix corresponding to the scaling factor vector.

6. The processor of claim 1 or more, wherein the first input matrix comprises multiple parts, and multiplying the one or more elements of the second input matrix by the one or more corresponding scaling factors comprises: Multiply one or more elements of each part of the first input matrix by the one or more corresponding scaling factors.

7. One or more processors according to claim 6, wherein the number of portions of the plurality of portions of the first input matrix is ​​indicated by the input of the instructions.

8. A method comprising: Receive an instruction call for performing a matrix multiplication-accumulation MMA operation, the instruction call having operands indicating one or more scaling data storage locations in a first input matrix, a second input matrix, and a tensor memory for storing scaling data; In response to the instruction call, a scaling matrix is ​​selected from the one or more scaling data storage locations based on the scaling vector size input and the scaling data selector operand according to the instruction call; Identify one or more scaling vectors in the scaling matrix, indicated by the scaling vector size input; Generate multiple scaled values ​​corresponding to the elements of the first input matrix, wherein each scaled value is the product of the corresponding matrix element in the first input matrix and the scaled element in one or more identified scaling vectors; as well as The MMA operation is performed to generate multiple elements of the result matrix based at least in part on the multiple scaled values ​​corresponding to the elements of the first input matrix and one or more elements of the second input matrix.

9. The method of claim 8, wherein the input of the instruction call indicates that the elements of the first input matrix are stored in 4 bits of tensor memory.

10. The method of claim 9, wherein the elements of the first input matrix are stored in a container represented by a 4-bit tensor memory, wherein the container is a storage unit representing the elements of the first input matrix.

11. The method of claim 8, wherein the position of the scaling element in the one or more scaling vectors is determined at least in part based on the position of the corresponding matrix element in the first input matrix.

12. The method according to claim 8, further comprising: Identify one or more blocks of the first input matrix, wherein each block includes one or more first input vectors of the first input matrix, and the number of blocks is based at least in part on the scaling vector size input of the instruction.

13. The method of claim 8, wherein the one or more scaling vectors are identified in the scaling matrix based at least in part on the byte selector operands of the instruction.

14. The method of claim 8, further comprising: Generate one or more scaled input vectors, each scaled input vector including one or more scaled values, said one or more scaled values ​​being generated by calculating one or more products of one or more elements of the corresponding input vector of the first input matrix and one or more corresponding elements of the scaled vector of the scaling matrix.

15. The method of claim 14, wherein the corresponding input vector is located in a block of the first input matrix, and the scaling vector corresponds to the block of the first input matrix.

16. A system comprising: One or more processors are configured to execute instructions to perform one or more matrix multiplication-accumulation (MMA) operations to multiply a first input matrix by a scaled matrix generated by multiplying one or more elements of a second input matrix by one or more corresponding scaling factors stored in a tensor memory.

17. The system of claim 16, wherein the one or more corresponding scaling factors are retrieved from scaling data indicated by the scaling data operand of the instruction, at least in part based on the scaling data selection operand of the instruction.

18. The system of claim 16, wherein the input of the instruction indicates the size of each element in the second input matrix, and the size of each element is 4 bits, 6 bits, or 8 bits.

19. The system of claim 16, wherein the one or more scaling factors are arranged as a plurality of vectors indicated by the input of the instruction, and one or more scaling factors of each vector are multiplied by one or more elements of each row or each column of a portion of the second input matrix corresponding to the scaling factor vector.

20. The system of claim 16, wherein the first input matrix comprises a plurality of parts, and multiplying the one or more elements of the second input matrix by the one or more corresponding scaling factors comprises: Multiply one or more elements of each part of the first input matrix by the one or more corresponding scaling factors.

Images