Data compression for neural networks

Abstract

The present application provides data compression for neural networks. Systems and methods for generating representative values for a data set determine a first common value by first compressing a portion of values in the data set and a second common value by further compressing a subset of the portion of values. The representative value is generated by taking a difference between the first common value and the second common value, where the representative value corresponds to a mathematical relationship between the first and second common values and each value within the subset of the portion of values. The representative value requires less storage than the first and second common values.

Description

Background Art

[0001] Being able to compress and decompress data and further reduce the number of bits when storing compressed data to achieve the desired storage metrics can involve various technical challenges. Algorithms associated with determining compressed numerical representations of compressed data elements and further storing these compressed numerical representations (often represented by a large number of bits) can exceed or otherwise overwhelm the storage capacity associated with the storage device. Moreover, due to the size of these numerical representations, the computational costs associated with training the neural network can also be affected. Furthermore, the need for numerical representations with a large number of bits often leads to complex issues associated with transmitting data within a device or between different devices given bandwidth limitations and / or constraints in the computing environment. BRIEF DESCRIPTION OF THE DRAWINGS

[0002] Various techniques are described with reference to the accompanying drawings, in which:

[0003] Figure 1 According to an embodiment, a computing environment is described in which a Block Floating Point (BFP) encoder is executed to generate one or more compressed numerical representations;

[0004] Figure 2 A diagram of a matrix including a set of values ​​represented by an exponent, a sign, and a mantissa (e.g., according to a block floating point format) and another diagram of a matrix including a set of values ​​represented by an exponent, a packed numeric representation, a sign, and a mantissa (according to a block floating point format with a packed numeric representation) are illustrated according to an embodiment;

[0005] Figure 3 Diagrams illustrating exemplary matrices, sub-matrices (or blocks), sub-matrices (or quads) of blocks, and scalars according to embodiments;

[0006] Figure 4 A diagram is described in which public exponent values ​​for blocks and quads are determined according to an embodiment;

[0007] Figure 5 A diagram is described in which a compressed numerical representation (or indicator) corresponding to a mathematical relationship between values ​​in a block and in a quad is determined according to an embodiment;

[0008] Figure 6 An illustrative example of a process for generating a compressed numerical representation (or indicator) corresponding to a mathematical relationship between index values ​​in a data set is shown according to an embodiment;

[0009] Figure 7 Another illustrative example of a process for generating a compressed numerical representation (or indicator) corresponding to a mathematical relationship between index values ​​in a data set is shown in accordance with an embodiment;

[0010] Figure 8 Examples of parallel processing units ("PPUs") are described according to embodiments;

[0011] Figure 9 An example of a general processing cluster ("GPC") is described according to an embodiment;

[0012] Figure 10 According to the embodiment, an example of a memory partition unit is described;

[0013] Figure 11 An example of a streaming multiprocessor is described according to an embodiment; and

[0014] Figure 12

[0066] A computer system is described in accordance with the embodiments in which various embodiments may be implemented. DETAILED DESCRIPTION

[0015] In one embodiment, systems and methods implemented in accordance with the present disclosure are used to generate compressed numerical representations in a dataset that are simultaneously well-suited for representing activations, weights, and gradients in deep neural networks and that can be efficiently implemented in hardware. In particular, in one embodiment, the compressed numerical representation is suitable for efficient implementation of matrix multiplications, as these multiplications primarily constitute a large portion of the computational cost of training neural networks. In one embodiment, the compressed numerical representation includes integers or other numerical values, but in some embodiments, it includes indicators of values, indicators determining the value, and / or indicators of mathematical relationships determining the value.

[0016] That is, in one embodiment, systems and methods implemented in accordance with the present disclosure are used to compress a data set in a multi-step process, where for each element in the data set, the entire data set is first compressed by finding a value that can be stored once rather than multiple times. In one embodiment, each subset of the data set is then compressed in the same manner to find (for each subset) a second value, and then, instead of storing that second value, an indicator of how to get from the first value to the second value is stored. This indicator typically takes up less space (fewer bits) than the second value itself and is therefore a more efficient way to store the values ​​of the data set.

[0017] In other words, in one embodiment, the multiple values ​​in the data set are received as a matrix with "m" rows and "n" columns, where each value (or scalar) of the matrix is ​​represented by a common floating-point number (or binary value). A scalar is a computer representation that approximates a real-valued value using a floating-point format. In one embodiment, the multiple values ​​in the data set are received as a multidimensional tensor (a four-dimensional tensor). In one embodiment, a tensor is a multidimensional array of scalars, including but not limited to a matrix of R×C elements. Each common floating-point number is represented by a sequence of bits, where a portion of the bits (typically a single bit) represents the sign, another portion of the bits represents the exponent, and another portion of the bits represents the mantissa. For example, in a 32-bit sequence, the sign is represented by a single bit "1", the exponent is represented by 8 bits "01111111", and the mantissa is represented by 23 bits "110000000000000000000000", which would generate a binary value of "1.11." In addition to floating-point numbers (such as logarithmic representations), other representations of different formats for values ​​can also be implemented, and floating-point numbers are just one example described herein.

[0018] In one embodiment, the matrix is ​​subdivided into a uniform network of non-overlapping tiles of N×M scalars (e.g., 4×4 scalars). In one embodiment, each non-overlapping tile that divides the matrix into a uniform grid is called a submatrix (or block). These tiles can be the same size or different sizes. The elements in the matrix are typically tiled to improve locality in intermediate cache storage. More specifically, the matrix is ​​divided by the weight matrix in tiles of a predefined size and the exponents are shared within the tile. The tile limits the number of values ​​of the shared exponents. In one embodiment, each block is compressed independently of each other by determining the information shared with the elements of the block, which stores the entries in a manner that uses fewer bits. This can be achieved by subtracting a value from the exponent of each element of the block, such as the maximum exponent of all elements in the block. Other methods can also be used to reduce the bits stored for each element or to achieve other advantages, such as extracting common bits, factoring, or using other methods. In one embodiment, the maximum or common exponent value is called "block_exp". As an illustrative example, in one embodiment, if each exponent in the block is at least n (n is a positive integer), n may be subtracted from each element in the block so that each element in the block may be stored using fewer bits.

[0019] Then, in one embodiment, each block is further subdivided into a uniform grid of smaller P×Q scalars (e.g., 2×2 scalars) of non-overlapping sub-blocks. In one embodiment, each of these uniform grids of non-overlapping sub-blocks is referred to as a "quad". In one embodiment, each quad in the block is compressed by determining a common value (or exponent value) for all elements in the quad, as previously described, for example, when combining blocks of compression matrices. In one embodiment, the common value (or common exponent value) of the quad is referred to as "quad_exp". In one embodiment, assume that each quad belongs to the same block and the value of quad_exp is less than or equal to block_exp. Furthermore, in one embodiment, as long as the block size of N×M is small enough, it is likely that the difference between block_exp and max_exp is small enough that additional compression can be performed.

[0020] In one embodiment, rather than having the system store quad_exp directly, a more efficient encoding is achieved by storing the difference between quad_exp and block_exp (or an indicator of the difference). In some cases, this difference may be referred to as a per-quad exponential increment, a compressed numerical representation, or a proxy (instead of quad_exp) that provides information allowing the system to determine quad_exp from block_exp. In one embodiment, this difference is represented as an unsigned integer with a predetermined number of bits. In one embodiment, if the actual difference is larger than what can be represented using the predetermined number of bits, the difference is saturated to the largest representative value. However, in some embodiments, rather than storing the difference or the difference value itself, an indicator of a mathematical relationship is determined and stored, and / or an indicator is determined and stored that determines the mathematical relationship between quad_exp and block_exp and the values ​​in all scalars of the quad.

[0021] The technology described herein is used to generate a difference or indicator that provides information to indicate the mathematical relationship between the standard block exponents. The existence of these differences increases the dynamic range of the value that can be supported in the submatrix (or block), while still having a lower memory overhead than having a separate exponent for each scalar. Larger dynamic range means conversely that the scalar in the submatrix (or block) can be represented by a mantissa that is less than the mantissa bits needed for numerical stability if there is no difference. Therefore, the storage space required for storing these differences is less than the amount of space required for the single exponent and separate exponent of each scalar.

[0022] In addition to providing a more efficient way to store exponents, the techniques described herein also conveniently allow for computing the transpose of matrices, since the regular structure of the submatrices (or blocks) and quads means that the transpose can be computed without further rounding or quantization errors. Furthermore, by implementing the techniques described herein and the resulting space savings, it will approximately double (2x) the end-to-end training software configuration across a wide variety of neural networks and / or various graphics processing units (GPUs).

[0023] In the preceding and following descriptions, various techniques are described. For illustrative purposes, specific configurations and details are presented to provide a comprehensive understanding of possible methods for implementing these techniques. However, it will be apparent that the techniques described below can be practiced in various configurations without these specific details. Furthermore, well-known features may be omitted or simplified to avoid obscuring the described techniques.

[0024] Figure 1 A computing environment 100 is illustrated in which a block floating point (BFP) encoder 104 is implemented in a computing device 130 along with a front-end service 102, a data store 106, a compressor 108, and a decompressor 110 to generate and store one or more compressed numerical representations. The one or more compressed numerical representations may be referred to as representative values, common index values, reference index values, or difference values. In one embodiment, the one or more compressed numerical representations are simply referred to as indicators 140 of a data set, depending on the embodiment. In one embodiment, the computing environment 100 is a representation associated with a neural network, wherein the neural network uses the compressed numerical representation to reduce the cost of computations performed within the neural network.

[0025] In one embodiment, the computing device 130 includes and all descriptions provided herein are associated with a graphics processing unit (GPU). In one embodiment, the computing device 130 is configured to receive or obtain input data 120 and generate a difference value, or at least an indicator 140, which includes information representing a mathematical relationship between block index values ​​of a data set of the input data 120. The computing device 130 can be connected to another computing device, computing service, virtual machine, virtual service (not in the network) outside or inside the computing environment via a network 180 (wireless network or wired network). Figure 1 ) receives input data 120.

[0026] In one embodiment, the computing device 130 includes a front-end service 102. The front-end service 102 can be a computer system or component thereof configured to receive or acquire input data 120. For example, the front-end service 102 can be an analog-to-digital (ADC) converter. In one embodiment, the ADC is a computing system capable of receiving or acquiring analog signals as input and converting these analog signals into digital signals. The analog signal can be a signal provided to the computing device 130, picked up by a microphone, or light entering a digital camera. In one embodiment, the ADC is then configured to convert the analog signal into a digital signal, which is then a real number represented in bits to form a data set or matrix (composed of rows and columns). As described herein with respect to Figure 1 As described, the ADC is an illustrative embodiment of a front-end receiving service / computing device that may be implemented in a computing device 130 that takes, receives, or acquires input data 120. In another embodiment, the input data 120 is already digital signals and these signals are converted to values ​​in a matrix before reaching the front-end service 102 (e.g., the ADC). Therefore, the ADC may be bypassed or subtracted from the computing device 130. In one embodiment, the front-end service 102 is not an ADC that receives signals, but rather a computing service / device, such as a server computer system, a server cluster, a virtual computer system, a virtual server, a virtual runtime environment, a container environment, a serverless execution environment, a service hosting system, or any suitable computing entity that uses and is associated with the computing device 130 to receive the input data 120.

[0027] In one embodiment, once a signal has been converted into a digital signal or has been verified as a digital signal, the digital signal is represented as a value in a data set. In one embodiment, these values ​​are stored as input data 120 in a matrix. In one embodiment, a computing device 130 includes a BFP encoder 104. In one embodiment, the BFP encoder 104 is referred to as a BFP generator. The BFP encoder 104 can be configured to receive or obtain input data 120 sets of values ​​from the front-end service 102. In other words, the input data 120 can use floating-point representation to represent real numbers. Therefore, as described herein, block floating-point or BFP encoder 104 is used to provide an arithmetic algorithm that assigns an exponent to each data block from the input data 120 sets. Like this, each individual exponent of each element in the block does not have to be stored.

[0028] Therefore, in one embodiment, the BFP encoder 104 performs a first step operation on the set of input data 120 to identify the common (or maximum) exponent value among all values ​​in the block of the set of input data 120. Specifically, in one embodiment, if the block of the set of input data 120 is a number in the format of (sign)a*2^n (e.g., a multiple of 2 to the power of n, which can be represented by storing bits representing the sign, a, and n), and each element has an "n" of at least "100", then "100" can be identified as common. This operation is performed because the space required to store 2^(n-100) for each element is less. For example, examples of previous solutions attempt to generate a common or maximum exponent value for a data set, such as BF16 format floating point numbers (bfloat16) (also known as e8m7), and various forms of block exponent representation.

[0029] For example, in another embodiment, if the entries in the block of the set of input data 120 are a*2^70, b*2^59, c*2^65, and d*2^73, the common element will be 73 (because the exponents of all entries are at most 73). This will result in using less storage space associated with the entries by storing 73, a*2^-3, b*2^-14, c*2^-8, and d*2^0 because storing the common exponent 73 and right-shifting each mantissa according to each exponent difference takes up less space than storing the exponent of each element separately.

[0030] In yet another embodiment, if the entries in a block of the set of input data 120 are a*2^70, a*2^59, a*2^65, and a*2^73, the common value may be "a" because "a" is present in each element; the common value may be a*2^59 because a*2^59 can be factored out of each element; and the common value may be the vector (a, 59) because each element has a coefficient a and an exponent of at least 59. In practice, by factoring out the common values ​​between entries of the data set, less storage space may be required.

[0031] In addition, in one embodiment, the BFP encoder 104 further performs a second step operation to identify a second common exponent value in the subset 150 of the values ​​of the block of the set of input data 120. That is, in one embodiment, in a subsequent step after identifying the common exponent value from the values ​​of the data set as described above, the BFP encoder 104 again performs the same operation as in the first step on each subset of the block. That is, following the example above, the BFP encoder 104 can identify "95" as common to each element of the subset 150 of the block (because each number in the subset is of the form a*2^n, and n is at most "95").

[0032] Since two public exponent values ​​are identified (one from the block of the set of input data 120 and the other from the subset 150 of the blocks of the set of input data 120), a difference value can be determined or calculated. The difference value can be stored as an indicator 140, where the indicator 140 will provide information to the system on how to obtain the public exponent value of the subset 150 of the data set from the public exponent value of the block of the set of input data 120. Therefore, the difference value can be simply stored as "5" (because it indicates that the system will need to subtract "5" from "100", the public exponent value of the entire block identified in the first step). Therefore, when instructions or information are provided to the BFP encoder 104 so that the indicator 140 is determined and ultimately stored with the value "5", greater compression can be achieved because storing "5" only once for the subset 150 takes up less space than storing "95".

[0033] In one embodiment, the compressor 108 provides information or instructions to the BFP encoder 104 on how to perform a first step of compressing blocks of the set of input data 120 and further how to perform a second step of compressing respective subsets of blocks of the set of input data 120. The compressor may provide instructions so that the BFP encoder 104 selects a first number of elements in the set of input data 120 as blocks to be compressed first, and further selects another number of elements (which are smaller than the blocks to be compressed again). The compressor 108 may identify and provide instructions to the BFP encoder 104 based at least in part on user input and / or system policy to determine how to compress the data set and its subsets to ultimately determine the difference. In one embodiment, as Figure 1 As shown, the compressor 108 is not a separate device or service from the BFP encoder 104, but is merged with the BFP encoder 104 to perform the above-described operations.

[0034] Specifically, as an exemplary embodiment, the BFP encoder 104 determines that a block from a set of input data 120 (e.g., a matrix) can have a size of 4×4 scalars. Once the block is compressed, a first common exponent or maximum exponent can be determined. Subsequently, instructions from the compressor 108 to the BFP encoder 104 can indicate a subset of the block to be compressed. In one embodiment, the subset or quad of blocks has a size of 2×2 scalars. After compressing the quad, a second common or maximum exponent value is determined. The difference between the first common exponent value and the second common exponent value is calculated and further stored. In other words, the difference itself can be a stored indicator, allowing the system to determine the second common exponent value and the value in the subset of blocks based on the first common exponent value. However, in one embodiment, the difference is not stored, but rather several indicators of the difference are stored. That is, after determining the difference, the system can simply store an indicator of the mathematical relationship between the first and second common exponent values ​​and the value of the subset, without having to store the second common exponent value itself. Specifically, in one embodiment, the indicator is a function that can be used to derive the first and second common exponent values.

[0035] In one embodiment, computing device 130 includes data store 106. Data store 106 is configured as a storage device for receiving and storing public index values ​​for a data set. Specifically, in one embodiment, data store 106 is configured to store difference values, indicators 140 corresponding to the mathematical relationship between the first and second public index values ​​and each value in one or more subsets of values, and / or instructions indicating how the mathematical relationship between the first and second public index values ​​and each value in the one or more subsets of values ​​is determined. The storage requirements for storing the difference values, indicators, and / or instructions for the mathematical relationship between the first and second public index values ​​are less than the storage requirements for separately storing each index value for an entry or element of a matrix.

[0036] In one embodiment, computing device 130 includes decompressor 110. Decompressor 110 can be configured to receive compressed floating point numbers with public exponent values ​​and decompose each floating point number into a byte array to form a byte matrix. The compressed floating point number can be received by decompressor 110 from BFP encoder 104, and / or it can come from data storage 106. Then, in one embodiment, decompressor 110 transposes the matrix to generate a new matrix, wherein the first row shows low entropy (exponent) and the last row has high entropy (last mantissa bit). In other words, decompressor 110 receives a compressed data set and stored difference and / or indicator, and is configured to decompress the data set to generate data output. That is, decompressor 110 can be configured to receive a set of instructions and / or algorithms that, when executed, cause decompressor 110 to expand the compressed data set back to its original form (e.g., before compressing the value, for example, by regenerating the value in matrix 120 through lossy compression technology or lossless compression technology). In one embodiment, the decompressor 110 decompresses the compressed data set after performing additional mathematical operations on the compressed data set and thus generates data output that is different from the original values ​​in the matrix 120 .

[0037] Figure 2 A diagram 200 of a matrix according to an embodiment is shown, which includes a set of values ​​represented by a public exponent 202, a sign 204, and a mantissa 206 when filled. In one embodiment, diagram 200 is a representation of a matrix having a data set using floating-point representation to represent real numbers. In one embodiment, floating-point numbers are represented in bits, where, for example, a subsequence of bits represents the exponent of the number "2". Specifically, a number can be represented as (sign)*a*2^n, and their floating-point representation has the bits of sign (sign), n, and a. In one embodiment, a value in diagram 200 is represented by a series of bits, where the public exponent 202 is "1000001", the sign 204 is "0", and the mantissa 206 is "01100000000000000000000". The result of this value is the decimal number 1.0011x2^4=22.

[0038] In one embodiment, if the index of each element in the 4x4 block of the matrix is ​​at most "k", the system performs the above steps for Figure 1 The first operation described above stores "k" once, which will be stored as the public index 202. Furthermore, in one embodiment, the system stores each value of the matrix with an index of "kn," which takes up fewer bits than storing it with an index of "n."

[0039] also, Figure 2Another diagram 210 of a matrix is ​​shown, according to an embodiment, that includes a set of values ​​represented by an exponent 202, a packed numerical representation 208, a sign 204, and a mantissa 206 (or block floating point with a packed numerical representation). In one embodiment, Figure 2 The second figure 210 of is subdivided into 2x2 submatrices, and after performing the same process as shown above, determines the second public index for the set that is subdivided.The size of the set that is subdivided can be predetermined, arbitrary, random and / or determined by user input.In one embodiment, the second public index of each 2x2 submatrix (or sub-block) is not stored.On the contrary, for each submatrix in these 2x2 submatrices, compression numerical value is calculated and represented 208.Compression numerical value represents 208 and can be calculated by subtracting the second public index from public index 202, and is stored as difference (or incremental value).The size of compression numerical value represents 208 can be the predetermined bit amount represented by digital bit Q, which is less than the size of the second public index calculated for each 2x2 submatrix (or sub-block).In other words, the size of predetermined bit amount will be no more than the bit of the second public index.

[0040] In other words, from each of these 2x2 sub-matrices, the same process of extracting the public exponent value 202 from the block is performed, but instead of storing the exponent that can be pulled out, a compressed numerical representation 208 between the exponent that has been extracted and the first public exponent value from the first step is stored. In one embodiment, this compressed numerical representation 208 is not an integer representing each of the 2x2 sub-matrices, but rather some indicator that determines the compressed numerical representation 208. The compressed numerical representation 208 is typically small and therefore requires fewer bits to store. In one embodiment, the compressed numerical representation 208 is simply an indicator of the stored function to enable the derivation of the second public exponent value.

[0041] Figure 3 An exemplary matrix 302 (similar to Figure 1 302 ), a submatrix (or block) 304, a submatrix (or quad) 306 of the block, and a scalar 308. In one embodiment, the data set is represented using floating point representation to represent real numbers as elements of the matrix 302. That is, in one embodiment, the matrix 302 has 24 rows and 16 columns. In one embodiment, the matrix 302 is configured with a different number of rows and columns.

[0042] In one embodiment, submatrix (or block) 304 is a block of 16 elements (e.g., 4x4 scalars) of matrix 302. The block size can be determined randomly or predetermined by user input. In addition, the block size can also be determined based at least in part on user input or metadata associated with the matrix.

[0043] In one embodiment, a quad 306 of size 2x2 scalar is a sub-portion of a sub-matrix (or block) 304. The size of the quad can be determined randomly or predetermined by user input. Additionally, the size of the quad can be determined based at least in part on user input or metadata associated with the matrix.

[0044] In one embodiment, scalar 308 is a sub-portion, an element, or a single element of matrix 302. Each scalar 308 may be a floating point number represented in bits.

[0045] Figure 4 According to an embodiment, a diagram 400 is shown in which common index values ​​are determined for sub-matrices of a matrix. In one embodiment, a sub-matrix (or block) 402 is obtained from Figure 3 The matrix 302 in the extract is a data set of scalar values ​​of size 4x4.

[0046] In one embodiment, as described above with respect to Figure 1 As described, block 402 and the size of block 402 are determined by the BFP encoder. In one embodiment, block 402 and its size are determined by user input and / or system policy. A first common or maximum exponent value 406 can be determined among all values ​​in block 402. In one embodiment, the first common or maximum exponent value 406 has a size of 10 bits. The first common exponent value 406 can represent a sign and mantissa 404 that at least shares the first common exponent value or the maximum exponent value. In this way, the common exponent value 406 is only stored once for block 402, without having to store each exponent value of each element in block 402 separately. In other words, 16 exponent values ​​for 16 scalars in block 402 are not stored, for example, but only a single shared exponent value among all elements in block 402 is stored.

[0047] In one embodiment, quad 408 is determined by subdividing block 402 into 2x2 scalars. As described above, quad 408 is a subset of block 402, and the size of the quad 408 can be predetermined or selected so that there is a high probability that the public exponent value 412 of quad 408 is at least the same as or less than the public exponent value 412 of block 402. Once quad 408 is determined, the next step will perform the same compression operation (the same compression operation performed on block 402) to determine the public exponent value 412 of quad 408. The public exponent value 412 of quad 408 will be the exponent value of the sign and mantissa 410 of quad 408.

[0048] As mentioned above, determine the difference between the public exponent value 406 of block 402 and the public exponent value 412 of quadruple 408.Then, this difference is used as the relation between two public exponent values ​​406,412 or the indicator of mathematical relationship.Different is usually very little (less than the size of the public exponent value 406 of block 402 and / or the size of the public exponent value 412 of quadruple 408), and can be assigned as being expressed as " q position " integer.User input and / or system policy can predetermine the size or quantity of " q position ".In one embodiment, perform this iteration of determining difference continuously, till all differences of determining matrix and submatrix.Iteration number may depend on the size of initial matrix.

[0049] Figure 5 According to an embodiment, a diagram 500 is shown in which a difference value (or indicator 516) corresponding to a mathematical relationship between the common values ​​506, 512 in the block and the quaternion value is determined. Figure 4 , after a block (submatrix of a matrix) 502 extracts or determines a public exponent value 506 (e.g., a first public value) of its sign and mantissa 504, the public exponent value 506 may be stored in a data store ( Figure 5 Not shown, but Figure 1 ). Subsequently, a submatrix (e.g., a quad) 508 of the block can be determined such that another public exponent value 512 (e.g., a second public value) is determined for the sign and mantissa 510 of the quad 508. In one embodiment, the other public exponent value 512 is not stored in the data store. Instead, a difference (e.g., a third value) between the public exponent value 506 and the other public exponent value 512 is determined and stored in the data store. The difference can be an integer, which is typically smaller than the other public exponent value 512. In one embodiment, the size of the difference (or indicator 516) is 6 bits. The size of the difference (or indicator 516) can vary, as the 6-bit size described here is only an example, and differences (or indicators 516) of other possible sizes can be stored. In one embodiment, the difference is not an integer, but rather an indication of how the difference is determined. Therefore, the difference (or indicator 516) is merely a proxy for the system to determine another public exponent value 512, when necessary, by using a calculation method 514 or a mathematical calculation based at least in part on the previously generated public exponent value 506.

[0050] Figure 6An illustrative example of a process 600 for generating a difference value (or indicator) is shown according to an embodiment. In one embodiment, a system has at least one or more processors (or one or more arithmetic logic unit (ALU) processors) configured to include logic 602 for receiving or acquiring a data set consisting of values. In addition, the one or more processors also include logic 604 for determining a first common value (e.g., a first common exponent value or a first maximum exponent value) within each value of a first set of values. In one embodiment, the one or more processors are further configured to include logic 606 for determining a second common value (e.g., a second common exponent value or a second maximum exponent value) within each value of one or more value subsets of the first set of values. In addition, the one or more processors are further configured to include logic for storing a third value corresponding to a mathematical relationship between the first and second common values ​​and each value within the one or more value subsets, wherein the third value requires less storage than the first and second common values ​​608. In one embodiment, the third value is an integer of a predetermined size. In one embodiment, the third value is simply an indicator or function that can be used to derive the first and second common values.

[0051] Figure 7 Another example of a process for generating a difference value or indicator is shown according to an embodiment. In one embodiment, the system has at least one or more processors configured to include logic for receiving or acquiring a data set. The data set can be compressed to determine a first value common to all values ​​of the data set 702. In addition, the processor may also include logic 704 for compressing a subset of the data set to determine a second value common to the values ​​of the subset. Once the first value and the second value are determined, the difference 706 between the two values ​​can be determined and stored. The difference value can be an indicator of the mathematical relationship between the first value and the second value and each value of the subset. In other words, the difference value can be a proxy (e.g., in place of the second value) that, when used with the first value, provides information to the system to determine the second value and the value of the subset. This difference value requires less storage space than storing the first and / or second value.

[0052] Figure 8A parallel processing unit ("PPU") 800 is shown according to one embodiment. In one embodiment, the PPU 800 is configured using machine-readable code that, if executed by the PPU, causes the PPU to perform some or all of the processes and techniques described throughout this disclosure. In one embodiment, the PPU 800 is a multi-threaded processor that includes one or more arithmetic logic units (ALUs) implemented on one or more integrated circuit devices and is designed to process computer-readable instructions (also referred to as machine-readable instructions or simply instructions) on multiple threads in parallel, utilizing multithreading as a latency hiding technique. In one embodiment, a thread refers to a thread of execution and is an instance of a group of instructions configured to be executed by the PPU 800. In one embodiment, the PPU 800 is a graphics processing unit ("GPU") configured to implement a graphics rendering pipeline for processing three-dimensional ("3D") graphics data to generate two-dimensional ("2D") image data for display on a display device, such as a liquid crystal display (LCD) device. In one embodiment, the PPU 800 is used to perform computations such as linear algebra operations and machine learning operations. Figure 8 An exemplary parallel processor is shown for purposes of illustration only and should be construed as a non-limiting example of a processor architecture contemplated within the scope of the present disclosure, and any suitable processor may be used in addition to and / or in place of the processor.

[0053] In one embodiment, one or more PPUs 300 are configured to accelerate high-performance computing (HPC), data center, and machine learning applications. In one embodiment, the PPU 800 is configured to accelerate deep learning systems and applications, including the following non-limiting examples: autonomous vehicle platforms, deep learning, high-precision speech, image, and text recognition systems, intelligent video analysis, molecular simulation, drug discovery, disease diagnosis, weather forecasting, big data analysis, astronomy, molecular dynamics simulation, financial modeling, robotics, factory automation, real-time language translation, online search optimization, and personalized user recommendations.

[0054] In one embodiment, the PPU 800 includes an input / output ("I / O") unit 805, a front-end unit 810, a scheduler unit 812, a work distribution unit 814, a hub 816, a crossbar switch ("Xbar") 820, one or more general processing clusters ("GPCs") 818, and one or more partition units 822. In one embodiment, the PPU 800 is connected to a host processor or other PPUs 800 via one or more high-speed GPU interconnects 808. In one embodiment, the PPU 800 is connected to a host processor or other peripheral devices via an interconnect 802. In one embodiment, the PPU 800 is connected to a local memory including one or more memory devices 804. In one embodiment, the local memory includes one or more dynamic random access memory ("DRAM") devices. In one embodiment, the one or more DRAM devices are configured and / or configurable as a high-bandwidth memory ("HBM") subsystem, where multiple DRAM dies are stacked within each device.

[0055] The high-speed GPU interconnect 808 may refer to a wire-based, multi-lane communication link used by a system to extend and include one or more PPUs 800 in conjunction with one or more CPUs, supporting cache coherency between the PPU 800 and the CPU, and CPU mastering. In one embodiment, data and / or commands are sent to / from other units of the PPU 800, such as one or more copy engines, video encoders, video decoders, power management units, and other components in the PPU 800, via the high-speed GPU interconnect 808 via the hub 816. Figure 8 Components that may not be explicitly shown.

[0056] In one embodiment, the I / O unit 805 is configured to receive data from the host processor ( Figure 8 In one embodiment, the I / O unit 805 communicates with the host processor directly via the system bus 802 or through one or more intermediate devices (e.g., a memory bridge). In one embodiment, the I / O unit 805 can communicate with one or more other processors (e.g., one or more PPUs 800) via the system bus 802. In one embodiment, the I / O unit 805 implements a Peripheral Component Interconnect Express ("PCIe") interface for communicating over the PCIe bus. In one embodiment, the I / O unit 805 implements an interface for communicating with external devices.

[0057] In one embodiment, the I / O unit 805 decodes packets received via the system bus 802. In one embodiment, at least some of the packets represent commands configured to cause the PPU 800 to perform various operations. In one embodiment, the I / O unit 805 sends the decoded commands to various other units of the PPU 800 specified by the commands. In one embodiment, the commands are sent to the front end unit 810 and / or to the hub 816 or other units of the PPU 800, such as one or more copy engines, video encoders, video decoders, power management units, etc. Figure 8 In one embodiment, I / O unit 805 is configured to route communications between the various logical units of PPU 800.

[0058] In one embodiment, a program executed by a host processor encodes a command stream in a buffer, which provides a workload to the PPU 800 for processing. In one embodiment, the workload includes instructions and data to be processed by these instructions. In one embodiment, the buffer is an area of ​​memory that is accessible (e.g., read / write) by the host processor and the PPU 800. The host interface unit can be configured to access the buffer in system memory connected to the system bus 802 via a memory request sent by the I / O unit 805 via the system bus 802. In one embodiment, the host processor writes a command stream to the buffer and then sends a pointer to the beginning of the command stream to the PPU 800, so that the front end unit 810 receives pointers to one or more command streams and manages the one or more command streams, reading commands from the streams and forwarding the commands to the various units of the PPU 800.

[0059] In one embodiment, the front end unit 810 is coupled to a scheduler unit 812, which configures the various GPCs 818 to process tasks defined by one or more streams. In one embodiment, the scheduler unit 812 is configured to track state information related to the various tasks managed by the scheduler unit 812, where the state information may indicate which GPC 818 the task is assigned to, whether the task is active or inactive, a priority associated with the task, and the like. In one embodiment, the scheduler unit 812 manages the execution of multiple tasks on one or more GPCs 818.

[0060] In one embodiment, the scheduler unit 812 is coupled to a work distribution unit 814, which is configured to dispatch tasks for execution on the GPCs 818. In one embodiment, the work distribution unit 814 tracks a plurality of scheduled tasks received from the scheduler unit 812 and manages a pending task pool and an active task pool for each GPC 818. In one embodiment, the pending task pool may include a number of time slots (e.g., 32 time slots) containing tasks assigned to be processed by a particular GPC 818; the active task pool may include a number of time slots (e.g., 4 time slots) for tasks being actively processed by the GPC 818, with the intention that when the GPC 818 completes execution of a task, the task is evicted from the active task pool of the GPC 818, and one of the other tasks from the pending task pool is selected and scheduled for execution on the GPC 818. In one embodiment, if an active task is idle on a GPC 818 , e.g., while waiting for a data dependency to be resolved, the active task is evicted from the GPC 818 and returned to the pending task pool, while another task in the pending task pool is selected and scheduled for execution on the GPC 818 .

[0061] In one embodiment, work distribution unit 814 communicates with one or more GPCs 818 via XBar 820. In one embodiment, XBar 820 is an interconnect network that couples many units of PPU 800 to other units of PPU 800 and can be configured to couple work distribution unit 814 to a specific GPC 818. Although not explicitly shown, one or more other units of PPU 800 can also be connected to XBar 820 via hub 816.

[0062] These tasks are managed by the scheduler unit 812 and dispatched to the GPCs 818 by the work distribution unit 814. The GPCs 818 are configured to process tasks and generate results. The results can be consumed by other tasks within the GPC 818, routed to different GPCs 818 via the XBar 820, or stored in the memory 804. The results can be written to the memory 804 via the partition unit 822, which implements a memory interface for reading and writing data from / to the memory 804. The results can be sent to another PPU 804 or the CPU via the high-speed GPU interconnect 808. In one embodiment, the PPU 800 includes a number U of partition units 822, which is equal to the number of independent and different memory devices 804 coupled to the PPU 800. Figure 10 The partition unit 822 is described in more detail.

[0063] In one embodiment, the host processor executes a driver kernel that implements an application programming interface ("API") that enables one or more applications executing on the host processor to schedule operations to be performed on the PPU 800. In one embodiment, the PPU 800 executes multiple computing applications simultaneously, and the PPU 800 provides isolation, quality of service ("QoS"), and independent address spaces for the multiple computing applications. In one embodiment, the application generates instructions (e.g., in the form of API calls) that cause the driver kernel to generate one or more tasks for execution by the PPU 800, and the driver kernel outputs the tasks to one or more streams being processed by the PPU 800. In one embodiment, each task includes one or more groups of related threads, which may be referred to as a warp. In one embodiment, a warp includes multiple related threads (e.g., 32 threads) that may be executed in parallel. In one embodiment, a cooperating thread may refer to multiple threads that include instructions to perform tasks and instructions to exchange data through shared memory. According to an embodiment, in combination Figure 10 Describes threads and cooperative threads in more detail.

[0064] Figure 9 According to one embodiment, a GPC 900 is shown, for example Figure 8 800 of the GPC shown. In one embodiment, each GPC 900 includes multiple hardware units for processing tasks, and each GPC 900 includes a pipeline manager 902, a pre-raster operation unit ("PROP") 904, a raster engine 908, a work distribution crossbar unit ("WDX") 916, a memory management unit ("MMU") 918, one or more data processing clusters ("DPC") 906, and any suitable combination of parts. It should be understood that Figure 9 The GPC 900 may include instead of or in addition to Figure 9 Other hardware units of the unit shown.

[0065] In one embodiment, the operation of GPC 900 is controlled by pipeline manager 902. Pipeline manager 902 manages the configuration of one or more DPCs 906 for processing tasks assigned to GPC 900. In one embodiment, pipeline manager 902 configures at least one of one or more DPCs 906 to implement at least a portion of a graphics rendering pipeline. In one embodiment, DPC 906 is configured to execute vertex shading programs on a programmable streaming multiprocessor ("SM") 914. In one embodiment, pipeline manager 902 is configured to route packets received from work assignments to appropriate logic units within GPC 900. Some packets may be routed to fixed-function hardware units in PROP 904 and / or raster engine 908, while other packets may be routed to DPC 906 for processing by primitive engine 912 or SM 914. In one embodiment, pipeline manager 902 configures at least one of one or more DPCs 906 to implement a neural network model and / or a compute pipeline.

[0066] In one embodiment, the PROP unit 904 is configured to route data generated by the raster engine 908 and the DPC 906 to a raster operation ("ROP") unit in the memory partition unit, as described in detail above. In one embodiment, the PROP unit 904 is configured to perform optimizations for color blending, organize pixel data, perform address translation, and the like. In one embodiment, the raster engine 908 includes several fixed-function hardware units configured to perform various raster operations, and the raster engine 908 includes a setup engine, a coarse raster engine, a culling engine, a clipping engine, a fine raster engine, and a tile aggregation engine, and any suitable combination of the above engines. In one embodiment, the setup engine receives transformed vertices and generates a plane equation associated with the geometric primitives defined by the vertices, which is sent to the coarse raster engine to generate coverage information for the primitives (e.g., an x, y coverage mask for the tile); the output of the coarse raster engine is sent to the culling engine, where fragments associated with primitives that fail the z-test are culled, and to the clipping engine, where fragments outside the view frustum are clipped. In one embodiment, those fragments that remain after clipping and culling can be passed to a fine raster engine to generate pixel fragment attributes based on the plane equation generated by the setup engine. In one embodiment, the output of the raster engine 908 includes fragments processed by any suitable entity, such as fragments processed by a fragment shader implemented in the DPC 906.

[0067] In one embodiment, each DPC 906 included in a GPC 900 includes an M-pipeline controller ("MPC") 910, a primitive engine 912, one or more SMs 914, and any suitable combination thereof. In one embodiment, the MPC 910 controls the operation of the DPC 906, routing packets received from the pipeline manager 902 to appropriate units within the DPC 906. In one embodiment, packets associated with vertices are routed to the primitive engine 912, which is configured to retrieve vertex attributes associated with the vertices from memory; conversely, packets associated with shader programs may be sent to the SM 914.

[0068] In one embodiment, SM 914 includes a programmable streaming processor configured to process tasks represented by multiple threads. In one embodiment, SM 914 is multithreaded and configured to simultaneously execute multiple threads (e.g., 32 threads) from a specific thread group and implement a SIMD (single instruction, multiple data) architecture, in which each thread in a group of threads (warp) is configured to process a different data set based on the same set of instructions. In one embodiment, all threads in a thread group execute the same instructions. In one embodiment, SM 914 implements a SIMT (single instruction, multiple thread) architecture, in which each thread in a group of threads is configured to process a different data set based on the same set of instructions, but individual threads in a thread group are allowed to diverge during execution. In one embodiment, a program counter, call stack, and execution state are maintained for each warp, enabling concurrency between warps and serial execution within warps when threads within a warp diverge. In another embodiment, a program counter, call stack, and execution state are maintained for each individual thread, enabling equal concurrency between all threads within and between warps. In one embodiment, execution state is maintained for each individual thread, allowing threads executing the same instructions to converge and execute in parallel for maximum efficiency. In one embodiment, SM 914 is described in more detail below.

[0069] In one embodiment, MMU 918 provides an interface between GPC 900 and the memory partition unit. MMU 918 provides virtual-to-physical address translation, memory protection, and arbitration of memory requests. In one embodiment, MMU 918 provides one or more translation lookaside buffers ("TLBs") for performing translation of virtual addresses into physical addresses in memory.

[0070] Figure 10A memory partition unit of a PPU is shown according to one embodiment. In one embodiment, the memory partition unit 1000 includes a raster operation ("ROP") unit 1002, a level 2 ("L2") cache 1004, a memory interface 1006, and any suitable combination thereof. The memory interface 1006 is coupled to a memory. The memory interface 1006 can implement a 32-, 64-, 128-, 1024-bit data bus, etc., for high-speed data transfer. In one embodiment, the PPU includes U memory interfaces 1006, one memory interface 1006 for each pair of partition units 1000, where each pair of partition units 1000 is connected to a corresponding memory device. For example, the PPU can be connected to up to Y memory devices, such as a high-bandwidth memory stack or graphics double data rate, version 5, synchronous dynamic random access memory ("GDDR5 SDRAM").

[0071] In one embodiment, memory interface 1006 implements an HBM2 memory interface, and Y is equal to half U. In one embodiment, the HBM2 memory stack is located on the same physical package as the PPU, providing significant power and area savings compared to traditional GDDR5 SDRAM systems. In one embodiment, each HBM2 stack includes four memory dies and Y is equal to 4, where the HBM2 stack includes two 128-bit channels per die, for a total of 8 channels and a data bus width of 1024 bits.

[0072] In one embodiment, the memory supports single error correction, double error detection ("SECDED") error correction code ("ECC") to protect data. ECC provides increased reliability for computing applications sensitive to data corruption. Reliability is particularly important in large-scale cluster computing environments where the PPU processes very large data sets and / or long-running applications.

[0073] In one embodiment, the PPU implements a multi-level memory hierarchy. In one embodiment, the memory partition unit 1000 supports unified memory to provide a single unified virtual address space for CPU and PPU memory, thereby enabling data sharing between virtual memory systems. In one embodiment, the frequency of PPU accesses to memory located on other processors is tracked to ensure that memory pages are moved to the physical memory of the PPU that accesses the pages more frequently. In one embodiment, the high-speed GPU interconnect 808 supports address translation services, allowing the PPU to directly access the CPU's page tables and providing full access to CPU memory by the PPU.

[0074] In one embodiment, the copy engine transfers data between multiple PPUs or between a PPU and a CPU. In one embodiment, the copy engine can generate a page fault for an address that is not mapped in the page table, and the memory partition unit 1000 then services the page fault, maps the address into the page table, and then the copy engine performs the transfer. In one embodiment, multiple copy engines operate on fixed memory (i.e., non-pageable) between multiple processors, substantially reducing the available memory. In one embodiment, following a hardware page fault, the address can be passed to the copy engine regardless of whether the memory page is resident, and the copy process is transparent.

[0075] According to one embodiment, from Figure 8 Data from the memory or other system memory is retrieved by the memory partition unit 1000 and stored in the L2 cache 1004 located on the chip and shared between the various GPCs. In one embodiment, each memory partition unit 1000 includes at least a portion of the L2 cache 960 associated with the corresponding memory device. In one embodiment, the lower level cache is implemented in each unit within the GPC. In one embodiment, each SM 1040 can implement a level 1 ("L1") cache, where the L1 cache is a dedicated memory dedicated to a specific SM 1040, and data from the L2 cache 1004 is retrieved and stored in each L1 cache for processing in the functional units of the SM 1040. In one embodiment, the L2 cache 1004 is coupled to the memory interface 1006 and the XBar 820.

[0076] In one embodiment, ROP unit 1002 performs graphics raster operations related to pixel color, such as color compression, pixel blending, and the like. In one embodiment, ROP unit 1002 implements depth testing in conjunction with raster engine 1025, receiving the depth of a sample location associated with a pixel fragment from the culling engine of raster engine 1025. In one embodiment, the depth of the sample location associated with the fragment is tested against the corresponding depth in the depth buffer. In one embodiment, if the fragment passes the depth test for the sample location, ROP unit 1002 updates the depth buffer and sends the result of the depth test to raster engine 1025. It should be understood that the number of partition units 1000 can be different than the number of GPCs, and thus, in one embodiment, each ROP unit 1002 can be coupled to each GPC. In one embodiment, ROP unit 1002 tracks packets received from different GPCs and determines to which GPC the results generated by ROP unit 1002 are routed via the Xbar.

[0077] Figure 11 According to one embodiment, a streaming multiprocessor is shown, for example Figure 9Streaming multiprocessor. In one embodiment, SM 1100 includes: an instruction cache 1102; one or more scheduler units 1104; a register file 1108; one or more processing cores 1110; one or more special function units ("SFUs") 1112; one or more load / store units ("LSUs") 1114; an interconnect network 1116; a shared memory / L1 cache 1118, and any suitable combination thereof. In one embodiment, a work distribution unit dispatches tasks for execution on the GPCs of the PPU, with each task assigned to a specific DPC within the GPC and, if the task is associated with a shader program, to SM 1100. In one embodiment, scheduler unit 1104 receives tasks from the work distribution unit and manages instruction scheduling for one or more thread blocks assigned to SM 1100. In one embodiment, scheduler unit 1104 schedules thread blocks for execution into warps of parallel threads, where each thread block is assigned at least one warp. In one embodiment, each warp executes a thread. In one embodiment, the scheduler unit 1104 manages multiple different thread blocks, assigns warps to different thread blocks, and then dispatches instructions from multiple different cooperation groups to various functional units (e.g., core 1110, SFU 1112, and LSU 1114) in each clock cycle.

[0078] Cooperative groups can refer to a programming model for organizing groups of communicating threads that allows developers to express the granularity at which threads are communicating, enabling the expression of richer and more efficient parallel decompositions. In one embodiment, a cooperative launch API supports synchronization between thread blocks to execute parallel algorithms. In one embodiment, applications of conventional programming models provide a single, simple structure for synchronizing cooperative threads: a barrier across all threads of a thread block (e.g., a syncthreads() function). However, programmers often want to define thread groups at a granularity smaller than the thread block granularity and synchronize within the defined group, enabling higher performance, design flexibility, and software reuse in the form of a collective group-wide function interface. Cooperative groups enable programmers to explicitly define thread groups at sub-block (i.e., as small as a single thread) and multi-block granularity, and perform collective operations (such as synchronization) on threads in the cooperative group. The programming model supports clean composition across software boundaries, so libraries and utility functions can safely synchronize in their local context without making assumptions about convergence. The cooperative group primitive supports new patterns of cooperative parallelism, including producer-consumer parallelism, opportunistic parallelism, and global synchronization across the entire grid of thread blocks.

[0079] In one embodiment, the dispatch unit 1106 is configured to send instructions to one or more functional units, and the scheduler unit 1104 includes two dispatch units 1106, so that two different instructions from the same warp can be dispatched during each clock cycle. In one embodiment, each scheduler unit 1104 includes a single dispatch unit 1106 or additional dispatch units 1106.

[0080] In one embodiment, each SM 1100 includes a register file 1108 that provides a set of registers for the functional units of SM 1100. In one embodiment, register file 1108 is divided between each functional unit, so that each functional unit is allocated a dedicated portion of register file 1108. In one embodiment, register file 1108 is divided between the different thread warps being executed by SM 1100, and register file 1108 provides temporary storage for operands connected to the data paths of the functional units. In one embodiment, each SM 1100 includes a plurality of L processing cores 1110. In one embodiment, SM 1100 includes a large number (e.g., 128 or more) of different processing cores 1110. In one embodiment, each core 1110 includes a fully pipelined, single-precision, double-precision, and / or mixed-precision processing unit, which includes a floating-point arithmetic logic unit and an integer arithmetic logic unit. In one embodiment, the floating-point arithmetic logic unit implements the IEEE 754-2008 standard for floating-point operations. In one embodiment, core 1110 includes 64 single-precision (32-bit) floating point cores, 64 integer cores, 32 double-precision (64-bit) floating point cores, and 8 tensor cores.

[0081] According to an embodiment, the tensor core is configured to perform matrix operations. In one embodiment, one or more tensor cores are included in core 1110. In one embodiment, the tensor core is configured to perform deep learning matrix algorithms, such as convolution operations for neural network training and inference. In one embodiment, each tensor core operates on 4x4 matrices and performs a matrix multiplication and accumulation operation D=A×B+C, where A, B, C, and D are 4x4 matrices.

[0082] In one embodiment, the matrix multiplication inputs A and B are 16-bit floating point matrices, and the accumulation matrices C and D are 16-bit floating point or 32-bit floating point matrices. In one embodiment, the tensor cores operate on 16-bit floating point input data and 32-bit floating point accumulations. In one embodiment, the 16-bit floating point multiplication requires 64 operations and produces a full-precision product, which is then accumulated with other intermediate products for the 4x4x4 matrix multiplication using 32-bit floating point addition. In one embodiment, the tensor cores are used to perform larger two-dimensional or higher dimensional matrix operations built from these smaller elements. In one embodiment, the API (such as the CUDA9 C++ API) exposes specialized matrix loads, matrix multiplication and accumulation operations, and matrix store operations to efficiently use the tensor cores from CUDA-C++ programs. In one embodiment, at the CUDA level, the warp-level interface assumes that the 16×16 size matrix spans all 32 threads of the warp.

[0083] In one embodiment, each SM 1100 includes M SFUs 1112 that perform special functions (e.g., attribute evaluation, reciprocal square root, etc.). In one embodiment, the SFUs 1112 include tree traversal units configured to traverse a hierarchical tree data structure. In one embodiment, the SFUs 1112 include texture units configured to perform texture map filtering operations. In one embodiment, the texture units are configured to load texture maps (e.g., 2D texture arrays) from memory and sample the texture maps to generate sampled texture values ​​for use in shader programs executed by the SMs 1100. In one embodiment, the texture maps are stored in shared memory / L1 cache. According to one embodiment, the texture units implement texture operations, such as filtering operations using mip maps (e.g., texture maps of different levels of detail). In one embodiment, each SM 1100 includes two texture units.

[0084] In one embodiment, each SM 1100 includes N LSUs 1054 that implement load and store operations between the shared memory / L1 cache 1006 and the register file 1108. In one embodiment, each SM 1100 includes an interconnect network 1116 that connects each functional unit to the register file 1108 and connects the LSUs 1114 to the register file 1108 and the shared memory / L1 cache 1006. In one embodiment, the interconnect network 1116 is a crossbar switch that can be configured to connect any functional unit to any register in the register file 1108 and to connect the LSUs 1114 to memory locations in the register file and the shared memory / L1 cache 1115.

[0085] In one embodiment, shared memory / L1 cache 1118 is an on-chip memory array that allows data storage and communication between SM 1100 and primitive engines, as well as between threads within SM 1100. In one embodiment, shared memory / L1 cache 1118 includes 128KB of storage capacity and is located in the path from SM 1100 to the partition unit. In one embodiment, shared memory / L1 cache 1118 is used to cache reads and writes. One or more of shared memory / L1 cache 1118, L2 cache, and memory are backing stores.

[0086] In one embodiment, data cache and shared memory functionality are combined into a single memory block to provide improved performance for both types of memory accesses. In one embodiment, this capacity can be used by the program as a cache that is not using shared memory or is available. For example, if the shared memory is configured to use half of its capacity, texture and load / store operations can use the remaining capacity. According to an embodiment, the integration within the shared memory / L1 cache 1118 enables the shared memory / L1 cache 1118 to operate as a high-throughput pipeline for streaming data, while providing high bandwidth and low latency access to frequently reused data. When configured for general-purpose parallel computing, a simpler configuration can be used compared to graphics processing. In one embodiment, the fixed-function graphics processing unit is bypassed, creating a simpler programming model. In one embodiment, in a general-purpose parallel computing configuration, the work distribution unit assigns and allocates thread blocks directly to the DPC. According to one embodiment, threads in a block execute the same program, use unique thread IDs in computations to ensure each thread generates unique results, use SM 1100 to execute the program and perform computations, use shared memory / L1 cache 1118 to communicate between threads, and read and write global memory through shared memory / L1 cache 1118 and a memory partitioning unit using LSU 1114. In one embodiment, when configured for general-purpose parallel computation, SM 1100 writes commands that the scheduler unit can use to start new work on a DPC.

[0087] In one embodiment, the PPU is included in or coupled to a desktop computer, laptop computer, tablet computer, server, supercomputer, smartphone (e.g., wireless, handheld device), personal digital assistant ("PDA"), digital camera, vehicle, head-mounted display, handheld electronic device, etc. In one embodiment, the PPU is disposed on a single semiconductor substrate. In one embodiment, the PPU is included in a system-on-chip ("SoC") along with one or more other devices, such as an additional PPU, memory, a reduced instruction set computer ("RISC") CPU, a memory management unit ("MMU"), a digital-to-analog converter ("DAC"), etc.

[0088] In one embodiment, the PPU may be included on a graphics card that includes one or more memory devices. The graphics card may be configured to interface with a PCIe slot on a motherboard of a desktop computer. In another embodiment, the PPU may be an integrated graphics processing unit ("iGPU") included in a chipset of the motherboard.

[0089] Figure 12 According to one embodiment, a computer system 1200 is shown upon which various architectures and / or functions may be implemented. In one embodiment, the computer system 1200 is configured to implement the various processes and methods described throughout this disclosure.

[0090] In one embodiment, computer system 1200 includes at least one central processing unit 1202 connected to a communication bus 1210 implemented using any suitable protocol, such as PCI (Peripheral Component Interconnect), PCI-Express, AGP (Accelerated Graphics Port), HyperTransport, or any other bus or point-to-point communication protocol. In one embodiment, computer system 1200 includes main memory 1204, which may take the form of random access memory ("RAM"), control logic (e.g., implemented in hardware, software, or a combination thereof), and data stored on main memory 1204. In one embodiment, network interface subsystem 1222 provides an interface to other computing devices and networks for receiving data from computer system 1200 and sending data to other systems.

[0091] In one embodiment, computer system 1200 includes an input device 1208, a parallel processing system 1212, and a display device 1206, which can be implemented using a conventional CRT (cathode ray tube), LCD (liquid crystal display), LED (light emitting diode), plasma display, or other suitable display technology. In one embodiment, user input is received from input device 1208 (such as a keyboard, mouse, touchpad, microphone, etc.). In one embodiment, each of the above modules can be located on a single semiconductor platform to form a processing system.

[0092] In this specification, a single semiconductor platform may refer to an integrated circuit or chip based on a single semiconductor. It should be noted that the term single semiconductor platform may also refer to a multi-chip module with increased connectivity that simulates on-chip operations and provides substantial improvements over traditional central processing units ("CPUs") and bus implementations. Of course, the modules may also be placed separately or in various combinations of semiconductor platforms, depending on the needs of the user.

[0093] In one embodiment, a computer program in the form of machine-readable executable code or computer control logic algorithms is stored in the main memory 1204 and / or the secondary memory. According to one embodiment, if the computer program is executed by one or more processors, it enables the system 1200 to perform various functions. The memory 1204, memory, and / or any other memory are possible examples of computer-readable media. The secondary memory can refer to any suitable storage device or system, such as a hard drive and / or a removable storage drive, representing a floppy disk drive, a tape drive, an optical drive, a digital versatile disk ("DVD") drive, a recording device, or a universal serial bus ("USB") flash memory.

[0094] In one embodiment, the architecture and / or functionality of the various prior figures may be implemented in the context of a central processing unit 1202, a parallel processing system 1212, an integrated circuit capable of simultaneously having at least a portion of the capabilities of the central processing unit 1202 and the parallel processing system 1212, a chipset (e.g., a group of integrated circuits designed to operate and be sold as a unit that performs related functions, etc.), and any suitable combination of integrated circuits.

[0095] In one embodiment, the architecture and / or functionality of the various preceding figures may be implemented in the context of a general purpose computer system, a circuit board system, a game console system dedicated for entertainment purposes, a dedicated system, etc. In one embodiment, the computer system 1200 may take the form of a desktop computer, a laptop computer, a tablet computer, a server, a supercomputer, a smartphone (e.g., wireless, handheld device), a personal digital assistant ("PDA"), a digital camera, a vehicle, a head mounted display, a handheld electronic device, a mobile telephone device, a television, a workstation, a game console, an embedded system, and / or any other type of logic.

[0096] In one embodiment, the parallel processing system 1212 includes multiple PPUs 1214 and associated memory 1216. In one embodiment, the PPUs are connected to a host processor or other peripheral devices via an interconnect 1218 and a switch 1220 or multiplexer. In one embodiment, the parallel processing system 1212 distributes computational tasks across parallelizable PPUs 1214, for example, as part of distributing computational tasks across multiple GPU thread blocks. In one embodiment, memory is shared and accessible (e.g., for read and / or write access) across some or all of the PPUs 1214, although such shared memory may incur a performance penalty due to the use of local memory and registers resident on the PPUs. In one embodiment, the operations of the PPUs 1214 are synchronized using a command such as __syncthreads(), which requires all threads in a block (e.g., executing across multiple PPUs 1214) to reach a certain execution point in the code before continuing.

[0097] Accordingly, the specification and drawings are to be regarded in an illustrative rather than a restrictive sense. It will, however, be evident that various modifications and changes may be made thereto without departing from the broader spirit and scope of the invention as set forth in the claims.

[0098] Other variations are within the spirit of the present disclosure. Thus, while the disclosed technology is susceptible to various modifications and alternative constructions, certain illustrated embodiments thereof are shown in the drawings and have been described above in detail. However, it should be understood that there is no intention to limit the invention to one or more specific forms disclosed, but on the contrary, it is intended to cover all modifications, alternative constructions, and equivalents that fall within the spirit and scope of the invention as defined by the appended claims.

[0099] In the context of describing the disclosed embodiments (particularly in the context of the following claims), the use of the terms "a" and "an" and "said" and similar references should be interpreted as including the singular and the plural, unless otherwise stated herein or the context clearly contradicts. Unless otherwise stated, the terms "comprise", "have", "include" and "include" should be interpreted as open terms (i.e., meaning "including but not limited to"). The term "connected", when unmodified and referring to a physical connection, should be interpreted as being partially or completely contained in, attached to or connected together, even if something intervenes. Unless otherwise stated herein, the enumeration of ranges of values ​​herein is merely a shorthand expression for each individual value within the range as a separate reference, and each individual value is incorporated into the specification as if it were separately listed herein. Unless otherwise stated or contradicted by the context, the use of the terms "set" (such as "set of items") or "subset" should be interpreted as a non-empty set containing one or more members. In addition, unless otherwise stated or contradicted by the context, the term "subset" of a corresponding set does not necessarily mean a proper subset of the corresponding set, but the subset and the corresponding set may be equal.

[0100] Conjunctions (such as phrases of the form "at least one of A, B, and C" or "at least one of A, B, C," etc.), unless expressly stated otherwise or clearly contradicted by context, should be understood to generally indicate that the context of an item, term, etc. can be A, B, or C, or any non-empty subset of the set of A, B, and C. For example, in the example of a set with three members, the conjunctions "at least one of A, B, and C" and "at least one of A, B, and C" refer to any one of the following sets: {A}, {B}, {C}, {A, B}, {A, C}, {B, C}, {A, B, C}. Thus, such conjunctions do not generally imply that certain embodiments require the presence of at least one of A, at least one of B, and at least one of C. Furthermore, unless otherwise indicated or contradicted by context, the term "plurality" denotes plurality (e.g., "plurality of items" means a plurality of items). The number of items in a plurality of items is at least two, but can be more when expressly stated or indicated by context. Further, the phrase "based on" means "based at least in part on" rather than "based solely on" unless specified otherwise or the context clearly dictates otherwise.

[0101] Unless otherwise indicated herein or clearly contradicted by the context, the operations of the processes described herein may be performed in any suitable order. In one embodiment, processes such as those described herein (or variations and / or combinations thereof) are performed under the control of one or more computer systems configured with executable instructions and are implemented as code (e.g., executable instructions, one or more computer programs, or one or more applications) that is executed collectively on one or more processors by hardware or a combination thereof. In one embodiment, the code is stored on a computer-readable storage medium, for example, in the form of a computer program that includes a plurality of instructions that can be executed by one or more processors. In one embodiment, the machine-readable medium or computer-readable storage medium is a non-transitory machine-readable medium or computer-readable storage medium that does not include transient signals (e.g., propagated instantaneous electrical or electromagnetic transmissions), but includes non-transitory data storage circuitry (e.g., buffers, cache memories, and queues) within a transceiver of transient signals. In one embodiment, code (e.g., executable code or source code) is stored on a combination of one or more non-transitory computer-readable storage media, having executable instructions stored thereon (or other memory storing executable instructions) that, when executed by (e.g., as a result of being executed by) one or more processors of a computer system, cause the computer system to perform the operations described herein. In one embodiment, the combination of one or more non-transitory computer-readable storage media or machine-readable media includes information that, if used by one or more computer instructions, configures one or more processors of the computer system and further causes the one or more processors to perform the operations described herein. In one embodiment, the combination of non-transitory computer-readable storage media includes multiple non-transitory computer-readable storage media and one or more individual non-transitory storage media of the multiple non-transitory computer-readable storage media lacks all code, while the multiple non-transitory computer-readable storage media collectively store all code. In one embodiment, execution of the executable instructions causes different instructions to be executed by different processors, for example, a non-transitory computer-readable storage medium stores instructions and a main CPU executes some instructions while a graphics processor unit executes other instructions. In one embodiment, different components of the computer system have separate processors, and the different processors execute different subsets of the instructions.

[0102] Thus, in one embodiment, a computer system is configured to implement one or more services that individually or collectively perform the operations of the processes described herein, and these computer systems are configured with applicable hardware and / or software capable of performing the operations. Furthermore, a computer system implementing an embodiment of the present disclosure is a single device, and in another embodiment is a distributed computer system comprising multiple devices that operate differently, such that the distributed computer system performs the operations described herein and such that no single device performs all of the operations.

[0103] The use of any and all examples or exemplary language (e.g., "such as") provided herein is intended merely to better illuminate embodiments of the invention and does not limit the scope of the invention unless otherwise claimed. No language in the specification should be construed as indicating any non-claimed element essential to the practice of the invention.

[0104] The embodiments of the present disclosure described herein include the best mode known to the inventors for implementing the invention. Upon reading the above description, variations of these embodiments will be apparent to those of ordinary skill in the art. The inventors expect that skilled technicians will appropriately adopt such variations, and the inventors intend to implement the embodiments of the present disclosure in a manner different from that specifically described herein. Therefore, the scope of the present disclosure includes all modifications and equivalents of the subject matter permitted by applicable law as described in the appended claims. In addition, unless otherwise specified herein or the context clearly contradicts otherwise, any combination of the above elements in all their possible variations is included within the scope of the present disclosure.

[0105] All references (including publications, patent applications, and patents) cited herein are hereby incorporated by reference to the same extent as if each reference were individually and specifically indicated to be incorporated by reference and were set forth in its entirety herein.

[0106] In the description and claims, the terms "coupled" and "connected" and their derivatives may be used. It should be understood that these terms may not be synonyms for each other. On the contrary, in specific examples,

[0107] "Connected" or "coupled" may be used to indicate that two or more elements are in direct or indirect physical or electrical contact with each other. "Coupled" may also mean that two or more elements are not in direct contact with each other, but still cooperate or interact with each other.

[0108] Unless otherwise specified, it is understood that throughout the specification, terms such as "process," "calculate," "measure," "determine," etc., refer to actions and / or processes of a computer or computing system or similar electronic computing device that processes and / or converts physical quantities (such as electronic quantities) in registers and / or memories of the computing system into other data similarly represented as physical quantities in memories, registers, or other such information storage, transmission, or display devices of the computing system.

[0109] In a similar manner, the term "processor" may refer to any device or part of a device that processes electronic data from registers and / or memory and converts the electronic data into other electronic data that can be stored in registers and / or memory. As non-limiting examples, a "processor" may be a central processing unit (CPU) or a graphics processing unit (GPU). A "computing platform" may include one or more processors. As used herein, a "software" process may include, for example, software and / or hardware entities that perform work over time, such as tasks, threads, and intelligent agents. In addition, each process may refer to multiple processes for executing instructions sequentially or in parallel, continuously or intermittently. The terms "system" and "method" are used interchangeably herein so long as a system can embody one or more methods and these methods can be considered to be a system.

[0110] In this document, reference may be made to obtaining, acquiring, receiving or inputting analog or digital data into a subsystem, computer system or computer-implemented machine. The process of obtaining, acquiring, receiving or inputting analog or digital data can be accomplished in a variety of ways, such as by receiving data as a parameter to a function call or an application programming interface call. In some implementations, the process of obtaining, acquiring, receiving or inputting analog or digital data can be accomplished by transmitting data via a serial or parallel interface. In another implementation, the process of obtaining, acquiring, receiving or inputting analog or digital data can be accomplished by transmitting data from a providing entity to an acquiring entity via a computer network. Reference may also be made to providing, outputting, sending, issuing or presenting analog or digital data. In various examples, the process of providing, outputting, sending, issuing or presenting analog or digital data can be accomplished by using the data as an input or output parameter to a function call, a parameter to an application programming interface or an inter-process communication mechanism.

[0111] While the above discussion illustrates exemplary implementations of the technology, other architectures may be used to implement the functionality described and are intended to be within the scope of this disclosure. Furthermore, while specific allocations of responsibilities have been defined above for discussion purposes, the various functions and responsibilities may be allocated and divided in different ways depending on the specific circumstances.

[0112] Furthermore, although the subject matter has been described in language specific to structural features and / or methodological acts, it is to be understood that the subject matter defined in the appended claims is not necessarily limited to the specific features or acts described. Rather, the specific features and acts disclosed are example forms of implementing the claims.

Claims

1. A processor, comprising: One or more circuits for compressing information, wherein the one or more circuits are configured to: receiving input data comprising a first set of values; identifying, by one or more encoders, a first common value within each value of the first set of values; identifying, by the one or more encoders, a second common value within each value of one or more subset values ​​of the first set of values; and Compressed information including a third value corresponding to a mathematical relationship between the first and second common values ​​and each of the one or more subset values ​​is stored. 2 . The processor of claim 1 , wherein the first set of values ​​is a uniform grid of non-overlapping tiles from a matrix of floating point numbers. 3 . The processor of claim 2 , wherein each subset of the one or more subset values ​​is a smaller uniform grid of non-overlapping tiles from the first set of values. 4 . The processor of claim 1 , wherein the first common value is identified from a compressed set of the first set values, and the second common value is identified from a second compressed set of the one or more subset values.

5. The processor of claim 4, wherein the third value comprises information for the processor to decompress the compressed set of the first set of values. The processor of claim 1 , wherein the third value is stored in a storage device without storing the second common value. 7 . The processor of claim 6 , wherein the third value is stored in the memory device as an unsigned integer not exceeding a predetermined number of bits.

8. A computer processing method comprising: receiving input data comprising a first set of values; identifying, by one or more encoders, a first common value within each value of the first set of values; identifying, by the one or more encoders, a second common value within each value of one or more subset values ​​of the first set of values; as well as Compressed information including a third value is stored, the third value corresponding to an indication of how to determine a mathematical relationship between the first and second common values ​​and each of the one or more subset values.

9. The method of claim 8, wherein the first common value is identified from a compressed set of the first set values, and the second common value is identified from a second compressed set of the one or more subset values, wherein the third value and the compressed set of the first set values ​​are used to train a neural network.

10. The method of claim 8, wherein the indication is a difference between the second common value and the first common value.

11. The method of claim 8, wherein the first set of values ​​is a data set of elements of a matrix of floating point numbers, wherein a bit of each of the floating point numbers represents an exponent value. 12 . The method of claim 11 , wherein identifying the first common value further comprises identifying a first maximum exponent value common to the first set of values.

13. The method of claim 12, wherein after identifying the first maximum exponent value, the method further comprises: The second common value is identified by identifying a second largest exponent value common to the one or more subset values.

14. A machine-readable medium having stored thereon a set of instructions that, if executed by one or more processors, cause the one or more processors to at least: receiving input data comprising a first set of values; identifying, by one or more encoders, a first common value within each value of the first set of values; identifying, by the one or more encoders, a second common value within each value of one or more subset values ​​of the first set of values; as well as Compressed information including a third value corresponding to a mathematical relationship between the first and second common values ​​and each value within the one or more subset values ​​is stored.

15. The machine-readable medium of claim 14, wherein the first common value is stored and the second common value is not stored.

16. The machine-readable medium of claim 14, wherein the third value is a difference between the first common value and the second common value.

17. The machine-readable medium of claim 16, wherein the difference value is a proxy for identifying the second common value from the first common value.

18. The machine-readable medium of claim 16, wherein the difference value has at least the same or fewer bits than the second common value.

19. The machine-readable medium of claim 14, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to compress the first set of values ​​to identify the first common value.

20. The machine-readable medium of claim 14, wherein the set of instructions further comprises instructions that, if executed by the one or more processors, cause the one or more processors to compress the one or more subset values ​​to identify the second common value.

21. A machine-readable medium having an instruction set stored thereon, which, when executed by one or more processors, causes the one or more processors to: receive a data set, compress the data set using one or more encoders to determine a first value, and then compress a subset of the data set by replacing elements of the subset with smaller elements using the one or more encoders and store a second value, wherein the second value indicates how to calculate the replaced elements based at least in part on the smaller elements and the first value.

22. The machine-readable medium of claim 21, wherein the data set comprises a matrix of floating point numbers.

23. The machine-readable medium of claim 21, wherein the first value is determined by subtracting a common exponent value associated with all values ​​of a first set of values ​​in the data set from an exponent of each value of the first set of values.

24. The machine-readable medium of claim 21, wherein the second value is determined by subtracting a common exponent value associated with all values ​​of the subset of the data set from an exponent of each value of the subset.

25. The machine-readable medium of claim 21, wherein the second value uses less storage than the first value.

Images