Method for optimizing implementation of convolution operator for tensor computing unit

Abstract

A kind of tensor computing unit convolution operator optimization implementation method faces, convolution operator is represented by the DSL of deep learning compiler, implicit general matrix multiplication is obtained by coordinate transformation to convolution calculation;Then after scheduling optimization is carried out to the convolution operator, the optimal search parameter is obtained by searching and the CUDA C code is generated through the back end of deep learning compiler, then the generated CUDA C code is integrated into neural network, the inference speed of convolutional neural network on NVIDIAGPU platform is improved.The application can improve the performance of automatic code generation of convolution operator in half-precision calculation, and provide guarantee for the performance of automatic code generation of fusion operator in neural network inference calculation.

Description

Technical Field

[0001] This invention relates to a technology in the field of artificial intelligence, specifically an optimized implementation method for the convolution operator of Tensor Cores in NVIDIA GPUs. Background Technology

[0002] Deep learning compiler technology has been proposed to improve the development efficiency of deep learning operators. Among them, TVM compilation optimization technology can effectively solve the problem of automatic code generation for fusion operators. However, the performance of operators generated by TVM is highly dependent on the development of operator scheduling strategies and the design of scheduling space. Real-world deep learning business scenarios require that the performance of operators generated by TVM can reach the ultimate performance of the hardware, just like manual optimization.

[0003] Existing Generalized Matrix Multiplication (GEMM) acceleration technologies have a narrow range of applications, and may not be able to perform convolution operator calculations or run fusion operators. Existing improved GEMM processing technologies are based on assembly instruction sets, which are difficult to program and do not have cross-platform versatility. Summary of the Invention

[0004] To address the aforementioned shortcomings of existing technologies, this invention proposes an optimized implementation method for convolution operators oriented towards tensor computation units. This method involves writing convolution operator computation and tensor computation unit scheduling code in a domain-specific language (DSL), and then generating convolution computation code through automatic tuning techniques. This invention improves the performance of automatic code generation for convolution operators in half-precision computation and ensures the performance of automatic code generation for fusion operators in neural network inference computation.

[0005] This invention is achieved through the following technical solution:

[0006] This invention relates to an optimized implementation method for convolution operators oriented towards tensor computation units. The method uses a deep learning compiler's DSL to represent the convolution operator, performs coordinate transformation on the convolution computation to obtain a computational representation of implicit general matrix multiplication, optimizes the scheduling of the convolution operator to obtain a scheduling template, searches for the optimal search parameters, generates CUDA C code through the backend of the deep learning compiler, and then integrates the generated CUDA C code into the neural network to improve the inference speed of the convolutional neural network on the NVIDIA GPU platform.

[0007] The aforementioned scheduling optimization of the convolution operator for Tensor Cores, resulting in a scheduling template, specifically includes:

[0008] Step 1) Cache the input data, weight data, and output results in local memory (shared memory) and tensor register (wmma fragment), respectively.

[0009] Step 2) Utilize the small storage space of half-precision floating-point types to optimize the cache read and write steps with double buffering.

[0010] Step 3) Utilize the advantage of high read / write bandwidth of half-precision floating-point type to optimize the vectorized scheduling of cache read / write steps.

[0011] Step 4) Divide the computational dimensions, that is, divide the GEMM_M dimension into four dimensions: bm, tm, om, and im; divide the GEMM_N dimension into four dimensions: bn, tn, on, and in; and divide the GEMM_K dimension into two dimensions: ok and ik according to the division parameters im and in.

[0012] Step 5) Perform GPU thread block binding and thread binding, that is, bind the split dimension bm to blockIdx.y, tm to threadIdx.y, bn to blockIdx.x, and in to threadIdx.x respectively.

[0013] Step 6) Map the unbound computation dimension to a GEMM computation represented by wmma::mma_sync; after binding, each thread block will compute om*on GEMMs, and the result size of the computation is im*in.

[0014] The segmentation parameters are searchable parameter templates, and the optimal segmentation parameters are obtained through automatic search in the deep learning compiler.

[0015] To improve the efficiency of parameter search, this invention specifically designs the search space and search method for segmenting parameters im and in, taking into account the characteristics of computational size and hardware architecture. Specifically, this includes:

[0016] ① When GEMM_M is a multiple of 32, set the search space of im to [8,16,32] and go to step ④; otherwise, go to step ②.

[0017] ② When GEMM_M is a multiple of 16, set the search space of im to [8,16] and go to step ④; otherwise, go to step ③.

[0018] ③ Set the search space of im to [8].

[0019] ④ Perform parameter search and optimization on IM.

[0020] ⑤ Determine the value of in based on the search results of im: when the search space of im is 8, 16 or 32, in corresponds to 32, 16 or 8.

[0021] This invention relates to a system for implementing the above-described method, comprising: an operator representation unit, a scheduling optimization unit, a parameter search unit, a code generation unit, and an inference optimization unit, wherein: the operator representation unit expresses computational logic using a deep learning compiler's DSL based on user-given computational semantics; the scheduling optimization unit performs scheduling optimization processing on the operator based on the operator's computational logic and the architecture characteristics of the hardware platform to obtain a scheduling optimization template for the operator; the parameter search unit performs heuristic search based on the scheduling optimization parameter template information to obtain the optimal splitting parameters; the code generation unit generates efficient CUDA C code based on the optimal splitting parameters; and the inference optimization unit compiles and connects the generated CUDA C code with the neural network on the NVIDIA GPU platform to accelerate the inference of the convolutional neural network.

[0022] Technical effect

[0023] This invention utilizes the coordinate transformation of the Implicit GEMM algorithm to optimize the scheduling of convolution operators, and designs specific parameter search methods tailored to the computational size and hardware platform architecture characteristics.

[0024] Compared with existing technologies, this invention significantly improves the computational performance and inference speed of convolution operators on NVIDIA GPU platforms. Attached Figure Description

[0025] Figure 1 This is a schematic diagram of the process of the present invention;

[0026] Figure 2 This is a schematic diagram of the scheduling process for convolution operators oriented towards tensor computation units;

[0027] Figure 3 A schematic diagram of the automatic search process for partitioning parameters for tensor computational units;

[0028] Figure 4 This is a schematic diagram of the original TVM convolution scheduling result;

[0029] Figure 5 This is a schematic diagram of the convolution scheduling results of this method. Detailed Implementation

[0030] like Figure 1 As shown, this embodiment relates to an optimized implementation method for a convolution operator oriented towards tensor computation units. The computational representation of the convolution operator is written using TVM's DSL, and after coordinate transformation based on the Implicit GEMM algorithm, it is implemented as follows: Figure 2The scheduling method shown performs Tensor Core-oriented scheduling optimizations for convolution computation. After scheduling optimization, the computation graph representation of the convolution operator is as follows: Figure 5 As shown, then use as Figure 3 The automatic tuning process shown searches for and determines the parameters for im and in, and finally uses the TVM backend to generate efficient CUDA C code based on the search results.

[0031] The scheduling optimization described above is written using TVM's scheduling primitives such as cache_read, cache_write, double_buffer, vectorize, split, and tensorize.

[0032] In the scheduling optimization process, this embodiment sets the block parameters as searchable parameters.

[0033] The automatic tuning process is implemented using the AutoTVM algorithm, which searches for the best-performing parameters in the search space through heuristic search and a cost model based on XGBoost.

[0034] Through specific practical experiments, tests were conducted on a platform with an NVIDIA RTX 2070 GPU. The GPU computation used CUDA framework version 11.0 and TVM version 0.8. Figure 4 Compared to the original TVM convolution operator scheduling optimization shown, this embodiment comprehensively considers the coordinate transformation method of Implicit GEMM, the use of tensor computation units, memory architecture, and the characteristics of half-precision floating-point types during the code generation stage, and designs a specific parameter search space according to the characteristics of the selected GPU architecture, thereby improving the performance of automatic generation of convolution operators.

[0035] Based on this embodiment, the correctness of the invention was verified through convolutional neural networks such as ResNet50, MobileNetv1, MobileNetv2, and SqueezeNet. At the same time, the invention has improved the inference speed of different networks to varying degrees compared with the community open-source TVM implementation.

[0036] The performance comparison of the convolution operator implemented in this invention with the convolution operator of the open-source deep learning compiler TVM is shown in the table below:

[0037]

[0038] The following table compares the inference speed of the convolutional neural network implemented using this invention with that of an open-source deep learning compiler:

[0039] Neural network (Batchsize=1) This invention TVM ResNet18 0.65ms 1.11ms ResNet50 1.26ms 2.14ms SqeeuzeNet1.0 0.28ms 0.41ms

[0040] Compared with existing technologies, the convolution operator obtained by this invention achieves a maximum speedup of 15.4 times compared with the open-source deep learning compiler TVM; the inference performance of the convolutional neural network achieves a maximum speedup of 1.70 times compared with TVM.

[0041] The above-described specific implementations can be partially adjusted by those skilled in the art in different ways without departing from the principles and purpose of the present invention. The scope of protection of the present invention is defined by the claims and is not limited to the above-described specific implementations. All implementation schemes within the scope of the claims are bound by the present invention.

Claims

1. A method for optimizing the implementation of convolution operators for tensor computation units, characterized in that, The convolution operator is represented by the DSL of the deep learning compiler. The implicit general matrix multiplication representation is obtained by performing coordinate transformation on the convolution calculation. Then, the convolution operator is scheduled and optimized to obtain the scheduling template. The optimal search parameters are obtained through search and CUDA C code is generated through the backend of the deep learning compiler. The generated CUDA C code is then integrated into the neural network to improve the inference speed of the convolutional neural network on the NVIDIA GPU platform. The aforementioned scheduling optimization of the convolution operator to obtain the scheduling template specifically includes: Step 1) Cache the input data, weight data, and output results in local memory and tensor registers respectively; Step 2) Utilize the small storage space of half-precision floating-point types to optimize cache read and write steps using double buffering; Step 3) Utilize the high read / write bandwidth of half-precision floating-point types to optimize the vectorized scheduling of cache read / write steps; Step 4) Divide the computational dimensions, that is, divide the GEMM_M dimension into four dimensions: bm, tm, om, and im; divide the GEMM_N dimension into four dimensions: bn, tn, on, and in; and divide the GEMM_K dimension into two dimensions: ok and ik according to the division parameters im and in. Step 5) Perform GPU thread block binding and thread binding, that is, bind the split dimension bm to blockIdx.y, tm to threadIdx.y, bn to blockIdx.x, and in to threadIdx.x respectively; Step 6) Map the unbound computation dimension to a GEMM computation represented by wmma::mma_sync; after binding, each thread block will compute om*on GEMMs, and the result size of the computation is im*in.

2. The method for optimizing the implementation of the convolution operator for tensor computation units according to claim 1, characterized in that, The segmentation parameters are searchable parameter templates, and the optimal segmentation parameters are obtained through automatic search in the deep learning compiler.

3. The method for optimizing the implementation of the convolution operator for tensor computation units according to claim 1 or 2, characterized in that, The aforementioned segmentation parameters are obtained in the following way: ① When GEMM_M is a multiple of 32, set the search space of im to [8, 16, 32] and go to step ④; otherwise, go to step ②. ② When GEMM_M is a multiple of 16, set the search space of im to [8, 16] and proceed to step ④; Otherwise, proceed to step ③; ③ Set the search space of im to [8]; ④ Perform parameter search and optimization on IM; ⑤ Determine the value of in based on the search results of im: when the search space of im is 8, 16 or 32, in corresponds to 32, 16 or 8.

4. A system for implementing the method of any one of claims 1 to 3, characterized in that, include: The system comprises an operator representation unit, a scheduling optimization unit, a parameter search unit, a code generation unit, and an inference optimization unit. Specifically, the operator representation unit expresses the computational logic using the deep learning compiler's DSL based on the computational semantics given by the user; the scheduling optimization unit performs scheduling optimization on the operator based on the computational logic of the operator and the architectural characteristics of the hardware platform to obtain the operator's scheduling optimization template. The parameter search unit performs a heuristic search based on the scheduling optimization parameter template information to obtain the optimal splitting parameters; the code generation unit generates efficient CUDA C code based on the optimal splitting parameters; the inference optimization unit compiles and connects the generated CUDA C code with the neural network on the NVIDIA GPU platform to accelerate the inference of the convolutional neural network.

Images