Convolution weight gradient calculation method, device, medium and equipment

Abstract

The application discloses a convolution weight gradient calculation method and device, a medium and equipment, and the method comprises the following steps: obtaining the mth gradient part and the corresponding mth feature part; all feature data rows of the mth feature part are loaded into S V sub-areas of a shared memory, and the row numbers of the feature data rows in each sub-area form an arithmetic sequence with a tolerance of S V ; S V is equal to a vertical step length; threads access different sub-areas in turn and read continuous memory blocks as needed; each thread performs sliding traversal and point multiplication operation on the loaded gradient data row and feature data block, and accumulates each point multiplication result to the corresponding weight gradient position; when the traversal calculation of the gradient data row and the corresponding feature data block of all gradient parts in the output gradient graph is completed, the weight gradient matrix of the convolution kernel is obtained. The application can convert the cross-step access into continuous access, significantly reduce the loading amount of redundant data, and further improve the memory access efficiency.

Description

TECHNICAL FIELD

[0001] The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. BACKGROUND

[0002] In a traditional convolution weight gradient calculation method, a weight gradient matrix is obtained by performing convolution calculation on the basis of an input feature map stored in a forward calculation process and an output gradient map generated in a backward propagation process. In the process of calculating the weight gradient, when the stride of the convolution kernel is greater than 1, the data access mode is a stride access mode. For an artificial intelligence processor with memory arranged continuously, in order to obtain target data distributed sparsely, a thread needs to load an entire block of continuous memory data and then filter the target data from the data. This approach results in the thread actually loading much more data than the effective calculation demand, which not only significantly reduces the memory access efficiency, but also easily causes a bandwidth bottleneck. SUMMARY

[0003] The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. V The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. V The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. V The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device.

[0004] The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device.

[0005] The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device.

[0006] The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. V The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. V The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. V The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device.

[0007] The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device. T The present application relates to the technical field of artificial intelligence, and in particular to a convolution weight gradient calculation method and device, a computer readable storage medium and an electronic device.T threads, so that the i-th thread obtains the i-th row of gradient data row; wherein, N T is the total number of rows of the m-th gradient part;

[0008] sequentially cyclically access S V sub-regions, and each time access the N T consecutive feature data rows in the current sub-region to the N T threads, so that the i-th thread obtains the feature data block composed of the feature data rows with the row numbers (i-1)×S V +1 to (i-1)×S V +K h ; wherein, 1≤i≤N T ;

[0009] perform sliding traversal and point multiplication operation in the gradient data row and the feature data block loaded by each thread, and accumulate each point multiplication 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 is the horizontal step; K h and K w are the height and width of the convolution kernel respectively;

[0010] perform traversal calculation on the gradient data row and the corresponding feature data block of all gradient parts in the output gradient map, to obtain the weight gradient matrix of the convolution kernel.

[0011] Optionally, based on the size of the convolution kernel, the vertical step and the horizontal step, determine the mapping region and the corresponding boundary expansion region of the m-th gradient part in the input feature map, to be the m-th feature part together, comprising:

[0012] when the vertical step and the horizontal step are both 1, directly determine the mapping region according to the position region of the m-th gradient part in the output gradient map;

[0013] when at least one of the vertical step and the horizontal step is greater than 1, up-sample the m-th gradient part in the sliding direction greater than 1, and determine the mapping region according to the up-sampled result;

[0014] determine the boundary expansion region of the periphery of the mapping region based on the size of the convolution kernel;

[0015] take the mapping region and the boundary expansion region as the m-th feature part.

[0016] Optionally, the loading of all feature data rows of the m-th feature part to the S V sub-regions of the shared memory comprises:

[0017] taking the jth row of feature data as a starting row and reading feature data rows in the mth feature part according to a vertical step to obtain a jth data subset;

[0018] loading the jth data subset to a jth sub-region of the shared memory; wherein 1≤j≤S V .

[0019] Optionally, the sequentially cyclically accessing S V sub-regions each time accessing continuously loading N T rows of feature data rows in a current sub-region to N T threads by row, comprising:

[0020] cyclically accessing in an order from a 1st sub-region to an S V th sub-region and determining a total number of accesses according to a height of the convolution kernel;

[0021] in each access process, determining a starting row position of feature data to be loaded in a current sub-region according to a current number of accesses;

[0022] starting from the starting row position, continuously loading N T rows of feature data rows in the current sub-region to N T threads by row.

[0023] Optionally, a division number of gradient parts in the output gradient map is determined by a computing resource of the artificial intelligence processor.

[0024] Optionally, the N T threads load corresponding gradient data rows and feature data blocks in parallel through a specified batch loading instruction.

[0025] Optionally, the N T threads release cache positions through a register reuse mechanism to load corresponding gradient data rows and feature data blocks in a rolling manner.

[0026] Optionally, the gradient data rows and the feature data blocks loaded by each thread are slidingly traversed and multiplied, and each multiplication result is accumulated to a corresponding weight gradient position, comprising:

[0027] sliding a window of the feature data blocks in the thread by the horizontal step, and simultaneously sliding a window of the gradient data rows by a step of 1; wherein the sliding window of the gradient data rows is 1×1;

[0028] after each sliding step, locating two window data participating in current calculation and performing a multiplication operation to obtain a corresponding multiplication result; wherein a size of the multiplication result is K h ×Kw .

[0029] The second aspect embodiment of the present application provides a convolution weight gradient calculation method, applied to an artificial intelligence processor, and the method comprises:

[0030] obtaining an input feature tensor and a corresponding output gradient tensor; wherein the input feature tensor is composed of input feature maps of N C channels; the output gradient tensor is composed of output gradient maps of N C channels; N C >1;

[0031] respectively performing the convolution weight gradient calculation method of any one of the first aspect embodiments on the input feature map and the output gradient map corresponding to each channel to obtain a weight gradient matrix of the corresponding channel.

[0032] Optionally, the parallel calculation quantity of the weight gradient matrix is determined by thread resources enabled by the artificial intelligence processor.

[0033] The third aspect embodiment of the present application provides a convolution weight gradient calculation device, comprising:

[0034] a first blocking module, configured to load an mth gradient part in an output gradient map to a shared memory, and determine a mapping region of the mth gradient part in an input feature map and a corresponding boundary expansion region based on the size of a convolution kernel, a vertical step and a horizontal step, to be used as an mth feature part; wherein m≥1;

[0035] a second blocking module, configured to load all feature data rows of the mth feature part to S V sub-regions of the shared memory; wherein the row numbers of the feature data rows contained in each sub-region in the mth feature part form an arithmetic sequence with a tolerance of S V ; there is no same row number between the arithmetic sequences; S V is equal to the vertical step;

[0036] a first loading module, configured to load N T row gradient data rows of the mth gradient part to N T threads row by row, so that an ith thread obtains an ith row gradient data row; wherein N T is the total number of rows of the mth gradient part;

[0037] a second loading module, configured to sequentially and circularly access S V sub-regions, and load N T row feature data rows in a current sub-region to N Tthreads, so that the ith thread obtains the feature data block composed of the feature data rows with the row numbers of (i-1)XS V +1 to (i-1)XS V +K h ; wherein, 1≤i≤N T ;

[0038] the thread calculation module is configured to perform sliding traversal and point multiplication operation in the gradient data row obtained by each thread and the feature data block, and to accumulate each point multiplication 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 is the horizontal step; K h and K w are the height and width of the convolution kernel respectively;

[0039] the first result acquisition module is configured to perform traversal calculation on the gradient data row of all gradient parts in the output gradient graph and the corresponding feature data block, to obtain the weight gradient matrix of the convolution kernel.

[0040] The fourth aspect embodiment of the present application provides a convolution weight gradient calculation device, comprising:

[0041] the data acquisition module is configured to acquire an input feature tensor and a corresponding output gradient tensor; wherein, the input feature tensor is composed of N C channel input feature maps; the output gradient tensor is composed of N C channel output gradient graphs; N C >1;

[0042] the second result acquisition module is configured to execute the convolution weight gradient calculation method of any one of the first aspect embodiments on the input feature graph and the output gradient graph corresponding to each channel respectively, to obtain the weight gradient matrix of the corresponding channel.

[0043] The fifth aspect embodiment of the present application provides a computer readable storage medium, comprising a stored computer program; wherein, the computer program controls the device where the computer readable storage medium is located to execute the convolution weight gradient calculation method of any one of the above embodiments when running.

[0044] The sixth aspect embodiment of the present application provides an electronic device, comprising a processor, a memory and a computer program stored in the memory and configured to be executed by the processor, wherein the processor realizes the convolution weight gradient calculation method of any one of the above embodiments when executing the computer program.

[0045] Compared with the prior art, the embodiment of the present application provides a convolution weight gradient calculation method, device, computer readable storage medium and electronic equipment, the embodiment of the present application respectively obtains the mth gradient part and the corresponding mth feature part from the output gradient graph and the input feature graph, and stores the mth gradient part directly into the shared memory; the mth feature part is divided into S V group feature data according to the row number, and is respectively stored in S V sub-areas, and the row numbers of the feature data rows in each sub-area constitute an arithmetic sequence with a tolerance of S V , and are arranged continuously in the physical memory to eliminate the cross-step influence; then, threads access different sub-areas in turn and read the continuous memory blocks as needed; in the calculation process, each thread performs sliding traversal and point multiplication operation on the loaded gradient data row and feature data block, obtains the point multiplication result of the current sliding position, and accumulates it into the cache position of the weight gradient matrix; when all the gradient parts in the output gradient graph and all the feature parts in the input feature graph are calculated and traversed, the weight gradient matrix of the convolution kernel can be obtained. The present application can convert the originally sparse and discrete cross-step access mode into an efficient continuous access mode, significantly reduce the loading amount of redundant data, and further improve the memory access efficiency and overall calculation performance. BRIEF DESCRIPTION OF DRAWINGS

[0046] Figure 1 is a schematic diagram of an embodiment of the input feature graph and the output gradient graph provided by the present application;

[0047] Figure 2 is a flowchart of an embodiment of the convolution weight gradient calculation method provided by the present application;

[0048] Figure 3 is a schematic diagram of another embodiment of the input feature graph and the output gradient graph provided by the present application;

[0049] Figure 4 is a schematic diagram of an embodiment of the feature part sub-area storage and access provided by the present application;

[0050] Figure 5 is a schematic diagram of another embodiment of the feature part sub-area storage and access provided by the present application;

[0051] Figure 6 is a schematic diagram of an embodiment of the thread loading feature data block provided by the present application;

[0052] Figure 7 is a schematic diagram of an embodiment of the point multiplication result provided by the present application;

[0053] Figure 8is a flow diagram of another embodiment of the method for calculating the gradient of the convolution weight provided by the present application;

[0054] Figure 9 is a structural diagram of an embodiment of the device for calculating the gradient of the convolution weight provided by the present application;

[0055] Figure 10 is a structural diagram of another embodiment of the device for calculating the gradient of the convolution weight provided by the present application;

[0056] Figure 11 is a structural diagram of an embodiment of the electronic device provided by the present application;

[0057] Figure 12 is a structural diagram of an embodiment of the artificial intelligence processor provided by the present application. DETAILED DESCRIPTION

[0058] The technical solutions in the embodiments of the present application will be described clearly and completely below with reference to the drawings in the embodiments of the present application. Obviously, the described embodiments are only a part of the embodiments of the present application, rather than all the embodiments. Based on the embodiments in the present application, all other embodiments obtained by those skilled in the art without any creative work fall within the scope of protection of the present application.

[0059] The artificial intelligence processor involved in the present application can be any one of the following: CPU (Central Processing Unit, central processor), GPU (Graphics Processing Unit, graphics processor), TPU (Tensor Processing Unit, tensor processor), NPU (Neural network Processing Unit, neural network processor), DPU (Deeplearning Processing Unit, deep learning processor), APU (Accelerated Processing Unit, accelerated processor), and GPGPU (General-Purpose computing on Graphics Processing Unit, general-purpose graphics processor). The specific product or technology is determined when the embodiments of the present application are applied.

[0060] In addition, the line numbers involved in the present application are the line numbers of the feature data lines in the corresponding feature part, and the line number sequence starts from 1. Figure 1 、 4 The blue blocks in the line number sequence in FIGS. 1, 2, 3, 4 and 5 represent the line numbers of the padding lines.

[0061] 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.

[0062] 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.

[0063] 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.

[0064] 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.

[0065] The inventor further found that when using an artificial intelligence processor to perform the above calculation, there is a problem that the amount of data actually loaded by the thread far exceeds the effective calculation requirement. Specifically, the inventor arranges threads in the vertical direction (h direction) of the output gradient map (as shown in Figure 1 (b)), each thread needs to load a row of gradient data rows and a feature data block (containing K h =3 rows of feature data rows) participating in the calculation of the gradient data row, such as thread T0 which needs to load the first row of gradient data rows and the feature data block composed of feature data rows with row numbers 1 to 3. To load Figure 1 the second column data in (c), the target data loaded by threads T0 to T3 in the first loading is the feature data with row numbers 1, 3, 5, and 7 in the second column, but in the actual loading operation, threads T0 to T3 will load all the feature data in the second column (including padding data), and then filter out the feature data with row numbers 1, 3, 5, and 7; in the second loading, the target data loaded by threads T0 to T3 is the feature data with row numbers 2, 4, 6, and 8 in the second column, and in the actual operation, all the feature data in the second column will be loaded, and then the feature data with row numbers 2, 4, 6, and 8 will be filtered out; similarly, in the third loading, the target data loaded by threads T0 to T3 is the feature data with row numbers 3, 5, 7, and 9 in the second column, and in the actual operation, all the feature data in the second column will be loaded, and then the feature data with row numbers 3, 5, 7, and 9 will be filtered out. As can be seen, when the vertical step S V is 2 in the forward propagation process, the amount of data loaded each time is doubled, but the proportion of effective data is only 50%. Obviously, when the vertical step S V takes a larger value, the proportion of effective data loaded each time will be lower. This cross-step data loading method not only leads to a significant decrease in memory access efficiency, but also increases the number of instructions emitted, and is prone to bandwidth bottlenecks and bank conflict problems.

[0066] To solve the above problems, the following describes a convolution weight gradient calculation method, device, computer readable storage medium, and electronic equipment provided by embodiments of the present application, taking a GPU as an example.

[0067] Referring to Figure 2 , which is a flowchart of an embodiment of the convolution weight gradient calculation method provided by the present application.

[0068] The first aspect of the present application provides a convolution weight gradient calculation method applied to an artificial intelligence processor, which comprises steps S11 to S16, as follows:

[0069] Step S11: load the mth gradient part in the output gradient map to the shared memory, and determine the mapping region of the mth gradient part in the input feature map and the corresponding boundary expansion region based on the size of the convolution kernel, the vertical step and the horizontal step, to jointly serve as the mth feature part; wherein m≥1;

[0070] Step S12: load all feature data rows of the mth feature part to S V sub-areas of the shared memory; wherein the row numbers of the feature data rows contained in each sub-area in the mth feature part form an arithmetic progression with a tolerance of S V ; there is no same row number between the arithmetic progressions; S V is equal to the vertical step;

[0071] Step S13: load the N T row gradient data rows of the mth gradient part by rows to N T threads, so that the ith thread obtains the ith row gradient data row; wherein N T is the total number of rows of the mth gradient part;

[0072] Step S14: sequentially access the S V sub-areas, and load the continuous N T row feature data rows in the current sub-area by rows to N T threads each time, so that the ith thread obtains a feature data block composed of the feature data rows with the row numbers of (i-1)×S V +1 to (i-1)×S V +K h ; wherein 1≤i≤N T ;

[0073] Step S15: perform sliding traversal and point multiplication operation in the gradient data row and the feature data block obtained by each thread, and accumulate each point multiplication 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 is the horizontal step; K h and K w are the height and width of the convolution kernel respectively;

[0074] Step S16: perform traversal calculation on the gradient data rows and the corresponding feature data blocks of all gradient parts in the output gradient map, to obtain the weight gradient matrix of the convolution kernel.

[0075] 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".

[0076] 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).

[0077] 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 W This 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.

[0078] 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)).

[0079] 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 WWhen 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.

[0080] 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.

[0081] by Figure 1 Taking the characteristic part of (c) as an example, if S V =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).

[0082] by Figure 3The first feature part is taken as an example, and a data loading operation is executed by a tensor memory accelerator (TMA): according to instruction parameters of a vertical loading step of 2 and a horizontal loading step of 1, data is read twice from the global memory and stored in different sub-areas of the GSM: ① data is read from the coordinates (-1, -1) of the input feature map, and the read result (with a size of 34x18) is stored in sub-area 1 (an odd row group) of the GSM; wherein (x, y) represents a row coordinate x and a column coordinate y; ② according to the same instruction parameters, data is read from the coordinates (0, -1) of the input feature map, and the read result (with a size of 34x18) is stored in sub-area 2 (an even row group) of the GSM. As shown in Figure 4 FIG. 1 is a schematic diagram of an embodiment of feature part sub-area storage and access provided by the application; Figure 4 The darker blue blocks and the darker green blocks in FIG. 1 represent padding. In order to make the diagram simple and avoid visual clutter, the padding is not shown in the diagram. Figure 4 The window size corresponding to the feature data blocks loaded by sub-area 1 and sub-area 2, T0 and T31 is scaled in FIG. 1, and is not strictly drawn according to the actual size. For example, the display size of sub-area 1 is 10x9, but the actual size is 34x18.

[0083] It is worth noting that the size of the first feature part corresponding to the first gradient part (with a size of 32x8) is 65x18 (the column coordinate range is -1~64). In theory, the read result size stored by sub-area 1 should be 33x18; the read result size stored by sub-area 2 should be 32x18, but in actual operation, the read result size of the two sub-areas can be uniformly set to 34x18 in order to simplify the loading process and improve the execution efficiency of the instructions, although it is slightly larger than the actual feature part, but does not affect the correctness of the subsequent thread calculation result.

[0084] In step S13, N T row gradient data contained in the mth gradient part is allocated to N T threads, that is, the ith thread loads and exclusively occupies the ith row gradient data row, so as to realize the binding of the thread and the gradient data row. As shown in FIG. 1(b), the threads are arranged in the vertical direction of the mth gradient part. Figure 1

[0085] In step S14, the N T threads are controlled to sequentially and circularly access S V sub-areas, and each time one sub-area is accessed, the N T threads all obtain a corresponding row of feature data; the total access times are K h , so that the ith thread obtains a row of feature data with a row number of (i-1)xs 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.

[0086] 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 a 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.

[0087] 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.

[0088] like Figure 1As shown in (b) and (c), thread T0 obtains the first row of gradient data and a feature data block composed of feature data rows with row numbers 1-3. When T0 initially calculates, the window of the row of gradient data is the window corresponding to B0, and the window of the feature data block is the purple frame area (3x3 in size). The data in the B0 window is multiplied by the data in the purple frame area to obtain the first point multiplication result and accumulate it in the cache location of the weight gradient. In the next step, the window is slid, the window of the row of gradient data is updated to the window corresponding to B1 (the sliding step is 1), and the window of the feature data block is the black frame area (3x3 in size, the sliding step is S W =2). The data in the B1 window is multiplied by the data in the black frame area to obtain the second point multiplication result and accumulate it in the cache location of the weight gradient. In this way, thread T0 completes the sliding calculation of the gradient data row it holds. The remaining threads T1-T3 also perform sliding traversal and point multiplication operation according to the above operation, and finally obtain a 3x3 weight gradient matrix at the cache location for accumulation.

[0089] In step S16, steps S11-S15 are repeatedly executed to perform traversal calculation on the gradient data rows of all gradient parts in the output gradient map and the corresponding feature data blocks. When all gradient data processing is completed, all point multiplication results are accumulated in the cache location of the weight gradient, and finally a K h x K w weight gradient matrix is obtained.

[0090] As described above, “each gradient data is multiplied by a corresponding input region with a size equal to the size of the convolution kernel”, but the embodiment of the present application divides the calculation task into two levels to realize parallel calculation: the first level (task division) is to divide the output gradient map into M gradient parts and divide the input feature map into M feature parts, so as to divide the global calculation task into M groups of “gradient part-feature part” parallel calculation; wherein M≥1; the second level (thread-level parallelism) is to assign the corresponding gradient data row and feature data block in each group of “gradient part-feature part” according to the thread dimension, and to complete the point multiplication operation and accumulation operation through thread-level parallelism.

[0091] In summary, the embodiment of the present application pre-organizes the stride data of the feature part into S V continuous data blocks and stores them separately (i.e. rearranges them in the GSM according to the calculation requirements), so that the row numbers of the feature data rows in each sub-region form an arithmetic sequence with a tolerance of the vertical step S V , and the data inside each sub-region is arranged continuously in the physical memory, thereby converting the originally sparse and discrete stride access mode into an efficient continuous access mode, significantly reducing the loading amount of redundant data, and further improving the memory access efficiency and overall calculation performance.

[0092] It should be noted that the embodiments of the present application are applicable to any output gradient map and corresponding input feature map of any channel layer in any batch training.

[0093] In an optional embodiment, the mapping region of the mth gradient part in the input feature map and the corresponding boundary expansion region are determined based on the size of the convolution kernel, the vertical step size and the horizontal step size, to jointly serve as the mth feature part, comprising:

[0094] When the vertical step size and the horizontal step size are both 1, the mapping region is directly determined according to the position region of the mth gradient part in the output gradient map;

[0095] When at least one of the vertical step size and the horizontal step size is greater than 1, the mth gradient part is up-sampled in the sliding direction with a step size greater than 1, and the mapping region is determined according to the up-sampled result;

[0096] The boundary expansion region of the periphery of the mapping region is determined based on the size of the convolution kernel;

[0097] The mapping region and the boundary expansion region serve as the mth feature part.

[0098] It should be noted that, for different values of the vertical step size S V and the horizontal step size S W , the embodiments of the present application adopt different strategies to determine the mapping region:

[0099] ① When the vertical step size and the horizontal step size are both 1 (i.e., S V =1 and S W =1), the size of the mapping region is consistent with the size of the mth gradient part, and no additional interpolation or up-sampling operation is needed. At this time, the corresponding mapping region of the mth gradient part in the input feature map is directly determined according to the position region of the mth gradient part in the output gradient map through one-to-one mapping.

[0100] ② 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), due to the size compression of the output gradient map relative to the input feature map caused by the stride convolution operation, up-sampling operation is needed for mapping. Specifically, for the sliding direction with a step size greater than 1, if S V >1, up-sampling operation is performed in the vertical direction, i.e., S V -1 all-zero rows are inserted between adjacent gradient data rows; if S W >1, up-sampling operation is performed in the horizontal direction, i.e., S W -1 all-zero columns are inserted between adjacent gradient data columns, such as Figure 1As 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).

[0101] 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.

[0102] 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:

[0103] 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;

[0104] Load the j-th data subset into the j-th sub-region of shared memory; where 1≤j≤S V .

[0105] 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:

[0106] 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.

[0107] 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;

[0108] 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.

[0109] 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.

[0110] For example, such asFigure 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:

[0111] 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.

[0112] 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;

[0113] 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;

[0114] 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.

[0115] 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.

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

[0117] In an optional embodiment, the number of gradient part divisions in the output gradient map is determined by the computing resources of the artificial intelligence processor.

[0118] It should be noted that the computing resources of the artificial intelligence processor include but are not limited to the number of available threads, register resources, and shared memory capacity. In the initial stage of the task, the number of gradient part divisions can be reasonably divided according to the current computing resources of the processor, so as to realize optimal allocation and efficient use of resources, and avoid the situation of idle thread resources or memory overflow.

[0119] For example, if the number of available threads of the processor is large and the shared memory capacity is sufficient, the size of the gradient part can be appropriately increased, and the number of divisions can be reduced. In particular, in the case of extremely abundant computing resources, the number of gradient parts can even be 1 (i.e. no division), so as to minimize the data loading times of the global memory and improve the data access efficiency. On the contrary, if limited by factors such as shared memory capacity, thread number, or register resources, the size of the gradient part can be correspondingly reduced, and the number of divisions can be increased, so as to avoid resource competition and overflow, thereby ensuring efficient completion of the calculation of each gradient part.

[0120] In an optional embodiment, the N T threads load the corresponding gradient data rows and feature data blocks in parallel through the specified batch loading instruction.

[0121] As Figure 6 shown in FIG. 1, which is a schematic diagram of one embodiment of the thread loading feature data block provided by the present application, is for the data loading process of the mth feature part (size 65x18) corresponding to the mth gradient part (size 32x8) in Figure 3 ; wherein m≥1; the convolution kernel size is 3x3, the vertical step size S V =2, and the horizontal step size S W =2. When the mth feature part is loaded into the GSM, it is stored in two sub-areas, i.e. sub-area 1 (odd row group, size 33x18) and sub-area 2 (even row group, size 32x18).

[0122] Figure 6 As shown in the left subgraph of Part 1, 32 threads are positioned to sub-area 1 in the first access, and 32 rows of feature data rows are continuously read from the initial row position in sub-area 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 of the 3x3 weight gradient matrix (as shown in the right subgraph of Part 1).

[0123] The middle subgraph of Part 1 shows the specific implementation of the above data loading: 32 threads perform efficient loading through the ldmatrix32r strider18 instruction, which has the function of: taking 32 threads as a group, and each thread is responsible for loading 1 data, a total of 32 threads are executed at the same time, and the address access of adjacent threads spans 18 coordinates (in the embodiment of the present application, it is equivalent to that the access step length of each thread to the memory is 18 data points, and the actual physical address interval is 18 times the number of bytes of each data point). Specifically, first, the starting point of loading is located at the initial row position (row number 1) in the sub-region 1, and through the ldmatrix32r strider18 instruction, the first to the thirty-second data points in the first column of the sub-region are loaded into 32 threads in parallel (corresponding to the green column in Part 1); continue to use the ldmatrix32r strider18 instruction to load the first to the thirty-second data points in the second column into 32 threads (corresponding to the yellow column in Part 1). In this way, the loading of all columns (18 columns) is completed, so that 32 threads obtain the first row of feature data rows in the corresponding feature data block.

[0124] The left subgraph of Part 2 shows that when 32 threads access for the second time, they are positioned to the sub-region 2, and 32 rows of feature data rows are continuously read from the initial row position (row number 2) so that each thread obtains the second row of feature data rows in the corresponding feature data block for calculating the weight gradient in the second row of the 3*3 weight gradient matrix (as shown in the right subgraph of Part 2); its loading process is consistent with that of Part 1: taking the initial row position (row number 2) of the sub-region 2 as the starting point, the ldmatrix32r strider18 instruction is used to realize the parallel loading of 32 threads on the second row of feature data of the feature data block (as shown in the middle subgraph of Part 2).

[0125] The left subgraph of Part 3 shows that when 32 threads access for the third time, they are repositioned to the sub-region 1, and 32 rows of feature data rows are continuously read from the feature data row with row number 3 so that each thread obtains the third row of feature data rows in the corresponding feature data block for calculating the weight gradient in the third row of the 3*3 weight gradient matrix (as shown in the right subgraph of Part 3); its loading process is consistent with that of Part 1: taking the feature data row with row number 3 in the sub-region 1 as the starting point, the ldmatrix32r strider18 instruction is used to realize the parallel loading of 32 threads on the third row of feature data in the feature data block (as shown in the middle subgraph of Part 3).

[0126] Similarly, when 32 threads load corresponding gradient data rows from the m-th gradient section (size 32x8), they only need to set the starting point to the 1st row of the gradient section, and then use the ldmatrix32r strider8 instruction (adjacent thread address access spans 8 coordinates) to achieve the goal.

[0127] Referring to Figure 7 , is a schematic diagram of an embodiment of obtaining a dot product result provided by the present application.

[0128] In an optional embodiment, the sliding traversal and dot product operation in the gradient data row and the feature data block loaded by each thread, and the accumulation of each dot product result to the corresponding weight gradient position, include:

[0129] Sliding the window of the feature data block in the thread by the horizontal step length, and sliding the window of the gradient data row by a step length of 1; wherein the sliding window of the gradient data row is 1x1;

[0130] After each sliding step, locate the two window data participating in the current calculation, and perform a dot product operation to obtain a corresponding dot product result; wherein the size of the dot product result is K h x K w .

[0131] Exemplarily, in combination with Figure 6 and Figure 7 , Figure 7 shows the data that a certain thread can hold (cached in registers), such as R100~R107 storing a gradient data row (size 1x8) in the gradient section, and registers R0~R53 storing a corresponding feature data block (size 3x18). In Figure 7 the left gradient data row, the 1x1 yellow area indicates the data points that have been traversed, and the 1x1 pink area indicates the new data points that are currently traversed; in the right feature data block, the leftmost column indicates the data of the boundary expansion column, and the blue dashed box (size 1x3) indicates the sliding window on the feature data block, and the yellow area indicates the data that has been traversed, and the pink area is the new data that is currently traversed.

[0132] In the process of calculation, the thread traverses along the horizontal direction, i.e., the sliding window of the gradient data row is 1x1 and the sliding step length is 1, and the sliding window of the feature data block is 3x3 and the sliding step length is the horizontal step length 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.

[0133] Specifically, during the computation process, each thread will use a designated register (such as...) Figure 8 The 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.

[0134] 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.

[0135] 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 9For example, if the thread does not have 54 registers when processing the feature data block, a horizontal register rolling reuse mode can be used: after the thread completes the calculation of the current sliding window data, the occupied register space is released in time, and then the released registers are used to store the data that has not been loaded into the thread to ensure the complete loading of the feature data block. In particular, this register rolling reuse mode can be applied to the case where the input feature map and the output gradient map are not split horizontally.

[0136] Referring to Figure 10 is a flowchart of another embodiment of the convolution weight gradient calculation method provided by the present application.

[0137] The second aspect embodiment of the present application provides a convolution weight gradient calculation method applied to an artificial intelligence processor, which comprises steps S21 to S22, and the details are as follows:

[0138] Step S21: obtaining an input feature tensor and a corresponding output gradient tensor; wherein the input feature tensor is composed of N C channel input feature maps; the output gradient tensor is composed of N C channel output gradient maps; N C > 1;

[0139] Step S22: performing the convolution weight gradient calculation method of any one of the above first aspect embodiments on the input feature map and the output gradient map corresponding to each channel respectively to obtain a weight gradient matrix corresponding to the channel.

[0140] It should be noted that the embodiments of the present application can be applied to the weight gradient calculation of depth separable convolution. For example, in depth separable convolution, the size of the input feature tensor is 1x64x64x192 (i.e. batchxheightxwidthxchannel number), the size of the output gradient tensor is 1x32x32x192, and each channel is processed independently by the convolution weight gradient calculation method of any one of the above first aspect embodiments to obtain a weight gradient matrix corresponding to the channel (192 weight gradient matrices corresponding to 192 channels), which can greatly improve the parallelism and hardware resource utilization.

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

[0142] It should be noted that when the thread resources available to the processor are sufficient, the weight gradient matrices of multiple channels can be calculated in parallel to fully utilize the parallel computing capability of the thread and realize efficient synchronous calculation of the weight gradient matrices of multiple channels.

[0143] Referring to Figure 11FIG. 1 is a structural schematic diagram of an embodiment of a convolution weight gradient calculation device provided by the present application.

[0144] The third aspect embodiment of the present application provides a convolution weight gradient calculation device, comprising:

[0145] The first block module 11 is configured to load the mth gradient part in the output gradient map to the shared memory, and determine the mapping region of the mth gradient part in the input feature map and the corresponding boundary expansion region based on the size, vertical step and horizontal step of the convolution kernel, so as to jointly serve as the mth feature part; wherein m≥1;

[0146] The second block module 12 is configured to load all feature data rows of the mth feature part to S V sub-areas of the shared memory; wherein the row numbers of the feature data rows contained in each sub-area in the mth feature part form an arithmetic progression with a tolerance of S V ; there is no same row number between the arithmetic progressions; S V is equal to the vertical step;

[0147] The first loading module 13 is configured to load N T row gradient data rows of the mth gradient part to N T threads row by row, so that the ith thread obtains the ith row gradient data row; wherein N T is the total number of rows of the mth gradient part;

[0148] The second loading module 14 is configured to sequentially and circularly access S V sub-areas, and load N T row feature data rows in the current sub-area to N T threads row by row each time, so that the ith thread obtains a feature data block composed of feature data rows with row numbers of (i-1)×S V +1 to (i-1)×S V +K h ; wherein 1≤i≤N T ;

[0149] The thread calculation module 15 is configured to perform sliding traversal and point multiplication operation in the gradient data row and the feature data block obtained by each thread, and accumulate each point multiplication 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 is the horizontal step; K h and K w are the height and width of the convolution kernel respectively;

[0150] The first result obtaining module 16 is configured to calculate the gradient data rows of all gradient parts in the output gradient map and the corresponding feature data blocks to obtain the weight gradient matrix of the convolution kernel.

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

[0152] Referring to Figure 11 , it is a structural schematic diagram of another embodiment of the convolution weight gradient calculation device provided by the present application.

[0153] The fourth aspect of the present application provides a convolution weight gradient calculation device, comprising:

[0154] The data obtaining module 21 is configured to obtain an input feature tensor and a corresponding output gradient tensor; wherein the input feature tensor is composed of N C channel input feature maps; the output gradient tensor is composed of N C channel output gradient maps; N C > 1.

[0155] The second result obtaining module 22 is configured to execute the convolution weight gradient calculation method described in any one of the first aspect of the present application on the input feature map and the output gradient map corresponding to each channel respectively to obtain the weight gradient matrix of the corresponding channel.

[0156] It should be noted that the convolution weight gradient calculation device provided in the fourth aspect of the present application can realize all processes of the convolution weight gradient calculation method described in any one of the second aspect of the present application, and the functions and technical effects of each module and unit in the device are the same as those of the convolution weight gradient calculation method described in any one of the second aspect of the present application, which will not be repeated here.

[0157] The fifth aspect of the present application provides a computer readable storage medium, which comprises a stored computer program; wherein the computer program controls the device where the computer readable storage medium is located to execute the convolution weight gradient calculation method described in any one of the above embodiments when running.

[0158] Referring to Figure 12 , it is a structural schematic diagram of an embodiment of the electronic device provided by the present application.

[0159] The sixth aspect of the present application provides an electronic device, comprising 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 of any of the above embodiments when executing the computer program.

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

[0161] The processor 31 can be any one of a CPU (Central Processing Unit), a GPU (Graphics Processing Unit), a TPU (Tensor Processing Unit), a NPU (Neural network Processing Unit), a DPU (Deep learning Processing Unit), an APU (Accelerated Processing Unit), and a GPGPU (General-Purpose computing on Graphics Processing Unit). The processor 31 is the control center of the electronic device, which connects various parts of the electronic device through various interfaces and lines.

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

[0163] It should be noted that the above-mentioned electronic device can include, but is not limited to, a processor, a memory, and the like, which can be understood by those skilled in the art, ​The structural diagram shown is only an example of the structure of the electronic device, and does not constitute a limitation on the structure of the electronic device. The electronic device can include more or fewer components than shown, or combine certain components, or different components.

[0164] Referring to ​ is a structural schematic diagram of an embodiment of the artificial intelligence processor provided by the present application.

[0165] The artificial intelligence processor provided by the embodiment of the present application includes a plurality of computing units (Computing Unit, CU). The CU is a core component for performing computing tasks, and can run multiple threads (Threads) in parallel to perform the convolution weight gradient calculation method described in any of the above embodiments. The CU also includes other hardware resources, such as thread group level shared memory (Group Shared Memory, GSM), thread level register (Thread Level Register, TLR), etc.; wherein the GSM is used to store the gradient part extracted from the output gradient map, and the S V group feature data rows; the S V group feature data rows are stored in the GSM V TLR is used to cache the gradient data rows and corresponding feature data blocks loaded by the thread during the calculation process, and temporarily store the intermediate calculation results. In addition, the artificial intelligence processor also includes a global memory (Global Memory, GLM) for storing the input feature map and the output gradient map in any of the embodiments of the first aspect, or the input feature tensor and the output gradient tensor in any of the embodiments of the second aspect.

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

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