Method for accelerating random precision sparse matrix multiplication and addition operation based on tensor core

Abstract

The invention discloses a tensor core-based method for accelerating arbitrary precision sparse matrix multiplication and addition operation, which comprises the following steps of: aiming at an architecture of a tensor core on A100 and a programming model thereof, combining program characteristics of deep learning parallel processing sparse matrix, combining an arbitrary precision sparse matrix multiplication and addition operation operator with parallel computing acceleration equipment, and performing parallel computing acceleration on the sparse matrix. The method comprises the following steps: decomposing a general matrix into a sparse matrix in a set format, partitioning the sparse matrix and assigning a corresponding parallel computing hardware unit for the sparse matrix, reducing memory conflicts and access delay by using a shared memory, calling a high-level language API function library to execute loading, computing and storing processes of a computing data volume, and finally performing specification on output data of a function. Therefore, the final result of the random-precision sparse matrix multiplication and addition operation on the parallel computing acceleration equipment is obtained. According to the method, parallel computing resources such as the tensor core on the A100 can be fully utilized, the program calculation amount and the communication overhead are effectively reduced, and the program operation performance is greatly improved.

Description

Technical Field

[0001] The present invention belongs to the field of parallel acceleration in computer system structure, and in particular relates to a method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores. Background Art

[0002] A100 GPU is a special hardware developed by NVIDIA for deep learning computing acceleration. It is an upgraded version of V100 GPU. Its biggest upgrade is the adoption of Ampere architecture, adding a powerful third-generation tensor core and new sparse features. The peak performance of ordinary FP16 TC calculations has been increased from the original 125TFLOPS to 312TFLOPS, and the peak performance of sparse FP16 TC calculations has been increased to 624TFLOPS. It has now been widely used in deep learning training. A100 GPU includes 108 Streaming Processors, referred to as SMs. Each SM contains 4 tensor cores and 64 FP32CUDA cores. Each A100 contains 432 tensor cores. Its shared memory and first-level data cache size is 192KB, and the video memory size is 80GB. Among them, the tensor core hardware unit supports high-level language API function library calls, and also supports the more low-level assembly language PTX (Parallel Thread Execution) instruction calls.

[0003] Arbitrary precision sparse matrix multiplication and addition refers to an important matrix operation in the field of deep learning and numerical computing. Its core idea is to multiply two matrices and accumulate the result with another matrix. This operation can be performed with arbitrary precision, such as using 16-bit integers, 8-bit integers, or 4-bit integers to meet the needs of different application scenarios. In deep learning, arbitrary precision sparse matrix multiplication and addition is often used in matrix multiplication operations of models such as convolutional neural networks (CNN) and recurrent neural networks (RNN), which can effectively accelerate the training and reasoning process of the model and improve computing efficiency and performance. By combining the characteristics of sparse matrices, arbitrary precision sparse matrix multiplication and addition can further reduce computing and storage overhead, which is of great significance for processing large-scale data and sparse matrices.

[0004] Since tensor cores are dedicated hardware units designed for parallel computing acceleration of deep learning, and their computing performance, communication latency, and memory access latency vary greatly, abstract allocation and division of computing data in tensor cores and pre-storage and pre-fetching become key steps in optimizing the performance of the computing process. How to combine the architecture and programming model of the tensor cores in A100 so that the parallel computing programs on them can fully utilize the powerful parallel processing capabilities of the tensor cores and perfectly maintain load balance between the tensor core computing capabilities and communication bandwidth has always been one of the main challenges in the field of parallel computing acceleration research. The implementation method of arbitrary precision sparse matrix multiplication and addition operations has certain guiding significance for the performance optimization of parallel computing acceleration programs on tensor cores. Summary of the invention

[0005] The purpose of the present invention is to provide a method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores, which fully utilizes the low-precision parallel computing resources on A100 by performing reasonable abstract division, alignment filling and pre-storage and pre-fetching of the computing data volume, and solves the problem of mismatch between the computing power of the tensor cores in A100 and the memory communication bandwidth between the host and the GPU.

[0006] To achieve the above object, the present invention adopts the following technical solutions:

[0007] A method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores, comprising the following steps:

[0008] Step 1, preprocess the original matrix data to generate a sparse matrix W2 in SB-BCSR format and a dense matrix A2 in bit continuous format;

[0009] Step 2, performing matrix operations based on the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit continuous format generated in step 1, obtaining a single-bit matrix operation result stored in the tensor core register, and performing matrix reduction according to the value of the single-bit matrix operation result to obtain a multi-bit matrix operation result;

[0010] Step 3: Based on the multi-bit matrix operation result obtained in step 2, for all threads in a block, each thread calculates the corresponding loading address of the multi-bit operation result, and copies the multi-bit operation result to the pre-specified video memory space in parallel.

[0011] A further improvement of the present invention is that the specific implementation method of step 1 is as follows:

[0012] Step 1.1, obtain the matrix data of the original matrices W and A in dense format on the CPU, and save the matrix data in the host memory. At the same time, preprocess the original matrix W, select the set vector length, and convert it into a sparse matrix W1 in a one-dimensional BCSR format suitable for tensor core calculation;

[0013] In step 1.2, according to the interface calculation operator of the tensor core and the number of bits of the sparse matrix W1 in the one-dimensional BCSR format, the one-dimensional BCSR format sparse matrix W1 is bit-decomposed and filled and aligned to generate a sparse matrix W2 in the SB-BCSR format; at the same time, according to the number of bits of the original matrix A, it is first bit-compressed to generate a dense matrix A1 in a bit-tight format, and then the dense matrix A1 in the bit-tight format is bit-decomposed to generate a dense matrix A2 in a bit-continuous format.

[0014] A further improvement of the present invention is that, in step 1.1, the dense format refers to a matrix that is stored in a video memory directly in the order of the matrix elements without compression or special encoding. In this format, each element occupies a storage unit in the memory, regardless of whether it is zero or non-zero; the BCSR format refers to a data format for storing sparse matrices. In the BCSR format, the matrix is ​​divided into rectangular blocks of fixed size, and each block is represented by a triple containing a row index, a column offset and a data value list. Specifically, the position of each block is represented by a row index and a column offset, and the content is represented by a data value list; the vector length refers to the block dimension of each block in the BCSR format in different directions. For the one-dimensional BCSR format, the vector length is a scalar, and for the two-dimensional BCSR format, the vector length is a vector. Specifically, the vector length determines the length of the data list in a triple.

[0015] A further improvement of the present invention is that in step 1.2, the tensor core refers to the tensor core in the A100 series GPU of NVIDIA Corporation;

[0016] The interface computing operator of the tensor core refers to the function interface provided by NVIDIA to call the tensor core hardware to perform matrix operations;

[0017] Bit decomposition is an operation that breaks a number into binary bits, which is used to encode or compress numbers in computer systems. In bit decomposition, a number is broken down into a series of binary bits, each bit representing a position value in the number. For a number N, R i is the ith binary value of N, and n is the number of binary digits of N; then N is expressed as:

[0018]

[0019] For m numbers N j , YesN j The i-th binary value of k ; The decomposition process is expressed as:

[0020]

[0021]

[0022]

[0023] Padding alignment refers to the operation of adding additional space or elements in the data structure or memory layout to meet the requirements of the set hardware or software, so that the storage location or layout of the data meets the set alignment requirements;

[0024] The SB-BCSR format refers to an improved format specifically designed to accommodate Tensor Core hardware and associated with the BCSR format;

[0025] A dense matrix in bit-continuous format refers to an equivalent matrix formed by bit-decomposing a dense matrix with continuous data. It continuously stores the binary bits of the same position of different data in the memory, and the binary bits of different positions of the same data are stored in a scattered manner. For a row L0 in the matrix, it can be regarded as an ordered set {N 0 , N 1 ,…,N n-1}, according to the decomposition process, the row is transformed into L1 after bit decomposition, and the ordered set is represented as {S 0 , S 1 , …, S m-1}.

[0026] A further improvement of the present invention is that the specific implementation method of step 2 is as follows:

[0027] Step 2.1, allocate space for the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit-continuous format on the GPU memory, and use the corresponding CUDA function to copy the matrix data from the host to the GPU; divide the matrix data into multiple blocks, and all threads in a block jointly and concurrently copy the corresponding matrix data that need to be calculated by themselves from the GPU memory to the shared memory of the SM; create and initialize member objects a, b and c associated with the registers in the tensor core on the GPU, which are used to store the sparse matrix W2 in SB-BCSR format, the dense matrix A2 in bit-continuous format and the single-bit operation results respectively;

[0028] Step 2.2, for all threads in a block, calculate the storage address of the matrix data that needs to be copied, call the load_matrix_sync() interface function, and load the matrix data from the SM's shared memory to the registers of the tensor core through the member objects a and b; specify the operation type as binary bitwise AND operation, call the CUDA operation function, and perform thread synchronization at the end of the program flow in a block operation; repeat the above process until the operation tasks of all blocks are completed and all single-bit results are stored in the registers of the tensor core;

[0029] Step 2.3: For all threads in a block, each thread calculates the loading address of the corresponding single-bit operation result, calls the store_matrix_sync() interface function, and uses the member object c to load the single-bit operation result from the register of the tensor core to the shared memory of the SM in parallel; the single-bit operation result in the shared memory is reduced, and the multi-bit operation result is generated by combining according to the bit number of the original matrices W and A.

[0030] A further improvement of the present invention is that, in step 2.1, the corresponding CUDA function refers to the interface function provided by NVIDIA for allocating address space and performing data migration on the GPU, including two key interface functions of cudaMalloc() and cudaMemcpy(); dividing the matrix data into multiple blocks means abstractly allocating the matrix data to different blocks according to the CUDA programming model, and each block performs the actual calculation operation; the member object associated with the register in the tensor core refers to the abstract data type provided by NVIDIA for the convenience of programmers to call the tensor core hardware, which includes three types, the matrix_a type indicates that the object is the first multiplicand, the matrix_b type indicates that the object is the second multiplicand, and the accumulator type indicates that the object is a source or target accumulator. When declaring this type of object, the calculation dimension and storage layout format are specified, and the object is modified, exported or performs operations through the four functions of load_matrix_sync(), store_matrix_sync(), fill_fragment() and mma_sync().

[0031] A further improvement of the present invention is that, in step 2.2, the address of the matrix data to be copied is calculated by the following formula, where A load Indicates the address of the matrix data to be copied, A shmemIndicates the starting address of the shared memory, warpId indicates the id number of the warp in the CUDA programming model, M indicates the height of the first matrix, Type indicates the storage type of the data, and the sizeof() function is used to calculate the space size of the data type in bytes:

[0032]

[0033] subjecttoi=1,2

[0034] Specifying the calculation type as binary bit AND calculation means setting the precision to b1 and the bit operation type to bmmaBitOpAND when calling the mma_sync() function for calculation; performing thread synchronization means calling the __syncthreads() interface function to synchronize all threads in the same block after completing a round of calculation to ensure that all calculation tasks in the previous stage have been completed.

[0035] A further improvement of the present invention is that, in step 2.3, the address for loading the single-bit operation result is calculated as follows, where A Sstore Indicates the address of the single-bit calculation result to be loaded, A shmem Indicates the starting address of the shared memory, warpId indicates the id number of the warp in the CUDA programming model, laneId indicates the id number of the thread in a warp, Type indicates the storage type of the data, and the sizeof() function is used to calculate the space size of the data type in bytes:

[0036]

[0037]

[0038] Combining and generating multi-bit operation results refers to combining the single-bit operation results output by the tensor core according to their bit positions to generate the final multi-bit operation results. The combination generation process is the inverse process of bit decomposition.

[0039] A further improvement of the present invention is that, in step 3, the pre-designated video memory space refers to the video memory space allocated in advance for storing the result during initialization, and the size of the video memory space can accommodate the final multi-bit operation result.

[0040] Compared with the prior art, the present invention has the following technical effects:

[0041] The present invention provides a method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores. The method aims at the architecture and programming model of tensor cores on A100, combines the program characteristics of deep learning parallel processing of sparse matrices, combines arbitrary precision sparse matrix multiplication and addition operators with parallel computing acceleration devices, decomposes general matrices into sparse matrices of set formats, divides the sparse matrix into blocks and specifies corresponding parallel computing hardware units for them, uses shared memory to reduce memory conflicts and access delays, calls high-level language API function libraries to perform the loading, calculation and storage processes of the calculation data volume, and finally reduces the output data of the function, thereby obtaining the final result of arbitrary precision sparse matrix multiplication and addition operations on the parallel computing acceleration device. This method provides a unique implementation idea and method for developers who develop, transplant or optimize parallel computing programs for parallel sparse matrix operation operators on tensor cores, can make full use of parallel computing resources such as tensor cores on A100, effectively reduce the program calculation amount and communication overhead, and greatly improve the program running performance. BRIEF DESCRIPTION OF THE DRAWINGS

[0042] Figure 1 It is a flow chart of the present invention. DETAILED DESCRIPTION

[0043] The present invention is further described below in conjunction with the accompanying drawings:

[0044] The present invention provides a method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores. The method targets the architecture of tensor cores on A100 and its programming model, combines the program characteristics of deep learning parallel processing of sparse matrices, combines arbitrary precision sparse matrix multiplication and addition operators with parallel computing acceleration devices, decomposes general matrices into sparse matrices of a set format, divides the sparse matrices into blocks and specifies corresponding parallel computing hardware units for them, uses shared memory to reduce memory conflicts and access delays, calls a high-level language API function library to execute the loading, calculation and storage processes of the calculation data volume, and finally reduces the output data of the function, thereby obtaining the final result of the arbitrary precision sparse matrix multiplication and addition operations on the parallel computing acceleration device.

[0045] See also Figure 1 The present invention provides a method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores, comprising the following steps:

[0046] Step 1.1, obtain the original matrices W and A in dense format, and convert them into a sparse matrix W1 in one-dimensional BCSR format suitable for tensor core calculation. W should be a matrix that meets certain sparse format and sparsity requirements through a special pruning strategy, and the vector length should be 2 n , and finally output three sets of data, including row index, column offset and data value.

[0047] Step 1.2, further preprocess the sparse format matrix W1 obtained in step 1.1 into a data format more suitable for tensor core calculations, where tensor core refers to a hardware unit used to perform matrix multiplication and accumulation operations, usually used for matrix operations in deep learning and numerical calculations. Here, it specifically refers to the tensor core in the A100 series GPU of NVIDIA. The architecture of A100 is a parallel computing acceleration device based on the Ampere architecture, including a total of 108 SMs, each of which contains 4 tensor cores and 64 FP32 CUDA cores. Each A100 contains 432 tensor cores, and its shared memory and first-level data cache size is 192KB, and the video memory size is 80GB.

[0048] There are three things you should include when actually testing on Tensor Cores:

[0049] 1) For hardware devices such as tensor cores, each group of computing data should be abstracted into a parallel computing framework related to the number of threads, and different tensor cores should be assigned to each group of computing data to maximize the computing efficiency of the device;

[0050] 2) Consider the impact of resource competition and use data padding, interleaved addressing, and pipelines to reduce addressing conflicts and communication delays;

[0051] 3) For each set of calculation data, multiple consecutive tests are required until the obtained running time meets the set accuracy requirements to eliminate the influence of randomness and show objective rules. For example, at a confidence level of 98%, the accuracy is ensured to reach 0.01.

[0052] The interface calculation operator of the tensor core refers to the function interface provided by NVIDIA to call the tensor core hardware to perform matrix operations. These operators are mainly used to accelerate matrix multiplication and accumulation operations in deep learning, and are optimized for common matrix operations in deep learning. They use the parallel computing capabilities and mixed precision computing characteristics of the tensor core hardware to effectively improve computing performance and efficiency. The interface calculation operator of the tensor core supports calculations of multiple precisions, including 8bit, 4bit, and 1bit. In matrix multiplication operations, the tensor core can perform multiple calculation operations at the same time and accumulate the results into the target matrix, which can significantly improve the computing throughput and parallelism.

[0053] According to the calculation operator of the tensor core and the number of bits of the W1 matrix, the W1 matrix is ​​first padded and aligned, and the blank space is padded with zeros. Then, the aligned W1 matrix is ​​bit-decomposed to convert it into a bit-continuous matrix, which is called the sparse matrix W2 in SB-BCSR format. In the bit decomposition, each bit represents a position or bit value in the number. For a number N, R iis the ith binary value of N, and n is the number of binary digits of N; then N can be expressed as

[0054]

[0055] For m numbers N j , YesN j The i-th binary value of k ; Then the bit decomposition process can be expressed as

[0056]

[0057]

[0058]

[0059] At the same time, according to the number of bits of matrix A, it is first bit compressed to generate a dense matrix A1 in a bit-tight format, and then the matrix A1 is bit-decomposed to generate a dense matrix A2 in a bit-continuous format.

[0060] Padding alignment refers to the operation of adding extra space or elements in the data structure or memory layout to meet the requirements of the set hardware or software, so that the storage location or layout of the data meets the set alignment requirements. The alignment here is to meet the bit alignment requirements of the binary and operation calculation operator of the tensor core. The alignment length is 128 bits, and the empty space is filled with zero elements;

[0061] The SB-BCSR format is an improved format specifically designed to adapt to the Tensor Core hardware and associated with the BCSR format. The main improvement is to perform bit decomposition and padding alignment operations on the data in each block after the block operation. It is a bit-continuous sparse matrix storage format.

[0062] A dense matrix in bit-continuous format refers to an equivalent matrix formed by bit-decomposing a dense matrix with continuous data. It continuously stores the binary bits of the same position of different data in the memory, and the binary bits of different positions of the same data are stored in a scattered manner. For a row L0 in the matrix, it can be regarded as an ordered set {N 0 , N 1 ,…,N n-1}, according to the above decomposition process, after bit decomposition, the row is transformed into L1, which can be represented by an ordered set {S 0 , S 1 , …, S m-1}.

[0063] Step 2.1, use the sparse matrix W2 and dense matrix A2 obtained in step 1.2 to initialize the calculation process, allocate space for the W2 and A2 matrices on the GPU's video memory, and use key interface functions such as cudaMalloc() and cudaMemcpy() to copy the matrix data from the host to the GPU; divide the matrix data into multiple blocks according to the CUDA programming model, and different blocks complete the actual calculation operations in different areas. The threads in the same block jointly and parallelly copy the corresponding matrix data that needs to be calculated from the GPU video memory to the shared memory of the SM; create and initialize member objects a, b, and c associated with the registers in the tensor core on the GPU, which are used to store W2, A2, and single-bit operation results, respectively.

[0064] Step 2.2, use the member objects a and b associated with the registers in the tensor core obtained in step 2.1 to start the tensor core to execute the calculation process. For all threads in a block, calculate the address of the matrix data to be copied according to the following formula, where A load Indicates the address of the matrix data to be copied, A shmem Indicates the starting address of the shared memory, warpId indicates the id number of the warp in the CUDA programming model, M indicates the height of the first matrix, Type indicates the storage type of the data, and the sizeof() function is used to calculate the space size of the data type in bytes.

[0065]

[0066] subjecttoi=1,2

[0067] Call the load_matrix_sync() interface function to load the matrix data from the shared memory of the SM into the a and b member objects associated with the registers in the tensor core; specify the calculation type as binary bit AND calculation, set the precision to b1, set the bit operation type to bmmaBitOpAND, and call the CUDA calculation function to perform the calculation; call the __syncthreads() interface function to synchronize all threads in the same block to achieve thread synchronization and ensure that all calculation tasks in the previous stage have been completed; repeat the above process until the calculation tasks in all areas of the block have been completed. At this time, the single-bit calculation result is stored in the object c, and the data type is a 32-bit integer.

[0068] Step 2.3, use step 2.2 to obtain the object c that stores the single-bit operation result, and reduce it to a multi-bit operation result. For all threads in a block, calculate the address for loading the single-bit operation result according to the following formula, where A SstoreIndicates the address of the single-bit operation result to be loaded, A shmem Indicates the starting address of the shared memory, warpId indicates the id number of the warp in the CUDA programming model, laneId indicates the id number of the thread in a warp, Type indicates the storage type of the data, and the sizeof() function is used to calculate the space size of the data type in bytes.

[0069]

[0070] Call the store_matrix_sync() interface function to load the operation result from the c object associated with the register in the tensor core to the shared memory of the SM; reduce the single-bit operation result in the shared memory, combine and generate multi-bit operation results according to the bit number of the W and A matrices, and store them in the register of the CUDA core. The combination generation process is the inverse process of the bit decomposition operation.

[0071] Step 3, for all threads in a block, calculate the address for loading the multi-bit operation result, and copy the multi-bit operation result to a pre-specified video memory space. The size of the video memory space should be allocated in advance during initialization and be able to accommodate the final multi-bit operation result.

[0072] After completing the above steps, the result of the specified precision sparse matrix multiplication and addition operation is obtained. Later, according to the application requirements, the input auxiliary program can be added or the result can be output to other operators and operator fusion can be performed. So far, the implementation method of accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores has been completed.

[0073] Although the present invention has been described in detail above with general descriptions and specific embodiments, it is obvious to those skilled in the art that some modifications or improvements may be made thereto based on the present invention. Therefore, these modifications or improvements made without departing from the spirit of the present invention all fall within the scope of protection claimed by the present invention.

Claims

1. A method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores, characterized in that: The following steps are involved: Step 1, preprocess the original matrix data to generate a sparse matrix W2 in SB-BCSR format and a dense matrix A2 in bit continuous format; Step 2, performing matrix operations based on the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit continuous format generated in step 1, obtaining a single-bit matrix operation result stored in the tensor core register, and performing matrix reduction according to the value of the single-bit matrix operation result to obtain a multi-bit matrix operation result; Step 3: Based on the multi-bit matrix operation result obtained in step 2, for all threads in a block, each thread calculates the corresponding loading address of the multi-bit operation result, and copies the multi-bit operation result to the pre-specified video memory space in parallel.

2. A method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 1, characterized in that: The specific implementation method of step 1 is as follows: Step 1.1, obtain the matrix data of the original matrices W and A in dense format on the CPU, and save the matrix data in the host memory. At the same time, preprocess the original matrix W, select the set vector length, and convert it into a sparse matrix W1 in a one-dimensional BCSR format suitable for tensor core calculation; In step 1.2, according to the interface calculation operator of the tensor core and the number of bits of the sparse matrix W1 in the one-dimensional BCSR format, the one-dimensional BCSR format sparse matrix W1 is bit-decomposed and filled and aligned to generate a sparse matrix W2 in the SB-BCSR format; at the same time, according to the number of bits of the original matrix A, it is first bit-compressed to generate a dense matrix A1 in a bit-tight format, and then the dense matrix A1 in the bit-tight format is bit-decomposed to generate a dense matrix A2 in a bit-continuous format.

3. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 2, characterized in that: In step 1.1, dense format means that a matrix is ​​stored in the video memory in the order of matrix elements without compression or special encoding. In this format, each element occupies a storage unit in the memory, regardless of whether it is zero or non-zero; The BCSR format refers to a data format for storing sparse matrices. In the BCSR format, the matrix is ​​divided into rectangular blocks of fixed size, and each block is represented by a triple containing row index, column offset and data value list. Specifically, the position of each block is represented by row index and column offset, and the content is represented by data value list. The vector length refers to the block dimension of each block in different directions in the BCSR format. For the one-dimensional BCSR format, the vector length is a scalar, and for the two-dimensional BCSR format, the vector length is a vector. Specifically, the vector length determines the length of the data list in a triple.

4. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 2, characterized in that: In step 1.2, tensor core refers to the tensor core in the A100 series GPU of NVIDIA; The interface computing operator of the tensor core refers to the function interface provided by NVIDIA to call the tensor core hardware to perform matrix operations; Bit decomposition is an operation that breaks a number into binary bits, which is used to encode or compress numbers in computer systems. In bit decomposition, a number is broken down into a series of binary bits, each bit representing a position value in the number. For a number N, R i is the ith binary value of N, and n is the number of binary digits of N; then N is expressed as: For m numbers N j , YesN j The i-th binary value of k ; The decomposition process is expressed as: Padding alignment refers to the operation of adding additional space or elements in the data structure or memory layout to meet the requirements of the set hardware or software, so that the storage location or layout of the data meets the set alignment requirements; The SB-BCSR format refers to an improved format specifically designed to accommodate Tensor Core hardware and associated with the BCSR format; A dense matrix in bit-continuous format refers to an equivalent matrix formed by bit-decomposing a dense matrix with continuous data. It continuously stores the binary bits of the same position of different data in the memory, and the binary bits of different positions of the same data are stored in a scattered manner. For a row L0 in the matrix, it can be regarded as an ordered set {N 0 , N 1 ,…,N n-1 }, according to the decomposition process, the row is transformed into L1 after bit decomposition, and is represented by an ordered set as {S 0 , S 1 , …, S m-1 }.

5. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 2, characterized in that: The specific implementation method of step 2 is as follows: Step 2.1, allocate space for the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit-continuous format on the GPU memory, and use the corresponding CUDA function to copy the matrix data from the host to the GPU; divide the matrix data into multiple blocks, and all threads in a block jointly and concurrently copy the corresponding matrix data that need to be calculated by themselves from the GPU memory to the shared memory of the SM; create and initialize member objects a, b and c associated with the registers in the tensor core on the GPU, which are used to store the sparse matrix W2 in SB-BCSR format, the dense matrix A2 in bit-continuous format and the single-bit operation results respectively; Step 2.2, for all threads in a block, calculate the storage address of the matrix data that needs to be copied, call the load_matrix_sync() interface function, and load the matrix data from the SM's shared memory to the registers of the tensor core through the member objects a and b; specify the operation type as binary bitwise AND operation, call the CUDA operation function, and perform thread synchronization at the end of the program flow in a block operation; repeat the above process until the operation tasks of all blocks are completed and all single-bit results are stored in the registers of the tensor core; Step 2.3: For all threads in a block, each thread calculates the loading address of the corresponding single-bit operation result, calls the store_matrix_sync() interface function, and uses the member object c to load the single-bit operation result from the register of the tensor core to the shared memory of the SM in parallel; the single-bit operation result in the shared memory is reduced, and the multi-bit operation result is generated by combining according to the bit number of the original matrices W and A.

6. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 5, characterized in that: In step 2.1, the corresponding CUDA function refers to the interface function provided by NVIDIA for allocating address space and performing data migration on the GPU, including the two key interface functions cudaMalloc() and cudaMemcpy(); dividing the matrix data into multiple blocks means abstractly allocating the matrix data to different blocks according to the CUDA programming model, and each block performs the actual calculation operation; the member object associated with the register in the tensor core refers to the abstract data type provided by NVIDIA for the convenience of programmers to call the tensor core hardware, which contains three types, the matrix_a type indicates that the object is the first multiplicand, the matrix_b type indicates that the object is the second multiplicand, and the accumulator type indicates that the object is the source or target accumulator. When declaring this type of object, the calculation dimension and storage layout format are specified, and the object is modified, exported or executed through the four functions of load_matrix_sync(), store_matrix_sync(), fill_fragment() and mma_sync().

7. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 6, characterized in that: In step 2.2, the address of the matrix data to be copied is calculated as follows, where A load Indicates the address of the matrix data to be copied, A shmem Indicates the starting address of the shared memory, warpId indicates the id number of the warp in the CUDA programming model, M indicates the height of the first matrix, Type indicates the storage type of the data, and the sizeof() function is used to calculate the space size of the data type in bytes: subjecttoi=1,2 Specifying the calculation type as binary bit AND calculation means setting the precision to b1 and the bit operation type to bmmaBitOpAND when calling the mma_sync() function for calculation; performing thread synchronization means calling the __syncthreads() interface function to synchronize all threads in the same block after completing a round of calculation to ensure that all calculation tasks in the previous stage have been completed.

8. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 7, characterized in that: In step 2.3, the address for loading the single-bit operation result is calculated as follows, where A Sstore Indicates the address of the single-bit calculation result to be loaded, A shmem Indicates the starting address of the shared memory, warpId indicates the id number of the warp in the CUDA programming model, laneId indicates the id number of the thread in a warp, Type indicates the storage type of the data, and the sizeof() function is used to calculate the space size of the data type in bytes: Combining and generating multi-bit operation results refers to combining the single-bit operation results output by the tensor core according to their bit positions to generate the final multi-bit operation results. The combination generation process is the inverse process of bit decomposition.

9. The method for accelerating arbitrary precision sparse matrix multiplication and addition operations based on tensor cores according to claim 1, characterized in that: In step 3, the pre-designated video memory space refers to the video memory space allocated in advance for storing the result during initialization, and the size of the video memory space can accommodate the final multi-bit operation result.

Images