A template computing code generation method and system for tensor accelerator

Abstract

The present invention relates to a method and system for generating code for stencil calculations for a tensor accelerator. The method comprises the following steps: S1: converting a stencil calculation into a domain-specific programming language designed by the present invention; S2: splitting the weight matrix of the stencil calculation into rows and columns; S3: stacking and rearranging the split weight rows and columns to form a weight matrix suitable for matrix multiplication; S4: starting an automatic tuning code generator, forming an optimized configuration for the matrix multiplication loop through evolutionary search, and screening candidate configurations using a cost model; S5: applying the screened candidate configurations to actual machine performance measurements, and using the measured results to update the cost model; S6: checking whether the number of evolutionary search iterations reaches an iteration threshold; if so, terminating the program and selecting the optimal configuration; otherwise, proceeding to S4. The present invention constructs a multi-stage soft pipeline for memory access and computation, which better conceals memory access overhead.

Description

Technical Field

[0001] The present invention belongs to the field of high-performance stencil computing, and in particular relates to a stencil computing code generation method and system for a tensor accelerator. Background Art

[0002] In the natural sciences, partial differential equations (PDEs) describe the laws governing natural phenomena by expressing relationships between variables and their partial derivatives, becoming an indispensable part of modern scientific applications. For example, PDEs are widely used in fields such as climate prediction, earthquake simulation, and fluid dynamics. Numerical methods are often used to solve the PDEs in these applications, with the finite difference method (FDM) being a common choice. Template calculation is one of the key computational modes within the FDM, based on uniform grids.

[0003] Template calculation considers the values ​​of its own grid point and its neighboring grid points in the spatial dimension, and determines the updated value of the grid point through a certain calculation model. In partial differential equations, the template calculations corresponding to different partial derivative operations may be different. These template calculations can be classified according to multiple characteristics, including dimension, shape, order, type of coefficients in the calculation, boundary processing method, etc. Figure 1 The figure shows the order R of the template calculation, which is equal to the distance from the grid point to be updated to the farthest neighbor grid point. From a mathematical point of view, the order of the template calculation is equivalent to the order of the Taylor expansion in the finite difference method.

[0004] Due to the high complexity of scientific simulations, scientific applications involve a variety of template calculations of varying orders, dimensions, and shapes. In complex scientific simulations, a single iteration may require mesh updates based on multiple template types. For example, COSMO, a widely used scientific software in meteorological simulation, has a complex execution process involving 30 template calculations of varying dimensions and orders. These template calculations cover dimensions from two to three dimensions, and orders from 2 to 16. Due to the diversity of different template calculation modes and memory access patterns, optimizing computational performance requires not only general template calculation optimization across a variety of computing hardware platforms, but also specialized optimization and performance tuning for each type of template calculation.

[0005] Most existing template computation code generation methods are based on the GPU's general-purpose computing units. For example, Artemis abstracts loop optimization into scheduling primitives and provides an automatic tuning algorithm to automatically generate high-performance template computation code. However, Artemis still performs computational optimization based on the GPU's CUDA Core, making it difficult to utilize the computing power of the tensor computing units and missing the opportunity for performance optimization. TCStencil, on the other hand, utilizes computational mapping to convert template computations into a series of reduction and summation operations, thereby achieving the mapping of template computations to tensor computing units. However, current work still relies on manual implementation of CUDA kernels, and computational mapping of three-dimensional and high-order two-dimensional templates is difficult to achieve. Existing research has established domain-specific programming language compilers and proposed various template computation parallel optimization methods. Although these methods fully exploit the parallel performance of CUDA Cores, they are difficult to fully utilize with the latest tensor accelerators. Furthermore, existing work has achieved performance acceleration by converting template computations into GEMM computations, thereby scheduling tensor computing units. However, this work is still limited to manual optimization and difficult to achieve automated code generation. Summary of the Invention

[0006] In order to solve the above technical problems, the present invention provides a template calculation code generation method and system for tensor accelerators, designs a code generation method for GPU tensor accelerators that can automatically generate high-performance template calculation code, and introduces shared memory asynchronous reading to make the memory access part of template calculation more efficient.

[0007] To achieve the above object, the technical solution adopted by the present invention is:

[0008] In one aspect, the present invention provides a method for generating template computing code for a tensor accelerator, comprising:

[0009] Step S1: converting the template calculation into a domain-specific programming language, wherein the domain-specific programming language includes: a basic formula of the template calculation and a loop optimization configuration specified by a user; setting the total number of iterations of the template calculation code generation program to m times;

[0010] Step S2: According to the template calculation mode expressed in the domain-specific programming language, the weight matrix of the template calculation is split into rows and columns, and non-zero elements in the template calculation weight matrix are extracted. If the non-zero element is located in a row, row splitting is adopted; if it is located in a column, column splitting is adopted;

[0011] Step S3: stacking and rearranging the split weight rows or columns to form a weight matrix suitable for matrix multiplication, selecting a matrix multiplication shape suitable for the tensor calculation unit, determining whether the weight rows or columns can be stacked and rearranged within the range of the matrix multiplication shape, and performing a rearrangement transformation on them;

[0012] Step S4: Start the automatic tuning code generator, set the number of iterations i=0, start iterative search for the optimal configuration of the template calculation code loop optimization, perform loop optimization on the rearranged and converted template calculation, generate a matrix multiplication loop optimization configuration through evolutionary search, build a candidate configuration queue, and use the cost model to screen the candidate configurations;

[0013] Step S5: The selected candidate configurations are used for actual device testing. After obtaining the measured performance results, the optimal configuration is updated. The measured results are used to update the cost model, and the number of iterations i is increased by 1.

[0014] Step S6: Check whether the number of evolutionary search iterations reaches the total number of iterations. If i=m at this time, the program ends and returns the current optimal configuration to obtain the optimal code generation sample. Otherwise, go to step S4.

[0015] On the other hand, the present invention provides a template computing code generation system for a tensor accelerator, comprising:

[0016] A template calculation intermediate representation module is used to convert the template calculation into a domain-specific programming language, wherein the domain-specific programming language includes: a basic formula of the template calculation and a loop optimization configuration specified by the user; and sets the total number of iterations of the template calculation code generation program to m times;

[0017] The template calculation splitting module is used to split the weight matrix of the template calculation into rows and columns according to the template calculation mode expressed in the domain-specific programming language, extract the non-zero elements in the template calculation weight matrix, and use row splitting when the non-zero elements are in a row, and use column splitting when they are in a column; stack and rearrange the split weight rows or columns to form a weight matrix suitable for matrix multiplication, select a matrix multiplication shape suitable for the tensor calculation unit, determine whether the weight rows or columns can be stacked and rearranged within the range of the matrix multiplication shape, and perform rearrangement conversion on them;

[0018] a loop optimization module for starting an automatic tuning code generator, setting the number of iterations i=0, and starting an iterative search for an optimal configuration for loop optimization of the template calculation code, loop optimizing the rearranged and transformed template calculation, generating a matrix multiplication loop optimization configuration through evolutionary search, building a candidate configuration queue, and screening the candidate configurations using a cost model;

[0019] The evolutionary search module is used to apply the selected candidate configurations to real-machine testing. After obtaining the measured performance results, the optimal configuration is updated and the measured results are used to update the cost model. The number of iterations i is incremented by 1.

[0020] The code generation module is used to check whether the number of evolutionary search iterations reaches the total number of iterations. If i=m at this time, the program ends and returns the current optimal configuration to obtain the optimal code generation sample. Otherwise, it goes to the loop optimization module.

[0021] Compared with the prior art, the present invention has the following advantages:

[0022] The present invention introduces cache read and cache write-back primitives to realize asynchronous reading of shared memory, constructs a multi-level soft pipeline for memory access and calculation, so that the memory access overhead can be better concealed. BRIEF DESCRIPTION OF THE DRAWINGS

[0023] Figure 1 This is a schematic diagram of a two-dimensional Star-type template calculation in an embodiment of the present invention;

[0024] Figure 2 This is a flow chart of a method for generating template calculation code for a tensor accelerator according to the present invention;

[0025] Figure 3 A schematic diagram of a three-dimensional Star-type template calculation and segmentation method according to an embodiment of the present invention;

[0026] Figure 4 Schematic diagram of stacking and rearranging weighted rows in an embodiment of the present invention;

[0027] Figure 5 Schematic diagram of stacking and rearranging weight columns in an embodiment of the present invention;

[0028] Figure 6 Schematic diagram of the automatic code generation principle of evolutionary search combined with real machine testing in an embodiment of the present invention;

[0029] Figure 7 This is a structural block diagram of a template computing code generation system for a tensor accelerator according to the present invention. DETAILED DESCRIPTION

[0030] In order to make the purpose, technical solutions and advantages of the present invention more clear, the present invention is further described in detail below through specific implementation and in conjunction with the accompanying drawings.

[0031] like Figure 2 As shown, an embodiment of the present invention provides a method for generating template calculation code for a tensor accelerator, comprising the following steps:

[0032] Step S1: converting the template calculation into a domain-specific programming language designed by the present invention, wherein the domain-specific programming language includes two parts: a basic formula for the template calculation and a loop optimization configuration specified by the user; setting the total number of iterations of the template calculation code generation program to m times;

[0033] Step S2: According to the template calculation mode expressed in the domain-specific programming language, the weight matrix of the template calculation is split into rows and columns, and non-zero elements in the template calculation weight matrix are extracted. If the non-zero element is located in a row, row splitting is adopted; if it is located in a column, column splitting is adopted;

[0034] Step S3: stacking and rearranging the split weight rows or columns to form a weight matrix suitable for matrix multiplication, selecting a matrix multiplication shape suitable for the GPU tensor computing unit, determining whether the weight rows / columns can be stacked and rearranged within the range of the matrix multiplication shape, and performing a rearrangement conversion on them;

[0035] Step S4: Start the automatic tuning code generator, set the number of iterations i = 0, and begin iteratively searching for the optimal configuration for loop optimization of the template calculation code. The rearranged and transformed template calculation is loop optimized, and a matrix multiplication loop optimization configuration is generated through evolutionary search. A candidate configuration queue is constructed, and the cost model is used to screen the candidate configurations.

[0036] Step S5: The selected candidate configurations are used for actual device testing. After obtaining the measured performance results, the optimal configuration is updated. The measured results are used to update the cost model, and the number of iterations i is increased by 1.

[0037] Step S6: Check whether the number of evolutionary search iterations reaches the total number of iterations. If i=m at this time, the program ends and returns the current optimal configuration to obtain the optimal code generation sample. Otherwise, go to step S4.

[0038] In one embodiment, the above step S1: converting the template calculation into a domain-specific programming language designed by the present invention, wherein the domain-specific programming language includes two parts: a basic formula for the template calculation and a user-specified loop optimization configuration; setting the total number of program iterations to m times, specifically includes:

[0039] First, the template calculation is expressed in the form of approximate mathematical calculation. The template calculation area is a two-dimensional grid or a three-dimensional grid. The size of the two-dimensional grid is dimX × dimY, and the size of the three-dimensional grid is dimX × dimY × dimZ. The weight array contains the offset and weight value corresponding to the template calculation. Taking a two-dimensional five-point template as an example, the hotspot diffusion speed can be expressed as The two-dimensional five-point template contains five weight values, where i and j represent the output grid point position of the current calculation, and offset x 、offset y Indicates the row and column offset of the input grid corresponding to the currently traversed weight k. In the embodiment of the present invention, the template calculation is expressed in an approximate mathematical form to further simplify the user programming process;

[0040] Secondly, the data required for the program initialization is input by the user, including input and output matrices and weight arrays. In addition, a loop optimization interface is provided for the user, allowing the user to specify the block size of each loop layer, select the loop body for loop fusion, and specify the vectorization length of the logical statements within the loop. Specifically, according to the GPU hardware storage hierarchy, the loop blocks correspond to global video memory, shared memory, and registers respectively. The user can guide the selection of loop block size based on the storage capacity of the current hardware. For vectorization length, it is necessary to select an appropriate vectorization length based on the hardware vector instructions. User guidance can make automatic code generation more efficient. In addition, the present invention also supports high-performance code generation without user guidance through random generation and evolutionary search.

[0041] At the same time, the total number of iterative searches of the program needs to be initialized and set to m times.

[0042] In one embodiment, step S2: the weight matrix calculated by the template is split into rows and columns, and non-zero elements in the weight matrix calculated by the template are extracted. When the non-zero element is in a row, row splitting is adopted; when it is in a column, column splitting is adopted, such as Figure 3 As shown, specifically including:

[0043] Step S21: defining the weight matrix for the template calculation. First, based on the radius r of the template, a weight matrix of size (2r+1)×(2r+1) is generated. The initial element values ​​of the weight matrix are all 0. Then, based on the weight array input by the user, the weight elements are assigned to the weight matrix.

[0044] Step S22: extracting weight rows / columns in the weight matrix using any of the following methods:

[0045] For two-dimensional template calculation, the non-zero elements in the weight matrix must be arranged in a row or a column, so there are two ways to extract the weight rows / columns in the weight matrix.

[0046] Method 1: extract the non-zero elements in the same row of the weight matrix, record the offset of their starting address relative to the starting address of the weight matrix, and use it as the weight row. The extracted weight elements in the weight matrix are set to 0.

[0047] Method 2: Extract the non-zero elements in the same column of the weight matrix, record the offset of their starting address relative to the starting address of the weight matrix, and use it as the weight column. Set the extracted weight elements in the weight matrix to 0;

[0048] In one embodiment, the above step S3: stacking and rearranging the split weight rows or columns to form a weight matrix suitable for matrix multiplication, selecting a matrix multiplication shape suitable for the GPU tensor computing unit, determining whether the weight rows or columns can be stacked and rearranged within the shape range, and performing a rearrangement conversion on them, specifically includes:

[0049] Step S31: Since the tensor computing units in different GPU architectures have different matrix multiplication shapes, for example, in the Volta architecture, the minimum-scale matrix multiplication shape is m8n8k4, and in the Amper architecture, the minimum-scale matrix multiplication shape is m16n8k8. The present invention constructs a matrix multiplication shape suitable for different inputs based on the minimum-scale matrix multiplication shape adapted to different hardware architectures. For example, for the Volta architecture, a matrix multiplication of m8n8k12 can be constructed, and for the Amper architecture, a matrix multiplication shape of m16n16k24 can be constructed. In an embodiment of the present invention, for different template calculation instances, all suitable matrix multiplication shape pools are constructed according to the tensor computing units in the hardware architecture used, and the shape with the highest non-zero element utilization is selected from the matrix multiplication shape pool. First, the length len of the weight row is defined, and the weight matrix shape of the matrix multiplication is k×n, such as Figure 4 As shown, taking the weight rows [a1, a2, a3, a4, a5] or [b1, b2, b3, b4, b5] as an example, the weight rows are stacked and rearranged to form a weight matrix of matrix multiplication. For a specific weight row length len and a specific matrix multiplication weight matrix shape k×n, the non-zero element coverage of the matrix multiplication weight matrix formed by stacking and rearranging the weight rows can be calculated. , where usage_ratio represents the non-zero element coverage, k represents the k-dimensional size of the matrix multiplication shape, len represents the weight row length, and n represents the n-dimensional size of the matrix multiplication shape. The matrix multiplication shape with the largest non-zero element coverage is selected to achieve the optimal GPU tensor computing unit utilization;

[0050] Step S32: In this embodiment of the present invention, weight rows or columns are extracted from the weight matrix. For the optimal matrix multiplication shape selected in step S31, the matrix multiplication calculation formula is expressed as C=A×B. The weight rows or columns extracted from the weight matrix are stacked and rearranged in the following two ways:

[0051] Method 1: The weight row represents the multiplication reduction sum of a row of the input grid. According to the computational semantics of matrix multiplication, the multiplication reduction sum is calculated by the rows of the input matrix and the columns of the weight matrix. Therefore, the input grid is used as the input matrix of the matrix multiplication, and the weight rows are stacked and rearranged to form the weight matrix, such as Figure 5As shown, taking the weight rows [A0, A1, A2, A3, A4] as an example, the weight rows are transposed into column tuples and arranged from left to right in the columns of the matrix B of the matrix multiplication. Each time a column is moved to the right, an element is shifted downward, and the vacant position is filled with 0 until the rightmost column is reached;

[0052] Method 2: The weight column represents the multiplication reduction sum of a column of the input grid. According to the computational semantics of matrix multiplication, the multiplication reduction sum is calculated by the rows of the input matrix and the columns of the weight matrix. Therefore, the input grid is used as the weight matrix of the matrix multiplication, and the weight columns are stacked and rearranged to form the input matrix of the matrix multiplication, such as Figure 5 As shown, taking the weight column [B0, B1, B2, B3] as an example, the weight column is transposed into a row tuple and arranged from top to bottom in the rows of the matrix A of the matrix multiplication. Each time it goes down a row, it shifts one element to the right and fills the vacant position with 0 until it reaches the last row;

[0053] In one embodiment, the above step S4 starts the automatic tuning code generator, performs loop optimization on the rearranged and converted template calculation, forms a matrix multiplication loop optimization configuration through evolutionary search, forms a candidate configuration queue, and uses a cost model to screen the candidate configurations, specifically including:

[0054] Step S41: First, define the loop optimization scheduling primitives used by the automatic tuning code generator. The scheduling primitives include loop fusion, cache read, tensor quantization, cache write back, loop blocking, and MMA instruction optimization:

[0055] __loop_fuse(a, b): fuses the first loop nest a and the second loop nest b into one loop nest;

[0056] __cache_read(a): Changes the data loading in the first loop nest a from global memory to registers to first loading from global memory to shared memory, and then loading from shared memory to registers;

[0057] __tensorize(a): rewrites the matrix multiplication loop nest in the first loop nest a into the constructed mma instruction that adapts to the shape of the GPU tensor computing unit;

[0058] __cache_write(a): Changes the data in the first loop nest a from registers to global memory to first write it back to shared memory, and then write it back to global memory from shared memory;

[0059] __loop_tile_for_block(a, tile_size): tiles the outermost loop of the first loop nest a using tile_size and maps it to GPU thread blocks;

[0060] __loop_tile_for_warp(a, tile_size): tiles the inner loop of the first loop nest a using tile_size and maps it to the GPU thread warp;

[0061] __mma_optimize(a, use_mma_shape): Use use_mma_shape to further divide the abstract MMA instructions in the first loop nest a and map them to the specific MMA instructions of the tensor computation unit;

[0062] Step S42: Start the automatic tuning code generator. For the template calculation loop nest generated in step S3, first apply loop_fuse. Step S3 splits the weight matrix of the template calculation into weight rows or columns. Each weight row or column will be stacked and rearranged to form a separate template calculation loop nest. The loop contains three elements: lower_bound, upper_bound and step. The loop with exactly the same three elements is loop fused through the loop_fuse primitive; initialize the initial population and high-quality population of the evolutionary search. The high-quality population is initially empty. Each sample in the population contains the tile_size and loop_tile_for_war used by loop_tile_for_block. The tile_size used by p and the basic hardware MMA instructions used by mma_optimize are used. 2048 samples are randomly selected from the initial population as the candidate set. The samples in the candidate set are hybridized with the samples in the high-quality population. The specific hybridization form is to exchange the tile_size used by the scheduling primitive or the basic MMA instructions used by mma_optimize to form the final candidate set. For each sample in the candidate set, cache_read, tensorize, cache_write, loop_tile_for_block, loop_tile_for_warp, and mma_optimize primitives are applied in sequence to complete the loop transformation.

[0063] Step S43: Apply the cache_read primitive, scan the template to calculate the memref.LoadOp in the loop nest, determine the size of the data block to be read based on the operator's loadIndices, construct a new memref.LoadOp in the loop, set the source address of the data to be read to global memory, set the target address to shared memory, change the original memref.LoadOp to vector.LoadOp, and read the vector array from the shared memory into the register;

[0064] Step S44: Apply the tensorize primitive to traverse the matrix multiplication loop nesting in the template calculation, determine the mma shape of the tensor calculation unit according to the three-layer loop of the matrix multiplication loop nesting, and define the three-layer loop of the matrix multiplication loop nesting as loop i, loop j, loop k, , }, construct the tensor computing unit abstract mma instruction with the same mma_shape. The abstract mma instruction is a virtual mapping of the tensor computing unit. It is necessary to use the mma_optimize primitive to determine the specific mma_shape to be converted into a specific mma instruction mapping.

[0065] Step S45: Apply the cache_write primitive, scan the template to calculate the memref.StoreOp in the loop nest, determine the size of the data block to be written back based on the storeIndices of the operator, change the memref.StoreOp to vector.StoreOp, write the vector array in the register back to the shared memory, and create a new memref.StoreOp to write the data block in the shared memory back to the global memory;

[0066] Step S46: Apply the loop_tile_for_block primitive. Based on the input loop nest a and tile size tile_size, select the outermost loop of loop nest a, define the loop lower bound as lower_bound, the upper bound as upper_bound, and the iteration step as step, and generate a new inner loop with a lower bound of 0, an upper bound of tile_size, and an iteration step of 1. The original lower and upper bounds of the outermost loop remain unchanged, and the iteration step is changed to step × tile_size. The loop variables of the outermost loop are bound to the GPU blockIdx.X, blockIdx.Y, and blockIdx.Z.

[0067] Step S47: Apply the loop_tile_for_warp primitive. Based on the input loop nest a and tile size tile_size, select the inner loop of loop nest a, define the loop lower bound as lower_bound, the upper bound as upper_bound, and the iteration step as step, and generate a new inner loop with a lower bound of 0, an upper bound of tile_size, and an iteration step of 1. The original inner loop lower and upper bounds remain unchanged, and the iteration step is changed to step × tile_size. The loop variables of the inner loop are bound to GPU threadIdx.Y and threadIdx.Z. ThreadIdx.X is fixed to 32, indicating a single thread warp.

[0068] Step S48: Apply the mma_optimize primitive to map the abstract mma instruction to the mma instruction of the specific hardware. In this embodiment of the present invention, according to the supported GPU hardware architecture, namely the Volta architecture and the Amper architecture, it is divided into mma_optmize_for_volta and mma_optimize_for_amper interfaces. For the mma_optimize_for_volta interface, the unique m8n8k4 hardware instruction under the Volta architecture is used for implementation. Taking the abstract m8n8k12 instruction as an example, the m8n8k4 instruction is looped three times in the K dimension to implement the abstract m8n8k12 instruction. For the mma_optimize_for_amper interface, the unique m16n8k8 and m16n8k16 hardware instructions under the Amper architecture are used for implementation. Taking the abstract m16n16k24 as an example, the m16n8k8 instruction is looped twice in the N dimension and three times in the K dimension to implement the abstract m16n16k24 instruction.

[0069] Step S49: Use xgboost to define a cost model. The input received by the cost model is the code sample generated after the cyclic transformation, and the output of the cost model is the predicted performance of the code sample. All code samples generated by the candidate set are input into the cost model, and the candidate set is ranked by prediction performance. By default, 64 candidate samples with priority performance are selected. In addition, samples from the remaining candidate set are randomly selected with a probability of 0.05 as the candidate set that needs to be tested on the real machine;

[0070] In one embodiment, Figure 6 As shown, after the template calculation is split, candidate samples will be generated through evolutionary search, and code features will be extracted for the candidate samples. These samples will be used to screen high-quality sample codes for the cost model. After the above step S5: the screened candidate configurations will be used for real machine testing. After obtaining the measured performance results, the optimal configuration will be updated, and the measured results will be used to update the cost model. The number of iterations i will be increased by 1.

[0071] Step S51: For the candidate set generated in step S4 that needs to be tested on a real machine, use the LLVM compiler to compile each sample in the candidate set to generate a binary executable file, construct a real input dataset as input for the binary executable file, execute it on real hardware, and use nsight-compute to collect program performance;

[0072] Step S52: Use the program performance of each sample in the candidate set in step S51 as the sample label and the samples in the candidate set as training data to construct a cost model training set, input the training set into the fit interface of xgboost, and update the current cost model;

[0073] Step S53: adding the three samples with the best performance in the candidate set measured on the real machine in step S51 to the high-quality population of evolutionary search, and performing hybridization with the randomly selected candidate set in the initial population during the evolutionary search process, ultimately completing one search iteration, and the number of iterations i is increased by 1;

[0074] In one embodiment, the above step S6: checks whether the number of evolutionary search iterations reaches the iteration number threshold. If i=m at this time, the program ends, returns the current optimal configuration, and saves the current cost model. Otherwise, go to step S4;

[0075] First, check whether the current number of iterations has reached the predetermined total number of iterations m. If it has reached m, the program terminates and the search ends. If the number of iterations has not reached the predetermined total number of iterations m at this time, go to step S4, continue the iterative search after performing evolutionary search, and execute subsequent steps.

[0076] For example, assuming m=100, each time an evolutionary search is completed, a search count check will be started. Since the iterative search process will continuously form high-quality code samples and add them to the high-quality population of evolutionary search, it is necessary to synchronize the optimal code transformation in the high-quality population to the search results. When the iterative search times reaches i=100, the current search result will be returned to obtain the optimal code generation sample.

[0077] like Figure 7 As shown, an embodiment of the present invention provides a template calculation code generation system for a tensor accelerator, including the following modules:

[0078] The template calculation intermediate representation module M1 is used to abstract the calculation area and calculation mode into the MLIR intermediate representation. It also supports the user to indicate the loop optimization configuration and provides manual experience guidance for loop optimization. The total number of iterations of the template calculation code generation program is set to n times.

[0079] The template calculation splitting module M2 is used to split the weight matrix of the template calculation into rows and columns, and extract the non-zero elements in the template calculation weight matrix. When the non-zero element is located in a row, row splitting is adopted; when it is located in a column, column splitting is adopted;

[0080] The loop optimization module M3 is used to abstract the loop optimization pass into scheduling primitives, including loop blocking, cache read, cache write back, and vectorization. The loop blocking size and vectorization length are exposed as tunable interfaces to the evolutionary search module.

[0081] The evolutionary search module M4 is used to initialize the loop optimization configuration, iteratively search for the optimal loop configuration, generate candidate loop configurations during the search process, and use the cost model to filter the candidate set. The final candidate set is used for real-machine testing, and the measured performance data is used to update the cost model and the evolutionary population. The number of iterations i is incremented by 1.

[0082] The code generation module M5 is used to detect that if i=n at this time, the program ends and the optimal loop configuration finally generated is used for code generation. Otherwise, it continues to the evolutionary search module for the next round of iterative search.

[0083] The specific embodiments described above further illustrate the objectives, technical solutions and beneficial effects of the present invention in detail. It should be understood that the above are only specific embodiments of the present invention and are 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 in the scope of protection of the present invention.

Claims

1. A method for generating template computing code for a tensor accelerator, characterized in that: include: Step S1: converting the template calculation into a domain-specific programming language, wherein the domain-specific programming language includes: a basic formula of the template calculation and a loop optimization configuration specified by a user; setting the total number of iterations of the template calculation code generation program to m times; Step S2: According to the template calculation mode expressed in the domain-specific programming language, the weight matrix of the template calculation is split into rows and columns, and non-zero elements in the template calculation weight matrix are extracted. If the non-zero element is located in a row, row splitting is adopted; if it is located in a column, column splitting is adopted; Step S3: stacking and rearranging the split weight rows or columns to form a weight matrix suitable for matrix multiplication, selecting a matrix multiplication shape suitable for the tensor calculation unit, determining whether the weight rows or columns can be stacked and rearranged within the range of the matrix multiplication shape, and performing a rearrangement transformation on them; Step S4: Start the automatic tuning code generator, set the number of iterations i=0, start iterative search for the optimal configuration of the template calculation code loop optimization, perform loop optimization on the rearranged and converted template calculation, generate a matrix multiplication loop optimization configuration through evolutionary search, build a candidate configuration queue, and use the cost model to screen the candidate configurations; Step S5: The selected candidate configurations are used for actual device testing. After obtaining the measured performance results, the optimal configuration is updated. The measured results are used to update the cost model, and the number of iterations i is increased by 1. Step S6: Check whether the number of evolutionary search iterations reaches the total number of iterations. If i=m at this time, the program ends and returns the current optimal configuration to obtain the optimal code generation sample. Otherwise, go to step S4.

2. A method for generating template calculation code for a tensor accelerator according to claim 1, characterized in that: The step S2 comprises: Step S21: defining a weight matrix for template calculation, marking the positions of non-zero elements therein; wherein the radius of template calculation is r, and the size of the weight matrix is ​​(2r+1)×(2r+1); Step S22: extracting weight rows / columns in the weight matrix using any of the following methods: Method 1: extract the non-zero elements in the same row of the weight matrix, and split the row where the current non-zero element is located from the weight matrix as the weight row; Method 2: extract the non-zero elements in the same column of the weight matrix, and split the column where the current non-zero element is located from the weight matrix as the weight column.

3. A method for generating template calculation code for a tensor accelerator according to claim 1, characterized in that: The step S3 comprises: Step S31: constructing a suitable matrix multiplication shape according to the GPU tensor computing unit in the hardware architecture used; Step S32: For the matrix multiplication calculation formula C=A×B, the weight rows or columns extracted from the weight matrix are stacked and rearranged in the following two ways: Method 1: Transpose the weight rows into column tuples and place them from left to right in the columns of the matrix B of the matrix multiplication. Each time you move to the right, shift one element downward and fill the vacant position with 0. Method 2: Transpose the weight column into a row tuple and arrange it from top to bottom in the rows of the matrix A of the matrix multiplication. Each time you go down a row, shift one element to the right and fill the vacant position with 0.

4. A method for generating template calculation code for a tensor accelerator according to claim 1, characterized in that: The step S4 forms a matrix multiplication loop optimization configuration through evolutionary search, specifically including: Step S41: Optimizing the rearranged and converted template calculation loop into loop scheduling primitives, including loop blocking, cache read, cache write-back, and vectorization; the evolutionary search is used to generate loop blocking size and vectorization length; Step S42: Start the automatic tuning code generator to calculate the loop nesting for the rearranged and transformed template generated in step S3; initialize the initial population and high-quality population of the evolutionary search, and select M samples from the initial population as candidate sets during the evolutionary search process; Step S43: hybridize the samples in the candidate set with the samples in the high-quality population to form a final candidate set. Apply the loop optimization scheduling primitive to each sample in the final candidate set to complete the loop transformation. Step S44: Construct a cost model to evaluate the performance of the loop configurations in the final candidate set, select the best N loop configurations, and randomly select a loop configuration from the remaining loop configurations with a probability of 0.05 to prevent the cost model screening from falling into a local optimum.

5. A template computing code generation system for a tensor accelerator, characterized in that: Includes the following modules: A template calculation intermediate representation module is used to convert the template calculation into a domain-specific programming language, wherein the domain-specific programming language includes: a basic formula of the template calculation and a loop optimization configuration specified by the user; and sets the total number of iterations of the template calculation code generation program to m times; The template calculation splitting module is used to split the weight matrix of the template calculation into rows and columns according to the template calculation mode expressed in the domain-specific programming language, extract the non-zero elements in the template calculation weight matrix, and use row splitting when the non-zero elements are in a row, and use column splitting when they are in a column; stack and rearrange the split weight rows or columns to form a weight matrix suitable for matrix multiplication, select a matrix multiplication shape suitable for the tensor calculation unit, determine whether the weight rows or columns can be stacked and rearranged within the range of the matrix multiplication shape, and perform rearrangement conversion on them; a loop optimization module for starting an automatic tuning code generator, setting the number of iterations i=0, and starting an iterative search for an optimal configuration for loop optimization of the template calculation code, loop optimizing the rearranged and transformed template calculation, generating a matrix multiplication loop optimization configuration through evolutionary search, building a candidate configuration queue, and screening the candidate configurations using a cost model; The evolutionary search module is used to apply the selected candidate configurations to real-machine testing. After obtaining the measured performance results, the optimal configuration is updated and the measured results are used to update the cost model. The number of iterations i is incremented by 1. The code generation module is used to check whether the number of evolutionary search iterations reaches the total number of iterations. If i=m at this time, the program ends and returns the current optimal configuration to obtain the optimal code generation sample. Otherwise, it goes to the loop optimization module.

Images