A method for accelerating arbitrary precision sparse matrix multiplication based on tensor cores

Abstract

The application discloses a method for accelerating arbitrary precision sparse matrix multiplication and addition operation based on a tensor core, which is in combination with deep learning and the program characteristics of parallel processing of sparse matrices, and is in combination of an arbitrary precision sparse matrix multiplication and addition operation operator and a parallel computing acceleration device, decomposes a general matrix into a sparse matrix of a set format, blocks the sparse matrix and specifies a corresponding parallel computing hardware unit for the sparse matrix, uses shared memory to reduce memory conflict and access delay, calls a high-level language API function library to execute the loading, calculation and storage process of the calculation data, and finally reduces the output data of the function, so that the final result of the arbitrary precision sparse matrix multiplication and addition operation on the parallel computing acceleration device is obtained. The application can fully utilize the parallel computing resources such as the tensor core on the A100, effectively reduce the program calculation amount and communication overhead, and greatly improve the program running performance.

Description

TECHNICAL FIELD

[0001] The present application belongs to the field of parallelization acceleration in computer system structure, and particularly relates to a method for accelerating arbitrary precision sparse matrix multiplication based on tensor core. BACKGROUND

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

[0003] Arbitrary precision sparse matrix multiplication refers to an important matrix operation in the field of deep learning and numerical computing. The core idea is to perform multiplication on two matrices and accumulate the result with another matrix. This operation can be performed at 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 is commonly used in matrix multiplication of convolutional neural network (CNN) and recurrent neural network (RNN) models, which can effectively accelerate the training and inference process of the model, improve the computing efficiency and performance. By combining the characteristics of sparse matrices, arbitrary precision sparse matrix multiplication can further reduce the computational and storage overhead, which is of great significance for processing large-scale data and sparse matrices.

[0004] Due to the tensor core is a special hardware unit designed for deep learning parallel computing acceleration, and its computing performance, communication delay and memory access delay have great performance difference, so the abstract allocation division and prefetch of the computing data volume in the tensor core becomes the key step of performance optimization of the computing process. How to combine the architecture and programming model of the tensor core in A100 to make the parallel computing program on it fully exert the powerful parallel processing capability of the tensor core and perfectly maintain the load balance between the computing capability of the tensor core and the communication bandwidth has been one of the main challenges in the field of parallel computing acceleration. The implementation method of the arbitrary precision sparse matrix multiplication and addition operation has certain guiding significance for the performance optimization of the parallel computing acceleration program on the tensor core. SUMMARY

[0005] The purpose of the present application is to provide a method for accelerating arbitrary precision sparse matrix multiplication and addition operation based on tensor core, which reasonably abstracts and divides the computing data volume, aligns and fills, and pre-stores and prefetches, fully utilizes the low-precision parallel computing resources on A100, and solves the problem of mismatch between the computing capability of the tensor core in A100 and the memory communication bandwidth between the host and GPU.

[0006] To achieve the above purpose, the technical scheme adopted by the present application is as follows:

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

[0008] Step 1, preprocessing 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 operation based on the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit continuous format generated in step 1 to obtain single-bit matrix operation results stored in the tensor core register, performing matrix reduction according to the value of the single-bit matrix operation results to obtain multi-bit matrix operation results;

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

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

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

[0013] Step 1.2, according to the interface calculation operator of the tensor core and the bit number of the one-dimensional BCSR format sparse matrix W1, perform bit decomposition and padding alignment operation on the one-dimensional BCSR format sparse matrix W1 to generate the SB-BCSR format sparse matrix W2; at the same time, according to the bit number of the original matrix A, first perform bit compression to generate the bit dense format dense matrix A1, and then perform bit decomposition on the bit dense format dense matrix A1 to generate the bit continuous format dense matrix A2.

[0014] The further improvement of the application is that in step 1.1, the dense format means that a matrix is stored in a memory in the order of matrix elements without compression or special encoding, and in this format, each element occupies a memory unit, whether it is zero or nonzero; the BCSR format means a data format for storing sparse matrices, in which the matrix is divided into fixed-size rectangular blocks, 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 means the block dimension in different directions in BCSR format, for one-dimensional BCSR format, the vector length is a scalar, for 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] The further improvement of the application is that in step 1.2, the tensor core refers to the tensor core in the A100 series GPU of NVIDIA company;

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

[0017] Bit decomposition refers to an operation of splitting a number into binary bits, which is used for encoding or compressing numbers in computer systems; in bit decomposition, a number is decomposed into a series of binary bits, each bit representing a bit value in the number; for a number N, R i is the i-th binary value of N, n is the binary bit number of N; then N is represented as:

[0018]

[0019] For the number N j , is the i-th bit binary value of N j , for the decomposed value S k ; the decomposition process is represented as:

[0020]

[0021]

[0022]

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

[0024] SB-BCSR format refers to an improved format specially designed for tensor core hardware and associated with BCSR format;

[0025] The bit-continuous format dense matrix refers to an equivalent matrix formed by bit decomposition of a data-continuous dense matrix, which stores the binary bits of different data in the same position in memory continuously, and the binary bits of the same data in different positions are stored dispersedly, for a certain 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 as an ordered set {S 0 , S 1 , …, S m-1}.

[0026] The further improvement of the application is that the specific implementation method of step 2 is as follows:

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

[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, load the matrix data from the shared memory of the SM into the registers of the tensor core through the member objects a and b; specify the operation type as binary bit AND operation, call the CUDA operation function, and perform thread synchronization at the end of the program stream 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, calculate the loading address of the corresponding single-bit operation result respectively, call the store_matrix_sync() interface function, and parallelly load the single-bit operation result from the registers of the tensor core into the shared memory of the SM through the member object c; reduce the single-bit operation result in the shared memory, and generate a multi-bit operation result according to the bit number of the original matrix W and A.

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

[0031] Further improvement of the application is that, in step 2.2, the address of the matrix data that needs to be copied refers to the following formula, wherein A load represents the address of the matrix data that needs to be copied, A shmemrepresents the start address of the shared memory, warpId represents the id number of the warp in the CUDA programming model, M represents the height of the first matrix, Type represents 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] subject to i = 1,2

[0034] The calculation type is specified as a binary bit AND calculation, which means that when the mma_sync() function is called for operation, the precision is set to b1 and the bit operation type is set to bmmaBitOpAND; thread synchronization is performed, which means that after completing a round of calculation, the __syncthreads() interface function is called to synchronize all threads in the same block, and it is ensured that all calculation tasks in the previous stage have been completed.

[0035] The further improvement of the application is that in step 2.3, the address for loading the single-bit operation result is calculated according to the following formula, wherein A Sstore represents the address of the single-bit operation result to be loaded, A shmem represents the start address of the shared memory, warpId represents the id number of the warp in the CUDA programming model, laneId represents the id number of the thread in a warp, Type represents 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] The combination of the single-bit operation results output by the tensor core generates the final multi-bit operation result according to the bit position of the single-bit operation result, and the combination generation process is the inverse process of the bit decomposition.

[0039] The further improvement of the application is that in step 3, the pre-designated video memory space is a video memory space allocated in advance for storing the result when the initialization is performed, and the size of the video memory space can accommodate the final multi-bit operation result.

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

[0041] The application provides a method for accelerating arbitrary precision sparse matrix multiplication operation based on a tensor core, which is in view of the architecture and programming model of the tensor core on the A100, in combination with the program characteristics of deep learning parallel processing of sparse matrices, combines the arbitrary precision sparse matrix multiplication operator with a parallel computing acceleration device, decomposes a general matrix into a sparse matrix of a set format, blocks the sparse matrix and specifies a corresponding parallel computing hardware unit for the sparse matrix, uses shared memory to reduce memory conflicts and access latency, calls a high-level language API function library to perform the loading, calculation and storage process of the calculation data, and finally reduces the output data of the function, so as to obtain the final result of the arbitrary precision sparse matrix multiplication operation on the parallel computing acceleration device. BRIEF DESCRIPTION OF DRAWINGS

[0042] Figure 1 The flowchart of the application. DETAILED DESCRIPTION

[0043] The application is further described below in combination with the drawings:

[0044] The application provides a method for accelerating arbitrary precision sparse matrix multiplication operation based on a tensor core, which is in view of the architecture and programming model of the tensor core on the A100, in combination with the program characteristics of deep learning parallel processing of sparse matrices, combines the arbitrary precision sparse matrix multiplication operator with a parallel computing acceleration device, decomposes a general matrix into a sparse matrix of a set format, blocks the sparse matrix and specifies a corresponding parallel computing hardware unit for the sparse matrix, uses shared memory to reduce memory conflicts and access latency, calls a high-level language API function library to perform the loading, calculation and storage process of the calculation data, and finally reduces the output data of the function, so as to obtain the final result of the arbitrary precision sparse matrix multiplication operation on the parallel computing acceleration device.

[0045] Please refer to Figure 1 The application provides a method for accelerating arbitrary precision sparse matrix multiplication operation based on a tensor core, which comprises the following steps:

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

[0047] Step 1.2, the sparse format matrix W1 obtained in step 1.1 is further preprocessed into a data format more suitable for tensor core calculation, where tensor core refers to a hardware unit used for matrix multiplication operations and accumulation operations, commonly used in matrix operations in deep learning and numerical calculations, here 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 Ampere architecture, including 108 SMs, each SM contains 4 tensor cores and 64 FP32 CUDA cores, each A100 contains 432 tensor cores, with shared memory and level 1 data cache size of 192KB, and memory size of 80GB.

[0048] Three points should be included when actually testing on the tensor core:

[0049] 1) For this hardware device, the tensor core, the amount of computation data for each group should be abstracted from the parallel computing framework related to the number of threads, and different tensor cores should be allocated for each group of computation data to maximize the computing efficiency of the device;

[0050] 2) Considering the impact of resource competition, data padding, interleaved addressing and pipelining are used to reduce addressing conflicts and communication delays;

[0051] 3) For each group of computation data, continuous multiple tests need to be conducted until the running time obtained meets the set accuracy requirement, to exclude the influence of randomness and show the objective law, for example, at a confidence level of 98%, ensure the accuracy to 0.01.

[0052] 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 operation. 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. It uses the parallel computing capability and mixed precision computing characteristics of the tensor core hardware to effectively improve the computing performance and efficiency. The interface computing operator of the tensor core supports multiple precision calculations, including 8bit, 4bit and 1bit, etc. In matrix multiplication operations, the tensor core can simultaneously perform multiple computing operations and accumulate the results into the target matrix, which can significantly improve the computing throughput and parallelism.

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

[0054]

[0055] For m numbers N j , is the i-th bit of N j , for n decomposed value S k ; the bit decomposition process can be expressed as

[0056]

[0057]

[0058]

[0059] At the same time, according to the bit number of the A matrix, it is first 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 alignment requirements set, and here the alignment is to meet the bit number alignment requirements of the binary and operation calculation operator of the tensor core, and the alignment length is 128 bits, and the spare space is filled with zero elements.

[0061] SB-BCSR format refers to a specially designed and associated improved format of BCSR format to adapt to the tensor core hardware, and the main improvement measure is to perform bit decomposition and padding alignment operation on the data in each block after block operation, which is a bit-continuous sparse matrix storage format.

[0062] The bit-continuous format dense matrix refers to an equivalent matrix formed by bit-decomposing a continuous dense matrix, which continuously stores the binary bits of different data in the same position in the memory, and the binary bits of the same data in different positions are stored dispersedly. For a certain 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, the row is transformed into L1 after bit decomposition, which can be expressed as an ordered set {S 0 , S 1 , …, S m-1}.

[0063] Step 2.1, using the sparse matrix W2 and the dense matrix A2 obtained in step 1.2 to initialize the calculation process, allocating space for W2 and A2 matrices on the GPU's video memory, using key interface functions such as cudaMalloc(), cudaMemcpy(), etc. to copy the matrix data from the host to the GPU; according to the CUDA programming model, the matrix data is divided into multiple blocks, and different blocks perform different actual calculation operations. Threads in the same block copy the corresponding matrix data needed for their own calculations from the GPU's video memory to the SM's shared memory in parallel. Create and initialize member objects a, b and c associated with the registers in the tensor core on the GPU, respectively for storing W2, A2 and single-bit operation results.

[0064] Step 2.2, using the member objects a, 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, the address of the matrix data that needs to be copied is calculated according to the following formula: A load A shmem represents the starting address of the shared memory, warpId represents the id number of the warp in the CUDA programming model, M represents the height of the first matrix, Type represents 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] subject toi = 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 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, implement thread synchronization, and ensure that all calculation tasks in the previous stage have been completed; repeat the above process until all calculation tasks in the block have been completed. At this time, the single-bit operation result is stored in object c, and the data type is 32-bit integer.

[0068] Step 2.3, using the object c obtained in step 2.2 to store the single-bit operation result, and reducing it to a multi-bit operation result. For all threads in a block, the address of the single-bit operation result is calculated according to the following formula: A SstoreA represents the address of the single-bit operation result to be loaded shmem warpId represents the id number of a warp in the CUDA programming model, laneId represents the id number of a thread in a warp, Type represents the storage type of data, and the sizeof() function is used to calculate the space size of the data type in bytes

[0069]

[0070] The store_matrix_sync() interface function is called to load the operation result from the c object associated with the register in the tensor core to the shared memory of the SM; the single-bit operation result in the shared memory is reduced, and according to the bit number of the W and A matrices, a multi-bit operation result is generated by combination and stored in the register of the CUDA core, and the combination generation process is the inverse process of the bit decomposition operation.

[0071] Step 3: For all threads in a block, the address of the multi-bit operation result is calculated, and the multi-bit operation result is copied to the pre-designated memory space, the size of which should be allocated in advance at the initialization and can accommodate the final multi-bit operation result.

[0072] After the above steps are completed, the result of the specified precision sparse matrix multiplication and addition operation is obtained, and subsequent input auxiliary programs can be added or the result can be output to other operators and fused. Thus, the implementation method of the tensor core acceleration of the arbitrary precision sparse matrix multiplication and addition operation is ended.

[0073] Although the present application has been described in detail in the foregoing description and specific embodiments, modifications and improvements can be made to the present application on the basis of the present application, which will be apparent to those skilled in the art. Therefore, these modifications or improvements made on the basis of not deviating from the spirit of the present application, all belong to the scope of protection claimed by the present application.

Claims

1. A method for accelerating arbitrary-precision sparse matrix multiplication and addition operations based on tensor kernels, characterized in that, Includes the following steps: 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; the specific implementation method is as follows: Step 1.1: Obtain the matrix data of the original dense matrix W and A 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 one-dimensional BCSR format suitable for tensor core calculation. Step 1.2: Based on the Tensor Core interface calculation operator and the bit count of the one-dimensional BCSR format sparse matrix W1, perform bit decomposition and padding alignment operations on the one-dimensional BCSR format sparse matrix W1 to generate the SB-BCSR format sparse matrix W2; simultaneously, based on the bit count of the original matrix A, first perform bit compression to generate a bit-tight format dense matrix A1, and then perform bit decomposition on the bit-tight format dense matrix A1 to generate a bit-continuous format dense matrix A2; where Tensor Core refers to the Tensor Core in NVIDIA's A100 series GPUs. The interface computation operators of TensorCore refer to the function interfaces provided by NVIDIA that call TensorCore hardware to perform matrix operations. Bit decomposition is an operation that breaks down a number into binary bits, used for encoding or compressing numbers in computer systems. In bit decomposition, a number is broken down into a series of binary bits, each bit representing a value within the number. For a number N, R... i Let be the i-th binary value of N, and n be the number of binary digits in N; then N can be represented as: For m numbers N j , It is N j The i A binary value, for the value S after decomposing n. k The decomposition process can be represented as: Padding alignment refers to the operation of adding extra space or elements to a data structure or memory layout to meet the requirements of a set hardware or software, so that the storage location or layout of the data conforms to the set alignment requirements. The SB-BCSR format is an improved format specifically designed for 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. In memory, it stores the binary bits at the same position for different data consecutively, and the binary bits at different positions for the same data are stored separately. 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, after bit decomposition, this row is transformed into L1, which is represented by an ordered set {S}. 0 S 1 S m-1 }; Step 2: Perform matrix operations on the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit-continuous format generated in Step 1 to obtain the single-bit matrix operation result stored in the tensor core register. Perform matrix reduction based on the value of the single-bit matrix operation result to obtain the multi-bit matrix operation result. Step 3: Based on the multi-bit matrix operation results obtained in Step 2, for all threads in a block, calculate the load address of the corresponding multi-bit operation result and copy the multi-bit operation result to the pre-specified video memory space in parallel.

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

3. The method for accelerating arbitrary-precision sparse matrix multiplication and addition based on tensor kernels according to claim 1, characterized in that, The specific implementation method of step 2 is as follows: Step 2.1: Allocate space on the GPU's video memory for the sparse matrix W2 in SB-BCSR format and the dense matrix A2 in bit-continuous format, and use the corresponding CUDA functions to copy the matrix data from the host to the GPU; divide the matrix data into multiple blocks, and all threads in a block copy the corresponding matrix data that needs to be calculated by itself from the GPU video memory to the shared memory of SM in parallel; create and initialize member objects a, b, and c on the GPU that are associated with the registers in the tensor core, and use them 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 to be copied, call the load_matrix_sync() interface function, and load the matrix data from the shared memory of SM into the registers of the Tensor Core through 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, calculate the load address of the corresponding single-bit operation result, call the store_matrix_sync() interface function, and use member object c to load the single-bit operation result from the register of the tensor core into the shared memory of SM in parallel; reduce the single-bit operation result in the shared memory, and combine it to generate a multi-bit operation result according to the number of bits of the original matrix W and A.

4. The method for accelerating arbitrary-precision sparse matrix multiplication and addition based on tensor kernels according to claim 3, characterized in that, In step 2.1, the corresponding CUDA functions refer to the interface functions provided by NVIDIA for allocating address space and performing data migration on the GPU, including the two key interface functions cudaMalloc() and cudaMemcpy(). Dividing matrix data into multiple blocks means that, according to the CUDA programming model, the matrix data is abstractly allocated to different blocks, and each block performs the actual computation operation. The member objects associated with the registers in the Tensor Core refer to the abstract data types provided by NVIDIA to facilitate programmers to call the Tensor Core hardware. These include three types: matrix_a, which indicates that the object is the first multiplicand; matrix_b, which indicates that the object is the second multiplicand; and accumulator, which indicates that the object is the source or destination accumulator. When declaring this type of object, the computation dimension and storage layout format are specified, and the object is modified, exported, or executed through the four functions load_matrix_sync(), store_matrix_sync(), fill_fragment(), and mma_sync().

5. The method for accelerating arbitrary-precision sparse matrix multiplication and addition based on tensor kernels according to claim 4, characterized in that, In step 2.2, the address of the matrix data to be copied is calculated using the following formula, where A load A represents the address of the matrix data to be copied. shmem The starting address of the shared memory is indicated by `warpId`, the warp ID in the CUDA programming model is indicated by `M`, the height of the first matrix is ​​indicated by `Type`, and the `sizeof()` function is used to calculate the size of the data in bytes. Specifying the computation type as binary bitwise AND means setting the precision to b1 and the bitwise operation type to bmmaBitOpAND when calling the mma_sync() function; performing thread synchronization means calling the __syncthreads() interface function to synchronize all threads in the same block after completing a round of computation, ensuring that all computation tasks in the previous stage have been completed.

6. The method for accelerating arbitrary-precision sparse matrix multiplication and addition based on tensor kernels according to claim 5, characterized in that, In step 2.3, the address for calculating the result of the single-bit operation is determined by the following formula, where A Sstore A represents the address of the single-bit computation result that needs to be loaded. shmem The `warpId` represents the starting address of the shared memory, `laneId` represents the warp ID in the CUDA programming model, `type` represents the thread ID within a warp, `type` represents the data storage type, and the `sizeof()` function is used to calculate the space size of the data type in bytes. Combining to generate multi-bit operation results refers to combining the single-bit operation results output by the tensor kernel according to their bit positions to generate the final multi-bit operation result. The combination generation process is the inverse process of bit decomposition.

7. The method for accelerating arbitrary-precision sparse matrix multiplication and addition based on tensor kernels according to claim 1, characterized in that, In step 3, the pre-specified video memory space refers to the video memory space allocated in advance during initialization to store the results. The size of this video memory space is sufficient to accommodate the final multi-bit operation results.

Images