Template calculation code generation method and system for tensor accelerator

Abstract

The invention relates to a tensor accelerator-oriented template calculation code generation method and system. The method comprises the following steps of: S1, calculating and converting a template into a field specific programming language designed by the tensor accelerator-oriented template calculation code generation method and the tensor accelerator-oriented template calculation code generation system; s2, carrying out row segmentation and column segmentation on the weight matrix calculated by the template; s3, stacking and rearranging the segmented weight rows / columns to form a weight matrix suitable for matrix multiplication; s4, starting an automatic tuning code generator, forming matrix multiplication cycle optimization configuration through evolutionary search, and screening candidate configuration by adopting a cost model; s5, the screened candidate configuration is used for actual measurement performance of a real machine, and an actual measurement result is used for updating a cost model; and S6, checking whether the number of evolution search iterations reaches an iteration number threshold or not, if so, ending the program, and selecting the optimal configuration, otherwise, turning to S4. According to the method, the multi-stage soft flow for memory access and calculation is constructed, so that the memory access overhead can be better masked.

Description

Technical Field

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

[0002] In the field of natural science, partial differential equations describe the laws of natural phenomena by expressing the relationship between variables and partial derivatives of variables, and have become an indispensable part of modern scientific applications. For example, partial differential equations are widely used in climate prediction, earthquake simulation, fluid mechanics and other fields. In order to solve the partial differential equations in these applications, numerical methods are usually used, among which the finite difference method (FDM) is a common choice. In the finite difference method, template calculation is one of the important calculation modes based on the uniform grid finite difference method.

[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 features, including dimension, shape, order, type of coefficients in calculation, boundary processing method, etc. Figure 1 As shown in Figure 1, the order R of the template calculation 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 set of template calculations of various orders, dimensions, and shapes. In complex scientific simulations, one round of iterations may require grid updates based on multiple types of templates. For example, COSMO is one of the widely used scientific software in the field of meteorological simulation. Its execution process is quite complex and includes 30 template calculations of different dimensions and orders. The dimensions of these template calculations cover two-dimensional and three-dimensional dimensions, and the orders range from 2nd to 16th order. Due to the diversity of different template calculation modes and memory access modes, optimizing computing performance requires not only general template calculation optimization on various computing hardware platforms, but also specialized optimization and performance tuning for each type of template calculation.

[0005] Most of the existing template calculation code generation methods are based on the general computing unit of GPU. For example, Artemis abstracts loop optimization into scheduling primitives and provides an automatic tuning algorithm to automatically generate high-performance template calculation code. However, Artemis still performs calculation optimization based on the CUDA Core of GPU, which makes it difficult to utilize the computing power of tensor computing units and loses the opportunity for performance optimization. TCStencil uses calculation mapping to convert template calculations into a series of reduction and summation operations, thereby realizing the mapping of template calculations to tensor computing units. However, current work still remains at the manual implementation of CUDA Kernel, and it is difficult to realize the calculation mapping of three-dimensional templates and high-order two-dimensional templates. Existing research has fully exploited the parallel performance of CUDA Cores by establishing domain-specific programming language compilers and proposing a variety of template calculation parallel optimization methods, but it is difficult to fully utilize the latest tensor accelerators. In addition, existing work converts template calculations into GEMM calculations to schedule tensor computing units to achieve performance acceleration, but it is still limited to manual optimization and difficult to achieve automatic 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 a tensor accelerator, designs a code generation method for a GPU tensor accelerator that can automatically generate high-performance template calculation code, and introduces shared memory asynchronous reading to make the memory access part of the 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 for 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 by 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. When the non-zero element is located in a row, row splitting is adopted, and when 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 conversion 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 loop optimization of the template calculation code, 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 a cost model to screen the candidate configurations;

[0013] Step S5: Use the selected candidate configurations for real machine testing. After obtaining the measured performance results, update the optimal configuration, and use the measured results to update the cost model. 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] The 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 for the template calculation and a loop optimization configuration specified by a 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 perform row and column splitting on the weight matrix of the template calculation according to the template calculation mode expressed by the domain-specific programming language, extract the non-zero elements in the weight matrix of the template calculation, 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 is used to start the automatic tuning code generator, set the number of iterations i=0, start iterative search for the optimal configuration of the loop optimization of the template calculation code, loop optimize the template calculation after the rearrangement conversion, generate the matrix multiplication loop optimization configuration through evolutionary search, build a candidate configuration queue, and use the cost model to screen the candidate configuration;

[0019] The evolutionary search module is used to apply the selected candidate configurations to real machine testing, and after obtaining the measured performance results, update the optimal configuration and use the measured results to update the cost model. The number of iterations i is increased 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, and 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 A schematic diagram of a two-dimensional Star-type template calculation in an embodiment of the present invention;

[0024] Figure 2 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 in an embodiment of the present invention;

[0026] Figure 4 A schematic diagram of stacking and rearranging weight rows in an embodiment of the present invention;

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

[0028] Figure 6 It is a 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 implementations and in conjunction with the accompanying drawings.

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

[0032] Step S1: converting the template calculation into the 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 a 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 by 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. When the non-zero element is located in a row, row splitting is adopted, and when it is located in a column, column splitting is adopted;

[0034] Step S3: 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 GPU tensor computing unit, determine whether the weight rows / columns can be stacked and rearranged within the range of the matrix multiplication shape, and perform rearrangement conversion on them;

[0035] Step S4: Start the automatic tuning code generator, set the number of iterations i=0, and start iterative search for the optimal configuration of the template calculation code loop optimization. The rearranged and transformed template calculation is loop optimized, and the 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: Use the selected candidate configurations for real machine testing. After obtaining the measured performance results, update the optimal configuration, and use the measured results to update the cost model. 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 loop optimization configuration specified by a user; setting the total number of iterations of the program 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. The user can specify the block size of each layer of loop, select the loop body for loop fusion, and specify the vectorization length of the logic statement in the loop. Specifically, according to the GPU hardware storage hierarchy, the loop blocks correspond to the global video memory, shared memory and registers respectively. The user can guide the selection of the loop block size according to the storage capacity of the current hardware. For the vectorization length, it is necessary to select the appropriate vectorization length according to the hardware vector instruction. 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 initializer program needs to be 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, and 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 of the template calculation, firstly generating a weight matrix of size (2r+1)×(2r+1) according to the radius r of the template, wherein the element values ​​of the initially generated weight matrix are all 0, and then assigning the weight elements to the weight matrix according to the weight array input by the user;

[0044] Step S22: extracting weight rows / columns in the weight matrix by 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 the following two ways to extract the weight row / column in the weight matrix.

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

[0047] Method 2: extract the non-zero elements in the same column of the weight matrix, record the offset of its starting address relative to the starting address of the weight matrix, and use it as the weight column, and 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 a 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, 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 according to the minimum-scale matrix multiplication shape suitable for 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, according to the tensor computing units in the hardware architecture used, a pool of all suitable matrix multiplication shapes is constructed, 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 the 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 for 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 they are 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 for 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 it is 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 template calculation after the rearrangement conversion, 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 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): merges 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 loading from global memory to registers to first loading from global memory to shared memory, and then 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): Change the data in the first loop nest a from registers to global memory to first write from registers to shared memory, and then write from shared memory to global 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 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, and 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, mma_optimize primitives are applied in sequence to complete the loop transformation;

[0063] Step S43: Apply cache_read primitive, scan the template to calculate memref.LoadOp in the loop nest, determine the size of the data block to be read according to the loadIndices of the operator, build a new memref.LoadOp in the loop, set the source address of the data read to the global memory, set the target address to the 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 an abstract mma instruction of the tensor computing unit 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 according to 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 loop_tile_for_block primitive, select the outermost loop of loop nest a according to the input loop nest a and tile_size, define the loop lower bound as lower_bound, upper bound as upper_bound, and iteration step as step, generate a new inner loop with lower bound as 0, upper bound as tile_size, and iteration step as 1; the original lower bound and upper bound of the outermost loop remain unchanged, the iteration step becomes step×tile_size, and the loop variables of the outermost loop are bound to blockIdx.X, blockIdx.Y, and blockIdx.Z of the GPU;

[0067] Step S47: Apply loop_tile_for_warp primitive, select the inner loop of loop nest a according to the input loop nest a and tile size tile_size, define the loop lower bound as lower_bound, upper bound as upper_bound, iteration step as step, generate a new inner loop, the lower bound is 0, the upper bound is tile_size, iteration step is 1; the original inner loop lower bound and upper bound remain unchanged, the iteration step becomes step×tile_size, the loop variables of the inner loop are bound to GPU threadIdx.Y, threadIdx.Z, and 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 the 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 3 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 2 times in the N dimension and 3 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 loop transformation. The output of the cost model is the predicted performance of the code sample. All the code samples generated by the candidate set are input into the cost model. The predicted performance of the candidate set is sorted. By default, 64 candidate samples with performance priority are selected. In addition, samples of 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, which will be used for the cost model to screen high-quality sample codes. After the above step S5: the screened candidate configurations will be used for real machine testing, and after obtaining the measured performance results, the optimal configuration will be updated, and the measured results will be used to update the cost model, and the number of iterations i will increase 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, build a real input data set as the input of 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 in step S51 to the high-quality population of evolutionary search, and hybridizing them with the candidate set randomly selected from the initial population during the evolutionary search process, and finally completing a search iteration, and the number of iterations i is increased by 1;

[0074] In one embodiment, the above step S6: check whether the number of evolutionary search iterations reaches the iteration number threshold. If i=m at this time, the program ends, and the current optimal configuration is returned and the current cost model is saved. 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 times, 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 detection will be initiated. 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 computing code generation system for a tensor accelerator, comprising 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, and supports the user to indicate the loop optimization configuration and provide 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 perform row splitting and column splitting on the weight matrix of the template calculation, extract the non-zero elements in the weight matrix of the template calculation, and when the non-zero elements are in a row, row splitting is adopted, and when they are 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, vectorization, and expose the loop block size and vectorization length as tunable interfaces to the evolutionary search module;

[0081] Evolutionary search module M4 is used to initialize loop optimization configuration, iteratively search for the optimal loop configuration, generate candidate loop configurations during the search process, and use the cost model to screen 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 evolutionary population. The number of iterations i is increased 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 description is only a specific embodiment of the present invention and is not intended to limit the present invention. Any modifications, equivalent substitutions, improvements, etc. made within the spirit and principles of the present invention should be included 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 for 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 by 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. When the non-zero element is located in a row, row splitting is adopted, and when 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 conversion 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 loop optimization of the template calculation code, 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 a cost model to screen the candidate configurations; Step S5: Use the selected candidate configurations for real machine testing. After obtaining the measured performance results, update the optimal configuration, and use the measured results to update the cost model. 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 by any of the following methods: Method 1: extract the non-zero elements in the same row in 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 cut the column where the current non-zero element is located from the weight matrix as the weight column.

3. The 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 row into a column tuple, and place it from left to right in the columns of the matrix B of the matrix multiplication. Each time you go to the right, you shift one element downward and fill the vacant position with 0. Method 2: Transpose the weight column into a row tuple, and place 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. The 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, which specifically includes: Step S41: abstracting the rearranged converted template calculation loop optimization into loop scheduling primitives, including loop blocking, cache reading, cache writing back, and vectorization; the evolutionary search is used to generate loop blocking size and vectorization length; Step S42: start the automatic tuning code generator, 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; Step S43: hybridize the samples in the candidate set with the samples in the high-quality population to form a final candidate set, and 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 loop configurations from the remaining loop configurations with a probability of 0.05 to avoid the cost model screening falling into a local optimum.

5. A template computing code generation system for a tensor accelerator, characterized in that: Includes the following modules: The 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 for the template calculation and a loop optimization configuration specified by a 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 perform row and column splitting on the weight matrix of the template calculation according to the template calculation mode expressed by the domain-specific programming language, extract the non-zero elements in the weight matrix of the template calculation, 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 is used to start the automatic tuning code generator, set the number of iterations i=0, start iterative search for the optimal configuration of the loop optimization of the template calculation code, loop optimize the template calculation after the rearrangement conversion, generate the matrix multiplication loop optimization configuration through evolutionary search, build a candidate configuration queue, and use the cost model to screen the candidate configuration; The evolutionary search module is used to apply the selected candidate configurations to real machine testing, obtain the measured performance results, update the optimal configuration, and use the measured results to update the cost model. The number of iterations i is increased 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