A dual-core collaborative parallel computing method for irregular matrix multiplication

Abstract

This invention relates to the technical field of high-performance computing, specifically disclosing a collaborative parallel computing method for CUDA Cores and Tensor Cores in irregular matrix multiplication. This method utilizes a hybrid core fusion approach to collaboratively schedule CUDA Cores and Tensor Cores, and optimizes computational resource allocation based on a thread partitioning coefficient model, thereby improving the computational efficiency of irregular matrix multiplication. For the first time, it systematically implements collaborative parallel computing between CUDA Cores and Tensor Cores at the SM granularity. The hybrid core fusion algorithm eliminates alternating idle states between cores, significantly improving GPU hardware resource utilization. Through shared memory partitioning, double buffering mechanisms, and data layout optimization, it avoids memory conflicts and hides memory access latency. Task partitioning based on a performance model achieves thread-level load balancing, further improving overall computational efficiency.

Description

Technical Field

[0001] This invention relates to the field of high-performance computing, and more specifically, to a dual-core collaborative parallel computing method for irregular matrix multiplication. Background Technology

[0002] With the rapid development of information technology, global computing power demand is growing exponentially, and traditional CPU architectures can no longer meet the needs of high-performance computing scenarios. In fields such as artificial intelligence and scientific computing, Generalized Matrix Multiplication (GEMM) is a core operation, and its performance directly affects the release of GPU computing power and the improvement of application efficiency. To accelerate GEMM operations, NVIDIA has launched a dedicated matrix computing unit, Tensor Core, which, together with CUDA Core, forms a heterogeneous GPU computing architecture.

[0003] Currently, Tensor Core programming primarily relies on the native WMMA API or high-level optimization libraries such as cuBLAS. However, both have significant limitations: third-party libraries like cuBLAS sacrifice optimization flexibility and debugging convenience; the WMMA API has strict regularity requirements for the input matrix dimension and only supports matrix operations at specific data granularities. More importantly, in the existing computing model, Tensor Cores and CUDA Cores cannot work in parallel, resulting in one core being active while the other is idle, leading to low hardware resource utilization and energy efficiency.

[0004] Existing optimization research mainly focuses on tapping the potential of a single computing core or improving algorithms and automatically tuning Tensor Cores, failing to systematically address the collaborative problem between the two types of computing cores. When handling irregular matrix multiplication, traditional methods require padding operations to satisfy dimensional constraints, which not only introduces computational redundancy but also increases memory overhead, severely impacting computational efficiency. Therefore, an innovative collaborative parallel computing method is urgently needed to efficiently schedule CUDA Cores and Tensor Cores, fully unleash the heterogeneous computing potential of GPUs, and improve the processing performance of irregular matrix multiplication. Summary of the Invention

[0005] To address the aforementioned technical problems in related technologies, this invention provides a dual-core collaborative parallel computing method for irregular matrix multiplication, which can solve the above problems.

[0006] To achieve the above-mentioned technical objectives, the technical solution of the present invention is implemented as follows:

[0007] A dual-core collaborative parallel computing method for irregular matrix multiplication includes the following steps:

[0008] S1. Resource and memory access characteristics analysis: In the GPU hybrid computing mode, the system analyzes the hardware resource usage characteristics of CUDA Core and TensorCore and their global memory and shared memory access behavior, and quantitatively evaluates the resource idle rate caused by serial execution in the traditional computing mode.

[0009] S2. Hybrid Core Fusion Scheduling: Based on the analysis results of step S1, a hybrid core fusion algorithm is constructed. Through a unified thread block grid and memory access coordination mechanism, collaborative parallel task scheduling of CUDA Cores and Tensor Cores is achieved at the instruction level and thread block level, eliminating alternating idle states between cores.

[0010] S3. Dynamic Task Partitioning and Optimization: A dynamic optimization model based on thread partitioning coefficients is introduced. The thread partitioning coefficients are adjustable parameters. The value of the thread partitioning coefficients dynamically determines the size of the matrix row blocks or column blocks allocated to CUDA Cores for processing. This dynamic optimization model adaptively partitions and allocates computing tasks to CUDA Cores and Tensor Cores based on the actual irregularity of the input matrix and the real-time hardware resource status of the GPU, optimizing thread-level load balancing and memory access locality, and avoiding redundant computation and memory overhead caused by padding operations.

[0011] S4. Collaborative Computation Execution: After scheduling, the CUDA Cores and Tensor Cores collaboratively perform irregular matrix multiplication operations according to the divided tasks.

[0012] Furthermore, the hybrid core fusion algorithm described in step S2 adopts a two-level hybrid fusion mechanism. At the macro level, it performs inter-thread block fusion and aggregates the task thread block space to expand parallel resources. At the micro level, it adopts a fixed dimension and condition partitioning strategy within the thread block to ensure that all threads in the Warp have consistent execution paths and eliminate thread divergence.

[0013] Furthermore, the fixed dimension of the thread block is set to a 2D grid arrangement, where the x-axis contains 32 threads and the y-axis contains 32 threads. Task allocation within the thread block is achieved through conditional partitioning based on threadIdx.y.

[0014] Furthermore, the thread partitioning coefficient mentioned in step S3 is derived in the following way:

[0015] Establish a hybrid kernel total execution time model, whereby the total execution time includes CUDA Core execution time, TensorCore execution time, inter-thread synchronization overhead, and memory transfer time;

[0016] The execution time of CUDA Core and Tensor Core is modeled separately, and the relative computing power ratio is quantified by combining the hardware specifications and theoretical computing power of the two.

[0017] Construct a shared memory resource allocation constraint model to ensure that CUDA Cores and Tensor Cores do not conflict under limited shared memory capacity;

[0018] Based on load balancing conditions, ignoring startup latency, and considering the characteristic that Tensor Cores complete 256 floating-point operations per MMA operation, the thread partitioning coefficient and the number of rows that CUDA Cores should process are derived.

[0019] Furthermore, in the shared memory resource allocation constraint model, a double buffering mechanism is introduced for Tensor Cores to hide data loading latency. That is, double data buffers are designed for Tensor Cores to overlap computation and data loading, thus hiding memory access latency.

[0020] Furthermore, the dynamic task partitioning process in step S3 also includes: padding and aligning the input matrix so that both the width and height of the matrix are multiples of 16, in order to meet the data granularity and memory access patterns required for efficient execution of Tensor Core.

[0021] Furthermore, separate data buffers are allocated for CUDA Cores and Tensor Cores in shared memory to optimize data layout, avoid storage conflicts, and maximize memory bandwidth supply.

[0022] Furthermore, in step S4, a unified composite kernel is launched, and within the kernel, matrix row computation tasks are dynamically distributed to CUDA Cores or Tensor Cores at runtime based on thread indexes and pre-computed partitioning coefficients, thereby achieving fine-grained collaborative execution.

[0023] The beneficial effects of this invention are:

[0024] Achieving heterogeneous core collaborative parallelism: For the first time, collaborative parallel computing between CUDA Cores and Tensor Cores at the SM granularity has been systematically achieved. By eliminating the alternating idle states between cores through a hybrid core fusion algorithm, the utilization rate of GPU hardware resources has been significantly improved.

[0025] Efficiently handles irregular matrices: By using a dynamic optimization model based on thread partitioning coefficients, computational tasks are adaptively partitioned, avoiding computational redundancy and memory overhead caused by traditional filling operations. In small-to-medium scale and irregular matrix scenarios, the performance is significantly better than existing solutions.

[0026] High versatility and robustness: It does not rely on specific hardware secrets or complex tuning, and can automatically adapt to different GPU architectures (such as V100, RTX4090, A100) and matrix shapes. Compared with mainstream optimization libraries cuBLAS and CUTLASS, it exhibits better computing performance and scalability under various matrix configurations.

[0027] Memory access and load balancing optimization: By using shared memory partitioning, double buffering mechanism and data layout optimization, storage conflicts are avoided and memory access latency is hidden; task partitioning based on performance model realizes thread-level load balancing, further improving overall computing efficiency. Attached Figure Description

[0028] To more clearly illustrate the technical solutions in the embodiments of the present invention or the prior art, the drawings used in the embodiments will be briefly introduced below. Obviously, the drawings described below are only some embodiments of the present invention. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.

[0029] Figure 1a This is a schematic diagram of the existing computing methods in GPU computing cores;

[0030] Figure 1b This is a schematic diagram of the dual-core collaborative parallel computing method in a GPU as described in an embodiment of the present invention;

[0031] Figure 2 This is a general overview diagram of the collaborative parallel computing algorithm described in this embodiment of the invention on a GPU equipped with Tensor Cores;

[0032] Figure 3a This is the basic kernel fusion (internal thread fusion) method described in the embodiments of the present invention;

[0033] Figure 3b This is the basic kernel fusion (thread block internal fusion) method described in the embodiments of the present invention;

[0034] Figure 3c This is the basic kernel fusion (fusion between thread blocks) method described in the embodiments of the present invention;

[0035] Figure 4 This is a schematic diagram of the hybrid kernel fusion method according to an embodiment of the present invention;

[0036] Figure 5 This is a schematic diagram of shared memory allocation according to an embodiment of the present invention;

[0037] Figure 6aThe result is the performance evaluation of GEMM operations with matrix shapes N=W, M=W, K=W on the V100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0038] Figure 6b The result is the performance evaluation of GEMM operation with matrix shape N=W, M=2W, K=4W on the V100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0039] Figure 6c The result is the performance evaluation of GEMM operations with matrix shapes N=W, M=4W, K=W on the V100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0040] Figure 6d The result is the performance evaluation of GEMM operation with matrix shape N=4W, M=4W, K=4W on V100 GPU architecture according to the embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0041] Figure 7a The result is the performance evaluation of GEMM operation with matrix shape N=W, M=W, K=W on the RTX4090 GPU architecture according to the embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0042] Figure 7b The result is the performance evaluation of GEMM operation with matrix shape N=W, M=2W, K=4W on the RTX4090 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0043] Figure 7c The result is the performance evaluation of GEMM operation with matrix shape N=W, M=4W, K=W on the RTX4090 GPU architecture according to the embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0044] Figure 7d The result is the performance evaluation of GEMM operation with matrix shape N=4W, M=4W, K=4W on the RTX4090 GPU architecture according to the embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0045] Figure 8a The result is the performance evaluation of GEMM operations with matrix shapes N=W, M=W, K=W on the A100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0046] Figure 8b The result is the performance evaluation of GEMM operation with matrix shape N=W, M=2W, K=4W on the A100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0047] Figure 8c The result is the performance evaluation of GEMM operation with matrix shape N=W, M=4W, K=W on the A100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200.

[0048] Figure 8d This is the result of a performance evaluation of GEMM operations with matrix shapes N=4W, M=4W, K=4W on the A100 GPU architecture according to an embodiment of the present invention, wherein the parameter W ranges from 1 to 200. Detailed Implementation

[0049] The technical solutions of the embodiments of the present invention will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some embodiments of the present invention, and not all embodiments. All other embodiments obtained by those skilled in the art based on the embodiments of the present invention are within the scope of protection of the present invention.

[0050] This invention proposes a collaborative parallel computing method for irregular matrix multiplication using CUDA Cores and Tensor Cores. It coordinates the scheduling of CUDA Cores and Tensor Cores through a hybrid core fusion method and optimizes the allocation of computing resources based on a thread partitioning coefficient model, thereby improving the computational efficiency of irregular matrix multiplication.

[0051] As shown in Figure 1, in existing technologies, CUDA Cores and Tensor Cores cannot work in parallel, resulting in idle resources. In the figure, SM stands for Streaming Multiprocessor, and CU stands for CUDA Core. This invention, however, achieves parallel collaboration between the two within the streaming multiprocessor through an innovative cooperation mechanism, such as... Figure 2The diagram illustrates the overall workflow of the algorithm on a GPU equipped with Tensor Cores, where TC stands for Tensor Core. After the kernel function starts, multiple thread blocks are allocated to streaming multiprocessors for intelligent scheduling. This algorithm overcomes the limitations of traditional single-computation-unit architectures through the following collaborative mechanisms, achieving task-level fusion and parallel processing of two computational units: It employs a hybrid kernel fusion method, reconstructing and integrating the kernel function originally designed for a single computational unit to form a composite kernel function capable of simultaneously scheduling and managing two types of computational units; and it proposes an intelligent thread partitioning strategy based on matrix sparsity patterns and real-time hardware status, achieving load balancing through dynamic task allocation and fine-grained scheduling.

[0052] In a specific embodiment of this application, the application first needs to perform resource and memory access characteristic analysis, specifically: in GPU hybrid computing mode, the system analyzes the hardware resource usage characteristics of CUDA Core and Tensor Core and their global memory and shared memory access behavior, and quantitatively evaluates the resource idle rate caused by serial execution of the two in traditional computing mode.

[0053] like Figures 3a-3c As shown, basic kernel fusion methods are mainly divided into three categories. Figure 3a This demonstrates a thread-level internal fusion method that merges the computation of two kernels into a single thread, with the thread and block dimensions becoming... (sth is the final selected value of the "block matrix height" in shared memory, while sth1 and sth2 are the "thread dimensions" calculated separately for different computing cores (CUDA Core / Tensor Core)) and (sbk is the thread block dimension, and sbk1 and sbk2 are the thread block dimensions calculated for CUDA Core and Tensor Core respectively.) Theoretically, this method can achieve optimal load balancing, but due to a serious instruction serialization problem, it cannot utilize two computing cores at the same time. Figure 3b This demonstrates a thread-block-based fusion method, where different threads within the same block are assigned to different kernels based on their thread IDs, thus transforming the dimensions of threads and blocks into... and This method has serious thread branching issues, leading to thread pauses and insufficient resource utilization. Figure 3c This demonstrates a block-level fusion method, where computations from different kernels are distributed across different thread blocks, with the dimensions of threads and blocks being [missing information]. and This method is simple to implement, but it may lead to workload imbalance across thread blocks, causing workload scheduling problems between kernels.

[0054] Resource and memory access characteristic analysis clarifies: the resource idle ratio of the two types of cores during serial execution (e.g., a certain type of core has an idle rate of 40%); the differences in hardware resource consumption between the two types of cores under different matrix dimensions (e.g., non-multiples of 16 dimensions in irregular matrices) (e.g., register usage when CUDA Cores process small blocks, shared memory waste when Tensor Cores process irregular blocks); ultimately transforming the fuzzy problem of "resource idleness" into a quantifiable optimization target, providing "specific data on hardware bottlenecks to be addressed" for subsequent collaborative scheduling. It also clarifies the differences in memory access preferences between the two types of cores: for example, Tensor Cores need to access 16×16 matrix blocks continuously (matching their data granularity), while CUDA Cores may have discrete memory accesses when processing irregular blocks; and the access latency and bandwidth consumption of different memory levels (global memory / shared memory): for example, global memory access latency is hundreds of times higher than shared memory, requiring data reuse through shared memory, but the shared memory usage requirements of the two types of cores (e.g., Tensor Cores) need to be clarified first (e.g., register usage when CUDA Cores process small blocks, shared memory waste). How much space is needed for Core double buffering? These analysis results provide direct basis (data support) for subsequent "memory coordination mechanism of hybrid core fusion algorithm" and "shared memory partitioning and double buffering design", avoiding performance loss caused by blind memory access design. Specifically, hybrid core fusion needs to achieve "cooperative parallelism between instruction level and thread block level", and the premise of parallel scheduling is to clarify "whether the thread configuration requirements of the two types of cores conflict" and "whether memory access can be coordinated" - through the analysis of "thread block resource occupation" and "memory access path", the dimension of the thread block after fusion (such as whether (32,32) is suitable for the two types of cores) and the synchronization time of memory access can be determined to avoid resource contention or thread divergence after fusion; dynamic task partitioning needs to adaptively allocate tasks based on "matrix irregularity" and "real-time hardware resource status", and "real-time hardware resource status" (such as the current CUDA) The quantitative standards for Core utilization and remaining shared memory capacity are predefined through the analysis of "hardware resource usage characteristics". At the same time, the analysis of "irregular matrix memory access overhead" also provides a basis for judging the dynamic adjustment of "thread partitioning coefficient" (e.g., when the degree of irregularity is high, more tasks should be allocated to CUDA Cores to reduce filling overhead).

[0055] Based on the above analysis, such as Figure 4As shown, this application constructs a hybrid core fusion algorithm. Through a unified thread block grid and memory access coordination mechanism, it achieves collaborative parallel task scheduling of CUDA Cores and Tensor Cores at the instruction and thread block levels, eliminating alternating idle states between cores. A two-level fusion mechanism is adopted: at the macro level, inter-block fusion is achieved by aggregating thread block spaces of different computing tasks, naturally expanding parallel computing resources; at the micro level, intra-block fusion is achieved using a carefully designed (32,32) thread block dimension and conditional partitioning based on threadIdx.y. This design ensures that all 32 threads within each thread bundle follow the same execution path, completely eliminating thread branches. The proposed hierarchical fusion retains the resource integration advantages of inter-block fusion while adapting to the special requirements of Tensor Cores for matrix computation through intra-block fusion, ensuring sufficient computing resource supply.

[0056] (32,32) The core meaning of thread block dimension refers to the fact that when the GPU kernel starts, each thread block is configured as a two-dimensional (2D) thread grid structure, with specific parameters clearly defined:

[0057] Horizontal axis (x-axis, corresponding to threadIdx.x): contains 32 threads, with index values ​​ranging from 0 to 31;

[0058] Vertical axis (y-axis, corresponding to threadIdx.y): contains 32 threads, with index values ​​ranging from 0 to 31;

[0059] The total number of threads in a single thread block is 32 × 32 = 1024. This value not only meets the maximum number of threads in a single thread block of NVIDIA GPUs, but is also one of the optimal values ​​to make full use of hardware resources.

[0060] This dimension was not chosen arbitrarily, but rather it was an optimized design specifically for the GPU hardware execution characteristics and the dual-core collaboration requirements, with two core objectives:

[0061] Matching the Warp thread bundle characteristic, thread divergence is avoided at its source: The smallest execution unit of NVIDIA GPUs is the Warp (thread bundle, 32 consecutive threads by default). Threads within the same Warp must execute the same instructions; otherwise, thread divergence will be triggered (some threads wait, resulting in a significant performance drop). In the (32,32) dimensional design, the 32 threadIdx.x threads (x-axis 0~31) corresponding to the same threadIdx.y form a complete Warp, laying the hardware foundation for subsequent task division along the y-axis and avoiding thread divergence.

[0062] It simultaneously adapts to the computing needs of both CUDA Core and Tensor Core: Tensor Core's efficient computing relies on 16×16 matrix blocks (MMA operation granularity), and (32,32) thread blocks can be perfectly split and adapted to this granularity; while when CUDA Core processes irregular matrix blocks, the 32-thread warp granularity can also ensure the parallel efficiency of general computing, achieving conflict-free collaboration between the two cores within the same thread block.

[0063] The conditional partitioning based on threadIdx.y is implemented within the (32,32) thread block. It is the core logic for allocating tasks between CUDA Cores and Tensor Cores. Essentially, it uses the y-axis index of the thread as the basis for judgment, dividing 1024 threads into two groups and binding them to two types of heterogeneous cores respectively.

[0064] In a specific embodiment of this application, the relative computing power ratio of CUDA Cores and Tensor Cores on the target GPU platform is quantified based on their hardware specifications and theoretical computing power. Based on this ratio and the total number of rows in the matrix to be computed, the optimal row partitioning point that should be processed by CUDA Cores is derived. A thread partitioning coefficient aimed at minimizing collaborative computing time is calculated, achieving the theoretically optimal allocation of computing tasks among heterogeneous cores. The specific steps are as follows:

[0065] (1) Establish a total execution time model (to provide a quantitative decision basis for task allocation between CUDA Core and Tensor Core, and to minimize the total computation time of the hybrid kernel by clarifying the time consumption composition and balancing the load, thereby fully releasing the computing power of the GPU heterogeneous resources): The total execution time of the hybrid kernel is , Primarily determined by CUDA Core execution time Tensor Core execution time Inter-thread synchronization overhead and memory transfer time Composition, model representation as The total execution time model breaks down the total execution time of the hybrid kernel into four quantifiable modules: CUDA Core execution time, Tensor Core execution time, thread synchronization overhead, and memory transfer time. This avoids the blindness of "fuzzy optimization"—allowing subsequent strategies to precisely target core time-consuming modules (especially the execution time of the two types of cores), rather than indiscriminately optimizing all aspects. The core optimization goal of the model is to "minimize the total execution time," and the key to achieving this goal is to balance the execution time of CUDA Cores and Tensor Cores (avoiding one type of core completing its task early and being idle while the other type of core is still computing). Through this model, combined with the hardware specifications of the two types of cores (such as the computing power ratio and the number of operations per cycle) and the matrix dimension (such as the total number of rows m), the "thread partitioning coefficient" and the optimal task allocation ratio of the two types of cores (such as the number of matrix rows that CUDA Cores should process) can be derived, ensuring that task partitioning is upgraded from "empirical value" to "theoretical optimal solution." The total execution time model is linked with the shared memory resource constraint model—when deriving the task partitioning ratio, the upper limit of shared memory usage by the two types of cores (such as Tensor Cores) is considered simultaneously. Core double buffering (memory requirements) avoids shared memory conflicts caused by unreasonable task allocation or additional data transfer time caused by insufficient memory, ensuring the optimization goals are achieved from a resource perspective.

[0066] (2) Perform execution time modeling for CUDA Core and Tensor Core respectively (to provide accurate quantitative basis for task allocation and load balancing of dual-core, clarify the performance difference between the two types of heterogeneous cores, and support the derivation of the optimal task partitioning strategy): For a Matrix A and Multiplication of matrix B. Matrix A contains m rows of data, which is the vertical dimension of matrix A. In this invention, the m rows will be dynamically divided into two parts, which are respectively assigned to CUDA Core and Tensor Core for processing (CUDA Core processes a portion of the rows, and Tensor Core processes the remaining rows). Here, k represents the number of columns in matrix A and the number of rows in matrix B, which is the "compatibility dimension" of matrix multiplication (the number of columns in matrix A must equal the number of rows in matrix B for multiplication to be performed); it represents the "intermediate dimension" when multiplying the two matrices, determining the number of matching elements in each row / column. Matrix B contains n columns of data, which is the horizontal dimension of matrix B; the final result matrix C of matrix A × B has a dimension of m × n (number of rows = number of rows in A, number of columns = number of columns in B).

[0067] Assumption of assignment If the line is processed by CUDA Core, then the execution time of CUDA Core... Modeling as a formula ,in , , , These represent the total number of CUDA cores, utilization, operations per cycle, and startup latency of a CUDA Core, respectively. Similarly, Tensor Cores process the remaining... Okay, its execution time Modeling as a formula ,in , , , These represent the total number of Tensor Cores, utilization, number of operations per cycle, and startup latency of a Tensor Core, respectively.

[0068] (3) Constructing a shared memory resource allocation constraint model: To improve computing performance, a data reuse strategy based on shared memory is adopted. The data reuse degree R is defined, and the shared memory usage of CUDA Cores and Tensor Cores is constrained by the following formulas: and formula To ensure that both have the total available capacity of the limited GPU shared memory. No conflict will occur, among which This represents the total capacity occupied by CUDA Cores in shared memory. This represents the total capacity occupied by Tensor Cores in shared memory. It represents the middle dimension of the block matrix (the block size corresponding to the middle dimension k of matrix multiplication). The height of the block matrix (the block size corresponding to the number of rows in matrix A). The width of the block matrix (the block size corresponding to the number of columns in matrix B), and a double buffering mechanism (i.e., double buffering) is introduced for the Tensor Core. This is to hide data loading delays.

[0069] The constraint model quantifies the memory usage of the two types of cores, specifying the "buffer capacity of CUDA Cores," the "double buffer capacity of Tensor Cores," and the upper limit of their sum. This ensures that the two types of cores have independent, non-overlapping spaces in shared memory, physically guaranteeing that parallel work does not interfere with each other. Simultaneously, by relating the "block size (sth, sbk)" to the "memory capacity," the constraint model ensures that the block size meets the computational needs of both types of cores (e.g., the 16×16 granularity of Tensor Cores) without causing memory shortages due to excessively large blocks. This achieves a balance between "maximizing data reuse" and "rationalizing memory usage," avoiding additional data transfer time caused by insufficient memory. Furthermore, when deriving the "thread partitioning coefficient" (task allocation ratio), the constraints of shared memory (such as the memory requirements of double buffering of Tensor Cores) should be considered simultaneously to avoid memory overflow due to excessive task allocation; the parameters of the constraint model (such as sth, sbk) will directly affect the execution time of the two types of cores (block size affects computational efficiency), thereby adjusting the optimization direction of the total execution time model in reverse to ensure that the goal of load balancing is achieved under limited resources.

[0070] (4) Derive the partition coefficient formula: To minimize the total time, the ideal state is to let (Load balancing conditions), ignoring startup latency, substitute the time model in (2) above into the balancing conditions, and combine the characteristic that Tensor Core completes 256 floating-point operations per MMA operation to derive the partitioning coefficient. and the number of lines that CUDA Core should process .

[0071] In a specific embodiment of this application, a shape adaptation and memory optimization strategy for irregular matrices is designed. The input matrix is ​​padded and aligned to meet the data granularity requirements of Tensor Core computation; independent data regions are designed for CUDA Cores and Tensor Cores in shared memory, and the data layout is optimized to avoid memory conflicts; based on the composite kernel and task partitioning results, unified kernel startup and dynamic task distribution at runtime are achieved. The specific steps are as follows:

[0072] (1) Data preparation and shape adaptive padding: Before computation begins, the dimensions of the input matrix are padded so that both its width and height are multiples of 16. This alignment is not arbitrary but is strictly designed to meet the specific data granularity (such as 16×16 matrix blocks) and memory access patterns required for efficient TensorCore execution, and is a key prerequisite for maximizing its peak performance. Padding ensures that the matrix blocks subsequently allocated to TensorCore for processing have a regular shape.

[0073] (2) Design of shared memory partitioning and double buffering mechanism: Under the constraint of limited shared memory capacity, separate data buffers are allocated for CUDACore and Tensor Core, with a double buffer designed for Tensor Core. Overlapped computation and data loading hide memory access latency. The total shared memory layout is based on the formula... Distribute, among which Indicates the total size of shared memory. As a synchronization flag, To conserve space, data layout is optimized to avoid storage conflicts, thereby minimizing access latency and maximizing memory bandwidth supply to both types of computing cores.

[0074] (3) Achieving unified kernel startup and dynamic task distribution at runtime: Based on the task partitioning results, a unified composite kernel is started. Within this kernel, matrix row computation tasks are dynamically distributed to CUDA Cores or Tensor Cores at runtime according to thread indices and pre-calculated partitioning coefficients. Specifically, through conditional judgment (such as based on threadIdx.y), threads within the same warp collaboratively process the data blocks allocated to their respective computational cores, achieving fine-grained collaborative execution. This mechanism ensures that two heterogeneous computing resources remain busy in the same computational context, maximizing the utilization of GPU's parallel hardware resources.

[0075] like Figure 5 As shown, this invention illustrates the shared memory allocation scheme. First, data preparation is performed by padding the width and height of the input matrix to multiples of 16 to meet the data granularity and memory access patterns required for efficient operation of 16×16 matrix blocks by the Tensor Cores. Then, the computational load is analyzed, and matrix row computations are strategically divided between general-purpose CUDA Cores and dedicated Tensor Cores, ensuring that the portion allocated to the Tensor Cores maintains critical 16-byte memory alignment. Based on the division results, the system intelligently selects the most suitable execution mode and allocates storage for the padded matrix structure. Finally, the shared memory configuration is optimized to minimize access latency, and the corresponding computational kernels are launched according to the selected mode.

[0076] In a specific embodiment of this application, such as Figures 6a-6dTables 7a-7d and 8a-8d respectively demonstrate the GEMM performance evaluation on four representative shape configurations on V100, RTX4090, and A100 GPU architectures. Experimental results show that the proposed method exhibits stable and excellent performance under various matrix configurations. Compared with mainstream computing libraries cuBLAS and CUTLASS, this implementation maintains lower computation time across different matrix dimensions, demonstrating excellent versatility and robustness. Particularly in small-scale square matrix configurations with shapes M=W, N=W, and K=W, its execution time is significantly lower than the comparative libraries, validating the effectiveness of our optimizations in memory access and parallel scheduling. As the matrix size increases, the execution time of all methods increases, but this method exhibits a more gradual growth trend, indicating better scalability.

[0077] The above description is only a preferred embodiment of the present invention and is not intended to limit the present invention. Any modifications, equivalent substitutions, improvements, etc., made within the spirit and principles of the present invention should be included within the protection scope of the present invention.

Claims

1. A dual-core collaborative parallel computing method for irregular matrix multiplication, characterized in that, Includes the following steps: S1. Resource and memory access characteristics analysis: In the GPU hybrid computing mode, the system analyzes the hardware resource usage characteristics of CUDA Core and Tensor Core and their global memory and shared memory access behavior, and quantitatively evaluates the resource idle rate caused by serial execution in the traditional computing mode. S2. Hybrid Core Fusion Scheduling: Based on the analysis results of step S1, a hybrid core fusion algorithm is constructed. Through a unified thread block grid and memory access coordination mechanism, collaborative parallel task scheduling of CUDA Cores and Tensor Cores is achieved at the instruction level and thread block level, eliminating alternating idle states between cores. S3. Dynamic Task Partitioning and Optimization: A dynamic optimization model based on thread partitioning coefficients is introduced. The thread partitioning coefficients are adjustable parameters. The value of the thread partitioning coefficients dynamically determines the size of the matrix row blocks or column blocks allocated to CUDA Cores for processing. This dynamic optimization model adaptively partitions and allocates computing tasks to CUDA Cores and Tensor Cores based on the actual irregularity of the input matrix and the real-time hardware resource status of the GPU, optimizing thread-level load balancing and memory access locality, and avoiding redundant computation and memory overhead caused by padding operations. S4. Collaborative Computation Execution: After scheduling, the CUDA Cores and Tensor Cores collaboratively perform irregular matrix multiplication operations according to the divided tasks.

2. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 1, characterized in that, The hybrid core fusion algorithm described in step S2 adopts a two-level hybrid fusion mechanism. At the macro level, it performs inter-thread block fusion and aggregates task thread block space to expand parallel resources. At the micro level, it adopts a fixed dimension and condition partitioning strategy within the thread block to ensure that all threads in the Warp have consistent execution paths and eliminate thread divergence.

3. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 2, characterized in that, The fixed dimension of the thread block is set to a 2D grid arrangement, where the x-axis contains 32 threads and the y-axis contains 32 threads. Task allocation within the thread block is achieved by dividing based on the condition threadIdx.y.

4. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 1, characterized in that, The thread partitioning coefficient mentioned in step S3 is derived in the following way: Establish a hybrid kernel total execution time model, whereby the total execution time includes CUDA Core execution time, TensorCore execution time, inter-thread synchronization overhead, and memory transfer time; The execution time of CUDA Core and Tensor Core is modeled separately, and the relative computing power ratio is quantified by combining the hardware specifications and theoretical computing power of the two. Construct a shared memory resource allocation constraint model to ensure that CUDA Cores and Tensor Cores do not conflict under limited shared memory capacity; Based on load balancing conditions, ignoring startup latency, and considering the characteristic that Tensor Cores complete 256 floating-point operations per MMA operation, the thread partitioning coefficient and the number of rows that CUDA Cores should process are derived.

5. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 4, characterized in that, In the shared memory resource allocation constraint model, a double buffering mechanism is introduced for Tensor Cores to hide data loading latency. That is, double data buffers are designed for Tensor Cores to overlap computation and data loading, thus hiding memory access latency.

6. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 1, characterized in that, The dynamic task partitioning process in step S3 also includes: padding and aligning the input matrix so that both the width and height of the matrix are multiples of 16, in order to meet the data granularity and memory access patterns required for efficient execution of Tensor Core.

7. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 6, characterized in that, Independent data buffers are allocated for CUDA Cores and Tensor Cores in shared memory, and the data layout is optimized to avoid storage conflicts and maximize memory bandwidth supply.

8. The dual-core collaborative parallel computing method for irregular matrix multiplication according to claim 1, characterized in that, In step S4, a unified composite kernel is started. Within the kernel, matrix row computation tasks are dynamically distributed to CUDA Cores or Tensor Cores at runtime based on thread indexes and pre-computed partitioning coefficients, thereby achieving fine-grained collaborative execution.

Images