Convolution operation method and device, electronic equipment and storage medium

Abstract

The invention relates to the technical field of artificial intelligence chips, and provides a convolution operation method and device, electronic equipment and a storage medium, and the method comprises the steps: traversing convolution kernel elements, and determining a current to-be-loaded data block based on the shape of a target data block and the coordinates of the traversed current convolution kernel element block; based on the current to-be-loaded data block, determining a to-be-covered data block and a newly-added data block, and loading the newly-added data block into the shared memory; based on the initial position, reading the input data block from the shared memory, applying the input data block and the current convolution kernel element block, performing matrix multiply-accumulate operation, and accumulating an operation result to an output result; and determining a convolution operation result based on an output result obtained by accumulation after traversal is completed. According to the method and the device, only the newly added data blocks are loaded into the shared memory, so that repeated data loading can be avoided, the data volume loaded each time is reduced, and the memory bandwidth is saved.

Description

Technical Field

[0001] The present invention relates to the technical field of artificial intelligence chips, and particularly to a convolution operation method, device, electronic device, and storage medium. Background Art

[0002] Convolution operation is one of the core operations in computer vision and deep learning, and is widely used in traditional image filtering (such as edge detection, blurring) and modern convolutional neural networks (CNNs). Its core advantages lie in local connection (only focusing on pixels within the local receptive field) and the convolutional kernel parameter sharing mechanism (that is, the same convolutional kernel is reused in the spatial dimension), which can efficiently extract local features of images while reducing the number of model parameters, enabling CNNs to dominate in tasks such as image recognition and image classification.

[0003] Currently, in order to adapt to parallel computing architectures such as Graphics Processing Unit (GPU), General-purpose computing on Graphics Processing Units (GPGPU), and Tensor Processing Unit (TPU), the method of converting convolution operation to General Matrix Multiplication (GEMM) operation by Image to Column (im2col) is usually adopted to accelerate the operation by using highly optimized matrix calculation libraries (such as BLAS, GEMM, etc.).

[0004] However, the above method needs to expand the input feature map data into a matrix, which will generate a lot of duplicate data, resulting in the same data being stored multiple times in memory, thus increasing memory occupancy and data loading volume. Summary of the Invention

[0005] The present invention provides a convolution operation method, device, electronic device, and storage medium to solve the defects of large memory space occupancy and large data loading volume existing when converting convolution operation to matrix multiplication operation in related technologies.

[0006] The present invention provides a convolution operation method, including: Traverse the convolution kernel elements, and determine the current data block to be loaded based on the target data block shape and the coordinates of the currently traversed convolution kernel element block, where the target data block shape is determined based on convolution parameters and input data shape; Based on the current data block to be loaded, determine the data block to be overwritten and the new data block, and load the new data block into the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory; Based on the starting position, read the input data block from the shared memory, apply the input data block and the current convolutional kernel element block to perform matrix multiplication and accumulation operations, and accumulate the operation results to the output result. The starting position is determined based on the coordinates of the current convolutional kernel element block; Based on the output result accumulated after traversal is completed, determine the convolution operation result.

[0007] According to a convolution operation method provided by the present invention, it further includes: Based on the target data block shape and the coordinates of the next convolutional kernel element block, determine the next data block to be loaded; Based on the next data block to be loaded, determine the new data block, and preload the new data block from the global memory to the cache. The preloading operation of the new data block is executed in parallel with the matrix multiplication and accumulation operation of the current convolutional kernel element block.

[0008] According to a convolution operation method provided by the present invention, the coordinates of the next convolutional kernel element block are determined based on the traversal order of the convolutional kernel elements and the coordinates of the current convolutional kernel element block, and the traversal order is any one of row-first traversal and column-first traversal.

[0009] According to a convolution operation method provided by the present invention, the determining the new data block based on the next data block to be loaded includes: In the case where both the row dimension index and the column dimension index of the next convolutional kernel element block change relative to the current convolutional kernel element block, use the data block to be loaded as the new data block; In the case where either the row dimension index or the column dimension index of the next convolutional kernel element block changes relative to the current convolutional kernel element block, determine the new data block from the next data block to be loaded.

[0010] According to a convolution operation method provided by the present invention, the determining the data block to be overwritten and the new data block based on the current data block to be loaded, and loading the new data block into the shared memory includes: In the case where the current convolutional kernel element block is the first convolutional kernel element block, determine that the data block to be overwritten is empty, and load the current data block to be loaded as the new data block from the global memory into the shared memory; When the current convolution kernel element block is not the first convolution kernel element block, based on the current data block to be loaded, determine the data block to be overwritten, and read the new data block from the cache to load the new data block from the cache into the shared memory.

[0011] According to a convolution operation method provided by the present invention, the loading the new data block into the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory includes: Based on the row dimension index or column dimension index of the data block to be overwritten, determine the storage location of the new data block in the shared memory, and load the new data block to the storage location.

[0012] According to a convolution operation method provided by the present invention, the shared memory is a circular buffer. Accordingly, the reading the input data block from the shared memory based on the starting position includes: Based on the starting position, apply a matrix loading instruction to circularly read the input data block from the shared memory row by row or column by column.

[0013] The present invention also provides a convolution operation device, including: A data block determination unit, configured to traverse the convolution kernel elements, and based on the target data block shape and the coordinates of the currently traversed convolution kernel element block, determine the current data block to be loaded, where the target data block shape is determined based on convolution parameters and input data shapes; A data loading unit, configured to determine the data block to be overwritten and the new data block based on the current data block to be loaded, and load the new data block into the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory; A matrix operation unit, configured to read the input data block from the shared memory based on the starting position, apply the input data block and the current convolution kernel element block to perform matrix multiplication and accumulation operations, and accumulate the operation results to the output result, where the starting position is determined based on the coordinates of the current convolution kernel element block; A result determination unit, configured to determine the convolution operation result based on the output result accumulated after the traversal is completed.

[0014] The present invention also provides an electronic device, including a memory, a processor, and a computer program stored on the memory and running on the processor, where when the processor executes the computer program, the convolution operation method described in any one of the above is implemented.

[0015] The present invention also provides a non-transitory computer-readable storage medium, on which a computer program is stored, and when the computer program is executed by a processor, the convolution operation method described in any one of the above is implemented.

[0016] The present invention also provides a computer program product, including a computer program which, when executed by a processor, implements the convolution operation method as described in any one of the above.

[0017] For the convolution operation method, device, electronic device and storage medium provided by the present invention, by traversing the convolution kernel elements, the current data block to be loaded can be dynamically determined according to the coordinates of the current convolution kernel element block. Based on the current data block to be loaded, the new data block can be directly determined, and only the new data block is loaded into the shared memory instead of loading the entire current data block to be loaded. Thus, the reuse of partial data in the shared memory can be realized, thereby avoiding repeated data loading, reducing the amount of data loaded each time, and further saving the memory bandwidth. When loading the new data block into the shared memory, it is loaded to the position of the data block to be overwritten in the shared memory so as to overwrite the data block to be overwritten. This mechanism ensures the efficient utilization of the shared memory and avoids unnecessary memory occupation and frequent data movement. In addition, compared with the traditional im2col method, the present invention does not need to first convert the input data into a matrix for storage. Therefore, no additional memory space is consumed to store the generated matrix, thereby reducing the occupation of memory space. BRIEF DESCRIPTION OF THE DRAWINGS

[0018] In order to more clearly illustrate the technical solutions in the present invention or related technologies, the following will briefly introduce the drawings required for use in the embodiments or related technology descriptions. Obviously, the drawings in the following description are some embodiments of the present invention. For those of ordinary skill in the art, other drawings can be obtained based on these drawings without creative efforts.

[0019] Figure 1 is a schematic diagram of the principle of converting a convolution operation into a matrix operation in related technologies; Figure 2 is a schematic diagram of data loading in the Implicit GEMM algorithm provided by the present invention; Figure 3 is a schematic diagram of the structure of the general image processor provided by the present invention; Figure 4 is a schematic flowchart of the convolution operation method provided by the present invention; Figure 5 is a schematic diagram of the coordinates of the convolution kernel element block provided by the present invention; Figure 6 is a schematic diagram of the shape of the input data after padding provided by the present invention; Figure 7 is one of the schematic diagrams of data block loading provided by the present invention; Figure 8It is a schematic diagram of parallel execution of data preloading and matrix multiplication and accumulation operations provided by the present invention; Figure 9 It is the second schematic diagram of data block loading provided by the present invention; Figure 10 It is the third schematic diagram of data block loading provided by the present invention; Figure 11 It is the fourth schematic diagram of data block loading provided by the present invention; Figure 12 It is a schematic structural diagram of a convolution operation device provided by the present invention; Figure 13 It is a schematic structural diagram of an electronic device provided by the present invention. Detailed implementation manners

[0020] To make the objectives, technical solutions and advantages of the present invention clearer, the technical solutions in the present invention will be clearly and completely described below with reference to the accompanying drawings in the present invention. Obviously, the described embodiments are some but not all of the embodiments of the present invention. All other embodiments obtained by those of ordinary skill in the art based on the embodiments in the present invention without making creative efforts shall fall within the protection scope of the present invention.

[0021] Convolutional neural networks have achieved remarkable results in computer vision tasks such as image classification, object detection, and video processing. The core computational workload mainly comes from the convolution operation process with multiple convolutional kernels and multiple channels in the convolutional layer. Convolution operation is a typical computationally intensive task. Since many hardware acceleration cards (such as GPUs, GPGPUs, TPUs, etc.) have been deeply optimized for general matrix multiplication (GEMM), in order to make full use of hardware resources, the most common practice in current convolutional operation implementations is to adopt the image-to-column matrix conversion strategy (i.e., the im2col algorithm), convert the convolution operation into matrix multiplication, and then use a highly optimized matrix calculation library to calculate and implement.

[0022] The core operation of convolution is sliding the window. The convolutional kernel slides on the input data (i.e., the convolution operation) to extract the local region features of the input data. The convolutional kernel is usually a small two-dimensional matrix (such as 3×3, 5×5), and for multi-channel inputs (such as color images), the convolutional kernel can be extended to a three-dimensional matrix (such as 3×3×3, the third dimension corresponds to the number of input channels). Each element in the convolutional kernel is a weight, which is used to multiply and accumulate the sum with the corresponding position of the input data. These weights determine the sensitivity of the convolutional kernel to the input data, and through training, the weights can be automatically adjusted to extract meaningful features.

[0023] Figure 1 It is a schematic diagram of the principle of converting a convolution operation into a matrix operation in the related art, such as Figure 1As shown, the shape of the input data is (N, H, W, C), and the shape of the convolutional kernel is (K, R, S, C). Here, N represents the batch size, that is, the number of data samples processed at one time. For example, N = 1 means there is only one sample; H represents the height of the input data, W represents the width of the input data; C represents the number of channels of the input data. For example, C = 2 means the input data has two channels. K represents the number of convolutional kernels, R represents the height of the convolutional kernel, S represents the width of the convolutional kernel, and C represents the number of channels of the convolutional kernel. This number of channels must match the number of channels of the input data for convolution operations. In the convolution operation, the convolutional kernel slides on the input data to calculate the weighted sum at each position, thereby generating the output data. The number of convolutional kernels (K) determines the number of channels of the output data. For a given input data shape (N, H, W, C) and convolutional kernel shape (K, R, S, C), the shape of the output data can be expressed as (N, P, Q, K), where P and Q represent the height and width of the output data respectively, and they depend on factors such as the height and width of the input data, the size of the convolutional kernel, the stride, and the padding. K is the number of convolutional kernels and also the number of channels of the output data. It should be understood that the stride refers to the number of pixels moved each time the convolutional kernel slides on the input data, that is, the relative offset of each convolution; padding refers to filling a certain number of pixel points at the boundary of the input data, and these pixel points are usually zeros.

[0024] Taking N = 1, H = 5, W = 5, C = 2, K = 2, R = 3, S = 3, stride = 1, padding = 1 as an example, in the convolution operation, it is usually to convert the convolution operation into a matrix operation through im2col, as Figure 1 shown. When converting the input data into the left matrix, each row of the left matrix is the input eigenvalue corresponding to one convolution operation. Taking the elements of the first row of the left matrix as an example, it is the input eigenvalue corresponding to the first convolution operation (that is, the first time the convolutional kernel slides on the input data). Figure 1 The colored block A shown in the upper left corner in is the input eigenvalue corresponding to the first time the convolutional kernel slides on the input data. When implementing the convolution operation using the im2col algorithm, it is necessary to expand the input eigenvalue corresponding to the colored block A into the elements of the first row of the left matrix. After expansion, each element in the first row of the left matrix corresponds one by one to the element in the corresponding position in the colored block A ( Figure 1 corresponding by the same color in ).

[0025] When converting the input data into the left matrix, each row of the left matrix has input eigenvalues. Since the stride = 1, the sliding window needs to be moved 25 times for convolution operations, so there are a total of OK, but due to the sliding window characteristic of convolution, a large amount of duplicate data will exist in the left matrix. The right matrix is a column matrix formed by expanding two convolutional kernels, and thus the corresponding output matrix can be obtained through matrix multiplication operation.

[0026] In the process of implementing the convolution operation as described above, the im2col algorithm will expand the input data to 9 times the original according to the size of the convolutional kernel, such as 3×3, and store it in the global memory. The amount of data expansion is proportional to the size of the convolutional kernel. For example, as Figure 1 shown, the amount of data of the original input data is , while the amount of data of the expanded left matrix is . Compared with the amount of data of the original input data, the amount of data after expansion increases by 8 times, and additional memory space is required to store it. In addition, there are many duplicate data in the expanded data. Because in the im2col algorithm, the input eigenvalue corresponding to each sliding window of the convolution operation will be converted into a row of the left matrix, and when the convolutional kernel slides on the input data, since the adjacent window positions will overlap, the elements in the corresponding left matrix of these windows will contain a lot of duplicate data, and these data need to be repeatedly loaded during calculation, wasting a large amount of memory bandwidth.

[0027] In response to this, a related technology has proposed another method for implementing convolution operation, namely the Implicit General Matrix Multiply (Implicit GEMM) algorithm. This algorithm does not need to explicitly construct the convolution matrix. Instead, when loading the input data from the global memory to the shared memory, it dynamically generates blocks of the convolution matrix through coordinate offset. Once the convolution matrix blocks are formed in the shared memory, the existing GEMM components can be used to accumulate the calculation results into the output tensor. Although this method does not require consuming memory space to store the additionally generated matrix, there are still many duplicate data when loading data from the global memory to the shared memory, wasting a large amount of memory bandwidth.

[0028] Figure 2 is a schematic diagram of data loading in the Implicit GEMM algorithm provided by the present invention. As Figure 2 shown, for a convolution operation with a convolutional kernel size of R S, the Implicit GEMM algorithm needs to perform R S times of data loading on the input data, and use the Matrix Multiply Accumulate (MMA) instruction to calculate the results. Finally, the results of R S times of MMA are accumulated together. Taking a 3×3 convolution (padding = 1, dilation = 1, stride = 1) as an example, the loading process of the input data is as Figure 2As shown here, dilation represents the dilation rate, which is used to increase the receptive field and enables the convolutional kernel to cover a larger input area. The blue area shown in the figure represents the valid area of the input data, the red dashed box represents the data area that needs to be loaded each time, the white area represents the loaded padding area, and r and s represent the coordinates of the corresponding convolutional kernel elements during each load. For a given input data shape (H, W, C), the data shape that needs to be loaded each time can be expressed as (P, Q, C), and its calculation formula is as follows: In the above formula, represents the dilation rate in the vertical direction, represents the dilation rate in the horizontal direction, , , and represent the padding values in the up, down, left, and right directions respectively, represents the stride in the vertical direction, represents the stride in the horizontal direction.

[0029] During the above process of loading input data, since it is necessary to perform R S times of data loading, and the amount of data loaded each time is P Q C, the total amount of data loaded is R S P Q C. It can be seen that for the case where the convolutional kernel size is greater than 1, there are still a lot of duplicate data when loading data from global memory to shared memory. For example, from the data areas loaded twice at r = 0, s = 0 and r = 0, s = 1, it can be seen that there is a large amount of duplicate data in these two loads, wasting a large amount of memory bandwidth.

[0030] Take the union of the data loaded R Figure 2 S times in , and the size of this union can be expressed as . When the stride stride = 1, and The calculation formulas are as follows: Therefore, the non - duplication rate of data loading is: Generally, padding is a small number such as 0 or 1, and dilation is a small number such as 1 or 2. When H and W are large, the above formula is approximately equal to , so the larger the shape of the convolution kernel, the greater the impact on data loading and the performance of the Implicit GEMM algorithm.

[0031] In response to this, the present invention provides a convolution operation method. By determining the data block to be overwritten and the newly added data block based on the currently to-be-loaded data block, and only loading the newly added data block into the shared memory, data duplication loading can be avoided, the amount of data loaded each time can be significantly reduced, thereby saving memory bandwidth, and further overcoming the above-mentioned defects.

[0032] It should be noted that convolution operation is a core operation in deep learning and is widely used in fields such as image processing, signal processing, and natural language processing. The method provided by the present invention is mainly used to implement the convolution operation between input data and a convolution kernel, where the form of the input data is diverse, specifically depending on the application scenario. For example, when the method provided by the present invention is applied to image processing tasks such as image classification, object detection, and image segmentation, the input data can be a single-channel grayscale image (height × width), such as a 224×224 grayscale image, and its shape can be expressed as (224, 224); the input data can also be a multi-channel color image (height × width × number of channels), such as a 224×224 RGB color image, and its shape can be expressed as (224, 224, 3). By sliding the convolution kernel on the image, features such as edges and textures can be extracted.

[0033] In addition, the execution subject of the convolution operation method provided by the present invention can be an artificial intelligence chip such as a GPU, GPGPU, or TPU. Such chips are highly optimized for matrix multiplication operations to accelerate calculations. By converting the convolution operation into the form of matrix multiplication, the present invention can efficiently execute it using the matrix multiplication operation unit on such chips, thereby optimizing the performance of the convolution operation. Taking the GPGPU as an example below, the structure of the execution subject of the present invention will be briefly introduced.

[0034] Figure 3 is a schematic structural diagram of the general image processor provided by the present invention, as Figure 3 shown, the general-purpose graphics processor (GPGPU) is actually an array of streaming processor clusters (abbreviated as SPC), for example, including Figure 3 the streaming processor cluster 1 shown in...,..., the streaming processor cluster M, and M is a positive integer greater than 1. In the GPGPU, 1 streaming processor cluster processes a computing task, or multiple streaming processor clusters process a computing task. Data sharing is performed between multiple streaming processor clusters through the global memory.

[0035] As shown Figure 3 in the figure, taking the streaming processor cluster 1 as an example, one streaming processor cluster includes multiple computing units, such as computing unit 1, computing unit 2, ……, computing unit N, where N is a positive integer. Each computing unit (Compute Unit, abbreviated as CU, for example, is a streaming processor) is used to perform arithmetic and logical operations other than matrix calculations such as matrix multiplication and convolution operations, such as accumulation, reduction, conventional addition, subtraction, multiplication, division, etc. A computing unit includes multiple cores (also called computing cores or computing kernels), and each computing core includes an arithmetic logic unit, a floating-point computing unit, etc., and the computing core is used to perform specific computing tasks. In addition, the computing unit also includes registers (such as Figure 3 the register bank shown in the figure), shared cache or shared memory, which are used to hierarchically store the source data and destination data related to the computing task, and the shared cache or shared memory in a computing unit is used to share data among the cores of the computing unit.

[0036] The streaming processor cluster 1 also includes a tensor operation unit, and the tensor operation unit is used to perform tensor calculations. For example, the tensor calculation may include matrix multiplication, convolution operation, etc. The tensor operation unit includes multiple tensor cores (i.e., TensorCore), which are used to perform specific computing tasks. An intermediate-level cache (such as L1 cache) can also be set in the streaming processor cluster 1, and the intermediate-level cache is used to share data among the computing units within a streaming processor cluster. The structures of other streaming processor clusters are the same as that of the streaming processor cluster 1 and will not be elaborated here. In addition, the GPGPU also includes an L2 cache, which is a high-speed cache located between the global memory and the intermediate-level cache (such as L1 cache), and is used to reduce the global memory access latency. The L2 cache is usually a shared resource of the entire GPGPU chip rather than private to each streaming processor cluster, and multiple streaming processor clusters can share data through the L2 cache.

[0037] Figure 4 is a schematic flowchart of the convolution operation method provided by the present invention. As shown Figure 4 in the figure, the method includes: Step 410, traverse the convolution kernel elements, and determine the current data block to be loaded based on the target data block shape and the coordinates of the currently traversed convolution kernel element block, where the target data block shape is determined based on the convolution parameters and the input data shape.

[0038] Specifically, traversing the convolution kernel elements means accessing all elements of the convolution kernel in a certain order. For example, the convolution kernel can be traversed by element-by-element traversal or block-by-block traversal. Here, element-by-element traversal means accessing each weight of the convolution kernel one by one; block-by-block traversal means dividing the convolution kernel into several sub-blocks and traversing them block by block.

[0039] The current convolutional kernel element block refers to the sub-region or element of the convolutional kernel that is currently being processed during traversal. The coordinates of the current convolutional kernel element block refer to the position index of the current convolutional kernel sub-block in the complete convolutional kernel. If the entire convolutional kernel is regarded as a grid, the coordinates (r, s) represent the row number r and column number s of the sub-block in the grid. Figure 5 It is a schematic diagram of the coordinates of the convolutional kernel element block provided by the present invention, as Figure 5 shown, for a convolutional kernel of size 3×3, if it is divided into 1×1 sub-blocks, the coordinates of each sub-block are Figure 5 shown. The traversal order can be (0, 0) → (0,1) → (0,2) → (1,0) → … → (2,2), processing one 1×1 weight block each time. This traversal order that first fixes the row index, traverses all columns, and then moves to the next row is called row-major traversal. In addition, the traversal order can also be (0, 0) →(1,0) → (2,0) → (0,1) → … → (2,2), also processing one 1×1 weight block each time. This traversal order that first fixes the column index, traverses all rows, and then moves to the next column is called column-major traversal. If it is divided into 2×2 sub-blocks (padding or adjustment is required), the coordinates of each convolutional kernel element block can be represented by the coordinate ranges of the starting element and the ending element, which will not be elaborated here.

[0040] The shape of the target data block refers to the shape of the data block loaded from the input data each time, which can be determined according to convolutional parameters and the shape of the input data, etc. Here, the convolutional parameters refer to the key parameters affecting the convolutional operation, specifically including convolutional kernel size, stride, padding, dilation rate, etc. Assuming the shape of the input data is (H, W, C), the shape of the target data block can be expressed as (P, Q, C), where the calculation formulas for P and Q can refer to the above Figure 2 related introduction and will not be elaborated here. It should be understood that in the embodiments of the present invention, by dynamically determining the shape of the target data block based on convolutional parameters and the shape of the input data, the present invention can flexibly adapt to different convolutional operation requirements and has wide applicability.

[0041] After determining the shape of the target data block, the traversal of the convolutional kernel elements can be started. According to the coordinates of the current convolutional kernel element block traversed and the shape of the target data block, the current data block to be loaded can be determined. Here, the current data block to be loaded refers to the data block that theoretically needs to be loaded from the input data to the shared memory for matrix multiply-accumulate (MMA) operation with the current convolutional kernel element block in the current calculation step. It is usually a local area of the input data, with the shape determined by the target data block and the position determined by the coordinates of the current convolutional kernel element block and the convolutional parameters.

[0042] Figure 6 It is a schematic diagram of the shape of the input data after padding provided by the present invention. As Figure 6 shown, taking the convolution operation with the shape of the input data being H = 6, W = 6, C = 1, the convolution kernel size being 3×3, the stride being stride = 1, the padding being padding = 1, and the dilation being dilation = 1 as an example, the shape of the complete input data after padding is as Figure 6 shown. According to the above parameters, it can be calculated that the height of the target data block is P = 6 and the width is Q = 6. When traversing the convolution kernel using a 1×1 block method, the starting addresses of the data blocks to be loaded each time during the process of traversing the convolution kernel elements are respectively located at Figure 6 the positions 0 to 8 marked in Figure 6 . Specifically, assuming row-major traversal, that is, traversing in the order of (0, 0) → (0,1) → (0,2) → (1,0) → … → (2,2), if the coordinates of the current convolution kernel element block are (0, 0), then the current data block to be loaded is Figure 6 the 6×6 data block corresponding to the starting address 0 in

[0043] Similarly, assuming column-major traversal, that is, traversing in the order of (0, 0) → (1,0) → (2,0) → (0,1) →… → (2,2), if the coordinates of the current convolution kernel element block are (0, 0), then the current data block to be loaded is Figure 6 the 6×6 data block corresponding to the starting address 0 in Figure 6 . By analogy. It can be seen that when traversing the convolution kernel elements in the order of column-major traversal, it slides along the row direction to determine the data block to be loaded corresponding to each convolution kernel element block.

[0044] Step 420: Based on the current data block to be loaded, determine the data block to be overwritten and the new data block, and load the new data block into the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory.

[0045] It should be noted that the core objective of step 420 is to optimize the utilization rate of shared memory through a dynamic coverage mechanism, reduce duplicate data loading, and improve the efficiency of convolution operations. This step specifically includes determining the data blocks to be covered and the newly added data blocks, and only loading the newly added data blocks into the shared memory to cover the data blocks to be covered.

[0046] Specifically, the data blocks to be covered refer to the data blocks that have been stored in the shared memory but are no longer needed in the current calculation step. They are used to mark the areas in the shared memory that can be covered by new data, avoiding redundant retention. The newly added data blocks are the data blocks that must be loaded in the current calculation step but have not been in the shared memory yet. They are used to supplement the new data required for the calculation and replace the data blocks to be covered. For example, assume that sub-blocks A, B, and C of the input data are stored in the shared memory, and the current calculation requires sub-blocks B, C, and D of the input data. Then, A can be determined as the data block to be covered, and D as the newly added data block. It is necessary to load data block D to the position of data block A in the shared memory to cover data block A.

[0047] It can be understood that since the traversal order of the convolution kernel elements is determined and the data in the input data is arranged in an orderly manner, during the traversal of the convolution kernel elements, for any two adjacent convolution kernel element blocks, the changes between their corresponding data blocks to be loaded are known. Therefore, for each current convolution kernel element block, based on its corresponding current data block to be loaded, the data blocks to be covered and the newly added data blocks can be directly determined without comparing the current data block to be loaded with the data blocks stored in the shared memory, which can simplify the process and improve the calculation efficiency.

[0048] After determining the data blocks to be covered and the newly added data blocks, first, according to the index of the data block to be covered, locate its address in the shared memory, and then directly write the newly added data block to the address where the data block to be covered is located in the shared memory, which can avoid additional memory operations.

[0049] Figure 7 is one of the schematic diagrams of data block loading provided by the present invention. As Figure 7 shown, in the column of original data on the left, the blue area represents the valid area of the input data, the red dashed box represents the data blocks that theoretically need to be loaded each time, and the white area represents the loaded padding area. The data blocks on the right are the data stored in the shared memory each time. Taking the traversal order of (0, 0) → (1,0) → (2,0) → (0,1) → … → (2,2) as an example, first, the coordinates of the convolution kernel element block traversed are r = 0, s = 0, and directly load the entire corresponding data block to be loaded (i.e., Figure 7Load the area represented by the red dashed box corresponding to r = 0 and s = 0 into the shared memory. Here, since the convolution kernel element block with r = 0 and s = 0 is the first convolution kernel element block traversed, in this case, the entire data block to be loaded corresponding to it is the newly added data block, and the data block to be overwritten is empty. It should be understood that the data block to be loaded is usually loaded from the global memory into the shared memory. To improve the data loading efficiency and reduce the number of global memory accesses, the data in the data block to be loaded can be loaded from the global memory into the shared memory in whole-block manner, or read and written into the shared memory in whole-row or whole-column manner. For example, when loading the data block to be loaded for the first time, the whole-block loading method can be adopted, and when loading the newly added data block subsequently, it can be loaded in whole-row or whole-column manner. Specifically, when traversing the convolution kernel elements in column-major order, since it slides along the row direction to determine the data block to be loaded corresponding to each convolution kernel element block, the data in the newly added data block can be loaded in whole-row manner; when traversing the convolution kernel elements in row-major order, since it slides along the column direction to determine the data block to be loaded corresponding to each convolution kernel element block, the data in the newly added data block can be loaded in whole-column manner.

[0050] Specifically, Figure 7 The traversal order shown in is column-major traversal, so the data can be loaded in whole-row manner. Before loading the data, indexes can be constructed for the padded input data in the row dimension and column dimension for subsequent data loading. For example, the row dimension index of the first row of data can be 0, the row dimension index of the second row of data is 1, and so on; similarly, the column dimension index of the first column of data can be 0, the column dimension index of the second column of data is 1, and so on. Figure 7 The numbers shown on the left side of a column of the original data in are the row dimension indexes corresponding to each row of data. For the convolution kernel element block with r = 0 and s = 0, since the shape of the target data block is 6×6, therefore, the range of the row dimension indexes of the data block to be loaded corresponding to this convolution kernel element block is 0 to 5. Since this data block to be loaded is the data block loaded for the first time, this data block to be loaded can be loaded from the global memory into the shared memory in whole-block manner. Figure 7 The first data block shown on the right side in is the data block loaded into the shared memory corresponding to the convolution kernel element block with r = 0 and s = 0, and matrix multiplication and accumulation operations will be performed between the two subsequently.

[0051] After completing the calculation of the convolution kernel element block with r = 0 and s = 0, then traverse the convolution kernel element block with coordinates r = 1 and s = 0. Since it slides in the row direction and the stride = 1, therefore, according to the current data block to be loaded (i.e., Figure 7(in the area represented by the red dashed box corresponding to r = 1, s = 0), the data in the 0th row (i.e., the data with a row dimension index of 0) can be directly determined as the data block to be overwritten, and the data in the 6th row (i.e., the data with a row dimension index of 6) as the newly added data block. Subsequently, in the way of loading the whole row, the newly added data block can be written into the 0th row of the shared memory, that is, Figure 7 the green data row corresponding to r = 1, s = 0 in overwrites the data in the 0th row of the shared memory to complete the loading of the newly added data block and overwrite the data block to be overwritten in the shared memory.

[0052] Similarly, when performing the matrix multiply-accumulate calculation of the convolution kernel element block corresponding to r = 2, s = 0, the same strategy as above is adopted for data loading, and it is only necessary to use the data of a new row (i.e., Figure 7 the data row marked in purple in to overwrite the data in the 1st row of the shared memory. It should be understood that the above data loading strategy also applies to sliding in the column direction, that is, it also applies to the scenario of row-major traversal.

[0053] In the embodiment of the present invention, by only loading the newly added data block into the shared memory and overwriting the data block to be overwritten, this overwriting mechanism can maximize the reuse of the loaded data, avoid repeated data loading, thereby reducing global memory access and saving memory bandwidth. In addition, due to reducing the overhead of data loading and storage, this operation mode can utilize computing resources more efficiently and improve the overall efficiency of convolution operations.

[0054] Step 430, based on the starting position, read the input data block from the shared memory, apply the input data block and the current convolution kernel element block, perform matrix multiply-accumulate operation, and accumulate the operation result to the output result. The starting position is determined based on the coordinates of the current convolution kernel element block.

[0055] Specifically, for each convolution kernel element block, after loading the corresponding newly added data block into the shared memory, the starting position for data reading can be determined according to the coordinates of the current convolution kernel element block, and then the input data block is read from the shared memory into the register according to this starting position for subsequent matrix multiply-accumulate (MMA) operation with the current convolution kernel element block. Here, the starting position refers to the starting address for reading the input data block from the shared memory, which is used to ensure the correct alignment of the read input data block with the convolution kernel element block to avoid invalid calculations or data misalignment.

[0056] Specifically, as Figure 7As shown, when traversing the convolution kernel elements in column - first order, the data is stored row - continuously in the shared memory. For the convolution kernel element block with coordinates r = 0, s = 0, since the data block to be loaded is directly stored into the shared memory row - by - row, its starting position in the shared memory can be determined as shm[0][0][0]. Here, the 0 in the first dimension represents the index in the row dimension (corresponding to the height dimension H), the 0 in the second dimension represents the index in the column dimension (corresponding to the width dimension W), and the 0 in the third dimension represents the index in the channel dimension C. When the channel dimension C = 1, the starting position can be simplified as shm[0][0].

[0057] After determining the starting position of data reading, the input data block can be read from the shared memory according to this starting position. For the convolution kernel element block with r = 0, s = 0, since the starting position of its corresponding input data block in the shared memory is shm[0][0], and the data is stored row - by - row in the shared memory, the corresponding data can be continuously read row - by - row from this starting position using the matrix load instruction (i.e., the ldmatrix instruction).

[0058] For the convolution kernel element block with coordinates r = 1, s = 0, since the new data block in the 6th row is written to the 0th row of the shared memory, its starting position in the shared memory can be determined as shm[1][0]. Subsequently, according to this starting address, the input data block can be read from the shared memory to the register using the strategy of circular row - by - row reading. Specifically, the data in the 1st row of the shared memory can be read first, and then the data in the 2nd, 3rd, 4th, 5th rows can be read in sequence, and finally the data in the 0th row can be read. Here, circular reading is a data access strategy that can regard the logical address space of the shared memory as a circular buffer with the head and tail connected through pointer wrapping. When accessing beyond the memory boundary, the address automatically wraps around to the starting position.

[0059] Similarly, for the convolution kernel element block with coordinates r = 2, s = 0, the corresponding input data block can be read from the shared memory using the above - mentioned circular reading strategy, which will not be elaborated here.

[0060] For each convolution kernel element block, after loading its corresponding input data block from the shared memory to the register, the MMA operation can be performed on the two. Taking the example in Figure 7 as an example, assuming that both the input data and the convolution kernel are single - channel, i.e., the channel dimension C = 1, each input data block loaded from the shared memory to the register is represented as Ai, and each convolution kernel element block corresponding to the input data block is represented as B[r][s] (if the channel dimension C is not 1, it can be represented as B[r][s][c]), and the matrix multiplication Ci = Ai Just use B[r][s]. Here, since the input data block Ai is a 6×6 (single-channel, equivalent to 6×6×1) matrix, and the convolutional kernel element block B[r][s] is a 1×1 (single-channel, equivalent to 1×1×1) weight block, B[r][s] can be regarded as a scalar and multiplied by each element in the input data block Ai. The result Ci is a 6×6 matrix of the same size as Ai, where each element is the product of the corresponding element in Ai and B[r][s]. It should be understood that before performing the MMA operation on each convolutional kernel element block, the convolutional kernel element block needs to be loaded from global memory to shared memory so that dedicated hardware such as TensorCore can read data from the shared memory for the MMA operation. If the Tensor Core is not used for the MMA operation subsequently and other computing units are used to implement it, the convolutional kernel element block needs to be loaded from the shared memory to the register so that the corresponding computing unit can read data from it for the operation. It should be noted that the storage areas of the convolutional kernel element block and the input data block in the shared memory are different, and the two are loaded into different registers for operation.

[0061] After each execution of the MMA operation and obtaining the corresponding operation result, the operation result can be accumulated on the dedicated accumulator to obtain the corresponding output result. For example, if the currently obtained operation result is C1, then each element on the C1 matrix is accumulated with the corresponding element on the C0 matrix to obtain the current output result. Here, the output result is the temporary storage in the accumulator. After all (r, s) traversals are completed, the accumulator result (i.e., the finally accumulated output result) is output to obtain the convolution operation result.

[0062] Step 440, determine the convolution operation result based on the output result accumulated after the traversal is completed.

[0063] Specifically, the output result accumulated after the traversal is completed refers to the final accumulated result obtained after all convolutional kernel element blocks (such as the 9 weight blocks of a 3×3 convolutional kernel) in Step 430 have been traversed and the corresponding MMA operations have been completed. The value at each output position in this output result is the accumulated sum of all relevant input data blocks and convolutional kernel element block operations. This output result is the convolution operation result. Here, the convolution operation result refers to the final output of the convolution operation of the input data and the convolutional kernel.

[0064] The method provided by the embodiments of the present invention can dynamically determine the currently to-be-loaded data block according to the coordinates of the current convolutional kernel element block by traversing the convolutional kernel elements. Based on the currently to-be-loaded data block, the newly added data block can be directly determined, and only the newly added data block is loaded into the shared memory instead of loading the entire currently to-be-loaded data block. Thus, the reuse of part of the data in the shared memory can be achieved, thereby avoiding repeated data loading, reducing the amount of data loaded each time, and further saving the memory bandwidth. When loading the newly added data block into the shared memory, it is loaded into the position of the data block to be overwritten in the shared memory to overwrite the data block to be overwritten. This mechanism ensures the efficient use of the shared memory and avoids unnecessary memory occupation and frequent data movement. In addition, compared with the traditional im2col method, the present invention does not need to first convert the input data into a matrix for storage. Therefore, there is no need to consume additional memory space to store the generated matrix, thereby reducing the occupation of memory space.

[0065] Based on the above embodiments, the method further includes: Step 450, determine the next to-be-loaded data block based on the target data block shape and the coordinates of the next convolutional kernel element block; Step 460, determine the newly added data block based on the next to-be-loaded data block, and preload the newly added data block from the global memory to the cache. The preloading operation of the newly added data block is executed in parallel with the matrix multiply-accumulate operation of the current convolutional kernel element block.

[0066] It should be noted that in traditional convolution operations, usually the currently required data is loaded from the global memory to the shared memory first, and then the data is loaded from the shared memory to the register for operation, that is, the data loading and data calculation are executed serially. However, when the time required to load data from the global memory to the shared memory is relatively long compared to the time required for data calculation, it will cause the computing unit to be idle waiting for data, resulting in low utilization.

[0067] In order to improve the utilization rate of the computing unit and further improve the convolution operation efficiency, in the embodiments of the present invention, while performing MMA operations on the current convolutional kernel element block and its corresponding input data block, the newly added data block corresponding to the next convolutional kernel element block is preloaded from the global memory to the cache (such as the L2 cache), realizing the overlap of calculation and data transfer, so that the data loading time is masked by the calculation task, thereby improving the utilization rate of the computing unit.

[0068] In addition, for the next convolution kernel element block, by preloading the corresponding newly added data block into the L2 cache, when performing MMA operations on the next convolution kernel element block, the corresponding newly added data block can be directly loaded from the L2 cache into the shared memory. Since the loading speed from the L2 cache to the shared memory is faster than that from the global memory to the shared memory, the data loading speed can be increased, thereby further improving the overall efficiency of the convolution operation.

[0069] Specifically, when traversing the convolution kernel elements, since the traversal order is known, the coordinates of the next convolution kernel element block can be determined according to the coordinates of the current convolution kernel element block and the traversal order. For example, as Figure 5 shown, for a 3×3 convolution kernel, using a 1×1 block division method and traversing in column-major order, if the coordinates of the current convolution kernel element block are (0, 0), the coordinates of the next convolution kernel element block are (1, 0); if the coordinates of the current convolution kernel element block are (1, 0), the coordinates of the next convolution kernel element block are (2, 0), and so on.

[0070] After determining the coordinates of the next convolution kernel element block, according to the coordinates and the shape of the target data block, the next data block to be loaded can be determined, and based on this data block to be loaded, the corresponding newly added data block can be determined. Here, the specific determination steps of the next data block to be loaded and the newly added data block can refer to the determination steps of the current data block to be loaded and its corresponding newly added data block in the above embodiments, which will not be elaborated here.

[0071] For the next convolution kernel element block, after determining its corresponding newly added data block, the newly added data block can be preloaded from the global memory into the L2 cache to enable the parallel execution of the preloading operation of the newly added data block and the MMA operation of the current convolution kernel element block.

[0072] Figure 8 is a schematic diagram of the parallel execution of data preloading and matrix multiply-accumulate operations provided by the present invention. As Figure 8As shown, for the convolutional kernel element block 1, after its corresponding newly added data block is loaded from the global memory to the shared memory, the corresponding input data block can be read from the shared memory into the register for MMA operations. While performing MMA operations on the convolutional kernel element block 1 and its corresponding input data block, steps 450 and 460 are executed to determine the newly added data block corresponding to the next convolutional kernel element block (i.e., convolutional kernel element block 2), and pre-load this newly added data block into the L2 cache. When traversing to convolutional kernel element block 2, its corresponding newly added data block can be loaded from the L2 cache into the shared memory, and then the corresponding input data block is read from the shared memory into the register to perform MMA operations on convolutional kernel element block 2 and this input data block. At the same time, pre-load the newly added data block corresponding to the next convolutional kernel element block into the L2 cache for data loading and MMA operations when traversing to this convolutional kernel element block later, and so on.

[0073] Based on any of the above embodiments, the coordinates of the next convolutional kernel element block are determined based on the traversal order of the convolutional kernel elements and the coordinates of the current convolutional kernel element block, and the traversal order is any one of row-major traversal and column-major traversal.

[0074] Specifically, when traversing the convolutional kernel elements, since the traversal order is known, the coordinates of the next convolutional kernel element block can be determined according to the coordinates of the current convolutional kernel element block and the traversal order of the convolutional kernel elements. Here, the traversal order of the convolutional kernel elements can be row-major traversal or column-major traversal, which can be specifically determined according to actual needs, and the embodiments of the present invention do not make specific limitations on this.

[0075] Exemplarily, as Figure 5 shown, for a 3×3 convolutional kernel, using a 1×1 block division method and traversing in column-major order, if the coordinates of the current convolutional kernel element block are (0, 0), then the coordinates of the next convolutional kernel element block are (1, 0); if the coordinates of the current convolutional kernel element block are (1, 0), then the coordinates of the next convolutional kernel element block are (2, 0), and so on.

[0076] Based on any of the above embodiments, in step 460, determining the newly added data block based on the next data block to be loaded includes: Step 461, in the case where both the row dimension index and the column dimension index of the next convolutional kernel element block change relative to the current convolutional kernel element block, using the data block to be loaded as the newly added data block; Step 462, in the case where either the row dimension index or the column dimension index of the next convolutional kernel element block changes relative to the current convolutional kernel element block, determining the newly added data block from the next data block to be loaded.

[0077] It should be noted that when the new data block is loaded into the shared memory using the overwrite mechanism, considering that when the row dimension index and column dimension index of the next convolutional kernel element block change relative to the current convolutional kernel element block, if only the new data block is still loaded into the shared memory, it may make the processes of data loading and subsequent data reading from the shared memory complex. To simplify the processes of data loading, overwriting, and subsequent data reading, in the embodiments of the present invention, the new data block is loaded into the shared memory only when the row dimension index changes or the column dimension index changes, and when both the row dimension index and column dimension index change, the entire data block to be loaded is loaded into the shared memory.

[0078] Specifically, for the next convolutional kernel element block, when determining its corresponding new data block, if both the row dimension index and column dimension index of the next convolutional kernel element block change relative to the current convolutional kernel element block, in this case, the entire data block to be loaded corresponding to the next convolutional kernel element block can be used as the new data block and pre-loaded into the L2 cache, so as to overwrite the entire input data block stored in the shared memory subsequently. If the row dimension index of the next convolutional kernel element block does not change or the column dimension index does not change relative to the current convolutional kernel element block, the new data block can be determined according to its corresponding data block to be loaded, and only the new data block is loaded into the L2 cache, so as to overwrite a certain row or a certain column of the data block to be overwritten in the shared memory subsequently.

[0079] Figure 9 is the second schematic diagram of data block loading provided by the present invention, as Figure 9 shown, when traversing the convolutional kernel elements in column-major order, that is, in the order of (0, 0) → (1,0) → (2,0) → (0,1) → … → (2,2), assuming that the coordinates of the convolutional kernel element block currently undergoing MMA operation are r = 2, s = 0, then the coordinates of the next convolutional kernel element block are r = 0, s = 1. Since the row dimension index of the next convolutional kernel element block changes from r = 2 to r = 0 and the column dimension index also changes from s = 0 to s = 1 relative to the current convolutional kernel element block, therefore, the entire data block to be loaded corresponding to the next convolutional kernel element block can be used as the new data block and pre-loaded into the L2 cache. When the convolutional kernel element block with coordinates r = 0, s = 1 is traversed subsequently, the new data block can be loaded from the L2 cache into the shared memory.

[0080] When performing MMA operations on the convolution kernel element block with r = 0 and s = 1, since the coordinates of the next convolution kernel element block are r = 1 and s = 1, relative to the current convolution kernel element block, only the row dimension index changes from r = 0 to r = 1, while the column dimension index remains unchanged. Therefore, the newly added data block (i.e., Figure 9 the green data row corresponding to r = 1 and s = 1 in

[0081] Based on any of the above embodiments, as Figure 9 shown, when the coordinates of the traversed convolution kernel element block are r = 0 and s = 1, the newly added data block in the L2 cache (i.e., the overall data block to be loaded corresponding to the convolution kernel element block with r = 0 and s = 1) needs to be loaded into the shared memory to overwrite the entire input data block stored in the shared memory. Although this data loading method can simplify the data loading and overwriting processes, there is still a problem of duplicate data loading, and the data stored in the shared memory is not reused. In response, the embodiments of the present invention provide another data loading method to avoid duplicate data loading.

[0082] Figure 10 FIG. Figure 10 is a third schematic diagram of data block loading provided by the present invention. As Figure 10 shown, for the convolution kernel element block with coordinates r = 0 and s = 1, its newly added data blocks are the data row with row dimension index 0 (i.e., Figure 10 the yellow data row marked in Figure 10 ), the data row with row dimension index 1 (i.e.,

[0083] the red data row marked in

[0084] Figure 10 ), and the data column with column dimension index 6 (i.e., Figure 10 the gray data column marked in Figure 10 ), while the data block to be overwritten is the data in the 0th row, 1st row, and 0th column in the shared memory. Subsequently, when loading the newly added data block, the newly added data block can be written to the position of the data block to be overwritten according to the position of the data block to be overwritten in the shared memory, thereby realizing the reuse of part of the data in the shared memory and avoiding duplicate data loading.

[0083] It can be understood that when subsequently reading the input data block from the shared memory to the register for MMA operations, for the convenience of reading, the data in the shared memory can be read element by element in the index order of each data element in the original input data. For example, for the data in the 0th row of the shared memory, it can be read in the order of indices 1 to 6.

[0084] Further, when the data volume is large, considering that reading element by element will reduce the data loading efficiency, to improve the loading efficiency, when both the row dimension index and the column dimension index of the convolutional kernel element block change relative to the previous convolutional kernel element block, a certain shift operation can be performed on the data in the shared memory when loading the new elements corresponding to the convolutional kernel element block into the shared memory. For example, as Figure 10 shown, for the convolutional kernel element block with r = 0 and s = 1, when loading the new data block into the shared memory, the yellow data row (except the last element) can be loaded into the 0th row of the shared memory first, the red data row (except the last element) can be loaded into the 1st row of the shared memory, then each column of data in the shared memory is shifted one column to the left, and finally the gray data column is loaded into the last column of the shared memory. After the data loading is completed, the input data block can be read from the shared memory row by row, thereby improving the data reading efficiency.

[0085] In addition, to address the problem of duplicate data loading existing in Figure 9 , the present invention also provides a new traversal order to solve the problem of duplicate data loading, and this new traversal order is called serpentine column-first traversal. It should be noted that the traditional column-first traversal means traversing column by column from top to bottom, and entering the next column after one column is completed. Different from the traditional column-first traversal, serpentine column-first traversal means that the column direction is the main direction (i.e., advancing column by column as a whole), but within each column, it alternates from top to bottom and from bottom to top (similar to a serpentine winding). For example, as Figure 5 shown, the traditional column-first traversal order is (0,0) → (1,0) → (2,0) → (0,1) → (1,1) → (2,1) → (0,2) → (1,2) → (2,2), while the serpentine column-first traversal order proposed in the embodiments of the present invention is (0,0) → (1,0) → (2,0) → (2,1) → (1,1) → (0,1) → (0,2) → (1,2) → (2,2).

[0086] Figure 11 is the fourth schematic diagram of data block loading provided by the present invention. As Figure 11 shown, when traversing the convolutional kernel elements in the order of serpentine column-first traversal, assuming the coordinates of the current convolutional kernel element block are r = 2 and s = 0, then the coordinates of the next convolutional kernel element block are r = 2 and s = 1. Since the row dimension index of the next convolutional kernel element block does not change relative to the current convolutional kernel element block, only the column dimension index changes, the new data block can be determined according to the data block to be loaded corresponding to the next convolutional kernel element block, and only the new data block is loaded. As can be seen from Figure 11 , for the convolutional kernel element block with coordinates r = 2 and s = 1, its new data block is only the data column with column dimension index 6 (i.e., Figure 11The gray data column marked in [Figure 0], and the data block to be overwritten is the data in the 0th column of the shared memory. Subsequently, when loading the newly added data block, the newly added data block can be written to the position of the data block to be overwritten according to the position of the data block to be overwritten in the shared memory, so as to realize the reuse of part of the data in the shared memory and avoid repeated data loading.

[0087] It can be understood that, as Figure 11 shown, assuming that the indexes of each data element in the newly added data block corresponding to the convolution kernel element block with r = 2 and s = 1 are 1 to 6 respectively. After loading the newly added data block into the shared memory, its index distribution in the shared memory is specifically as Figure 11 shown. Subsequently, when reading the input data block from the shared memory to the register for MMA operation, it can be realized by circularly reading by column in the column direction and adopting the strategy of circularly reading by row from the starting position in the row direction.

[0088] Similarly, for the problem of repeated data loading existing in row-major traversal, the snake-shaped row-major traversal order can also be adopted to solve it, which will not be elaborated here.

[0089] Based on any of the above embodiments, step 420 specifically includes: Step 421, when the current convolution kernel element block is the first convolution kernel element block, determine that the data block to be overwritten is empty, and load the current data block to be loaded as the newly added data block from the global memory into the shared memory; Step 422, when the current convolution kernel element block is not the first convolution kernel element block, based on the current data block to be loaded, determine the data block to be overwritten, and read the newly added data block from the cache to load the newly added data block from the cache into the shared memory.

[0090] Specifically, if the current convolution kernel element block is the first convolution kernel element block (i.e., the first convolution kernel element block traversed), it can be determined that the corresponding data block to be overwritten is empty, that is, there is no data in the shared memory by default in the initial state, and there is no need to overwrite. The current data block to be loaded can be directly used as the newly added data block and loaded from the global memory into the shared memory. Thus, the initialization of the shared memory is realized, providing the first input data block for subsequent calculations.

[0091] If the current convolution kernel element block is not the first convolution kernel element block (such as the second convolution kernel element block, the third convolution kernel element block, etc.), at this time, the input data block corresponding to the first convolution kernel element block has been stored in the shared memory. Therefore, it is necessary to determine the data block to be overwritten according to the current data block to be loaded, and load the newly added data block from the cache (such as L2 cache) into the shared memory to overwrite to the position of the data block to be overwritten.

[0092] Based on any of the above embodiments, in step 420, loading the newly added data block into the shared memory to overwrite the data block to be overwritten in the shared memory includes: Based on the row dimension index or column dimension index of the data block to be overwritten, determine the storage location of the newly added data block in the shared memory, and load the newly added data block to the storage location.

[0093] Specifically, the row dimension index or column dimension index of the data block to be overwritten refers to the position identifier of the row or column dimension of the data block to be overwritten in the shared memory, which is used to mark which memory areas can be overwritten by new data. The storage location of the newly added data block refers to the starting address and layout of the newly added data block in the shared memory, and needs to be aligned with the dimension of the data block to be overwritten to ensure data continuity.

[0094] Specifically, when loading data by row reading, the storage location of the newly added data block in the shared memory can be determined according to the row dimension index of the data block to be overwritten. For example, if the row dimension index of the data block to be overwritten is 0, the newly added data block needs to be written to row 0 of the shared memory. When loading data by column reading, the storage location of the newly added data block in the shared memory can be determined according to the column dimension index of the data block to be overwritten. For example, if the column dimension index of the data block to be overwritten is 1, the newly added data block needs to be written to column 1 of the shared memory.

[0095] Based on any of the above embodiments, the shared memory is a circular buffer. Correspondingly, in step 430, reading the input data block from the shared memory based on the starting position includes: Based on the starting position, apply the matrix loading instruction to circularly read the input data block from the shared memory by row or by column.

[0096] Specifically, the physical space of the shared memory is a limited linear address block, but through pointer wrapping or modulo operation, it can be logically regarded as "circular", that is, when accessing beyond the memory boundary, the address automatically wraps around to the starting position. Therefore, the shared memory can be regarded as a circular buffer. The starting position for reading data from this circular buffer is a dynamic value. When reading data by row, for different r values (i.e., row indices), the starting position can be expressed as shm[r][0][0]; when reading data by column, for different s values (i.e., column indices), the starting position can be expressed as shm[0][s][0].

[0097] When circularly reading data from the shared memory by row or by column according to the starting position, it can be implemented by using the matrix loading instruction. Here, the matrix loading instruction is a low-level operation for efficiently loading matrix data, which is used to optimize the loading process of matrix data from memory to registers to support high-performance matrix calculations.

[0098] Based on any of the above embodiments, an embodiment of the present invention provides an optimization method for convolution operations, and the method includes: Step S1, according to the shape of the input data (H, W, C) and the convolution parameters (including the convolution kernel size R, S, and padding, stride, dilation rate, etc.), determine the shape of the input data block to be loaded each time, that is, determine the target data block shape (P, Q, C).

[0099] Step S2, traverse the convolution kernel elements, and according to the coordinates of the current traversed convolution kernel element block and the target data block shape, determine the starting positions of the current input data block to be loaded and the shared memory.

[0100] Step S3, if the current traversed convolution kernel element block is the first convolution kernel element block, then load the current input data block to be loaded from the global memory to the shared memory, and then according to the starting position of the shared memory, read the input data block from the shared memory and load it into the register, and at the same time load the current convolution kernel element block into the shared memory or another register (specifically depending on the subsequent operation method) to perform MMA operations on the convolution kernel element block and its corresponding input data block, and accumulate the operation results to the dedicated accumulator. At the same time, according to the coordinates of the next convolution kernel element block and the target data block shape, determine the next input data block to be loaded, and accordingly determine its corresponding new data block, and then preload the new data block from the global memory to the L2 cache to realize parallel execution of data preloading and MMA operations.

[0101] Step S4, if the current traversed convolution kernel element block is not the first convolution kernel element block, such as the second convolution kernel element block, the third convolution kernel element block, etc., then determine the data block to be overwritten according to the current input data block to be loaded, read the new data block from the L2 cache, and then load the new data block from the L2 cache to the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory. The subsequent MMA operation process is similar to that described in Step S3. By analogy, after completing the traversal of all convolution kernel element blocks, the final accumulator result is the convolution operation result.

[0102] The convolution operation device provided by the present invention will be described below. The convolution operation device described below can be correspondingly referred to the convolution operation method described above.

[0103] Based on any of the above embodiments, Figure 12 is a schematic structural diagram of the convolution operation device provided by the present invention, as Figure 12 shown, the device includes: A data block determination unit 1210 is configured to traverse the convolutional kernel elements, and determine a current data block to be loaded based on the target data block shape and the coordinates of the currently traversed convolutional kernel element block, where the target data block shape is determined based on the convolutional parameters and the input data shape; A data loading unit 1220 is configured to determine a data block to be overwritten and a new data block based on the current data block to be loaded, and load the new data block into the shared memory, so that the new data block overwrites the data block to be overwritten in the shared memory; A matrix operation unit 1230 is configured to read an input data block from the shared memory based on a starting position, apply the input data block and the current convolutional kernel element block, perform a matrix multiplication and accumulation operation, and accumulate the operation result to the output result, where the starting position is determined based on the coordinates of the current convolutional kernel element block; A result determination unit 1240 is configured to determine the convolutional operation result based on the output result accumulated after the traversal is completed.

[0104] The device provided by the embodiment of the present invention can dynamically determine the current data block to be loaded according to the coordinates of the current convolutional kernel element block by traversing the convolutional kernel elements. Based on the current data block to be loaded, the new data block can be directly determined, and only the new data block is loaded into the shared memory, rather than loading the entire current data block to be loaded. Thus, the reuse of part of the data in the shared memory can be achieved, thereby avoiding repeated data loading, reducing the amount of data loaded each time, and further saving the memory bandwidth. When loading the new data block into the shared memory, it is loaded into the position of the data block to be overwritten in the shared memory so that it overwrites the data block to be overwritten, which ensures the efficient utilization of the shared memory and avoids unnecessary memory occupation and frequent data movement. In addition, compared with the traditional im2col method, the present invention does not need to first convert the input data into a matrix for storage. Therefore, no additional memory space is consumed to store the generated matrix, thereby reducing the occupation of memory space.

[0105] Based on any of the above embodiments, the device further includes a preloading unit, and the preloading unit is configured to: determine the next data block to be loaded based on the target data block shape and the coordinates of the next convolutional kernel element block; determine a new data block based on the next data block to be loaded, and preload the new data block from the global memory to the cache, and the preloading operation of the new data block is executed in parallel with the matrix multiplication and accumulation operation of the current convolutional kernel element block.

[0106] Based on any of the above embodiments, the coordinates of the next convolutional kernel element block are determined based on the traversal order of the convolutional kernel elements and the coordinates of the current convolutional kernel element block, and the traversal order is any one of row-first traversal and column-first traversal.

[0107] Based on any of the above embodiments, the preloading unit is specifically configured to: when both the row dimension index and the column dimension index of the next convolutional kernel element block change with respect to the current convolutional kernel element block, use the data block to be loaded as a new data block; when either the row dimension index or the column dimension index of the next convolutional kernel element block changes with respect to the current convolutional kernel element block, determine a new data block from the next data block to be loaded.

[0108] Based on any of the above embodiments, the data loading unit 1220 is specifically configured to: when the current convolutional kernel element block is the first convolutional kernel element block, determine that the data block to be overwritten is empty, and load the current data block to be loaded as a new data block from the global memory into the shared memory; when the current convolutional kernel element block is not the first convolutional kernel element block, determine the data block to be overwritten based on the current data block to be loaded, and read the new data block from the cache, so as to load the new data block from the cache into the shared memory.

[0109] Based on any of the above embodiments, the data loading unit 1220 is specifically configured to: determine the storage location of the new data block in the shared memory based on the row dimension index or the column dimension index of the data block to be overwritten, and load the new data block to the storage location.

[0110] Based on any of the above embodiments, the shared memory is a circular buffer. Accordingly, the matrix operation unit 1230 is specifically configured to: based on the starting position, apply a matrix loading instruction to circularly read the input data block from the shared memory row by row or column by column.

[0111] Figure 13 An entity structure diagram of an electronic device is exemplified. As Figure 13 shown, the electronic device may include: a processor 1310, a communication interface 1320, a memory 1330, and a communication bus 1340. Among them, the processor 1310, the communication interface 1320, and the memory 1330 complete communication with each other through the communication bus 1340. The processor 1310 may call logic instructions in the memory 1330 to execute a convolution operation method, which includes: traversing convolutional kernel elements, determining the current data block to be loaded based on the target data block shape and the coordinates of the currently traversed convolutional kernel element block, where the target data block shape is determined based on convolutional parameters and the input data shape; determining the data block to be overwritten and the new data block based on the current data block to be loaded, and loading the new data block into the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory; reading the input data block from the shared memory based on the starting position, applying the input data block and the current convolutional kernel element block to perform matrix multiply-accumulate operations, and accumulating the operation results to the output result, where the starting position is determined based on the coordinates of the current convolutional kernel element block; determining the convolution operation result based on the output result accumulated after the traversal is completed.

[0112] In addition, when the logical instructions in the above-mentioned memory 1330 are implemented in the form of software functional units and sold or used as independent products, they can be stored in a computer-readable storage medium. Based on such an understanding, the technical solution of the present invention, in essence, or the part that contributes to the related technology, or a part of the technical solution, can be embodied in the form of a software product. The computer software product is stored in a storage medium and includes several instructions for causing a computer device (which may be a personal computer, a server, or a network device, etc.) to execute all or part of the steps of the methods in various embodiments of the present invention. The aforementioned storage medium includes: various media such as USB flash drives, mobile hard disks, read-only memories (ROMs), random access memories (RAMs), magnetic disks, or optical discs that can store program codes.

[0113] On the other hand, the present invention also provides a computer program product. The computer program product includes a computer program that can be stored on a non-transitory computer-readable storage medium. When the computer program is executed by a processor, the computer can execute the convolution operation method provided by the above-mentioned various methods. The method includes: traversing the convolution kernel elements, determining the current data block to be loaded based on the target data block shape and the coordinates of the currently traversed convolution kernel element block, where the target data block shape is determined based on the convolution parameters and the input data shape; determining the data block to be overwritten and the new data block based on the current data block to be loaded, and loading the new data block into the shared memory so that the new data block overwrites the data block to be overwritten in the shared memory; reading the input data block from the shared memory based on the starting position, applying the input data block and the current convolution kernel element block, performing matrix multiplication and accumulation operations, and accumulating the operation results onto the output result, where the starting position is determined based on the coordinates of the current convolution kernel element block; determining the convolution operation result based on the output result accumulated after the traversal is completed.

[0114] In another aspect, the present invention also provides a non-transitory computer-readable storage medium, on which a computer program is stored. When the computer program is executed by a processor, it implements the convolution operation method provided by the above-mentioned various methods. The method includes: traversing the convolution kernel elements, determining the current data block to be loaded based on the target data block shape and the coordinates of the currently traversed convolution kernel element block, where the target data block shape is determined based on the convolution parameters and the input data shape; determining the data block to be overwritten and the newly added data block based on the current data block to be loaded, and loading the newly added data block into the shared memory so that the newly added data block overwrites the data block to be overwritten in the shared memory; reading the input data block from the shared memory based on the starting position, applying the input data block and the current convolution kernel element block, performing matrix multiplication and accumulation operations, and accumulating the operation results to the output result, where the starting position is determined based on the coordinates of the current convolution kernel element block; determining the convolution operation result based on the output result accumulated after the traversal is completed.

[0115] The device embodiments described above are merely illustrative. The units described as separate components may or may not be physically separated, and the components shown as units may or may not be physical units, that is, they may be located in one place or distributed to multiple network units. Some or all of the modules can be selected according to actual needs to achieve the purpose of the solution of this embodiment. Those of ordinary skill in the art can understand and implement it without creative efforts.

[0116] Through the description of the above embodiments, those skilled in the art can clearly understand that each embodiment can be implemented by means of software plus a necessary general hardware platform, and of course, it can also be implemented by hardware. Based on this understanding, the essence of the above technical solution, or the part that contributes to the related technology, can be embodied in the form of a software product. The computer software product can be stored in a computer-readable storage medium, such as ROM / RAM, magnetic disk, optical disc, etc., and includes several instructions for causing a computer device (which can be a personal computer, a server, or a network device, etc.) to execute the methods of each embodiment or some parts of the embodiments.

[0117] Finally, it should be noted that: the above embodiments are only used to illustrate the technical solutions of the present invention, rather than to limit them; although the present invention has been described in detail with reference to the foregoing embodiments, those of ordinary skill in the art should understand that they can still modify the technical solutions described in the foregoing embodiments, or perform equivalent replacements on some of the technical features; and these modifications or replacements do not make the essence of the corresponding technical solutions deviate from the spirit and scope of the technical solutions of each embodiment of the present invention.

Claims

1. A convolution operation method, characterized in that, Including: Traverse the convolutional kernel elements, and determine the currently to-be-loaded data block based on the target data block shape and the coordinates of the currently traversed convolutional kernel element block, where the target data block shape is determined based on convolutional parameters and the input data shape; Based on the currently to-be-loaded data block, determine the data block to be overwritten and the newly added data block, and load the newly added data block into the shared memory so that the newly added data block overwrites the data block to be overwritten in the shared memory; Based on the starting position, read the input data block from the shared memory, apply the input data block and the currently traversed convolutional kernel element block, perform matrix multiplication and accumulation operations, and accumulate the operation results onto the output result, where the starting position is determined based on the coordinates of the currently traversed convolutional kernel element block; Determine the convolutional operation result based on the output result accumulated after the traversal is completed.

2. The convolution operation method according to claim 1, wherein Also including: Determine the next to-be-loaded data block based on the target data block shape and the coordinates of the next convolutional kernel element block; Based on the next to-be-loaded data block, determine the newly added data block, and pre-load the newly added data block from the global memory into the cache, where the pre-loading operation of the newly added data block is executed in parallel with the matrix multiplication and accumulation operation of the currently traversed convolutional kernel element block.

3. The convolution operation method according to claim 2, wherein The coordinates of the next convolutional kernel element block are determined based on the traversal order of the convolutional kernel elements and the coordinates of the currently traversed convolutional kernel element block, and the traversal order is any one of row-major traversal and column-major traversal.

4. The convolution operation method according to claim 2, wherein The determining the newly added data block based on the next to-be-loaded data block includes: In the case where both the row dimension index and the column dimension index of the next convolutional kernel element block change relative to the currently traversed convolutional kernel element block, use the to-be-loaded data block as the newly added data block; In the case where either the row dimension index or the column dimension index of the next convolutional kernel element block changes relative to the currently traversed convolutional kernel element block, determine the newly added data block from the next to-be-loaded data block.

5. The convolution operation method according to claim 1, wherein The determining the data block to be overwritten and the newly added data block based on the currently to-be-loaded data block, and loading the newly added data block into the shared memory includes: In the case where the currently traversed convolutional kernel element block is the first convolutional kernel element block, determine that the data block to be overwritten is empty, and load the currently to-be-loaded data block as the newly added data block from the global memory into the shared memory; In the case where the currently traversed convolutional kernel element block is not the first convolutional kernel element block, determine the data block to be overwritten based on the currently to-be-loaded data block, and read the newly added data block from the cache to load the newly added data block from the cache into the shared memory.

6. The convolution operation method according to any one of claims 1 to 5, characterized in that The loading the newly added data block into the shared memory so that the newly added data block overwrites the data block to be overwritten in the shared memory includes: Based on the row dimension index or the column dimension index of the data block to be overwritten, determine the storage position of the newly added data block in the shared memory, and load the newly added data block to the storage position.

7. The convolution operation method according to any one of claims 1 to 5, characterized in that, The shared memory is a circular buffer. Correspondingly, the reading the input data block from the shared memory based on the starting position includes: Based on the starting position, apply a matrix loading instruction to circularly read the input data block from the shared memory row by row or column by column.

8. An electronic device, comprising a memory, a processor, and a computer program stored on the memory and running on the processor, characterized in that When the processor executes the computer program, it implements the convolution operation method according to any one of claims 1 to 7.

9. A non-transitory computer-readable storage medium having a computer program stored thereon, characterized in that, When the computer program is executed by the processor, it implements the convolution operation method according to any one of claims 1 to 7.

10. A computer program product, comprising a computer program, characterized in that, When the computer program is executed by the processor, it implements the convolution operation method according to any one of claims 1 to 7.

Images