Self-adaptive two-stage matrix division method based on GEMM (General Empirical Matrix Modulation)

Abstract

The invention discloses an adaptive two-stage matrix division method based on GEMM (General Empirical Matrix Modulation). The method comprises the following steps: determining the size of a block for the first time; carrying out first blocking processing on the input matrix from a global memory to a shared memory; loading the matrix partitioned for the first time to double buffer areas of a shared memory; determining the size of a secondary block; partitioning the partitioned matrix blocks from the shared memory to the register again; and loading the matrix blocks which are partitioned again to a buffer area of the register for calculation. The invention aims to provide a parallel optimization method based on hardware-software collaboration aiming at the bottlenecks of the current GPU-based GEMM in the aspects of data transmission efficiency, computing resource utilization rate, load balance and the like. By introducing an optimization scheme of a self-adaptive two-stage matrix partitioning mechanism in a matrix calculation process, the method aims to improve data locality, reduce memory bandwidth pressure and realize dynamic balance of loads among threads. According to the method, by reducing data transmission delay and optimizing allocation of computing resources, the efficiency of GEMM operation is improved from the source, and the efficient requirements of deep learning and large-scale scientific computing are met.

Description

Technical Field

[0001] The present invention mainly relates to the field of high performance computing and deep learning acceleration technology, and in particular to an adaptive two-level matrix partitioning method based on general matrix multiplication (GEMM). Technical Background

[0002] In the field of large-scale computing and deep learning, general matrix multiplication (GEMM) is one of the most commonly used computing operations. With the continuous growth of the scale of deep learning models and the increasing computing requirements, the demand for efficient matrix multiplication has become more urgent. As a highly parallel hardware platform, GPU has the characteristics of multi-threading, fast cache, and high-bandwidth memory access, so it has become the main hardware platform for accelerating GEMM operations. However, with the increase in computing density and model complexity, existing GEMM calculation methods still face multiple challenges on GPU.

[0003] First, because GEMM operations involve large-scale data transmission, existing methods still have bottlenecks in data access and transmission latency. Large-scale matrix operations lead to frequent global memory access, which significantly increases memory access latency. To this end, some solutions try to optimize data transmission through data segmentation, double buffering, and prefetching, but these methods often lack effective resource scheduling and parallel development methods when dealing with dynamic and complex data dependencies.

[0004] Secondly, in the existing GEMM implementation, the utilization and load balance of hardware resources are key issues. Since the parallel requirements of data access for each computing unit in the GPU architecture are different, how to reasonably divide the computing tasks and map them to different computing units has become an important factor affecting performance. The current mainstream methods are mostly based on strategies such as block partitioning and loop unrolling. However, these methods mainly use the fixed size of the matrix as the basis for block partitioning, which is difficult to adapt to the changes in GPU hardware and heterogeneous computing requirements, resulting in bottlenecks in hardware resource utilization and load balance.

[0005] In addition, existing GEMM optimization methods also have certain problems in the synchronization overhead and load balancing of parallel computing. Since large-scale parallel computing involves a large number of thread divisions and synchronization operations, the current thread division method often uses coarse-grained thread blocks as the minimum unit, failing to fully exploit the inherent data locality of each operation in the computing process. As the scale of computing increases, this division strategy is prone to introduce synchronization overhead and resource competition, reducing the overall computing efficiency.

[0006] Therefore, it is necessary to propose an adaptive two-level matrix partitioning method based on GEMM. By analyzing the runtime characteristics of GEMM, this paper designs an optimization method for the adaptive two-level matrix partitioning mechanism, aiming to improve the memory bandwidth utilization of GEMM, reduce data transmission delay, and realize dynamic load balancing of computing resources, thereby improving the overall computing efficiency on the GPU platform. Summary of the invention

[0007] The purpose of the present invention is to provide an adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM). The current GPU-based GEMM has problems in data transmission efficiency, computing resource utilization and load balancing.

[0008] The present invention is implemented in that an adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) comprises the following steps: Step 1: Determine the size of the first block; Step 2: The input matrix is ​​first divided into blocks from global memory to shared memory; Step 3: Load the first-blocked matrix into the double buffer of the shared memory; Step 4: Determine the size of the next block; Step 5: Move the divided matrix blocks from the shared memory to the register for further block processing; Step 6: Load the re-divided matrix blocks into the register buffer for calculation.

[0009] Furthermore, the step 1 evaluates the hardware resources to obtain the maximum capacity S of the shared memory. sm , the size of each matrix element S elem , the dimension T of the thread block x ×T y and make sure the chunk size meets the shared memory limit: According to the input matrices A (M×K) and B (K×N), the block size is selected so that b1 and b2 can divide M and N as much as possible. In order to optimize the access mode of the global memory, we choose b1, b2 and 32 as multiples to meet the memory alignment requirements. If M and N cannot divide b1 and b2, special processing is added to the boundary (such as padding the boundary part).

[0010] Furthermore, the step 2 divides the thread block tasks, each thread block is responsible for loading a b1×k and k×b2 sub-block of matrices A and B, each thread in the thread block is responsible for loading one or more matrix elements, and adjusting the thread index so that the global memory access satisfies the conditions for merged access. Then each thread loads a matrix element from the global memory and writes it to the corresponding position of the shared memory. Then the edge area of ​​the matrix is ​​checked to avoid out-of-bounds access. Finally, __syncthreads() is used to ensure that all threads complete data loading.

[0011] Furthermore, the step 3 described above loads the matrix after the first block division into the double buffer of the shared memory:

[0012] Step 1: Allocate two buffers sharedBuffer[2] in the shared memory, each with a size of b1×k and k×b2;

[0013] Step 2: The current buffer bufferIndex stores the data currently being calculated, and another buffer 1-bufferIndex1 loads the next block of data. To overlap data loading with calculation, the data in the current buffer is used for calculation, and the next block of data is loaded into the backup buffer at the same time, forming a pipeline operation. Finally, when the buffer is switched, __syncthreads() is used to ensure that data loading is completed.

[0014] Furthermore, the step 4 will evaluate the register capacity to ensure that the block sizes r1×k and k×r2 meet the register capacity constraints: Then, according to the number of registers and the number of parallel threads, the register space of each thread is reasonably allocated.

[0015] Furthermore, the step five described above will load data by thread, and each thread will load the r1×k or k×r2 data of the current block from the shared memory into the register, and use the register to store data during calculation to avoid frequent access to the shared memory. Finally, the register calculation of the current thread block and the data loading of the next thread block are performed simultaneously to form a pipeline operation.

[0016] Furthermore, the step 6 loads the re-divided matrix blocks into the buffer of the register for calculation:

[0017] Step 1: Each thread uses the data in the register to complete the matrix multiplication of the current block, and before writing the result back, it performs a boundary check to ensure that it does not cross the boundary of the global memory;

[0018] Step 2: Store the calculation results in global memory and continue switching to the next block until the entire matrix is ​​processed.

[0019] The beneficial effects of the present invention are as follows: (1) The present invention can effectively reduce the bottleneck of memory bandwidth through an adaptive two-level block method. The matrix blocks are gradually loaded from the global memory to the shared memory, and then to the hierarchical processing in the register, which greatly reduces the delay of data transmission and improves the utilization of parallel computing resources, thereby accelerating the GEMM operation. (2) The present invention adopts different-sized block strategies in different computing stages to achieve more flexible memory management. The adaptive adjustment of the first block and the second block size ensures the optimal configuration of memory resources, avoids memory overflow and redundant calculations, and improves the stability and resource utilization of the system. (3) Each computing stage of the present invention processes data in parallel in the buffer of the shared memory and register, which can effectively improve the throughput of the GPU. In addition, by dynamically adjusting the block strategy according to the resource usage of the thread block and the matrix size, the method has strong scalability and can adapt to computing tasks of different scales, especially for the processing of large-scale matrix operations. BRIEF DESCRIPTION OF THE DRAWINGS

[0020] Figure 1 It is a flow chart of an adaptive two-level matrix partitioning method based on general matrix multiplication (GEMM) provided by the present invention.

[0021] Figure 2 The invention provides an adaptive two-level matrix partitioning mechanism based on GPU hardware characteristics.

[0022] Figure 3 This is the first block division strategy provided by the implementation of the present invention.

[0023] Figure 4 This is the re-blocking strategy provided by the implementation of the present invention. DETAILED DESCRIPTION

[0024] In order to make the purpose, technical solution and advantages of the present invention more clearly understood, the present invention is further described in detail below in conjunction with the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are only used to explain the present invention and are not intended to limit the present invention.

[0025] Figure 1 The process of the adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) of the present invention is shown. As shown in the figure, the present invention is implemented in this way. An adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) includes the following steps: S101: Determine the size of the first block; S102: Perform the first block processing on the input matrix from the global memory to the shared memory; S103: Load the matrix after the first block division into the double buffer of the shared memory; S104: Determine the size of the next block; S105: transferring the divided matrix blocks from the shared memory to the register for further block processing; S106: Load the re-divided matrix blocks into the buffer of the register for calculation.

[0026] Furthermore, the step S101 evaluates the hardware resources to obtain the maximum capacity S of the shared memory. sm , the size of each matrix element S elem , the dimension T of the thread block x ×T y and make sure the chunk size meets the shared memory limit: According to the input matrices A (M×K) and B (K×N), the block size is selected so that b1 and b2 can divide M and N as much as possible. In order to optimize the access mode of the global memory, we choose b1, b2 and 32 as multiples to meet the memory alignment requirements. If M and N cannot divide b1 and b2, special processing is added to the boundary (such as padding the boundary part).

[0027] Furthermore, the S102 divides the thread block tasks, each thread block is responsible for loading a b1×k and k×b2 sub-block of matrices A and B, each thread in the thread block is responsible for loading one or more matrix elements, and adjusting the thread index so that the global memory access satisfies the conditions for merged access. Then each thread loads a matrix element from the global memory and writes it to the corresponding position of the shared memory. Then the edge area of ​​the matrix is ​​checked to avoid out-of-bounds access. Finally, __syncthreads() is used to ensure that all threads complete data loading.

[0028] Furthermore, the S103 loads the matrix after the first block division into the double buffer of the shared memory:

[0029] Step 1: Allocate two buffers sharedBuffer[2] in the shared memory, each with a size of b1×k and k×b2;

[0030] Step 2: The current buffer bufferIndex stores the data currently being calculated, and another buffer 1-bufferIndex1 loads the next block of data. To overlap data loading with calculation, the data in the current buffer is used for calculation, and the next block of data is loaded into the backup buffer at the same time, forming a pipeline operation. Finally, when the buffer is switched, __syncthreads() is used to ensure that data loading is completed.

[0031] Furthermore, the S104 will evaluate the register capacity to ensure that the block sizes r1×k and k×r2 meet the register capacity constraints: Then, according to the number of registers and the number of parallel threads, the register space of each thread is reasonably allocated.

[0032] Furthermore, the S105 described above will load data by thread, and each thread will load the r1×k or k×r2 data of the current block from the shared memory into the register, and use the register to store data during calculation to avoid frequent access to the shared memory. Finally, the register calculation of the current thread block and the data loading of the next thread block are performed simultaneously to form a pipeline operation.

[0033] Furthermore, the S106 loads the re-divided matrix blocks into the buffer of the register for calculation:

[0034] Step 1: Each thread uses the data in the register to complete the matrix multiplication of the current block, and before writing the result back, it performs a boundary check to ensure that it does not cross the boundary of the global memory;

[0035] Step 2: Store the calculation results in global memory and continue switching to the next block until the entire matrix is ​​processed.

[0036] Example: Assume that the shared memory of each thread is 16KB and the size of each register is 8KB. The size of matrix A is 1024×1024, and the size of matrix B is 1024×1024. Then we need to calculate the matrix multiplication C=A×B, where C is the final output layer matrix. If the traditional general matrix multiplication method is used for calculation, it will bring huge time overhead, but under the condition of matrix size and hardware resources, it can effectively verify the performance improvement brought by our proposed method. The steps of the adaptive two-level matrix partitioning method for implementing general matrix multiplication on convolutional neural network are as follows:

[0037] Step 1: Determine the size of the first block; The matrix size of the first block is calculated by formula 1. The matrix block size of the first block cannot exceed 8KB (i.e. half of 16KB). Therefore, the size of each matrix block must be controlled within 8KB. Assuming that each matrix element occupies 4 bytes, the size of each matrix block is: 8KB / 4 bytes = 2048 elements. Then the blocks of A and B are 1024 elements respectively. And because 32×32=1024 elements, the first block size is 32×32, matrix A is divided into bm×bk=32×32, and matrix B is divided into bk×bn=32×32.

[0038] Step 2: First block processing; Matrices A and B are 1024×1024 matrices. Each matrix is ​​divided into 32×32 sub-blocks, that is, A will be divided into 32 32×32 blocks, and B will also be divided into 32 32×32 blocks. Each thread block is responsible for calculating a 32×32 sub-matrix in matrix C.

[0039] Step 3: Load the first block matrix into shared memory; Load the first-blocked matrix blocks A1 and A2, B1 and B2 into the double buffer of the shared memory.

[0040] When matrix block A1 enters the calculation phase, it starts multiplying with B1. Then, it continues to load the next matrix blocks A3 and B3 into the shared memory, ensuring that when one matrix is ​​calculated, the calculation of the next matrix block starts immediately.

[0041] Step 4: Determine the size of the next block. The size of each register is 8KB. In order to optimize the use of registers, the size of the matrix block needs to adapt to the capacity of the register. The size of the block is obtained by Formula 2. As in step 1, we choose a 16×16 sub-block, which ensures that the register size used by each thread is within a reasonable range.

[0042] Step 5: Divide into blocks again. When A1 and B1 have been loaded into the shared memory, they are now further divided into smaller matrix blocks (rm×rk=16×16, rk×rn=16×16) according to the sizes determined in step 4. Each thread is responsible for calculating a 16×16 sub-matrix in the matrix C1.

[0043] Step 6: Load the re-divided matrix blocks into the register buffer and calculate. The matrix A will be divided into blocks again 11 and A 12 , B 11 and B 12 Load into double buffer of registers for calculations, avoiding memory waits.

[0044] In calculating A 11 and B 11 When the matrix block is completed, continue to load the next matrix block A 13 and B 13 To the double buffer of registers, maintain the continuity of calculation.

[0045] The flowchart of the adaptive two-level matrix partitioning method proposed above is as follows: Figure 2 As shown in the figure. Since this method divides the input matrix into smaller sub-blocks (such as bm, bn, bk) according to size by gradually dividing and transferring data between global memory, shared memory and registers, and performs block loading and calculation between different memory levels, it avoids the delay and bandwidth bottleneck caused by frequent access to global memory. This not only effectively utilizes the high-speed access advantages of shared memory and registers, but also realizes the parallelization of data transmission and calculation processes, thereby improving computing efficiency and reducing GPU resource waste, and ultimately significantly accelerating matrix multiplication operations.

[0046] Although the above describes the specific implementation mode of the present invention in conjunction with the accompanying drawings, it is not intended to limit the scope of protection of the present invention. Those skilled in the art should understand that various modifications or variations that can be made by those skilled in the art on the basis of the technical solution of the present invention without creative labor are still within the scope of protection of the present invention.

Claims

1. An adaptive two-level matrix partitioning method based on general matrix multiplication (GEMM), characterized in that: The adaptive two-level matrix partitioning method based on general matrix multiplication (GEMM) comprises the following steps: Step 1: Determine the size of the first block; Step 2: The input matrix is ​​first divided into blocks from global memory to shared memory; Step 3: Load the first-blocked matrix into the double buffer of the shared memory; Step 4: Determine the size of the next block; Step 5: Move the divided matrix blocks from the shared memory to the register for further block processing; Step 6: Load the re-divided matrix blocks into the register buffer for calculation.

2. The adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) as claimed in claim 1, characterized in that: Step 1 will evaluate the hardware resources and obtain the maximum capacity S of the shared memory. sm , the size of each matrix element S elem , the dimension T of the thread block x ×T y and make sure the chunk size meets the shared memory limit: According to the input matrices A (M×K) and B (K×N), the block size is selected so that b1 and b2 can divide M and N as much as possible. In order to optimize the access mode of the global memory, we choose b1, b2 and 32 as multiples to meet the memory alignment requirements. If M and N cannot divide b1 and b2, special processing is added to the boundary (such as padding the boundary part).

3. The adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) as claimed in claim 1, characterized in that: The step 2 described above divides the thread block tasks. Each thread block is responsible for loading a b1×k and k×b2 sub-block of matrices A and B. Each thread in the thread block is responsible for loading one or more matrix elements and adjusting the thread index so that the global memory access satisfies the conditions for merged access. Then each thread loads a matrix element from the global memory and writes it to the corresponding position of the shared memory. Then the edge area of ​​the matrix is ​​checked to avoid out-of-bounds access. Finally, __syncthreads() is used to ensure that all threads complete data loading.

4. The adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) as claimed in claim 1, characterized in that: The step 3 described above loads the matrix after the first block division into the double buffer of the shared memory, and the specific steps are as follows: Step 1: Allocate two buffers sharedBuffer[2] in the shared memory, each with a size of b1×k and k×b2; Step 2: The current buffer bufferIndex stores the data currently being calculated, and another buffer 1-bufferIndex1 loads the next block of data. To overlap data loading with calculation, the data in the current buffer is used for calculation, and the next block of data is loaded into the backup buffer at the same time, forming a pipeline operation. Finally, when the buffer is switched, __syncthreads() is used to ensure that data loading is completed.

5. The adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) as claimed in claim 1, characterized in that: Step 4 described above will evaluate the register capacity to ensure that the block sizes r1×k and k×r2 meet the register capacity constraints: Then, according to the number of registers and the number of parallel threads, the register space of each thread is reasonably allocated.

6. The adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) as claimed in claim 1, characterized in that: The step five described above will load data by thread. Each thread loads the r1×k or k×r2 data of the current block from the shared memory to the register. The register is used to store data during calculation to avoid frequent access to the shared memory. Finally, the register calculation of the current thread block and the data loading of the next thread block are performed simultaneously to form a pipeline operation.

7. The adaptive two-stage matrix partitioning method based on general matrix multiplication (GEMM) as claimed in claim 1, characterized in that: The step six loads the re-divided matrix blocks into the buffer of the register for calculation, and the specific steps are as follows: Step 1: Each thread uses the data in the register to complete the matrix multiplication of the current block, and before writing the result back, it performs a boundary check to ensure that it does not cross the boundary of the global memory; Step 2: Store the calculation results in global memory and continue switching to the next block until the entire matrix is ​​processed.

Images