Convolution weight gradient calculation method and device, medium and equipment

Abstract

The invention discloses a convolution weight gradient calculation method and device, a medium and equipment. The method comprises the following steps: acquiring an mth gradient part and a corresponding mth feature part; loading all the feature data rows of the mth feature part to an SV sub-region of the shared memory, wherein the row numbers of the feature data rows in each sub-region form an arithmetic progression of which the tolerance is SV; sV is equal to the vertical step length; the threads access different sub-regions in turn and read continuous memory blocks as required; each thread performs sliding traversal and point multiplication operation on the gradient data rows and the feature data blocks obtained by loading, and accumulates each point multiplication result to a corresponding weight gradient position; and after traversal calculation of gradient data rows of all gradient parts in the output gradient map and corresponding feature data blocks is completed, obtaining a weight gradient matrix of the convolution kernel. According to the method, step access can be converted into continuous access, the loading amount of redundant data is remarkably reduced, and then the memory access efficiency is improved.

Description

Technical Field

[0001] This invention relates to the field of artificial intelligence technology, and in particular to a method, apparatus, computer-readable storage medium, and electronic device for calculating convolution weight gradients. Background Technology

[0002] In traditional convolutional weight gradient calculation methods, the weight gradient matrix is ​​obtained by performing convolution calculations using the input feature map stored during the forward computation and the output gradient map generated during the backpropagation. During weight gradient calculation, when the stride of the convolutional kernel is greater than 1, the data access mode is a strafing access mode. For AI processors with contiguous memory, to obtain sparsely distributed target data, threads need to load the entire contiguous block of memory data and then filter out the required target data. This approach results in the actual amount of data loaded by the thread far exceeding the effective computational requirements, leading not only to a significant decrease in memory access efficiency but also easily causing bandwidth bottlenecks. Summary of the Invention

[0003] The purpose of this invention is to provide a method, apparatus, computer-readable storage medium, and electronic device for calculating convolutional weight gradients, which divides the input feature map into S rows according to row numbers. V Group feature data and store them separately in S V Within each sub-region, the row numbers contained in each sub-region constitute a tolerance of S. V The arithmetic sequence; during the thread loading process, the thread takes turns accessing different sub-regions and reads consecutive memory blocks as needed, thereby converting step-by-step access into continuous access to reduce redundant data loading.

[0004] A first aspect of the present invention provides a method for calculating convolution weight gradients, applied to an artificial intelligence processor, the method comprising: The m-th gradient part in the output gradient map is loaded into shared memory, and the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map are determined based on the size of the convolution kernel, the vertical stride and the horizontal stride, so as to be used as the m-th feature part; where m≥1; Load all feature data rows of the m-th feature part into the shared memory S. V There are several sub-regions; where the row numbers of the feature data rows in the m-th feature part of each sub-region constitute a common difference of S. V An arithmetic sequence; there are no identical row numbers among the arithmetic sequences; S V It is equal to the vertical step size; N of the m-th gradient part T Row gradient data is loaded row by row into N T N threads are used to ensure that the i-th thread obtains the i-th row of gradient data; where NT This represents the total number of rows in the m-th gradient part; Visit S sequentially in a loop. V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T There are 10 threads, such that the i-th thread obtains the result of the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block consists of the feature data rows; where 1≤i≤N T ; In each thread, a sliding traversal and dot product operation are performed on the gradient data rows and feature data blocks loaded, and each dot product result is accumulated to the corresponding weight gradient position; wherein, the sliding window of the feature data block is K. h ×K w And the sliding step size is a horizontal step size; K h and K w These are the height and width of the convolution kernel, respectively; The gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient graph are traversed and calculated to obtain the weight gradient matrix of the convolution kernel.

[0005] Optionally, determining the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map based on the kernel size, vertical stride, and horizontal stride, to collectively serve as the m-th feature part, includes: When both the vertical step size and the horizontal step size are 1, the mapping region is directly determined based on the position region of the m-th gradient part in the output gradient map. When at least one of the vertical step size and the horizontal step size is greater than 1, the m-th gradient part is upsampled in the sliding direction greater than 1, and the mapping region is determined based on the upsampled result. Based on the size of the convolution kernel, the boundary extension region surrounding the mapping region is determined; The mapped region and the boundary extension region are taken as the m-th feature part.

[0006] Optionally, the step of loading all feature data rows of the m-th feature portion into shared memory S V Each sub-region includes: Starting with the j-th row of feature data, and reading feature data rows in the m-th feature part according to the vertical step size, we obtain the j-th data subset; Load the j-th data subset into the j-th sub-region of shared memory; where 1≤j≤S V .

[0007] Optionally, the sequential cyclic access to S V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T One thread, including: According to the first sub-region to the Sth sub-region V The sub-regions are visited sequentially, and the total number of visits is determined based on the height of the convolution kernel. During each access, the starting row position of the feature data to be loaded in the current sub-region is determined based on the current access count; Starting from the starting row position, select N consecutive rows in the current sub-region. T Row feature data is loaded into N rows. T One thread.

[0008] Optionally, the number of divisions of the gradient portion in the output gradient graph is determined by the computing resources of the artificial intelligence processor.

[0009] Optionally, the N T Each thread loads the corresponding gradient data rows and feature data blocks in parallel using a specified batch loading instruction.

[0010] Optionally, the N T Each thread releases cache locations through a register reuse mechanism to allow for the rolling loading of corresponding gradient data rows and feature data blocks.

[0011] Optionally, the step of performing sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded in each thread, and accumulating each dot product result to the corresponding weight gradient position, includes: The window of the feature data block within the thread is slid with the horizontal step size, while the window of the gradient data row is slid with a step size of 1; wherein, the sliding window of the gradient data row is 1×1. After each sliding step, the two window data involved in the current calculation are located, and a dot product operation is performed to obtain the corresponding dot product result; wherein, the size of the dot product result is K. h ×K w .

[0012] A second aspect of the present invention provides a method for calculating convolution weight gradients, applied to an artificial intelligence processor, the method comprising: Obtain the input feature tensor and the corresponding output gradient tensor; wherein, the input feature tensor consists of N... C The output gradient tensor consists of N input feature maps of each channel; C The output gradient map consists of N channels; C >1; For each channel, the convolution weight gradient calculation method described in any embodiment of the first aspect is executed to obtain the weight gradient matrix of the corresponding channel.

[0013] Optionally, the number of parallel computations of the weight gradient matrix is ​​determined by the thread resources enabled by the artificial intelligence processor.

[0014] A third aspect of the present invention provides a convolution weight gradient calculation apparatus, comprising: The first block module is used to load the m-th gradient part of the output gradient map into shared memory, and determine the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map based on the size of the convolution kernel, the vertical stride and the horizontal stride, so as to jointly serve as the m-th feature part; where m≥1; The second block module is used to load all feature data rows of the m-th feature part into the shared memory S. V There are several sub-regions; where the row numbers of the feature data rows in the m-th feature part of each sub-region constitute a common difference of S. V An arithmetic sequence; there are no identical row numbers among the arithmetic sequences; S V It is equal to the vertical step size; The first loading module is used to load N of the m-th gradient part. T Row gradient data is loaded row by row into N T N threads are used to ensure that the i-th thread obtains the i-th row of gradient data; where N T This represents the total number of rows in the m-th gradient part; The second loading module is used to sequentially and cyclically access S. V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T There are 10 threads, such that the i-th thread obtains the result of the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block consists of the feature data rows; where 1≤i≤N T ; The thread computation module performs sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded by each thread, and accumulates each dot product result to the corresponding weight gradient position; wherein, the sliding window of the feature data block is K. h ×K w And the sliding step size is a horizontal step size; K h and K w These are the height and width of the convolution kernel, respectively; The first result acquisition module is used to traverse and calculate the gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient map to obtain the weight gradient matrix of the convolution kernel.

[0015] A fourth aspect of the present invention provides a convolution weight gradient calculation apparatus, comprising: The data acquisition module is used to acquire the input feature tensor and the corresponding output gradient tensor; wherein, the input feature tensor consists of N... C The output gradient tensor consists of N input feature maps of each channel; C The output gradient map consists of N channels; C >1; The second result acquisition module is used to perform the convolution weight gradient calculation method described in any embodiment of the first aspect on the input feature map and output gradient map corresponding to each channel respectively, so as to obtain the weight gradient matrix of the corresponding channel.

[0016] A fifth aspect of the present invention provides a computer-readable storage medium comprising a stored computer program; wherein, when the computer program is executed, it controls the device on which the computer-readable storage medium is located to perform the convolution weight gradient calculation method described in any of the above embodiments.

[0017] A sixth aspect of the present invention provides an electronic device including a processor, a memory, and a computer program stored in the memory and configured to be executed by the processor, wherein the processor, when executing the computer program, implements the convolution weight gradient calculation method described in any of the above embodiments.

[0018] Compared with the prior art, the embodiments of the present invention provide a method, apparatus, computer-readable storage medium, and electronic device for calculating convolution weight gradients. The embodiments of the present invention obtain the m-th gradient part and the corresponding m-th feature part from the output gradient map and the input feature map, respectively, and directly store the m-th gradient part into shared memory; the m-th feature part is divided into S according to row numbers. V Group feature data and store them separately in S V Within each sub-region, the row numbers of the feature data rows within each sub-region form a common difference of S. VThe algorithm constructs an arithmetic sequence and arranges it continuously in physical memory to eliminate the effects of stepping. Then, threads take turns accessing different sub-regions and reading contiguous memory blocks as needed. During computation, each thread performs sliding traversal and dot product operations on the loaded gradient data rows and feature data blocks to obtain the dot product result at the current sliding position, which is then accumulated in the cache location of the weight gradient matrix. Once all gradient parts in the output gradient map and all feature parts in the input feature map have been traversed and calculated, the weight gradient matrix of the convolution kernel is obtained. This invention can transform the originally sparse and discrete stepping access mode into an efficient continuous access mode, significantly reducing the amount of redundant data loaded, thereby improving memory access efficiency and overall computational performance. Attached Figure Description

[0019] Figure 1 This is a schematic diagram of an embodiment of the input feature map and output gradient map provided by the present invention; Figure 2 This is a flowchart illustrating an embodiment of the convolution weight gradient calculation method provided by the present invention; Figure 3 This is a schematic diagram of another embodiment of the input feature map and output gradient map provided by the present invention; Figure 4 This is a schematic diagram of an embodiment of the feature portion regional storage and access provided by the present invention; Figure 5 This is a schematic diagram of another embodiment of the feature portion regional storage and access provided by the present invention; Figure 6 This is a schematic diagram of an embodiment of the thread loading feature data block provided by the present invention; Figure 7 This is a schematic diagram of an embodiment of obtaining the dot product result provided by the present invention; Figure 8 This is a flowchart illustrating another embodiment of the convolution weight gradient calculation method provided by the present invention; Figure 9 This is a schematic diagram of the structure of an embodiment of the convolution weight gradient calculation device provided by the present invention; Figure 10 This is a schematic diagram of another embodiment of the convolution weight gradient calculation device provided by the present invention; Figure 11 This is a schematic diagram of the structure of an embodiment of the electronic device provided by the present invention; Figure 12 This is a schematic diagram of the structure of an embodiment of the artificial intelligence processor provided by the present invention. Detailed Implementation

[0020] The technical solutions of the embodiments of the present invention will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some embodiments of the present invention, and not all embodiments. Based on the embodiments of the present invention, all other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of the present invention.

[0021] The artificial intelligence processor involved in this invention can be any one of CPU (Central Processing Unit), GPU (Graphics Processing Unit), TPU (Tensor Processing Unit), NPU (Neural Network Processing Unit), DPU (Deep Learning Processing Unit), APU (Accelerated Processing Unit), and GPGPU (General-Purpose computing on Graphics Processing Unit), depending on its application to a specific product or technology in the embodiments of this invention.

[0022] Furthermore, the row numbers involved in this invention are all row numbers of the feature data rows in the corresponding feature parts, and the row number order starts from 1. Figure 1 , 4 The blue blocks in the number sequence 5 represent the row numbers of the padding rows.

[0023] It should be noted that in backpropagation, the calculation of the weight gradient is a convolution operation between the input feature map and the output gradient map, i.e. ;in, The weight gradient matrix; It is the input feature map (Activation) in the forward propagation; It is the output gradient map (ActivationGradient) in backpropagation. This is a convolution operation; This is the loss function. Depend on It consists of several weight gradients; Convolution weights gradient, This represents the vertical offset between the convolution weights and the center of the convolution kernel. This represents the horizontal offset between the convolution weights and the center of the convolution kernel. The height of the convolution kernel; is the width of the convolution kernel.

[0024] If the sliding step size (e.g., horizontal step size, vertical step size) of the convolution kernel is greater than 1 during forward propagation, the output gradient map needs to be upsampled, and then the upsampled output gradient map is convolved with the input feature map as described above.

[0025] like Figure 1 The diagram shown is a schematic representation of an embodiment of the input feature map and output gradient map provided by the present invention. Figure 1 In the middle (a), the input feature map is shown. Figure 1 Image (b) shows the output gradient. Taking a 3×3 convolution kernel as an example, its vertical stride S during forward propagation is... V For 2 and horizontal step size S W The parameter is 2, and the single-sided padding parameter P is 1. When calculating the weight gradient using related techniques, the output gradient map needs to be upsampled, and then the upsampled result is convolved with the input feature map. That is, the upsampled result (…). Figure 1 (d) requires input feature map ( Figure 1 Nine offset operations are performed on (a) to obtain nine weight gradients respectively.

[0026] The inventors discovered that each gradient data point is multiplied by a corresponding input region of size equal to the convolution kernel size. For example... Figure 1 As shown in (c), gradient data B0 is multiplied by the purple box region centered at A0, gradient data B1 is multiplied by the black box region centered at A2, gradient data B7 is multiplied by the red box region centered at A28, and so on. By summing the product results of all gradient data, a 3×3 weight gradient matrix can be obtained.

[0027] Further research by the inventors revealed that when using an AI processor for the aforementioned calculations, there was a problem where the actual amount of data loaded by the threads far exceeded the effective computational requirements. Specifically, the inventors addressed this by arranging the threads along the vertical direction (h direction) of the output gradient graph (e.g., Figure 1 In (b), each thread needs to load one row of gradient data and a block of feature data (containing K) to participate in the calculation of that gradient data row. h =3 rows of feature data), for example, thread T0 needs to load the first row of gradient data, and the feature data block consisting of feature data rows numbered 1 to 3. To load Figure 1Taking the data in column 2 of (c) as an example, during the first load, threads T0 to T3 load the feature data in column 2 with row numbers 1, 3, 5, and 7. However, in the actual loading operation, threads T0 to T3 load all feature data in column 2 (including padding data) and then filter out the feature data with row numbers 1, 3, 5, and 7. During the second load, threads T0 to T3 load the feature data in column 2 with row numbers 2, 4, 6, and 8. In the actual operation, all feature data in column 2 is loaded, and then the feature data with row numbers 2, 4, 6, and 8 is filtered out. Similarly, during the third load, threads T0 to T3 load the feature data in column 2 with row numbers 3, 5, 7, and 9. In the actual operation, all feature data in column 2 is loaded, and then the feature data with row numbers 3, 5, 7, and 9 is filtered out. Therefore, it can be seen that the vertical step size S during the forward propagation process... V When the vertical step size is 2, the amount of data loaded each time doubles, but the effective data ratio is only 50%. Clearly, when the vertical step size S... V When the value is larger, the percentage of effective data loaded each time will be lower. This step-by-step data loading method not only leads to a significant decrease in memory access efficiency, but also increases the number of instructions issued, and is very likely to cause bandwidth bottlenecks and bank conflicts.

[0028] To address the aforementioned issues, the following description uses a GPU as an example of an AI processor to illustrate the convolution weight gradient calculation method, apparatus, computer-readable storage medium, and electronic device provided in this embodiment of the invention.

[0029] See Figure 2 This is a flowchart illustrating an embodiment of the convolution weight gradient calculation method provided by the present invention.

[0030] A first aspect of the present invention provides a method for calculating convolution weight gradients, applied to an artificial intelligence processor. The method includes steps S11 to S16, as follows: Step S11: Load the m-th gradient part in the output gradient map into shared memory, and determine the mapping region and corresponding boundary expansion region of the m-th gradient part in the input feature map based on the size of the convolution kernel, the vertical stride and the horizontal stride, so as to jointly serve as the m-th feature part; where m≥1; Step S12: Load all feature data rows of the m-th feature part into shared memory S V There are several sub-regions; where the row numbers of the feature data rows in the m-th feature part of each sub-region constitute a common difference of S. V An arithmetic sequence; there are no identical row numbers among the arithmetic sequences; S V It is equal to the vertical step size; Step S13: Calculate N for the m-th gradient part. T Row gradient data is loaded row by row into N T N threads are used to ensure that the i-th thread obtains the i-th row of gradient data; where N T This represents the total number of rows in the m-th gradient part; Step S14: Visit S sequentially in a loop V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T There are 10 threads, such that the i-th thread obtains the result of the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block consists of the feature data rows; where 1≤i≤N T ; Step S15: Perform sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded by each thread, and accumulate each dot product result to the corresponding weight gradient position; wherein, the sliding window of the feature data block is K. h ×K w And the sliding step size is a horizontal step size; K h and K w These are the height and width of the convolution kernel, respectively; Step S16: Traverse and calculate the gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient map to obtain the weight gradient matrix of the convolution kernel.

[0031] It should be noted that step S16 is equivalent to "repeating steps S11 to S15 until all gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient map are traversed to obtain the weight gradient matrix of the convolution kernel".

[0032] In step S11, when the sizes of the input feature map and the output gradient map are large, the input feature map and the output gradient map can be divided according to the required computational size (e.g., ...). Figure 1 The division method shown by the red dashed line in (b) only requires ensuring that the divided gradient part and feature part satisfy the coverage constraint relationship of the convolution operation to achieve the parallel decomposition of computation tasks (such as dot product and accumulation calculation).

[0033] Specifically, the output gradient map is divided into M gradient parts, and the m-th gradient part is selected and loaded into shared memory (GSM); where M ≥ 1. Based on the vertical step size S... V and horizontal step size S WThis involves determining the mapping region (i.e., the effective data region) corresponding to the m-th gradient part in the input feature map; simultaneously, it considers the size of the convolution kernel (height K). h Width K w The boundary processing requirements of the convolution operation are determined, thus obtaining the boundary extension region outside the mapped region. The mapped region and the boundary extension region together constitute the m-th feature part, which covers all the input feature data required for calculation with the m-th gradient part.

[0034] like Figure 1 As shown, due to the small size of the output gradient map, no partitioning operation was performed, and it only contains one gradient part (M=1). Based on the vertical step size (S) V =2) and horizontal step size (S W =2), determining the mapping region corresponding to the gradient part in the input feature map as . Figure 1 The region enclosed by the blue dashed box in (c) is used to define the boundary extension region outside the mapped region. Combined with the kernel size (3×3), these two elements together constitute the complete feature part (i.e., Figure 1 (The entire area shown in (c)).

[0035] like Figure 3 The diagram shown is a schematic diagram of another embodiment of the input feature map and output gradient map provided by the present invention; Figure 3 (a) is the input feature map (size 64×64). Figure 3 (b) shows the output gradient map (32×32). The output gradient map is divided into four gradient parts (32×8). The convolution kernel size is 3×3, S... V =2 and S W When the value is 2, the first feature is the area enclosed by a purple dashed box (size 65×18), the second feature is the area enclosed by a blue dashed box, the third feature is the area enclosed by a red dashed box, and the fourth feature is the area enclosed by a green dashed box.

[0036] In step S12, within the GSM, the m-th feature part is rearranged according to computational requirements. Specifically, the m-th feature part is arranged by row number with a vertical step size S. V Divide the interval into S V The feature data is grouped, with each group stored in an independent sub-region, and the row numbers contained in each sub-region have a tolerance of S. V The arithmetic sequence consists of non-overlapping arithmetic sequences. In this embodiment of the invention, the row numbers refer to the row numbers of the feature data rows within the corresponding feature portions, and the row numbering sequence starts from 1.

[0037] by Figure 1 Taking the characteristic part of (c) as an example, if SV =2, then the data is divided into two sub-regions: sub-region 1 (odd-numbered rows) and sub-region 2 (even-numbered rows). Sub-region 1 contains the feature data rows with row numbers 1, 3, 5, 7, and 9; sub-region 2 contains the feature data rows with row numbers 2, 4, 6, and 8. (Loading...) Figure 1 Taking the second column of data in (c) as an example, during the first load, threads T0 to T3 continuously load N data in sub-region 1, starting from row number 1. T =4 feature data (the target data has row numbers 1, 3, 5, and 7); during the second loading, threads T0 to T3 continuously load N data points in sub-region 2, starting from row number 2. T =4 feature data (the target data has row numbers 2, 4, 6, and 8); during the third load, threads T0 to T3 continuously load N data points in sub-region 1, starting from row number 3. T =4 feature data (the row numbers of the target data are 3, 5, 7, and 9).

[0038] by Figure 3 Taking the first feature as an example, the data loading operation is performed through the Tensor Memory Accelerator (TMA): according to the instruction parameters of a vertical loading step size of 2 and a horizontal loading step size of 1, data is read from global memory twice and stored in different sub-regions of GSM: ① Data is read starting from the coordinates (-1, -1) of the input feature map, and the reading result (size 34×18) is stored in sub-region 1 (odd row group) of GSM; where (x, y) represents the row coordinate x and the column coordinate y; ② According to the same instruction parameters, data is read starting from the coordinates (0, -1) of the input feature map, and the reading result (size 34×18) is stored in sub-region 2 (even row group) of GSM. Figure 4 The diagram shown is a schematic representation of an embodiment of the feature portion regional storage and access provided by the present invention. Figure 4 The darker blue and green patches represent padding values. To keep the diagram simple and avoid visual clutter, [the following text is incomplete and likely refers to a different color scheme]. Figure 4 The window sizes corresponding to the feature data blocks loaded in sub-regions 1 and 2, T0 and T31 have been scaled and are not drawn strictly according to the actual size. For example, the display size of sub-region 1 is 10×9, but its actual size is 34×18.

[0039] It's worth noting that the size of the first feature part corresponding to the first gradient part (32×8) is 65×18 (column coordinates range from -1 to 64). Theoretically, the size of the read result stored in sub-region 1 should be 33×18; the size of the read result stored in sub-region 2 should be 32×18. However, to simplify the loading process and improve the execution efficiency of instructions, in practice, the size of the read result of both sub-regions can be uniformly set to 34×18. Although this slightly exceeds the actual feature part, it does not affect the correctness of the calculation results of subsequent threads.

[0040] In step S13, the N contained in the m-th gradient part is... T Row gradient data, respectively distributed to N T There is one thread, meaning the i-th thread loads and exclusively uses the i-th row of gradient data, thus binding the thread to the gradient data row. Figure 1 As shown in (b), the threads are arranged in the vertical direction of the m-th gradient part.

[0041] In step S14, control N T One thread sequentially and cyclically accesses S V There are N sub-regions, and each time a sub-region is accessed, N... T Each thread obtains a corresponding row of feature data; the total number of accesses is K. h This allows the i-th thread to obtain the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block is composed of feature data rows. For example... Figure 4 As shown, during the first access, 32 threads (T0~T31) access sub-region 1 and extract 32 consecutive rows of feature data starting from row number 1. At this time, thread T0 obtains the first row of feature data, thread T1 obtains the third row of feature data, and so on, until thread T31 obtains the 63rd row of feature data. During the second access, they access sub-region 2 and extract 32 consecutive rows of feature data starting from row number 2. At this time, thread T0 obtains the second row of feature data, thread T1 obtains the fourth row of feature data, and so on, until thread T31 obtains the 64th row of feature data. During the third access, they access sub-region 1 and extract 32 consecutive rows of feature data starting from row number 3. At this time, thread T0 obtains the third row of feature data, thread T1 obtains the fifth row of feature data, and so on, until thread T31 obtains the 65th row of feature data. In other words, through three access operations to the sub-region, T0 obtains a feature data block (3×18) consisting of feature data rows numbered 1 to 3; T1 obtains a feature data block consisting of feature data rows numbered 3 to 5, ..., and thread T31 obtains a feature data block consisting of feature data rows numbered 63 to 65.

[0042] In this embodiment of the invention, when the feature portion is loaded into GSM, the step data in the feature portion is pre-organized into S. V The data consists of several contiguous blocks stored separately; the characteristic data rows within each sub-region are numbered using the format S. V An arithmetic sequence with common difference (i.e., following the vertical step size S) V (Pre-sorted), such as 1, 3, 5, 7, 9, ..., 63, 65. These feature data rows are stored contiguously in physical memory (sub-regions), which can transform the originally sparse step-by-step access pattern (such as reading feature data rows with row numbers 1, 3, 5, 7, 9, ..., 63 in the feature part) into a contiguous access pattern within the sub-region, thereby significantly reducing the amount of redundant data loaded, and thus improving memory access efficiency and overall computing performance.

[0043] In step S15, each thread performs sliding traversal and dot product operations on the loaded gradient data rows and feature data blocks; wherein, the sliding window of the gradient data rows is 1×1 and the sliding step size is 1; the sliding window of the feature data blocks is K. h ×K w And the sliding step size is the horizontal step size S W Each slide performs a dot product operation on the gradient data and feature data within the window, and the corresponding product result is accumulated into the corresponding weight gradient position.

[0044] like Figure 1 As shown in (b) and (c), thread T0 obtains the first row of gradient data and a feature data block consisting of feature data rows numbered 1 to 3. During the initial calculation, the window for the gradient data rows is the window corresponding to B0, and the window for the feature data block is the purple box area (3×3 in size). The data within the B0 window is multiplied by the data within the purple box area to obtain the first dot product result, which is then accumulated in the weight gradient cache. The next step involves window sliding; the window for the gradient data rows is updated to the window corresponding to B1 (sliding step size of 1), and the window for the feature data block is the black box area (3×3 in size, sliding step size of S). W =2), multiply the data in window B1 with the data in the black box area to obtain the second dot product result and accumulate it in the weight gradient cache position; and so on, thread T0 will complete all sliding calculations related to the gradient data row it holds. The other threads T1~T3 perform sliding traversal and dot product operations in the same way as above, and finally obtain a 3×3 weight gradient matrix in the cache position used for accumulation.

[0045] In step S16, steps S11 to S15 are repeated to traverse and calculate the gradient data rows and their corresponding feature data blocks for all gradient parts in the output gradient graph. After all gradient data has been processed, all dot product results have been accumulated at the cache location of the weight gradient, ultimately yielding K.h ×K w The weight gradient matrix.

[0046] This is just as mentioned earlier, "each gradient data is multiplied by the corresponding input region with a size equal to the convolution kernel size." However, in this embodiment of the invention, the computation task is divided into the following two levels to achieve parallel computation: The first level (task division) is to divide the output gradient map into M gradient parts and the input feature map into M feature parts, thereby splitting the global computation task into M groups of parallel computation of "gradient part-feature part"; where M≥1; The second level (thread-level parallelism) is to allocate the corresponding gradient data rows and feature data blocks according to the thread dimension in each group of "gradient part-feature part", and complete the dot product operation and accumulation operation through thread-level parallelism.

[0047] In summary, the embodiments of the present invention pre-organize the step data of the feature portion into S V The data is stored in contiguous blocks and separated (i.e., rearranged in GSM according to computational needs), such that the row numbers of the characteristic data rows in each sub-region have a tolerance of vertical step size S. V The arithmetic sequence is such that the data within each sub-region is arranged continuously in physical memory, thereby transforming the original sparse and discrete step access mode into an efficient continuous access mode, significantly reducing the amount of redundant data loaded, and thus improving memory access efficiency and overall computing performance.

[0048] It is worth noting that the embodiments of the present invention are applicable to any output gradient map and corresponding input feature map of any channel layer in any batch training.

[0049] In an optional embodiment, determining the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map based on the size of the convolution kernel, the vertical stride, and the horizontal stride, to collectively serve as the m-th feature part, includes: When both the vertical step size and the horizontal step size are 1, the mapping region is directly determined based on the position region of the m-th gradient part in the output gradient map. When at least one of the vertical step size and the horizontal step size is greater than 1, the m-th gradient part is upsampled in the sliding direction greater than 1, and the mapping region is determined based on the upsampled result. Based on the size of the convolution kernel, the boundary extension region surrounding the mapping region is determined; The mapped region and the boundary extension region are taken as the m-th feature part.

[0050] It should be noted that, for the vertical step size S V and horizontal step size S WFor different values ​​of , different strategies are adopted in this embodiment of the invention to determine the mapping region: ① When both the vertical step size and the horizontal step size are 1 (i.e., S) V =1 and S W =1), the size of the mapped region is the same as the size of the m-th gradient part, requiring no additional interpolation or upsampling operations. At this point, the corresponding mapped region in the input feature map is determined directly based on the position region of the m-th gradient part in the output gradient map through a one-to-one mapping.

[0051] ② When at least one of the vertical step size and the horizontal step size is greater than 1 (i.e., S) V >1 and / or S W >1), because the strided convolution operation causes the output gradient map to be compressed relative to the input feature map, it needs to be mapped by upsampling. Specifically, for sliding directions with a stride greater than 1, if S V If the value is greater than 1, then an upsampling operation is performed in the vertical direction, that is, S is inserted between adjacent gradient data rows. V -1 rows are all zero rows; if S W If the value is greater than 1, then an upsampling operation is performed in the horizontal direction, that is, S is inserted between adjacent gradient data columns. W -1 column is all zeros, such as Figure 1 As shown in (d), the upsampled result is used to perform coordinate mapping on the input feature map to obtain the corresponding mapped region, as shown in (d). Figure 1 The area enclosed by the blue dashed box in (c).

[0052] Furthermore, based on the size of the convolution kernel, the required boundary extension region around the mapped region is determined. The mapped region and its corresponding boundary extension region together constitute the m-th feature part to ensure the completeness of the weight gradient calculation.

[0053] In an optional embodiment, the step of loading all feature data rows of the m-th feature portion into shared memory S V Each sub-region includes: Starting with the j-th row of feature data, and reading feature data rows in the m-th feature part according to the vertical step size, we obtain the j-th data subset; Load the j-th data subset into the j-th sub-region of shared memory; where 1≤j≤S V .

[0054] Furthermore, the sequential cyclic access to S V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T One thread, including: According to the first sub-region to the Sth sub-regionV The sub-regions are visited sequentially, and the total number of visits is determined based on the height of the convolution kernel. During each access, the starting row position of the feature data to be loaded in the current sub-region is determined based on the current access count; Starting from the starting row position, select N consecutive rows in the current sub-region. T Row feature data is loaded into N rows. T One thread.

[0055] It should be noted that the row number sequence (arithmetic sequence) of the feature data rows in the j-th data subset is [j, j+S]. V ,j+2S V [, ...], and there are no duplicate feature data rows between any two data subsets. This ensures that all feature data rows of the m-th feature part are completely divided into S. V A subset of data, and stored separately in GSM S V Each sub-region.

[0056] For example, such as Figure 5 The diagram shown is a schematic representation of another embodiment of the feature portion regional storage and access provided by the present invention. The embodiments of the present invention... Figure 5 The text only shows the vertical feature data distribution and access logic, where the convolution kernel size is 5×5 and the vertical stride is S. V The gradient part has a total of 3 rows, N. T =32 (corresponding to 32 threads being enabled, i.e., T0 to T31). Because S V If the value is greater than 1, the gradient needs to be upsampled in the vertical direction. After upsampling, the row number range corresponding to the mapped region is 3 to 96, and the row numbers occupied by the outer boundary expansion region are 1, 2, 97, and 98. The above mapped region and the boundary expansion region together constitute the feature part containing 98 rows of feature data. The first data subset is formed according to the row number sequence [1,4,7,...97] (arithmetic sequence) and stored in subregion 1 in GSM; the second data subset is formed according to the row number sequence [2,5,8,...98] (arithmetic sequence) and stored in subregion 2 in GSM; the third data subset is formed according to the row number sequence [3,6,9,...96] (arithmetic sequence) and stored in subregion 3 in GSM. During the data loading phase of 32 threads, a loop access mechanism is used to perform data reading operations: During the first access, the system locates sub-region 1 and reads 32 consecutive rows of feature data starting from row number 1, so that each thread can load the first row of feature data in the feature data block. On the second access, locate sub-region 2 and read 32 consecutive rows of feature data starting from row number 2, so that each thread can load the second row of feature data in the feature data block; On the third access, locate sub-region 3, and read 32 consecutive rows of feature data starting from row number 3, so that each thread can load the third row of feature data in the feature data block; On the 4th access, the system is repositioned to sub-region 1 and reads 32 consecutive rows of feature data starting from row number 4, so that each thread can load the 4th row of feature data in the feature data block.

[0057] On the 5th access, the system locates sub-region 2 and reads 32 consecutive rows of feature data starting from row number 5, so that each thread can load the 5th row of feature data from the feature data block.

[0058] It should be noted that when the same sub-region is accessed repeatedly, the starting row position is positioned within that sub-region by shifting down one row each time, to ensure that the feature data rows are loaded into the corresponding threads in order.

[0059] In an optional embodiment, the number of divisions of the gradient portion in the output gradient graph is determined by the computing resources of the artificial intelligence processor.

[0060] It should be noted that the computing resources of an AI processor include, but are not limited to, the number of available threads, register resources, and shared memory capacity. In the initial stage of a task, the amount of the gradient portion can be rationally allocated based on the current processor's computing resources, thereby achieving optimal resource allocation and efficient utilization, and avoiding situations where thread resources are idle or memory overflows.

[0061] For example, if the processor has a large number of available threads and sufficient shared memory capacity, the size of the gradient part can be appropriately increased, and the number of partitions can be reduced. Especially when computing resources are extremely abundant, the number of gradient parts can even be 1 (i.e., no partitioning), in order to minimize the number of data loads into global memory and improve data access efficiency. Conversely, if limited by factors such as shared memory capacity, number of threads, or register resources, the size of the gradient part should be reduced accordingly, and the number of partitions should be increased to avoid resource contention and overflow, thereby ensuring that the computation of each gradient part can be completed efficiently.

[0062] In an optional embodiment, the N T Each thread loads the corresponding gradient data rows and feature data blocks in parallel using a specified batch loading instruction.

[0063] like Figure 6 The diagram shown is a schematic representation of an embodiment of the thread loading feature data block provided by the present invention, which is for... Figure 3The data loading process for the m-th feature part (65×18) corresponding to the m-th gradient part (32×8); where m≥1; the convolution kernel size is 3×3, and the vertical stride is S. V =2, horizontal step size S W =2. When the m-th feature is loaded into GSM, it is stored in two sub-regions, namely sub-region 1 (odd-numbered row group, size 33×18) and sub-region 2 (even-numbered row group, size 32×18).

[0064] Figure 6 The left sub-graph of Part 1 shows that when the 32 threads access the region for the first time, they locate the sub-region 1 and read 32 rows of feature data continuously from the initial row position in the sub-region 1, so that each thread obtains the first row of feature data in the corresponding feature data block, which is used to calculate the weight gradient of the first row in the 3×3 weight gradient matrix (as shown in the right sub-graph of Part 1).

[0065] The intermediate sub-diagram in Part 1 illustrates the specific implementation of the aforementioned data loading: 32 threads efficiently load data using the `ldmatrix32r strider18` instruction. This instruction works by grouping the 32 threads together, with each thread responsible for loading one piece of data. All 32 threads execute simultaneously, and adjacent thread address accesses span 18 coordinates (in this embodiment, this is equivalent to each thread's memory access step being 18 data points; the actual physical address interval is 18 multiplied by the number of bytes per data point). Specifically, the loading starting point is first located at the initial row position (row number 1) in sub-region 1. Using the `ldmatrix32r strider18` instruction, data points 1 to 32 in the first column of the sub-region are loaded in parallel into the 32 threads (corresponding to the green column in Part 1). The `ldmatrix32r strider18` instruction is then used to load data points 1 to 32 in the second column into the 32 threads (corresponding to the yellow column in Part 1). By doing so, all columns (18 columns) are loaded, so that each of the 32 threads obtains the first row of the corresponding feature data block.

[0066] The left subplot of Part 2 shows that during the second access, the 32 threads locate subregion 2 and continuously read 32 rows of feature data starting from the initial row position (row number 2). This allows each thread to obtain the second row of feature data in the corresponding feature data block, which is used to calculate the weight gradient of the second row in the 3×3 weight gradient matrix (as shown in the right subplot of Part 2). The loading process is the same as in Part 1: starting from the initial row position (row number 2) of subregion 2, the parallel loading of the second row of feature data in the feature data block by the 32 threads is achieved through the ldmatrix32r strider18 instruction (as shown in the middle subplot of Part 2).

[0067] The left subplot of Part 3 shows that during the third access, the 32 threads reposition to sub-region 1 and continuously read 32 rows of feature data starting from the feature data row with row number 3. This allows each thread to obtain the 3rd row of feature data in the corresponding feature data block, which is used to calculate the weight gradient of the 3rd row in the 3×3 weight gradient matrix (as shown in the right subplot of Part 3). The loading process is the same as in Part 1: starting from the feature data row with row number 3 in sub-region 1, the 32 threads load the 3rd row of feature data in the feature data block in parallel using the ldmatrix32rstrider18 instruction (as shown in the middle subplot of Part 3).

[0068] Similarly, when 32 threads load the corresponding gradient data row from the m-th gradient part (32×8), they only need to position the loading starting point at the first row of the gradient part and use the ldmatrix32r strider8 instruction (adjacent thread address access spans 8 coordinates) to achieve this.

[0069] See Figure 7 This is a schematic diagram of an embodiment of obtaining the dot product result provided by the present invention.

[0070] In an optional embodiment, the step of performing sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded in each thread, and accumulating each dot product result to the corresponding weight gradient position, includes: The window of the feature data block within the thread is slid with the horizontal step size, while the window of the gradient data row is slid with a step size of 1; wherein, the sliding window of the gradient data row is 1×1. After each sliding step, the two window data involved in the current calculation are located, and a dot product operation is performed to obtain the corresponding dot product result; wherein, the size of the dot product result is K. h ×K w .

[0071] For example, combined Figure 6 and Figure 7 , Figure 7 This demonstrates the data that a particular thread can hold (cached in registers). For example, registers R100~R107 store a gradient data row (1×8) from the gradient part, and registers R0~R53 store the corresponding feature data block (3×18). Figure 7 In the gradient data row on the left, the 1×1 yellow area represents the data points that have been traversed, and the 1×1 pink area represents the new data points that have been traversed. In the feature data block on the right, the leftmost column represents the data of the boundary expansion column, the blue dashed box (size 1×3) represents the window that slides on the feature data block, the yellow area represents the data that has been traversed, and the pink area represents the new data that has been traversed.

[0072] During the computation, the thread traverses horizontally, meaning the sliding window for the gradient data rows is 1×1 with a step size of 1, and the sliding window for the feature data blocks is 3×3 with a horizontal step size of S. W =2. During each slide, a dot product is performed on the data at the current window, resulting in a 3×3 product. Each element of the product is then added to its corresponding weight gradient position. This process can be achieved using the FMADW instruction. The FMADW instruction performs both a multiplication and an addition operation simultaneously, mathematically expressed as: ;in, , and It is an operand (usually a floating-point number); This is the result. The accumulated result during the calculation process is stored in the corresponding temporary accumulation register (such as...). Figure 7 In R80~R88), the shared memory / global memory is written out only after the cumulative calculation is completed.

[0073] Specifically, during the computation process, each thread will use a designated register (such as...) Figure 7The R80~R88 caches partial summations of the gradients in the weight gradient matrix. After all gradient and feature parts have been traversed, the partial summations cached in each thread are merged to obtain the weight gradient matrix of the convolution kernel. If the splitting task is assigned to different computing units (i.e., corresponding thread groups) for parallel execution, with each computing unit responsible for its corresponding gradient-feature part, then the partial summations of each thread must be completed within all computing units before the partial summations of all threads are aggregated across units to obtain the weight gradient matrix of the convolution kernel. Therefore, in this embodiment, the dot product result can be first accumulated in the designated register of the corresponding thread. After all gradient and feature parts have been traversed, the partial summations in each designated register are merged into the cache location of the weight gradient matrix in shared memory. Alternatively, the intermediate partial summation result can be temporarily stored in a designated register, and the elements of the dot product result can be directly accumulated in real time to the cache location of the weight gradient matrix in shared memory. The specific scheme adopted can be flexibly determined according to the actual computing resources and performance requirements.

[0074] In an optional embodiment, the N T Each thread releases cache locations through a register reuse mechanism to allow for the rolling loading of corresponding gradient data rows and feature data blocks.

[0075] It should be noted that when thread register resources are limited, register reuse mechanisms can be used to dynamically manage cache locations, enabling rolling loading of required data. Figure 7 For example, if a thread does not have 54 registers when processing feature data blocks, a horizontal register rolling reuse method can be used: after completing the calculation of the current sliding window data, the thread promptly releases the occupied register space, and then uses these released registers to store data that has not yet been loaded into the thread, ensuring the complete loading of the feature data block. In particular, this register rolling reuse method can be applied to cases where the input feature map and the output gradient map are not split horizontally.

[0076] See Figure 8 This is a flowchart illustrating another embodiment of the convolution weight gradient calculation method provided by the present invention.

[0077] A second aspect of the present invention provides a method for calculating convolution weight gradients, applied to an artificial intelligence processor. The method includes steps S21 to S22, as follows: Step S21: Obtain the input feature tensor and the corresponding output gradient tensor; wherein, the input feature tensor is composed of N C The output gradient tensor consists of N input feature maps of each channel; C The output gradient map consists of N channels; C>1; Step S22: For each channel, perform the convolution weight gradient calculation method described in any of the embodiments of the first aspect to obtain the weight gradient matrix of the corresponding channel.

[0078] It should be noted that the embodiments of the present invention can be applied to the weight gradient calculation of depthwise separable convolution. For example, in depthwise separable convolution, the size of the input feature tensor is 1×64×64×192 (i.e., batch × height × width × number of channels), and the size of the output gradient tensor is 1×32×32×192. By using the convolution weight gradient calculation method described in any of the embodiments of the first aspect above, each channel is processed independently to obtain the weight gradient matrix of the corresponding channel (192 channels correspond to 192 weight gradient matrices), which can greatly improve parallelism and hardware resource utilization.

[0079] In an optional embodiment, the number of parallel computations of the weight gradient matrix is ​​determined by the thread resources enabled by the artificial intelligence processor.

[0080] It should be noted that when the processor has sufficient available thread resources, the weight gradient matrices of multiple channels can be calculated in parallel to fully utilize the parallel computing capabilities of threads and achieve efficient synchronous calculation of multi-channel weight gradient matrices.

[0081] See Figure 9 This is a schematic diagram of an embodiment of the convolution weight gradient calculation device provided by the present invention.

[0082] A third aspect of the present invention provides a convolution weight gradient calculation apparatus, comprising: The first block module 11 is used to load the m-th gradient part of the output gradient map into the shared memory, and determine the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map based on the size of the convolution kernel, the vertical stride and the horizontal stride, so as to jointly serve as the m-th feature part; where m≥1; The second block module 12 is used to load all feature data rows of the m-th feature part into the shared memory S. V There are several sub-regions; where the row numbers of the feature data rows in the m-th feature part of each sub-region constitute a common difference of S. V An arithmetic sequence; there are no identical row numbers among the arithmetic sequences; S V It is equal to the vertical step size; The first loading module 13 is used to load N of the m-th gradient part. T Row gradient data is loaded row by row into N T N threads are used to ensure that the i-th thread obtains the i-th row of gradient data; where N TThis represents the total number of rows in the m-th gradient part; The second loading module 14 is used to sequentially and cyclically access S. V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T There are 10 threads, such that the i-th thread obtains the result of the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block consists of the feature data rows; where 1≤i≤N T ; The thread computation module 15 is used to perform sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded by each thread, and to accumulate each dot product result to the corresponding weight gradient position; wherein, the sliding window of the feature data block is K. h ×K w And the sliding step size is a horizontal step size; K h and K w These are the height and width of the convolution kernel, respectively; The first result acquisition module 16 is used to traverse and calculate the gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient map to obtain the weight gradient matrix of the convolution kernel.

[0083] It should be noted that the convolution weight gradient calculation device provided in the third aspect embodiment of the present invention can realize all the processes of the convolution weight gradient calculation method described in any of the first aspect embodiments. The functions and technical effects of each module and unit in the device are the same as the functions and technical effects of the convolution weight gradient calculation method described in any of the first aspect embodiments, and will not be repeated here.

[0084] See Figure 10 This is a schematic diagram of another embodiment of the convolution weight gradient calculation device provided by the present invention.

[0085] A fourth aspect of the present invention provides a convolution weight gradient calculation apparatus, comprising: Data acquisition module 21 is used to acquire the input feature tensor and the corresponding output gradient tensor; wherein, the input feature tensor consists of N C The output gradient tensor consists of N input feature maps of each channel; C The output gradient map consists of N channels; C >1; The second result acquisition module 22 is used to perform the convolution weight gradient calculation method described in any embodiment of the first aspect on the input feature map and output gradient map corresponding to each channel respectively, so as to obtain the weight gradient matrix of the corresponding channel.

[0086] It should be noted that the convolution weight gradient calculation device provided in the fourth aspect embodiment of the present invention can realize all the processes of the convolution weight gradient calculation method described in any of the second aspects embodiments. The functions and technical effects of each module and unit in the device are the same as the functions and technical effects of the convolution weight gradient calculation method described in any of the second aspects embodiments, and will not be repeated here.

[0087] A fifth aspect of the present invention provides a computer-readable storage medium comprising a stored computer program; wherein, when the computer program is executed, it controls the device on which the computer-readable storage medium is located to perform the convolution weight gradient calculation method described in any of the above embodiments.

[0088] See Figure 11 This is a schematic diagram of an embodiment of the electronic device provided by the present invention.

[0089] A sixth aspect of the present invention provides an electronic device including a processor 31, a memory 32, and a computer program stored in the memory 32 and configured to be executed by the processor 31, wherein the processor 31 implements the convolution weight gradient calculation method described in any of the above embodiments when executing the computer program.

[0090] Preferably, the computer program can be divided into one or more modules / units (such as computer program one, computer program two, ...), and the one or more modules / units are stored in the memory 32 and executed by the processor 31 to complete the present invention. The one or more modules / units can be a series of computer program instruction segments capable of performing specific functions, and the instruction segments are used to describe the execution process of the computer program in the electronic device.

[0091] The processor 31 can be any one of a CPU (Central Processing Unit), GPU (Graphics Processing Unit), TPU (Tensor Processing Unit), NPU (Neural Network Processing Unit), DPU (Deep Learning Processing Unit), APU (Accelerated Processing Unit), and GPGPU (General-Purpose Computing on Graphics Processing Unit). The processor 31 is the control center of the electronic device, connecting various parts of the electronic device via various interfaces and lines.

[0092] The memory 32 mainly includes a program storage area and a data storage area. The program storage area can store the operating system, applications required for at least one function, etc., and the data storage area can store related data, etc. In addition, the memory 32 can be a high-speed random access memory, or a non-volatile memory, such as a plug-in hard disk, a smart media card (SMC), a secure digital (SD) card, and a flash card, etc., or the memory 32 can also be other volatile solid-state storage devices.

[0093] It should be noted that the aforementioned electronic devices may include, but are not limited to, processors and memory, as will be understood by those skilled in the art. Figure 11 The structural block diagram shown is merely a structural example of the above-described electronic device and does not constitute a limitation on the structure of the above-described electronic device. The above-described electronic device may include more or fewer components than shown, or combine certain components, or different components.

[0094] See Figure 12 This is a schematic diagram of an embodiment of the artificial intelligence processor provided by the present invention.

[0095] The artificial intelligence processor provided in this embodiment of the invention includes multiple computing units (CUs). The CU is the core component for performing computational tasks and can run multiple threads in parallel to execute the convolution weight gradient calculation method described in any of the above embodiments. The CU also includes other hardware resources, such as group-shared memory (GSM) and thread-level registers (TLRs); wherein, the GSM is used to store the gradient portion extracted from the output gradient map, and the S contained in the feature portion of the input feature map. V Group feature data rows; S V The group feature data rows are stored in the S of GSM respectively. V The TLR is used to cache the gradient data rows and corresponding feature data blocks loaded by the thread during the computation process, as well as to temporarily store intermediate computation results. In addition, the AI ​​processor also includes Global Memory (GLM) for storing the input feature map and output gradient map in any embodiment of the first aspect, or the input feature tensor and output gradient tensor in any embodiment of the second aspect.

[0096] The above description is only a preferred embodiment of the present invention. It should be noted that for those skilled in the art, several improvements and modifications can be made without departing from the technical principles of the present invention, and these improvements and modifications should also be considered within the scope of protection of the present invention.

Claims

1. A method for calculating the gradient of convolution weights, characterized in that, Applied to an artificial intelligence processor, the method includes: The m-th gradient part in the output gradient map is loaded into shared memory, and the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map are determined based on the size of the convolution kernel, the vertical stride and the horizontal stride, so as to be used as the m-th feature part; where m≥1; Load all feature data rows of the m-th feature part into the shared memory S. V There are several sub-regions; where the row numbers of the feature data rows in the m-th feature part of each sub-region constitute a common difference of S. V An arithmetic sequence; there are no identical row numbers among the arithmetic sequences; S V It is equal to the vertical step size; N of the m-th gradient part T Row gradient data is loaded row by row into N T N threads are used to ensure that the i-th thread obtains the i-th row of gradient data; where N T This represents the total number of rows in the m-th gradient part; Visit S sequentially in a loop. V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T There are 10 threads, such that the i-th thread obtains the result of the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block consists of the feature data rows; where 1≤i≤N T ; In each thread, a sliding traversal and dot product operation are performed on the gradient data rows and feature data blocks loaded, and each dot product result is accumulated to the corresponding weight gradient position; wherein, the sliding window of the feature data block is K. h ×K w And the sliding step size is a horizontal step size; K h and K w These are the height and width of the convolution kernel, respectively; The gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient graph are traversed and calculated to obtain the weight gradient matrix of the convolution kernel.

2. The convolution weight gradient calculation method as described in claim 1, characterized in that, The process of determining the mapping region and corresponding boundary expansion region of the m-th gradient part in the input feature map based on the size of the convolution kernel, the vertical stride, and the horizontal stride, to collectively constitute the m-th feature part, includes: When both the vertical step size and the horizontal step size are 1, the mapping region is directly determined based on the position region of the m-th gradient part in the output gradient map. When at least one of the vertical step size and the horizontal step size is greater than 1, the m-th gradient part is upsampled in the sliding direction greater than 1, and the mapping region is determined based on the upsampled result. Based on the size of the convolution kernel, the boundary extension region surrounding the mapping region is determined; The mapped region and the boundary extension region are taken as the m-th feature part.

3. The method for calculating convolution weight gradients as described in claim 1, characterized in that, The step of loading all feature data rows of the m-th feature part into shared memory S V Each sub-region includes: Starting with the j-th row of feature data, and reading feature data rows in the m-th feature part according to the vertical step size, we obtain the j-th data subset; Load the j-th data subset into the j-th sub-region of shared memory; where 1≤j≤S V .

4. The method for calculating convolution weight gradients as described in claim 1, characterized in that, The sequential cyclic access S V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T One thread, including: According to the first sub-region to the Sth sub-region V The sub-regions are visited sequentially, and the total number of visits is determined based on the height of the convolution kernel. During each access, the starting row position of the feature data to be loaded in the current sub-region is determined based on the current access count; Starting from the starting row position, select N consecutive rows in the current sub-region. T Row feature data is loaded into N rows. T One thread.

5. The method for calculating convolution weight gradients as described in claim 1, characterized in that, The number of gradient divisions in the output gradient graph is determined by the computing resources of the artificial intelligence processor.

6. The method for calculating convolution weight gradients as described in claim 1, characterized in that, The N T Each thread loads the corresponding gradient data rows and feature data blocks in parallel using a specified batch loading instruction.

7. The method for calculating convolution weight gradients as described in claim 1, characterized in that, The N T Each thread releases cache locations through a register reuse mechanism to allow for the rolling loading of corresponding gradient data rows and feature data blocks.

8. The method for calculating convolution weight gradients as described in claim 1, characterized in that, The process of performing sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded in each thread, and accumulating each dot product result to the corresponding weight gradient position, includes: The window of the feature data block within the thread is slid with the horizontal step size, while the window of the gradient data row is slid with a step size of 1; wherein, the sliding window of the gradient data row is 1×1. After each sliding step, the two window data involved in the current calculation are located, and a dot product operation is performed to obtain the corresponding dot product result; wherein, the size of the dot product result is K. h ×K w .

9. A method for calculating the gradient of convolution weights, characterized in that, Applied to an artificial intelligence processor, the method includes: Obtain the input feature tensor and the corresponding output gradient tensor; wherein, the input feature tensor consists of N... C The output gradient tensor consists of N input feature maps of each channel; C The output gradient map consists of N channels; C >1; For each channel, the convolution weight gradient calculation method as described in any one of claims 1 to 8 is executed to obtain the weight gradient matrix of the corresponding channel.

10. The method for calculating convolution weight gradients as described in claim 9, characterized in that, The number of parallel computations of the weight gradient matrix is ​​determined by the thread resources enabled by the artificial intelligence processor.

11. A convolution weight gradient calculation device, characterized in that, include: The first block module is used to load the m-th gradient part of the output gradient map into shared memory, and determine the mapping region and corresponding boundary extension region of the m-th gradient part in the input feature map based on the size of the convolution kernel, the vertical stride and the horizontal stride, so as to jointly serve as the m-th feature part; where m≥1; The second block module is used to load all feature data rows of the m-th feature part into the shared memory S. V There are several sub-regions; where the row numbers of the feature data rows in the m-th feature part of each sub-region constitute a common difference of S. V An arithmetic sequence; there are no identical row numbers among the arithmetic sequences; S V It is equal to the vertical step size; The first loading module is used to load N of the m-th gradient part. T Row gradient data is loaded row by row into N T N threads are used to ensure that the i-th thread obtains the i-th row of gradient data; where N T This represents the total number of rows in the m-th gradient part; The second loading module is used to sequentially and cyclically access S. V Each sub-region will be accessed, and N consecutive values ​​within the current sub-region will be retrieved on each visit. T Row feature data is loaded into N rows. T There are 10 threads, such that the i-th thread obtains the result of the row number (i-1)×S. V +1 to (i-1)×S V +K h The feature data block consists of the feature data rows; where 1≤i≤N T ; The thread computation module performs sliding traversal and dot product operations on the gradient data rows and feature data blocks loaded by each thread, and accumulates each dot product result to the corresponding weight gradient position; wherein, the sliding window of the feature data block is K. h ×K w And the sliding step size is a horizontal step size; K h and K w These are the height and width of the convolution kernel, respectively; The first result acquisition module is used to traverse and calculate the gradient data rows and corresponding feature data blocks of all gradient parts in the output gradient map to obtain the weight gradient matrix of the convolution kernel.

12. A convolution weight gradient calculation device, characterized in that, include: The data acquisition module is used to acquire the input feature tensor and the corresponding output gradient tensor; wherein, the input feature tensor consists of N... C The output gradient tensor consists of N input feature maps of each channel; C The output gradient map consists of N channels; C >1; The second result acquisition module is used to perform the convolution weight gradient calculation method as described in any one of claims 1 to 8 on the input feature map and output gradient map corresponding to each channel to obtain the weight gradient matrix of the corresponding channel.

13. A computer-readable storage medium, characterized in that, The computer-readable storage medium includes a stored computer program; wherein, when the computer program is executed, it controls the device on which the computer-readable storage medium is located to perform the convolution weight gradient calculation method as described in any one of claims 1 to 10.

14. An electronic device, characterized in that, The system includes a processor, a memory, and a computer program stored in the memory and configured to be executed by the processor, wherein the processor, when executing the computer program, implements the convolution weight gradient calculation method as described in any one of claims 1 to 10.

Images