Data processing method and electronic device

Abstract

The application discloses a data processing method and electronic equipment in the technical field of computer application, and discloses a data processing method and electronic equipment.In the application, for the target hardware, the corresponding hardware image can be acquired, so that the candidate sub-blocks are constructed based on the hardware image, when the request instruction is received, different use schemes of the sub-blocks matched with the instruction are generated, the optimal use scheme is selected based on the quantitative evaluation, and the resources of the target hardware are coordinately configured, so that the sub-blocks used on the target hardware are more consistent with the current hardware and computing requirements, the filling waste is effectively eliminated, the resource utilization is effectively improved, the hardware limitation faced by the fixed sub-block scheme is broken, the application can be applied to any type of hardware, and the technical effect of enhancing the universality of cross hardware platforms is achieved.

Description

Technical Field

[0001] This invention relates to the field of computer application technology, and in particular to a data processing method and an electronic device. Background Technology

[0002] In GEMM (General Matrix Multiply) optimization techniques, redundant computations often occur when dealing with non-standard matrix sizes, resulting in a lack of cross-hardware platform versatility. Furthermore, when faced with non-standard matrix sizes or novel computing architectures, manual adjustment of block parameters is required, and thread mapping rules can easily lead to shared memory bank conflicts.

[0003] Therefore, how to improve the optimization technology of GEMM to achieve cross-hardware platform versatility is a technical problem that urgently needs to be solved by those skilled in the art. Summary of the Invention

[0004] This invention provides a data processing method and an electronic device to improve the optimization technology of GEMM and achieve cross-hardware platform versatility.

[0005] This invention provides a data processing method, comprising: acquiring a hardware profile of target hardware; determining the atomic instruction specifications supported by the target hardware using the hardware profile, and constructing multiple candidate blocks that conform to hardware constraints based on the atomic instruction specifications; receiving a request instruction, and determining different usage schemes for the blocks using the instruction specifications of the request instruction and the candidate blocks; quantitatively evaluating the usage schemes, and determining a target scheme from the usage schemes using the quantitative evaluation results; coordinating and configuring the resources of the target hardware according to the target scheme, and performing matrix multiplication calculations using the blocks selected by the target scheme.

[0006] The present invention also provides a data processing apparatus, comprising: a hardware profiling module for acquiring a hardware profile of a target hardware; a block pre-construction module for determining the atomic instruction specifications supported by the target hardware using the hardware profile, and constructing multiple candidate blocks that conform to hardware constraints based on the atomic instruction specifications; a scheme generation module for receiving a request instruction, and determining different usage schemes for the blocks using the instruction specifications of the request instruction and in combination with the candidate blocks; a scheme optimization module for quantitatively evaluating the usage schemes, and determining a target scheme from the usage schemes using the quantitative evaluation results; and a data calculation module for coordinating and configuring the resources of the target hardware according to the target scheme, and then performing matrix multiplication calculations using the blocks selected by the target scheme.

[0007] The present invention also provides an electronic device, comprising: a memory for storing a computer program; and a processor for executing the computer program to implement the steps of any of the above-described data processing methods.

[0008] The present invention also provides a computer-readable storage medium storing a computer program, wherein the computer program, when executed by a processor, implements the steps of any of the above-described data processing methods.

[0009] The present invention also provides a computer program product, including a computer program that, when executed by a processor, implements the steps of any of the above-described data processing methods.

[0010] This invention improves GEMM's optimization technology and achieves cross-hardware platform versatility by employing a dynamic approach to the target hardware. Specifically, to ensure that the hardware block usage scheme conforms to the current hardware constraints, a hardware profile of the target hardware is first obtained, and multiple candidate blocks conforming to the hardware constraints are constructed based on the atomic instruction specifications determined by this hardware profile. Upon receiving a request instruction, different usage schemes for the blocks are determined using the instruction specifications of the request instruction and the candidate blocks. By evaluating the block usage schemes, the optimal target scheme can be identified for resource coordination and matrix multiplication calculations.

[0011] Because it targets specific hardware, its corresponding hardware profile can be obtained, allowing candidate blocks to be constructed based on this profile. Upon receiving a request instruction, different usage schemes are determined, and the optimal target scheme is selected based on quantitative evaluation. The resources of the target hardware are then coordinated and configured, ensuring that the final executed target scheme better meets the current hardware and computing requirements. This effectively eliminates filler waste and improves resource utilization, breaking through the hardware limitations of fixed block schemes. It can be applied to any type of hardware, enhancing the technical effect of cross-hardware platform versatility. Attached Figure Description

[0012] To more clearly illustrate the embodiments of the present invention, the accompanying 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.

[0013] Figure 1 A flowchart illustrating a data processing method provided in an embodiment of the present invention;

[0014] Figure 2 This is a schematic diagram illustrating an embodiment of a data processing method provided by the present invention.

[0015] Figure 3A detailed schematic diagram illustrating a data processing method provided in an embodiment of the present invention;

[0016] Figure 4 This is a schematic diagram of the structure of a data processing device provided in an embodiment of the present invention;

[0017] Figure 5 This is a schematic diagram of the structure of an electronic device provided in an embodiment of the present invention;

[0018] Figure 6 This is a schematic diagram of the specific structure of an electronic device provided in an embodiment of the present invention. Detailed Implementation

[0019] 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. Based on the embodiments of the present invention, all other embodiments obtained by those of ordinary skill in the art without creative effort are within the protection scope of the present invention.

[0020] It should be noted that, in the description of this invention, the terms "comprising," "including," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements includes not only those elements but also other elements not expressly listed, or elements inherent to such a process, method, article, or apparatus. The terms "first," "second," etc., used in this invention are used to distinguish similar objects and are not used to describe a specific order or sequence.

[0021] Due to limitations in hardware and storage resources, GEMM computation often requires block processing. Most GEMM optimization techniques employ fixed block strategies, pre-setting block sizes (e.g., 16x16x16 or 32x32x8) for specific GPU (Graphics Processing Unit) architectures, lacking cross-platform versatility. When faced with non-standard matrix sizes or new computing architectures, these solutions require manual adjustment of block parameters, and thread mapping rules (e.g., row-major indexing) can easily lead to shared memory bank conflicts. GPU shared memory is divided into multiple independent storage units called banks (typically 32). Each bank can read and write data simultaneously, but only one thread can access a single bank within the same clock cycle. When multiple threads access different addresses within the same bank simultaneously, these accesses are executed serially (queuing), resulting in performance degradation, known as bank conflicts. Furthermore, solutions supporting flexible configuration rely on developers manually adapting to hardware characteristics, leading to higher learning costs and maintenance difficulties.

[0022] For example, a fixed-block strategy can be used to optimize the 16x16x16 block computation mode specifically for its Tensor Cores (dedicated hardware computing units). For instance, when running FP16 matrix multiplication on a GPU, cuBLAS (a high-performance linear algebra library) strictly divides the input matrix into 16x16x16 blocks, matching hardware characteristics through hard-coded warp scheduling and shared memory allocation strategies. This approach is highly efficient on standard-sized matrices (e.g., 2048x2048), but it incurs approximately 12.5% ​​padding waste for matrix sizes that are not multiples of 16 (e.g., 2000x2000) and is incompatible with other vendors' hardware architectures.

[0023] Furthermore, the HIPBLAS library (an open-source linear algebra library) employs a 32x32x8 partitioning scheme specifically optimized for certain accelerator cards. When processing int8 matrix operations, this scheme requires the matrix dimension to be a multiple of 32; otherwise, zero-padding is necessary. This not only increases computational overhead but also necessitates maintaining two different sets of kernel code in hybrid computing environments (such as those using different hardware devices simultaneously), significantly increasing development and maintenance costs. These examples demonstrate that related technologies are limited by vendor-specific hardware designs, lacking cross-platform versatility and dynamic adaptability.

[0024] As can be seen, adopting a fixed block partitioning strategy, which strictly relies on the computing unit design of specific hardware, leads to three key problems: First, when the matrix size is not an integer multiple of the block size (e.g., a 2000x2000 matrix needs to be padded to 2048x2048), computational redundancy occurs; second, cross-platform compatibility is poor, requiring separate development and maintenance of independent code for different graphics cards; third, the thread scheduling rules are fixed, making it impossible to dynamically optimize shared memory access patterns for different hardware, which easily leads to Bank conflicts. These limitations significantly restrict the efficiency and versatility of related solutions in heterogeneous computing scenarios.

[0025] In order to solve the above problems, this invention proposes to perform dynamic block division before matrix multiplication calculation, so that the block division used on the target hardware is more in line with the current hardware and computing requirements. This can effectively improve resource utilization by eliminating fill waste and enhance cross-hardware platform versatility.

[0026] Specifically, in order to enable those skilled in the art to better understand the present invention, the present invention will be further described in detail below with reference to the accompanying drawings and specific embodiments.

[0027] Please refer to Figure 1 The data processing method provided in this embodiment of the invention includes the following steps.

[0028] S101. Obtain the hardware profile of the target hardware.

[0029] In order to ensure that dynamic partitioning conforms to hardware constraints, in this embodiment, the hardware profile corresponding to the target hardware performing matrix multiplication calculation is first obtained.

[0030] Hardware profiling refers to a comprehensive and accurate characterization and description of hardware formed by collecting and integrating multi-dimensional data information of hardware devices and using specific analysis methods.

[0031] In this embodiment, a hardware profile can be obtained by pre-collecting hardware data and storing it in a storage medium, then reading the data. Alternatively, a hardware profile can be obtained by scanning and exploring the target hardware using data acquisition tools. For example, key parameters of the target hardware can be obtained through the GPU driver interface, including: supported MMA instruction shapes (16x16x16, 32x32x8), number of registers, shared memory size and number of banks, number of SM cores, and maximum thread block size.

[0032] In one specific embodiment of the present invention, obtaining a hardware profile of the target hardware includes: using a hardware detection engine to detect the target hardware and obtain a hardware profile.

[0033] In this embodiment, a hardware detection engine can be used to detect the target hardware, thereby obtaining a hardware profile. The target hardware can be a GPU or other hardware capable of implementing GEMM. The target detection engine is an engine that can obtain relevant information about the target hardware (such as storage, supported request commands, parallel processing, etc.).

[0034] In one specific embodiment of the present invention, a hardware probing engine is used to probe the target hardware to obtain a hardware profile, including: probing the computation instruction analysis layer of the target hardware using the hardware probing engine to obtain the matrix computation instruction set specifications supported by the target hardware; probing the storage system parsing layer of the target hardware using the hardware probing engine to obtain storage structure information; probing the parallel architecture modeling layer of the target hardware using the hardware probing engine to obtain computing unit organization structure information; and constructing a hardware profile using the matrix computation instruction set specifications, storage structure information, and computing unit organization structure information.

[0035] In other words, the system can first activate the hardware detection engine to build a complete device profile through multi-level hardware characteristic scanning. Specifically, the computation instruction analysis layer accurately identifies the matrix computation instruction set specifications supported by the GPU, including core parameters such as instruction shape and data type support. These core parameters form the theoretical basis for subsequent block design. The storage system analysis layer maps in detail the capacity, bandwidth, and access characteristics of storage structures such as register files, shared memory, and caches, especially the bank distribution pattern of shared memory. This storage structure information provides a basis for storage access optimization. The parallel architecture modeling layer analyzes the organizational structure of computing units, including key parameters such as SM array size and thread scheduling granularity. These key parameters determine the design space of parallel strategies.

[0036] S102. Use hardware profiling to determine the atomic instruction specifications supported by the target hardware, and construct multiple candidate blocks that conform to hardware constraints based on the atomic instruction specifications.

[0037] After obtaining the hardware profile, the rules of instruction supported by the target hardware can be determined based on the hardware profile, and candidate blocks that conform to the hardware constraints can be constructed based on the atomic instruction specifications. In this embodiment, several candidate blocks can be constructed for subsequent selection.

[0038] Specifically, based on the common matrix input specifications, the underlying blocks can be initially combined. For example, most users' input matrices are 1024x1024 or 2048x2048, which are divisible by 32, 64, 128, etc. Based on the hardware-supported instructions, such as a computing card (target hardware) supporting 16x16 or 16x8 instructions, a preliminary combination can be performed. Sixteen 16x16 matrices can be combined into a 64x64 matrix (16x16 matrices are superimposed four times in the first dimension to become 64x16, and superimposed four times in the second dimension to become 64x64, i.e., 16 16x16 matrices = 64x64). This is the first sub-block combination, which is based on the instructions supported by the actual computing card. It divides the data into blocks of 64 and 128 sizes. It can also be divided according to the needs of the scenario. For example, if the input of a large model in a certain field is 240x160, some irregular blocks can be divided to deal with irregular blocks. For example, some 24x16 blocks can be divided first. In this way, the user's 240x160 input can be processed by 24x16 blocks. The 24x16 blocks are composed of a 16x16 block at the bottom layer plus an 8x16 block at the bottom layer (if the request instruction is available, it can be supported).

[0039] In one specific embodiment of the present invention, constructing multiple candidate blocks that conform to hardware constraints based on atomic instruction specifications includes: determining the atomic instruction specifications as basic computing units; and constructing candidate blocks of atomic instruction specifications whose size is an integer multiple of the basic computing units.

[0040] In the feasibility space construction, i.e., the construction of candidate blocks, candidate blocks that meet hardware constraints can be generated based on the instruction set specifications. For example, firstly, the atomic instruction specifications supported by the hardware (such as 16x16x16 MMA instructions) are used as the basic computational units, and then candidate blocks are generated according to the computational scale of the target matrix multiplication. Specifically, the macroscopic block task is decomposed into executable combination patterns of multiple atomic instructions. For example, if the specified specification of the multiplication matrix is ​​64x64x64, it can be decomposed into 4×4×4 candidate blocks of 16x16x16 atomic instructions, or into 4×8×4 candidate blocks composed of 16x8x16 sub-blocks (each 16x8x16 sub-block is implemented by 2 16x8x8 atomic instructions), ensuring that each candidate block strictly satisfies the integer multiple combination relationship of atomic instructions, thereby guaranteeing the hardware-level executability of all generated block usage schemes.

[0041] It is important to note that since the target hardware may support different atomic instructions, each candidate block strictly satisfies the integer multiple combination relationship of atomic instructions supported by the current hardware. This means that the specification of the candidate block is an integer multiple of the specifications of multiple atomic instructions. A candidate block may correspond to only one atomic instruction specification or multiple atomic instruction specifications.

[0042] S103. Receive the request instruction and determine different usage schemes for the block by using the instruction specification of the request instruction and combining it with the candidate blocks.

[0043] In this embodiment of the invention, upon receiving a request instruction, since the target hardware has already pre-divided the candidate blocks to match the current hardware, different usage schemes corresponding to different blocks and reuse times can be determined from the candidate blocks based on the instruction specifications of the request instruction. That is, each usage scheme can complete the matrix multiplication calculation corresponding to the request instruction, but due to the different blocks selected and the different block reuse methods, the technical effects of these usage schemes, such as latency, will differ. Different usage schemes of blocks may have completely different selected blocks or partially overlapping blocks, and the reuse times of different usage schemes may be the same or different.

[0044] In one specific embodiment of the present invention, different usage schemes of blocks are determined by using the instruction specification of the request instruction and combining it with the candidate blocks, including: obtaining the instruction specification of the request instruction and obtaining the block specification of the candidate blocks; using the instruction specification and the block specification, determining the blocks and reuse parameters corresponding to different usage schemes.

[0045] The instruction specification refers to the data size corresponding to the request instruction, such as 1024x1024. The candidate block specification refers to the data size supported by the candidate block, such as 64x64. The size of the selected candidate block needs to match the instruction specification of the request instruction; that is, the specifications of one or more selected candidate blocks, when superimposed and combined for multiplexing, should be greater than or equal to the instruction specification.

[0046] For example, an instruction specification of 1024x1024 can be composed of many combinations of candidate blocks. Larger candidate blocks mean fewer blocks are used (e.g., using all 32x32 pre-blocks in parallel for 32x32 iterations can complete the 1024x1024 calculation), which may lead to insufficient parallelism, low computational efficiency, and underutilization of hardware. However, using smaller candidate blocks results in a larger number of blocks, and excessive parallelism can cause resource queuing, putting too much pressure on the hardware and reducing efficiency (e.g., using all 32x8 pre-blocks in parallel for 32x32x4 iterations is needed to complete the 1024x1024 calculation). Therefore, in selecting candidate blocks, it is necessary to consider not only the instruction specification and block specifications but also load balancing strategies. For example, a numerical range for candidate blocks and a hardware utilization ratio range can be set, allowing the selection of candidate blocks that meet these requirements. In practical applications, candidate block selection strategies corresponding to some common instruction specifications can be directly specified. For example, for a 1024 x 1024 instruction specification, four 64*64 candidate blocks can be directly specified and reused four times.

[0047] For example, after hardware profiling, candidate blocks corresponding to the target solution are obtained. These candidate blocks include blocks with block sizes of 32x32 and 16x32. When user input (i.e., request commands and data) arrives, for example, a request command of 1024x1024, then a complete 1024x1024 can be calculated using 32x32 candidate blocks arranged in an 8x8 pattern. Here, 8x8 represents the selection and use of blocks based on user input (it can also be considered as further segmentation based on candidate blocks, i.e., dividing smaller blocks into larger-scale blocks). Of course, 32x32 candidate blocks can be combined to form 1024x1024, and 16x32 candidate blocks can also be combined to form 1024x1024. The specific combination method chosen depends on dynamic decision-making.

[0048] S104. Conduct a quantitative evaluation of the usage plan and use the results of the quantitative evaluation to determine the target plan from the usage plan.

[0049] After determining the usage schemes for the partitioned blocks, these schemes can be quantitatively evaluated. Based on the evaluation results, the optimal target scheme can be selected from the various usage schemes. The quantitative evaluation can assess different usage schemes from dimensions such as storage, resource utilization, and efficiency, thus selecting the scheme with the best evaluation across these dimensions as the target scheme. For example, the system's quantitative evaluation tool can be used to quantitatively evaluate the usage schemes; this tool can be any tool that can quantitatively evaluate different usage schemes from dimensions such as storage, resource utilization, and efficiency.

[0050] In one specific embodiment of the present invention, the quantitative evaluation of the usage scheme includes: determining the instruction coverage, storage access conflict rate, and resource utilization of the usage scheme; performing a weighted summation of the instruction coverage, storage access conflict rate, and resource utilization to obtain a quantitative value for each usage scheme; and determining the quantitative value as the quantitative evaluation result.

[0051] Among them, determining the target block scheme from the usage schemes using the quantitative evaluation results includes: sorting the quantitative values ​​to determine the usage scheme with the highest quantitative value; and determining the usage scheme with the highest quantitative value as the target scheme.

[0052] Multi-objective evaluation can be performed on various usage schemes for the partitions, i.e., a quantitative evaluation of each usage scheme, including: computational density evaluation (i.e., determining resource utilization): accurately calculating the occupancy rate of each partition on the computing unit (i.e., utilization); storage stress testing (i.e., determining storage access conflict rate): simulating and predicting the probability of shared memory access conflicts (i.e., storage access conflict rate); boundary effect analysis (i.e., determining instruction coverage): calculating the padding overhead of unaligned parts (i.e., instruction coverage). Then, a dynamic decision-making process is used to select the target scheme. That is, a weighted scoring model can be used, combined with the characteristics of the current computing task (such as matrix size), to select the optimal usage scheme. The weighted scoring model evaluates the partition usage schemes through three core dimensions: instruction coverage (weight 40%), storage access conflict rate (weight 35%), and resource utilization rate (weight 25%). In practice, a full-system performance analysis tool is used for each usage scheme to test and obtain quantitative indicators such as instruction coverage, storage access conflict rate (i.e., bank conflict rate), and resource utilization rate (i.e., register utilization rate). The total score is calculated by weighted summation, and the scheme with the highest score is selected. For example, if the 16x16x16 block partitioning scheme performs well in terms of computational efficiency (0.9 points) and storage access (0.7 points), its total score = 0.9×0.45+0.7×0.4+0.6×0.15=0.775, which is better than other partitioning schemes, and it can be used as the target scheme.

[0053] In one specific embodiment of this invention, usage schemes are output on a visual interface for users to choose from; the visual interface is monitored to obtain user-specified information; and a target scheme is determined from the usage schemes based on the user-specified information. That is, after determining the different usage schemes for each segment, they can be output externally to facilitate user selection, allowing users to choose and confirm according to their actual needs.

[0054] S105. After coordinating and configuring the resources of the target hardware according to the target scheme, matrix multiplication is performed using the blocks selected in the target scheme.

[0055] Once the target solution is obtained, the resources of the target hardware can be coordinated and configured according to that solution. That is, the resources are determined and allocated in blocks according to the target solution, thus enabling matrix multiplication calculations to be performed based on the target solution.

[0056] In one specific embodiment of this invention, the resources of the target hardware are coordinated and configured according to the target scheme, including: dividing the matrix into blocks based on the block size in the target scheme; using the computing unit organization structure information, setting the number of threads in the thread block to an integer multiple of the number of storage units, and allocating them cyclically so that each thread bundle processes the same number of matrix blocks, while keeping the data regions processed by adjacent thread bundles continuous. That is, after determining the target scheme, the thread block dimension and thread bundle allocation can be determined based on the block size and the computing unit organization structure information in the hardware profile, thereby ensuring full utilization of computing resources. Here, a thread bundle (Warps or warp) is the core execution unit in the GPU parallel computing architecture and a key design for achieving efficient parallel computing in GPUs; essentially, it is a group of threads simultaneously scheduled and executed by the GPU hardware.

[0057] Based on the block size and hardware parallelism capabilities, the thread block dimension and thread bundle allocation scheme can be accurately calculated to ensure full utilization of computing resources. A warp-level load balancing algorithm is adopted: the number of threads in each thread block is set to an integer multiple of 32 (the number of storage units in the normal case, but other values ​​are also possible) (e.g., 128 threads = 4 warps). A round-robin allocation strategy ensures that: each warp processes the same number of matrix blocks (e.g., each warp is responsible for 2 8×8 sub-blocks); and the data areas processed by adjacent warps remain contiguous, improving cache hit rate.

[0058] After the resource coordination and configuration are completed, the matrix multiplication calculation corresponding to the request instruction can be executed on the selected block of the target scheme to obtain the data processing result.

[0059] This means performing matrix multiplication calculations on the target hardware, i.e., performing operations on blocks determined by dynamic partitioning.

[0060] In one specific embodiment of the present invention, matrix multiplication calculation using the blocks selected by the target scheme includes: during the matrix multiplication calculation using the selected blocks, decomposing non-standard blocks into optimized combinations that match the blocks selected by the target scheme; and / or, during the matrix multiplication calculation using the selected blocks, if the data size is smaller than the selected blocks, zero-padding is performed on the regions that are not filled with data, and the invalid calculation of the zero-padding regions is skipped by masking the zero-padding regions.

[0061] That is, when executing a request instruction, the instruction can be reorganized, that is, non-standard blocks can be broken down into optimized combinations of request instructions, such as converting a 16x8x16 block into two 16x8x8 instruction sequences.

[0062] When executing request instructions, masks can be used to intelligently skip invalid calculations in padding regions while maintaining the continuity of the instruction pipeline. Specifically, masks can be used to obscure, hide, or replace invalid data in padding regions to avoid calculating invalid data, thereby reducing invalid computations. For example, a mask can be generated for each data element to identify whether it is valid data. When executing calculation instructions, the hardware or compiler dynamically skips the calculation of invalid elements based on the mask, without interrupting the instruction flow, keeping the pipeline fully loaded. If the target hardware natively supports 16×16×16 calculation blocks, but the actual data is 16×15×16, masking techniques can be used to treat this data as 16×16×16, allowing the hardware to run at full speed, while muting the data in the 16th column and excluding it from the result, thus maintaining full pipeline load while avoiding invalid computations.

[0063] Instruction reorganization and masking to skip invalid calculations can be used individually or in combination; the two are not mutually exclusive.

[0064] In one specific embodiment of this invention, matrix multiplication calculation is performed using the block partitioning selected in the target scheme. This includes: during shared memory access, adding a base offset in the row direction and applying a rotation step in the column direction to ensure that concurrent access request instructions are evenly distributed across different memory banks. This optimizes storage access. Specifically, a conflict-preventing access mode can be designed based on bank distribution characteristics. A dynamic offset strategy can be used to ensure that concurrent accesses are evenly distributed across different banks. For example, a composite addressing mode of row offset + column rotation can be used for shared memory access. In the row direction: a base offset of threadIdx.x % 32 is added; in the column direction: a rotation step of (threadIdx.y * 7) % 32 is used (7 and 32 are coprime). Here, threadIdx.x is a built-in variable for thread ID, used to identify the index of the current thread in the x-dimensional region within the thread block; threadIdx.y is one of the built-in variables, used to identify the index of the current thread in the y-dimensional region of its respective thread block.

[0065] In one specific embodiment of this invention, the method further includes: monitoring the registers of the target hardware; and adjusting the loop unrolling factor to reduce register usage when register demand reaches a threshold. Register usage can be optimized through activity analysis to avoid register overflow while ensuring performance. The loop unrolling factor is adjusted to reduce register usage when register demand is detected to be close to its upper limit. In loop unrolling optimization technology, the loop unrolling factor is the number of original loop iterations processed in a single iteration. It determines the batch processing scale of the code after loop unrolling, and its core function is to balance performance improvement and code overhead.

[0066] In one specific embodiment of this invention, continuous monitoring during system operation is also included: performance indicator collection: real-time acquisition of key indicators such as instruction throughput and storage bandwidth utilization; parameter fine-tuning: dynamic adjustment of execution parameters based on runtime data, such as appropriately adjusting the thread block size to adapt to the actual load; and anomaly handling: automatic triggering of the re-optimization process when a surge in bank conflicts is detected. That is, in this embodiment, after dynamic block partitioning, performance can be monitored during matrix multiplication calculations based on the dynamic block partitioning, and the block size can be adaptively adjusted based on the collected key indicators. Furthermore, re-optimization can be performed after a sudden increase in storage conflicts within a short period of time.

[0067] By applying the method provided in this embodiment of the invention, in order to improve the optimization technology of GEMM and achieve cross-hardware platform versatility, a dynamic scheme is adopted for the use of target hardware. Specifically, in order to ensure that the scheme for using hardware blocks conforms to the hardware constraints of the current hardware, a hardware profile of the target hardware is first obtained, and multiple candidate blocks that conform to the hardware constraints are constructed based on the atomic instruction specifications determined by the hardware profile. Upon receiving a request instruction, different usage schemes for the blocks are determined using the instruction specifications of the request instruction and the candidate blocks. By evaluating the usage schemes of the blocks, the optimal target scheme can be found for resource coordination and matrix multiplication calculation.

[0068] Because it targets specific hardware, a corresponding hardware profile can be obtained, allowing candidate blocks to be constructed based on this profile. Upon receiving a request instruction, multiple usage schemes are built, and the optimal target scheme is selected based on quantitative evaluation to coordinate and configure the resources of the target hardware. This ensures that the final executed target scheme better meets the current hardware and computing requirements, effectively eliminating filler waste and improving resource utilization. It breaks through the hardware limitations faced by fixed block schemes, making it applicable to any type of hardware and enhancing the technical effect of cross-hardware platform versatility.

[0069] To facilitate those skilled in the art to better understand and implement the data processing method provided in the embodiments of the present invention, the data processing method will be described in detail below with reference to specific application scenarios.

[0070] Please refer to Figure 2This method includes: hardware parameter acquisition: obtaining the shape of MMA (Matrix Multiply-Accumulate) instructions, such as the number of registers and the size of shared memory; dynamic block generation, including: candidate block generation and cost model evaluation; dynamic thread remapping, including: thread grouping, nonlinear mapping design and cross-architecture adaptation; instruction decomposition and padding: decomposing non-standard blocks and using masks to shield invalid calculations; kernel code generation and execution: automatically generating optimized GPU kernel code, dynamically loading and executing it.

[0071] Specifically, the first step is hardware parameter acquisition: Key parameters of the target hardware are obtained through the GPU driver interface, including: supported MMA instruction shapes (e.g., 16x16x16, 32x32x8), number of registers, shared memory size and number of banks, number of SM cores, and maximum thread block size.

[0072] Step 2, Dynamic Block Generation: 1. Candidate Block Generation: Based on hardware MMA instructions, enumerate all usable sub-block combinations (i.e., block usage schemes) (e.g., 16x8x16, 8x16x16, deriving 16x16x16, 16x32x16, etc.); 2. Cost Model Evaluation: The cost system calculates a comprehensive score (cost) for each usage scheme and prioritizes the scheme with the highest score as the target scheme for implementation. The weight parameters (e.g., computation priority or storage priority) can be automatically fine-tuned according to the target hardware and can be dynamically generated based on real-time hardware parameters and matrix features. For example, a model can be trained to output corresponding weight parameters based on real-time hardware parameters and matrix features. Alternatively, historical data can be used to find the correlation between real-time hardware parameters, matrix features, and the optimal configuration of weight parameters, thereby dynamically determining the weight parameters.

[0073] Step 3, Dynamic Thread Remapping: 1. Thread Grouping: Divide the threads into groups based on the block size (e.g., M_Tile × N_Tile) and the hardware warp size; 2. Non-linear Mapping: Design an index function with Bank offset; 3. Cross-Architecture Adaptation: Adjust the cooperation strategy for 32-thread or 64-thread systems.

[0074] Step 4, instruction decomposition and padding: A pair of non-standard blocks (such as 16x8x16) are decomposed into two 16x8x8 sub-blocks, and the 16x16x16 instructions are reused; zeros are padded to the insufficient parts, but invalid calculations are masked by a mask.

[0075] Step 5: Kernel code generation and execution: Automatically generate optimized GPU kernel code and dynamically load and execute it.

[0076] Please refer to Figure 3On GPU hardware that supports accelerated matrix computation, FP16 precision generalized matrix multiplication (GEMM) is computed with matrix sizes M=1024, N=768, and K=512 (all non-integer multiples of the standard block size). Fixed block strategies lead to significant fill waste and decreased computational efficiency. This invention addresses this problem through a hardware-aware dynamic optimization method.

[0077] In the hardware characteristic perception stage: the system first starts the hardware detection engine and establishes a complete device profile through multi-level hardware characteristic scanning.

[0078] For the computation instruction analysis layer: accurately identify the matrix computation instruction set specifications supported by the GPU, including core parameters such as instruction shape and data type support. This is the theoretical basis for subsequent block design.

[0079] For the storage architecture resolution layer: detailed mapping of the capacity, bandwidth, and access characteristics of storage structures such as register files, shared memory, and caches, especially the bank distribution pattern of shared memory. This provides a basis for storage access optimization.

[0080] For the parallel architecture modeling layer: analyze the organization of computing units, including key parameters such as SM array size and thread scheduling granularity. This determines the design space for parallel strategies.

[0081] In the intelligent block generation stage: Based on the hardware profile, the system enters a dynamic programming process, including: 1. Feasibility space construction: Based on the instruction set specifications, all candidate blocks that meet hardware constraints are generated. For example, firstly, based on the hardware-supported atomic instruction specifications (such as 16x16x16 MMA instructions) as the basic computational unit, and then for the computational scale of the target matrix multiplication, all possible combinations of atomic instructions are systematically generated, i.e., candidate blocks. Specifically, the macro-block task is decomposed into multiple executable combination patterns of atomic instructions. For example, the 64x64x64 matrix multiplication is decomposed into a regular combination of 4×4×4 16x16x16 atomic instructions, or into 4×8×4 combination schemes composed of 16x8x16 sub-blocks (each 16x8x16 sub-block is implemented by 2 16x8x8 atomic instructions), ensuring that each candidate block strictly satisfies the integer multiple combination relationship of atomic instructions, thereby guaranteeing the hardware-level executability of all subsequently generated block usage schemes. 2. Multi-objective evaluation: Quantitative evaluation of each usage scheme, including: Computational density evaluation: accurately calculating the utilization rate of each block on computing units; Storage stress testing: simulating and predicting the probability of shared memory access conflicts; Boundary effect analysis: calculating the padding overhead of unaligned parts. 3. Dynamic decision-making: Employing a weighted scoring model, combined with the characteristics of the current computing task (such as matrix size), to select the optimal overall target scheme. The weighted scoring model evaluates the block usage scheme through three core dimensions: computational efficiency (weight 40%), storage access (weight 35%), and resource utilization (weight 25%). In practice, a system-wide performance analysis tool is used for each block scheme to test and obtain quantitative indicators such as instruction coverage, bank conflict rate, and register utilization. The total score is calculated by weighted summation, and the scheme with the highest score is selected. For example, if a 16x16x16 block performs well in computational efficiency (0.9 points) and storage access (0.7 points), its total score = 0.9×0.45+0.7×0.4+0.6×0.15=0.775, which is better than other usage schemes.

[0082] During the resource coordination and optimization phase, after selecting the target solution, the system initiates comprehensive resource coordination, including: 1. Thread topology design: Based on the block size and hardware parallelism capabilities, accurately calculate the thread block dimension and thread bundle allocation scheme to ensure full utilization of computing resources. A warp-level load balancing algorithm is adopted: the number of threads in each thread block is set to an integer multiple of 32 (e.g., 128 threads = 4 warps). A round-robin allocation strategy ensures that each warp processes the same number of matrix blocks (e.g., each warp is responsible for 2 8×8 sub-blocks); the data areas processed by adjacent warps remain continuous, improving cache hit rate. 2. Storage access optimization: Based on bank distribution characteristics, a conflict-preventing access mode is designed. A dynamic offset strategy is adopted to evenly distribute concurrent access across different banks. A bank conflict elimination strategy is implemented: a composite addressing mode of "row offset + column rotation" is used for shared memory access: row direction: add a base offset of threadIdx.x % 32; column direction: use a rotation step of (threadIdx.y * 7) % 32 (7 and 32 are coprime). 3. Intelligent Register Allocation: Optimizes register usage through liveness analysis to prevent register overflow while maintaining performance. When register demand is detected to be approaching its limit, the loop unrolling factor is adjusted to reduce register usage.

[0083] During the instruction pipeline optimization phase, instruction reorganization is performed: non-standard blocks are broken down into optimized combinations that match the blocks selected in the target scheme, for example, converting a 16x8x16 block into two 16x8x8 instruction sequences; boundary handling is performed: masking is used to intelligently skip invalid calculations in padding regions while maintaining the continuity of the instruction pipeline; double buffering optimization is performed: overlapping computation and data transfer hides memory access latency. Specifically, memory access (especially from global memory, host memory, or the previous level cache) is often slower than computation. If this latency is not hidden, computation units will idle, severely degrading performance. Double buffering achieves overlapping computation and data transfer by prefetching and alternating between two buffers, allowing the current buffer to be used for computation while the next buffer asynchronously transfers data in the background, thus hiding the latency.

[0084] During the dynamic feedback optimization phase, the system is continuously monitored during runtime to collect performance metrics: real-time acquisition of key indicators such as instruction throughput and storage bandwidth utilization. Parameter fine-tuning is performed: execution parameters are dynamically adjusted based on runtime data, such as appropriately adjusting thread block sizes to adapt to the actual load. Anomaly handling is implemented: when a surge in bank conflicts is detected, a re-optimization process is automatically triggered.

[0085] Once the input size (e.g., 1024x1024) is determined, it is usually the same size provided subsequently. For example, in the common case of license plate recognition, Company A's license plate photos are all taken by fixed cameras on gate machines. The photo size only needs to be input once to determine the size. After dynamic optimization, it will always follow the selected block combination, resulting in the best performance and fastest recognition, and of course, no further optimization process is needed. However, if Company A temporarily uses a different camera to take photos, but the computing hardware remains unchanged, the execution parameters can be dynamically adjusted based on the data, without needing to resample the hardware profile or select and configure the block scheme.

[0086] Therefore, the data processing method provided by this invention can significantly improve computational efficiency and reduce resource waste. Through a hardware-aware dynamic partitioning strategy, this invention can automatically adapt to the characteristics of computing units on different GPU architectures, maximizing the utilization of matrix calculation instructions and significantly reducing filler waste caused by matrix size mismatch. It enables cross-platform general optimization, reducing development and maintenance costs. Related solutions require developing optimization code separately for different GPU architectures, while this invention, through dynamic hardware parameter acquisition and adaptive partitioning strategies, requires only one set of code to automatically adapt to multiple hardware platforms, greatly reducing manual tuning workload and improving the flexibility and portability of algorithm deployment. It can intelligently avoid storage access bottlenecks and improve parallel efficiency. Combining hardware storage hierarchy characteristics (such as shared memory bank distribution and register capacity), this invention dynamically optimizes thread mapping rules and data access patterns, effectively reducing bank conflicts and register overflow problems, ensuring continuous and efficient operation of computing units. Especially in non-standard matrix operation scenarios, it can still maintain performance close to the theoretical peak.

[0087] Through the above description of the embodiments, those skilled in the art can clearly understand that the methods according to the above embodiments can be implemented by means of software plus necessary general-purpose hardware platforms. Of course, they can also be implemented by hardware, but in many cases the former is a better implementation method.

[0088] Embodiments of the present invention also provide a data processing apparatus, which can be cross-referenced with the data processing method provided in the above embodiments.

[0089] Please refer to Figure 4 The device includes the following modules.

[0090] The hardware profiling module 101 is used to obtain a hardware profile of the target hardware.

[0091] The block pre-construction module 102 is used to determine the atomic instruction specifications supported by the target hardware using the hardware profile, and to construct candidate blocks that conform to hardware constraints based on the atomic instruction specifications.

[0092] The scheme generation module 103 is used to receive request instructions and determine different usage schemes for the blocks by using the instruction specifications of the request instructions and combining them with the candidate blocks.

[0093] The optimal solution module 104 is used to quantitatively evaluate the solutions used and determine the target solution from the solutions used based on the quantitative evaluation results.

[0094] The data calculation module 105 is used to coordinate and configure the resources of the target hardware according to the target scheme, and then perform matrix multiplication calculations using the blocks selected by the target scheme.

[0095] The apparatus provided in this invention improves GEMM optimization techniques and achieves cross-hardware platform versatility by employing a dynamic approach to the target hardware. Specifically, to ensure that the hardware block usage scheme conforms to the hardware constraints of the current hardware, a hardware profile of the target hardware is first obtained, and multiple candidate blocks conforming to the hardware constraints are constructed based on the atomic instruction specifications determined by the hardware profile. Upon receiving a request instruction, different usage schemes for the blocks are determined using the instruction specifications of the request instruction and the candidate blocks. By evaluating the block usage schemes, the optimal target scheme can be identified for resource coordination and matrix multiplication calculations.

[0096] Because it targets specific hardware, a corresponding hardware profile can be obtained, allowing candidate blocks to be constructed based on this profile. Upon receiving a request instruction, multiple usage schemes are built, and the optimal target scheme is selected based on quantitative evaluation to coordinate and configure the resources of the target hardware. This ensures that the final executed target scheme better meets the current hardware and computing requirements, effectively eliminating filler waste and improving resource utilization. It breaks through the hardware limitations faced by fixed block schemes, making it applicable to any type of hardware and enhancing the technical effect of cross-hardware platform versatility.

[0097] In one specific embodiment of the present invention, the quantitative evaluation of the usage scheme includes: determining the instruction coverage, storage access conflict rate, and resource utilization of the usage scheme; performing a weighted summation of the instruction coverage, storage access conflict rate, and resource utilization to obtain a quantitative value for each usage scheme; and determining the quantitative value as the quantitative evaluation result so that the usage scheme with the highest quantitative value is determined as the target scheme.

[0098] In one specific embodiment of the present invention, different usage schemes of blocks are determined by using the instruction specification of the request instruction and combining it with the candidate blocks, including: obtaining the instruction specification of the request instruction and obtaining the block specification of the candidate blocks; using the instruction specification and the block specification, determining the blocks and reuse parameters corresponding to different usage schemes.

[0099] In one specific embodiment of the present invention, matrix multiplication calculation is performed using the block partitioning selected by the target scheme, including: adding a reference offset to the row direction and using a rotation step size to the column direction during shared memory access, so that concurrent access request instructions are evenly distributed on different memory banks.

[0100] In one specific embodiment of the present invention, matrix multiplication calculation is performed using blocks selected by the target scheme, including: during the matrix multiplication calculation using the selected blocks, decomposing non-standard blocks into optimized combinations that match the blocks selected by the target scheme; and / or, during the matrix multiplication calculation using the selected blocks, if the data size is smaller than the selected blocks, zero-padding is performed on the regions that are not filled with data, and the invalid calculation of the zero-padding regions is skipped by masking the zero-padding regions.

[0101] In one specific embodiment of the present invention, constructing multiple candidate blocks that conform to hardware constraints based on atomic instruction specifications includes: determining the atomic instruction specifications as basic computing units; and constructing candidate blocks of atomic instruction specifications whose size is an integer multiple of the basic computing units.

[0102] In one specific embodiment of the present invention, the method further includes: monitoring the registers of the target hardware; and adjusting the loop unrolling factor to reduce register occupancy when the register demand is detected to reach a threshold.

[0103] In one specific embodiment of the present invention, obtaining a hardware profile of the target hardware includes: using a hardware detection engine to detect the computation instruction analysis layer of the target hardware to obtain the matrix computation instruction set specifications supported by the target hardware; using a hardware detection engine to detect the storage system parsing layer of the target hardware to obtain storage structure information; using a hardware detection engine to detect the parallel architecture modeling layer of the target hardware to obtain computation unit organization structure information; and using the matrix computation instruction set specifications, storage structure information, and computation unit organization structure information to construct a hardware profile.

[0104] In one specific embodiment of the present invention, the resources of the target hardware are coordinated and configured according to the target scheme, including: dividing the matrix into blocks based on the block size in the target scheme; using the computing unit organization structure information, setting the number of threads in the thread block to an integer multiple of the number of storage blocks, and cyclically allocating so that the thread bundles process the same number of matrix blocks, while keeping the data areas processed by adjacent thread bundles continuous.

[0105] For a description of the features in the embodiment corresponding to the data processing device, please refer to the relevant description in the embodiment corresponding to the data processing method, which will not be repeated here.

[0106] Corresponding to the above method embodiments, this invention also provides an electronic device. The electronic device described below and the data processing method described above can be referred to each other.

[0107] See Figure 5 As shown, the electronic device includes: a memory 332 for storing a computer program; and a processor 322 for executing the computer program to implement the steps of the data processing method described in the above method embodiment.

[0108] For details, please refer to Figure 6 , Figure 6 This is a schematic diagram of the specific structure of an electronic device provided in this embodiment. The electronic device can vary significantly due to differences in configuration or performance. It may include one or more central processing units (CPUs) (e.g., one or more processors) and a memory 332. The memory 332 stores one or more computer programs 342 or data 344. The memory 332 can be temporary or permanent storage. The program stored in the memory 332 may include one or more modules (not shown in the diagram), each module may include a series of instruction operations on the data processing device. Furthermore, the processor 322 may be configured to communicate with the memory 332 and execute the series of instruction operations stored in the memory 332 on the electronic device 301.

[0109] Electronic device 301 may also include one or more power supplies 326, one or more wired or wireless network interfaces 350, one or more input / output interfaces 358, and / or one or more operating systems 341.

[0110] The steps in the data processing method described above can be implemented by the structure of an electronic device.

[0111] In one specific embodiment of the present invention, a data processing method includes: acquiring a hardware profile of a target hardware; using the hardware profile to determine the atomic instruction specifications supported by the target hardware, and constructing multiple candidate blocks that conform to hardware constraints based on the atomic instruction specifications; receiving a request instruction, and using the instruction specifications of the request instruction in combination with the candidate blocks to determine different usage schemes for the blocks; performing a quantitative evaluation of the usage schemes, and using the quantitative evaluation results to determine a target scheme from the usage schemes; coordinating and configuring the resources of the target hardware according to the target scheme, and then performing matrix multiplication calculations using the blocks selected by the target scheme.

[0112] In one specific embodiment of the present invention, the quantitative evaluation of the usage scheme includes: determining the instruction coverage, storage access conflict rate, and resource utilization of the usage scheme; performing a weighted summation of the instruction coverage, storage access conflict rate, and resource utilization to obtain a quantitative value for each usage scheme; and determining the quantitative value as the quantitative evaluation result so that the usage scheme with the highest quantitative value is determined as the target scheme.

[0113] In one specific embodiment of the present invention, different usage schemes of blocks are determined by using the instruction specification of the request instruction and combining it with the candidate blocks, including: obtaining the instruction specification of the request instruction and obtaining the block specification of the candidate blocks; using the instruction specification and the block specification, determining the blocks and reuse parameters corresponding to different usage schemes.

[0114] In one specific embodiment of the present invention, matrix multiplication calculation is performed using the block partitioning selected by the target scheme, including: adding a reference offset to the row direction and using a rotation step size to the column direction during shared memory access, so that concurrent access request instructions are evenly distributed on different memory banks.

[0115] In one specific embodiment of the present invention, matrix multiplication calculation is performed using blocks selected by the target scheme, including: during the matrix multiplication calculation using the selected blocks, decomposing non-standard blocks into optimized combinations that match the blocks selected by the target scheme; and / or, during the matrix multiplication calculation using the selected blocks, if the data size is smaller than the selected blocks, zero-padding is performed on the regions that are not filled with data, and the invalid calculation of the zero-padding regions is skipped by masking the zero-padding regions.

[0116] In one specific embodiment of the present invention, constructing multiple candidate blocks that conform to hardware constraints based on atomic instruction specifications includes: determining the atomic instruction specifications as basic computing units; and constructing candidate blocks of atomic instruction specifications whose size is an integer multiple of the basic computing units.

[0117] In one specific embodiment of the present invention, the method further includes: monitoring the registers of the target hardware; and adjusting the loop unrolling factor to reduce register occupancy when the register demand is detected to reach a threshold.

[0118] In one specific embodiment of the present invention, obtaining a hardware profile of the target hardware includes: using a hardware detection engine to detect the computation instruction analysis layer of the target hardware to obtain the matrix computation instruction set specifications supported by the target hardware; using a hardware detection engine to detect the storage system parsing layer of the target hardware to obtain storage structure information; using a hardware detection engine to detect the parallel architecture modeling layer of the target hardware to obtain computation unit organization structure information; and using the matrix computation instruction set specifications, storage structure information, and computation unit organization structure information to construct a hardware profile.

[0119] In one specific embodiment of the present invention, the resources of the target hardware are coordinated and configured according to the target scheme, including: dividing the matrix into blocks based on the block size in the target scheme; using the computing unit organization structure information, setting the number of threads in the thread block to an integer multiple of the number of storage blocks, and cyclically allocating so that the thread bundles process the same number of matrix blocks, while keeping the data areas processed by adjacent thread bundles continuous.

[0120] Corresponding to the above method embodiments, this invention also provides a readable storage medium. The readable storage medium described below corresponds to the data processing method described above. This invention also provides a computer-readable storage medium storing a computer program configured to execute the steps of any of the above data processing method embodiments at runtime.

[0121] In one exemplary embodiment, the aforementioned computer-readable storage medium may include, but is not limited to, various media capable of storing computer programs, such as a USB flash drive, read-only memory (ROM), random access memory (RAM), portable hard disk, magnetic disk, or optical disk.

[0122] Embodiments of the present invention also provide a computer program product, which includes a computer program that, when executed by a processor, implements the steps in any of the above data processing method embodiments.

[0123] Embodiments of the present invention also provide another computer program product, including a non-volatile computer-readable storage medium storing a computer program, which, when executed by a processor, implements the steps in any of the above-described data processing method embodiments.

[0124] Those skilled in the art will further recognize that the units and algorithm steps of the various examples described in conjunction with the embodiments disclosed herein can be implemented in electronic hardware, computer software, or a combination of both. To clearly illustrate the interchangeability of hardware and software, the components and steps of the various examples have been generally described in terms of functionality in the foregoing description. Whether these functions are implemented in hardware or software depends on the specific application and design constraints of the technical solution. Those skilled in the art can use different methods to implement the described functions for each specific application, but such implementations should not be considered beyond the scope of the invention.

[0125] This article uses specific examples to illustrate the principles and implementation methods of the present invention. The descriptions of the above embodiments are only intended to help understand the method and core ideas of the present invention. It should be noted that those skilled in the art can make various improvements and modifications to the present invention without departing from its principles, and these improvements and modifications also fall within the protection scope of the present invention.

Claims

1. A data processing method, characterized in that, include: Obtain a hardware profile of the target hardware; The hardware profile is used to determine the atomic instruction specifications supported by the target hardware, and multiple candidate blocks that conform to hardware constraints are constructed based on the atomic instruction specifications. Receive a request instruction, and use the instruction specification of the request instruction and the candidate blocks to determine different usage schemes for the blocks; The proposed usage schemes are quantitatively evaluated, and the target scheme is determined from the proposed usage schemes using the quantitative evaluation results. After coordinating and configuring the resources of the target hardware according to the target scheme, matrix multiplication calculation is performed using the blocks selected by the target scheme. The process of obtaining a hardware profile of the target hardware includes: The hardware detection engine is used to detect the computation instruction analysis layer of the target hardware to obtain the matrix computation instruction set specifications supported by the target hardware. The hardware detection engine is used to detect the storage system parsing layer of the target hardware to obtain storage structure information; The hardware detection engine is used to detect the parallel architecture modeling layer of the target hardware to obtain the computing unit organization structure information; The hardware profile is constructed using the matrix calculation instruction set specification, the storage structure information, and the computing unit organization structure information. The coordinated configuration of resources for the target hardware according to the target scheme includes: The blocks are divided based on the block size partitioning matrix in the target scheme; Using the organizational structure information of the computing unit, the number of threads in the thread block is set to an integer multiple of the number of storage units, and the thread bundles are allocated in a cyclic manner so that they process the same number of matrix blocks, while keeping the data regions processed by adjacent thread bundles continuous.

2. The method according to claim 1, characterized in that, A quantitative evaluation of the aforementioned usage scheme includes: Determine the instruction coverage, storage access conflict rate, and resource utilization of the usage scheme; The instruction coverage, the storage access conflict rate, and the resource utilization are weighted and summed to obtain the quantitative value of each usage scheme; The quantified value is determined as the quantified evaluation result, so that the usage scheme with the highest quantified value is determined as the target scheme.

3. The method according to claim 1, characterized in that, Using the instruction specification of the request instruction and combining it with the candidate blocks, different usage schemes for the blocks are determined, including: Obtain the instruction specification of the request instruction, and obtain the block specification of the candidate block; Using the instruction specification and the block specification, determine the block and multiplexing parameters to be selected for different usage schemes.

4. The method according to claim 1, characterized in that, Matrix multiplication calculations are performed using the blocks selected in the target scheme, including: During shared memory access, a base offset is added in the row direction and a rotation step is used in the column direction to ensure that concurrent access requests are evenly distributed across different memory banks.

5. The method according to claim 1, characterized in that, Matrix multiplication calculations are performed using the blocks selected in the target scheme, including: In the process of performing the matrix multiplication calculation using the selected blocks, the non-standard blocks are decomposed into an optimized combination that matches the blocks selected by the target scheme; And / or, during the matrix multiplication calculation using the selected block, if the data size is smaller than the selected block, the unfilled areas are padded with zeros, and the invalid calculations of the zero-filled areas are skipped by masking the zero-filled areas.

6. The method according to claim 1, characterized in that, Based on the atomic instruction specification, multiple candidate blocks conforming to hardware constraints are constructed, including: The atomic instruction specification is defined as the basic computing unit; Based on the basic computing unit, candidate blocks of the atomic instruction specification with a size that is an integer multiple of the specified size are constructed.

7. The method according to claim 1, characterized in that, Also includes: Monitor the registers of the target hardware; If register demand is detected to reach a threshold, the loop unrolling factor is adjusted to reduce register usage.

8. An electronic device, characterized in that, include: Memory, used to store computer programs; A processor for executing the computer program to implement the steps of the data processing method as described in any one of claims 1 to 7.

Images