Gradient calculation method and device of convolution weight arranged continuously in vertical direction and medium

Abstract

The application discloses a convolution weight gradient calculation method and device arranged in a vertical direction continuously and a medium. The method comprises the following steps: role distribution and vertical continuous splicing are performed on N1 input feature maps and corresponding output gradient maps, an overall sliding map and an overall fixed map are obtained, and a sliding part and a fixed part of current calculation are extracted respectively according to the scale of thread granularity adaptation; a thread mask is introduced to activate corresponding calculation threads to skip the calculation of the boundary expansion row corresponding to each static fixed map; in each activated thread, synchronous sliding and point multiplication operation are performed, and the point multiplication result of each step of sliding is accumulated to the specified row position in the corresponding weight gradient matrix; after the fixed data row and the corresponding sliding data row in all covered areas of the overall fixed map are calculated, N1 weight gradient matrices are obtained. The application can reduce the number of idle threads and the calculation overhead of existing thread resources.

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 arranged in a vertical direction, a computer readable storage medium, an electronic device, and a computer program product. BACKGROUND

[0002] In a conventional convolution weight gradient calculation method, a weight gradient matrix is obtained by performing convolution calculation on the basis of input feature maps in a forward calculation process and output gradient maps in a backward propagation process. In the weight gradient calculation process, the size of the input feature maps and the output gradient maps is usually not a multiple of the thread granularity, which results in idle threads due to the lack of effective data processing, and thread resources are wasted. SUMMARY

[0003] The purpose of the embodiments of the present application is to provide a convolution weight gradient calculation method and device arranged in a vertical direction, a computer readable storage medium, an electronic device, and a computer program product. The input feature maps and the output gradient maps are respectively vertically and continuously spliced, and the sliding part and the fixed part of the current calculation are extracted according to the scale of the thread granularity, which can reduce the number of idle threads. In addition, the invalid calculation of the boundary expansion row of each static fixed map is skipped by using a thread mask, which can equivalently implement a 0 padding operation, so as to further reduce the calculation overhead of the existing thread resources.

[0004] The first embodiment of the present application provides a convolution weight gradient calculation method arranged in a vertical direction, which is applied to an artificial intelligence processor. The method comprises the following steps:

[0005] Role allocation is performed on N1 input feature maps and N1 output gradient maps to be processed, N1 dynamic sliding maps and corresponding N1 static fixed maps are divided, and logical splicing is performed in a vertical direction to obtain corresponding overall sliding maps and overall fixed maps; wherein N1≥1.

[0006] According to the thread granularity, the mth sliding part is extracted from the overall sliding map, and the corresponding region of the overall fixed map is read by boundary expansion to obtain the mth fixed part; wherein m≥1.

[0007] According to K h vertical offset relationships, K h covered regions in the mth fixed part and K h thread mask subsequences are determined; wherein each thread mask subsequence is used to control the activation state of a thread to shield the calculation of the corresponding boundary expansion row of the static fixed map in the corresponding covered region.

[0008] Load an i-th row of sliding data in an m-th sliding part and an i-th row of fixed data in a j-th covered area into an i-th thread; where i≥1; 1≤j≤K h ; K h is a height of the convolution kernel;

[0009] According to a thread mask sub-sequence corresponding to the j-th covered area, activate the corresponding thread to perform synchronous sliding and dot multiplication operation, and accumulate the dot multiplication result of each step of sliding to a j-th row position in the weight gradient matrix; where the fixed data row is slid with a 1×K w window; K w is a width of the convolution kernel;

[0010] Calculate the fixed data row and the corresponding sliding data row in all the covered areas of the whole fixed graph to obtain N1 weight gradient matrices.

[0011] Optionally, the execution of the synchronous sliding and the dot multiplication operation, and the accumulation of the dot multiplication result of each step of sliding to a j-th row position in the weight gradient matrix, comprises:

[0012] In each activated thread, according to a specified sliding step, synchronously slide the window of the sliding data row and the fixed data row loaded on the thread; where the sliding data row is slid with a 1×1 window;

[0013] After each step of sliding, locate the two window data participating in the current calculation, and perform dot multiplication operation to obtain the corresponding dot multiplication result; where the size of the dot multiplication result is 1×K w ;

[0014] Accumulate the dot multiplication result to a j-th row position in the weight gradient matrix of the corresponding dynamic sliding graph.

[0015] Optionally, the K h thread mask sub-sequences are obtained by the following steps:

[0016] S31: based on a k-th vertical offset relationship between the whole sliding graph and the whole fixed graph, locate the row sequence number in the whole sliding graph which needs to be multiplied by the corresponding boundary expansion row of the static fixed graph, and based on all the row sequence numbers, generate a k-th thread mask total sequence; where 1≤k≤K h ;

[0017] S32: repeat step S31 until K h thread mask total sequences are generated; where the length of each thread mask total sequence is equal to the number of rows of the whole sliding graph;

[0018] S33: according to the row sequence number interval of the mth sliding part in the overall sliding graph, extracting K h thread mask sub-sequences from the K h thread mask total sequence.

[0019] Optionally, the order of traversing the extracted sliding parts in the overall sliding graph is vertical direction first or horizontal direction first.

[0020] Optionally, the boundary expansion reading of the corresponding region of the overall fixed graph to obtain the mth fixed part comprises:

[0021] According to the first logical coordinates of the top-left data point in the mth sliding part in the overall sliding graph and the size of the convolution kernel, the second logical coordinates of the top-left data point in the mth fixed part in the overall fixed graph are calculated.

[0022] According to the size of the mth sliding part and the size of the convolution kernel, the size of the mth fixed part is determined, and the mth fixed part is read from the overall fixed graph based on the second logical coordinates.

[0023] Optionally, all the threads load the corresponding sliding data rows and fixed data rows in parallel through a specified batch loading instruction.

[0024] Optionally, each thread releases the cache location through a register multiplexing mechanism to load the data in the corresponding sliding data row and fixed data row in a rolling manner.

[0025] Optionally, the number of divisions of the sliding part in the overall sliding graph is determined by the computing resources of the artificial intelligence processor.

[0026] The second aspect embodiment of the present application provides a vertically continuous arrangement of convolution weight gradient calculation method, applied to an artificial intelligence processor, comprising:

[0027] The input feature tensor of N2 batches of samples and the corresponding output gradient tensor are role-assigned, and N2 dynamic sliding tensors and corresponding N2 static fixed tensors are divided; wherein each dynamic sliding tensor is composed of N C channels of dynamic sliding graphs; each static fixed tensor is composed of N C channels of static fixed graphs; N C > 1;

[0028] The vertically continuous arrangement of convolution weight gradient calculation method described in any of the first aspect embodiments is performed on the dynamic sliding graph and the corresponding static fixed graph of all the batches of samples in each channel, respectively, to obtain N2 weight gradient matrices of the corresponding channel; wherein N2=N1.

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

[0030] The third aspect of the present application provides a convolution weight gradient calculation device arranged in a vertical direction, comprising:

[0031] The merging and splicing module is configured to perform role allocation on the N1 input feature maps and the N1 output gradient maps, divide N1 dynamic sliding maps and corresponding N1 static fixed maps, and perform logical splicing in the vertical direction to obtain corresponding overall sliding maps and overall fixed maps; wherein N1≥1.

[0032] The first loading module is configured to extract an mth sliding part from the overall sliding map according to a thread granularity, and perform boundary extension reading on a corresponding region of the overall fixed map to obtain an mth fixed part; wherein m≥1.

[0033] The mask generation module is configured to determine K h offset relationships in the mth fixed part, K h covered regions in the mth fixed part, and K h thread mask subsequences; wherein each thread mask subsequence is used to control the activation state of a thread to shield the calculation of the corresponding boundary extension row of the static fixed map in the corresponding covered region.

[0034] The second loading module is configured to load an ith sliding data row in the mth sliding part and an ith fixed data row in the jth covered region to an ith thread; wherein i≥1; 1≤j≤K h ; K h is the height of the convolution kernel.

[0035] The thread calculation module is configured to activate the corresponding thread to perform synchronous sliding and dot multiplication operation according to the thread mask subsequence corresponding to the jth covered region, and accumulate the dot multiplication result of each step of sliding to the jth row position in the corresponding weight gradient matrix; wherein the fixed data row is slid with a window of 1×K w ; K w is the width of the convolution kernel.

[0036] The first result acquisition module is configured to perform traversal calculation on the fixed data rows and the corresponding sliding data rows in all the covered regions of the overall fixed map to obtain N1 weight gradient matrices.

[0037] The fourth aspect of the present application provides a convolution weight gradient calculation device arranged in a vertical direction, comprising:

[0038] a role assignment module configured to assign roles to input feature tensors and corresponding output gradient tensors of N2 batches of samples, and divide N2 dynamic sliding tensors and corresponding N2 static fixed tensors; wherein each dynamic sliding tensor is composed of dynamic sliding maps of N1 channels; each static fixed tensor is composed of static fixed maps of N1 channels; N1 > 1. C C C

[0039] a second result obtaining module configured to respectively perform the convolution weight gradient calculation method in the vertical direction for the dynamic sliding maps and the corresponding static fixed maps of all the batches of samples in each channel to obtain N2 weight gradient matrices of the corresponding channel; wherein N2 = N1.

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

[0041] The sixth aspect embodiment of the present application provides a computer program product, which comprises a computer program; the computer program realizes the convolution weight gradient calculation method in the vertical direction according to any one of the above embodiments when executed by a processor.

[0042] The seventh aspect embodiment of the present application provides an electronic device, which comprises a processor, a memory, and a computer program stored in the memory and configured to be executed by the processor; the processor realizes the convolution weight gradient calculation method in the vertical direction according to any one of the above embodiments when executing the computer program.

[0043] ​​​Compared with the prior art, the embodiment of the present application provides a convolution weight gradient calculation method, device, computer readable storage medium, electronic equipment and computer program product arranged in a vertical direction, the embodiment of the present application first assigns roles to N1 input feature maps and N1 output gradient maps to be processed, and respectively performs logical splicing in a vertical direction to obtain corresponding overall sliding maps and overall fixed maps; then, according to the size of thread granularity adaptation, the sliding part and the fixed part of the current calculation are extracted from the overall sliding map and the overall fixed map respectively; by introducing a thread mask, the corresponding calculation thread is activated to skip the invalid calculation of the boundary expansion row corresponding to each static fixed map; in each activated thread, synchronous sliding and point multiplication operations are performed, and the point multiplication result of each step of sliding is accumulated to the specified row position in the corresponding weight gradient matrix; when the fixed data row and the corresponding sliding data row in all covered areas of the overall fixed map are calculated, N1 weight gradient matrices are obtained. Through the division mechanism of adaptive thread granularity, the number of idle threads can be reduced, and the 0 padding operation is equivalent to be realized through the thread mask, further reducing the calculation overhead of the existing thread resources. BRIEF DESCRIPTION OF DRAWINGS

[0044] Figure 1 is an example diagram of an input feature map and an output gradient map provided by the present application;

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

[0046] Fig. 3(a) is a flowchart of an embodiment of obtaining the first row gradient in the weight gradient matrix provided by the present application;

[0047] Fig. 3(b) is a flowchart of an embodiment of obtaining the second row gradient in the weight gradient matrix provided by the present application;

[0048] Fig. 3(c) is a flowchart of an embodiment of obtaining the third row gradient in the weight gradient matrix provided by the present application;

[0049] Figure 4 is a flowchart of an embodiment of the convolution weight gradient calculation method arranged in a vertical direction provided by the present application;

[0050] Figure 5 is a schematic diagram of an embodiment of extracting the fixed part from the overall fixed map provided by the present application;

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

[0052] Figure 7is a schematic diagram of an embodiment of the partial thread mask sub-sequence provided by the present application;

[0053] Figure 8 is a schematic diagram of an embodiment of the total thread mask sequence provided by the present application;

[0054] Figure 9 is a schematic diagram of an embodiment of the dot product result acquisition provided by the present application;

[0055] Figure 10 is a flow schematic diagram of another embodiment of the vertically-sequentially-arranged convolution weight gradient calculation method provided by the present application;

[0056] Figure 11 is a structural schematic diagram of an embodiment of the vertically-sequentially-arranged convolution weight gradient calculation device provided by the present application;

[0057] Figure 12 is a structural schematic diagram of another embodiment of the vertically-sequentially-arranged convolution weight gradient calculation device provided by the present application;

[0058] Figure 13 is a structural schematic diagram of an embodiment of the electronic device provided by the present application;

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

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

[0061] The artificial intelligence processor involved in the present application can be any one of a CPU (Central Processing Unit), a GPU (Graphics Processing Unit), a TPU (Tensor Processing Unit), an NPU (Neural network Processing Unit), a DPU (Deeplearning Processing Unit), an APU (Accelerated Processing Unit), and a GPGPU (General-Purpose computing on Graphics Processing Unit), as determined when the embodiments of the present application are applied to specific products or technologies.

[0062] In addition, in the embodiments of the present application, a "dynamic sliding graph" is defined as a graph used to slide on a "static fixed graph" according to a specified sliding step to cover all possible positions in sequence during a convolution operation. Correspondingly, a "static fixed graph" is defined as a graph that remains in position during a convolution operation, used for the "dynamic sliding graph" to slide on the entire spatial dimension (including the width and height directions) of the "static fixed graph". Therefore, the "static fixed graph" needs to consider padding operation. In particular, each group of "dynamic sliding graph-static fixed graph" corresponds to a weight gradient matrix.

[0063] The inventor's basic inventive concept, i.e., the convolution weight gradient calculation method based on 1 input feature graph and 1 output gradient graph, will be described below to better understand the technical solutions in the embodiments of the present application:

[0064] In back propagation, the calculation of the weight gradient is a convolution operation between an input feature graph and an output gradient graph, i.e., ; wherein, is the weight gradient matrix; is the input feature graph (Activation) in forward propagation; is the output gradient graph (ActivationGradient) in back propagation; is the convolution operation; is the loss function. consists of weight gradients; is the gradient of the convolution weight , and is the position offset of the convolution weight and the center of the convolution kernel in the vertical direction, is the horizontal offset between the convolution weight and the convolution center; is the height of the convolution kernel; is the width of the convolution kernel.

[0065] See also Figure 1 , is an example diagram of an input feature map and an output gradient map provided by the present invention. Figure 1 In the figure, the yellow image represents the input feature map (as a dynamic sliding map), the blue image represents the output gradient map (as a static fixed map), and the corresponding weight gradient matrix has a size of 3×3.

[0066] If we use the related technology - Image to Column (Im2Col) method to calculate the above The weight gradient matrix of , then since the input feature map needs to perform 9 offset operations on the output gradient map in sequence to obtain 9 weight gradients respectively, that is Figure 1 The output gradient map in is read 9 times, so that its expanded matrix consists of 9 columns of data. Figure 2 The figure is a flow chart of an embodiment of the convolution weight gradient calculation method provided by the present invention. The inventors have found that in the expanded matrix of the output gradient map, each row corresponds to Figure 1 A local area of ​​the output gradient map; for example, data point A0 in the input feature map corresponds to the local area surrounded by the orange border in the output gradient map (3×3 matrix); data point A3 corresponds to the local area surrounded by the green border; data point A7 corresponds to the local area surrounded by the pink border; and so on, data point A63 corresponds to the local area surrounded by the purple border. If the input feature map has 64 data points, then it corresponds to 64 local areas, and through the dot multiplication operation, 64 local calculation results are obtained ( Matrix), then accumulate the elements of the corresponding positions in the 64 local calculation results (that is, add the corresponding elements of all matrices), and finally get a 3×3 weight gradient matrix .

[0067] After further research, the inventors found that each of the above 3×3 local areas can be split into three row matrices (each with a size of 1×3), so that the original traversal calculation task of the 3×3 local area can be decomposed into three traversal calculation tasks of 1×3 row matrices.

[0068] Taking a 3*3 weight gradient matrix as an example, as shown in FIG. 3(a), which is a flowchart of an embodiment provided by the present application for obtaining the gradient of the first row in the weight gradient matrix. In FIG. 3(a), the data points A0 to A63 in the input feature map are respectively multiplied with the first row in the corresponding 3*3 local region to generate 64 corresponding multiplication results (each with a size of 1*3), and then the elements in the corresponding positions of the 64 multiplication results are accumulated to obtain the gradient of the first row in the weight gradient matrix. As shown in FIG. 3(b), which is a flowchart of an embodiment provided by the present application for obtaining the gradient of the second row in the weight gradient matrix. In FIG. 3(b), the data points A0 to A63 in the input feature map are respectively multiplied with the second row in the corresponding 3*3 local region to generate 64 corresponding multiplication results, and then the elements in the corresponding positions of the 64 multiplication results are accumulated to obtain the gradient of the second row in the weight gradient matrix. As shown in FIG. 3(c), which is a flowchart of an embodiment provided by the present application for obtaining the gradient of the third row in the weight gradient matrix. In FIG. 3(c), the data points A0 to A63 in the input feature map are respectively multiplied with the third row in the corresponding 3*3 local region to generate 64 corresponding multiplication results, and then the elements in the corresponding positions of the 64 multiplication results are accumulated to obtain the gradient of the third row in the weight gradient matrix. By analogy, for a weight gradient matrix with a larger size, the local region thereof is the same size as the weight gradient matrix, and more row matrices can be split, and the weight gradient of the corresponding row position can be calculated by the multiplication operation and the accumulation operation.

[0069] The inventor also found that, when using an artificial intelligence processor to calculate the basic inventive concept proposed by the inventor, there is a problem of idle thread resources due to misalignment of thread granularity. Specifically, when only a single set of “dynamic sliding graph-static fixed graph” is calculated, the inventor arranges threads in the vertical direction (h direction) of the dynamic sliding graph to realize the calculation of the basic inventive concept; as shown in FIGS. 3(a) to 3(c), when the dynamic sliding graph has 8 rows of data, only 8 threads (T0 to T7) are needed to complete the calculation of the corresponding weight gradient matrix. However, thread scheduling is performed in units of thread bundles, and the thread granularity is usually 32 threads, so 24 threads are in an idle state, resulting in waste of thread resources.

[0070] To solve the above problem, the following takes a GPU as an example to illustrate the convolution weight gradient calculation method and device provided by the embodiment of the present application, computer readable storage medium, electronic equipment and computer program product.

[0071] Referring to Figure 4 , is an embodiment of the vertical direction continuous arrangement of convolution weight gradient calculation method provided by the application.

[0072] The first aspect embodiment of the application provides a vertical direction continuous arrangement of convolution weight gradient calculation method, applied to an artificial intelligence processor, the method comprises steps S11 to S16, as follows:

[0073] Step S11: role assignment is performed on N1 input feature maps and N1 output gradient maps to be processed, N1 dynamic sliding maps and corresponding N1 static fixed maps are divided, and logical splicing is performed in the vertical direction respectively to obtain corresponding overall sliding maps and overall fixed maps; wherein N1≥1;

[0074] Step S12: according to the thread granularity, the mth sliding part is extracted from the overall sliding map, and the corresponding area of the overall fixed map is read by boundary expansion to obtain the mth fixed part; wherein m≥1;

[0075] Step S13: according to K h vertical offset relations, K h covered areas and K h thread mask subsequences in the mth fixed part are determined; wherein each thread mask subsequence is used to control the activation state of the thread to shield the calculation of the corresponding boundary expansion row of the static fixed map in the corresponding covered area;

[0076] Step S14: the ith row of sliding data in the mth sliding part and the ith row of fixed data in the jth covered area are loaded into the ith thread; wherein i≥1;1≤j≤K h ; K h is the height of the convolution kernel;

[0077] Step S15: according to the thread mask subsequence corresponding to the jth covered area, the corresponding thread is activated to perform synchronous sliding and dot multiplication operation, and the dot multiplication result of each step of sliding is accumulated to the jth row position in the corresponding weight gradient matrix; wherein the fixed data row is slid with a window of 1×K w ; K w is the width of the convolution kernel;

[0078] Step S16: the fixed data row and the corresponding sliding data row in all the covered areas of the overall fixed map are calculated to obtain N1 weight gradient matrices.

[0079] It should be noted that step S16 corresponds to "repeatedly performing steps S12 to S15 until the fixed data rows and the corresponding sliding data rows in all the covered areas of the overall fixed graph are traversed and calculated to obtain N1 weight gradient matrices". In addition, for any one group of "input feature graphs - output gradient graphs ", since , the input feature graphs can be used to perform sliding calculation on the output gradient graphs to obtain , or the output gradient graphs can be used to perform sliding calculation on the input feature graphs to obtain . In other words, if the input feature graphs are dynamic sliding graphs, the output gradient graphs are static fixed graphs that are covered; if the output gradient graphs are dynamic sliding graphs, the input feature graphs are static fixed graphs that are covered. Obviously, the static fixed graph needs to consider the padding operation to ensure the effective calculation of the boundary data.

[0080] In step S11, first, the N1 input feature graphs to be processed and the N1 output gradient graphs are role-divided to form N1 groups of "dynamic sliding graphs-static fixed graphs", so as to obtain the corresponding N1 weight gradient matrices; then, the N1 dynamic sliding graphs are logically spliced in the vertical direction (h direction) to be combined into one overall sliding graph; similarly, the N1 static fixed graphs are logically spliced in the vertical direction (h direction) to be combined into one overall fixed graph. Through the splicing operation, the data row number in the vertical direction is expanded, which provides a data basis for subsequent matching of thread granularity.

[0081] In step S12, the overall sliding graph obtained by vertical splicing is usually large in size, and the overall sliding graph can be divided according to the required calculation size, and the overall fixed graph is synchronously divided in a manner suitable for the overall sliding graph, so as to ensure that the divided sliding part and fixed part meet the coverage constraint relationship of the convolution operation, thereby realizing parallel decomposition of the calculation task.

[0082] Specifically, the overall sliding graph can be divided in the horizontal direction (w direction) as shown by the red dashed line in Figure 1 , and in the vertical direction (h direction), the overall sliding graph is divided according to the thread granularity (such as a multiple of 32) to extract the mth sliding part (such as 32x16 in size) from the overall sliding graph to be assigned to the scheduled threads (32) for parallel calculation, avoiding the idle of thread resources. Since the calculation of the sliding part involves the boundary part of the corresponding area in the overall fixed graph, in order to ensure the completeness of the calculation, the corresponding boundary extension rows (the number of rows is determined by the height Kh decided), the corresponding boundary expansion columns (the number of columns is decided by the width K of the convolution kernel) on the left and right sides of the corresponding region in the w direction are read. w decided), the mth fixed part matching the mth sliding part is finally obtained (such as 34x18 in size).

[0083] As shown in Figure 5 , it is a schematic diagram of one embodiment of extracting a fixed part from an overall fixed graph provided by the application. In Figure 5 , N is the serial number of the static fixed graph (starting from 0), tileH and tileW are the height and width of the fixed part respectively. Since the mth sliding part is aligned with the thread granularity in the h direction, there may be a case of taking data across multiple dynamic sliding graphs; similarly, the mth fixed part also takes data across multiple static fixed graphs accordingly (such as Figure 5 , taking data involves 3 static fixed graphs). Of course, there may also be a case where the number of dynamic sliding graphs and static fixed graphs is greater than the specified extraction row number (such as 32 rows).

[0084] Generally, N1 dynamic sliding graphs (overall sliding graphs) and N1 static fixed graphs (overall fixed graphs) are stored in the global memory, which can be continuously read by the tensor memory accelerator (TMA) to obtain the mth sliding part and the mth fixed part, and then stored in the group shared memory (GSM) to provide the required calculation data for the thread.

[0085] In step S13, as shown in Figure 6 , it is a schematic diagram of one embodiment of thread loading fixed data rows provided by the application. Taking a 3x3 weight gradient matrix as an example, there are 3 offset relationships between the mth sliding part and the mth fixed part in the vertical direction, so there are also 3 regions covered by the mth sliding part in the mth fixed part (referred to as covered regions); as shown in the pink area in the left subgraph of Part 1 to Part 3 in Figure 6 , when the size of the mth sliding part is 32x16, the size of the covered region corresponding to the three vertical offset relationships is 32x18 (including the boundary expansion columns on the left and right sides), and each vertical offset relationship is used to calculate the gradient of the specified row position in the corresponding weight gradient matrix.

[0086] As shown in Figure 7 , it is a schematic diagram of one embodiment of a partial thread mask subsequence provided by the application. Figure 7The left, middle and right subgraphs of FIG. 3 (a) respectively show the data reading region under three vertical offset relationships, and each yellow box represents the data reading region of the corresponding static fixed graph (similar to the covered region), and each black box represents the real data region (similar to the sliding part), and H is the height of the static fixed graph / dynamic sliding graph. Because the mth sliding part and the mth fixed part involve the dynamic sliding graph and the static fixed graph sharing the same offset relationship, and the threads are arranged in rows in the h direction, for the same vertical offset relationship, the corresponding thread mask of each static fixed graph / dynamic sliding graph is the same under the same size. For example, logically splicing multiple dynamic sliding graphs similar to FIG. 3 (a) in the vertical direction, when 32 threads are scheduled, data is read across 4 dynamic sliding graphs, so T0 to T7 are used for the calculation of the 1st dynamic sliding graph, and T8 to T31 are used for the calculation of the subsequent 2nd to 4th dynamic sliding graphs; in addition, because the four dynamic sliding graphs share the same offset relationship, in any thread mask subsequence (length 32), the thread mask value corresponding to each dynamic sliding graph is the same, that is, T8 to T15, T16 to T23, T24 to T31 are the same as the thread mask value corresponding to T0 to T7.

[0087] When the mth fixed part is extracted from the whole fixed graph of the global memory to the GSM by the TMA, it is a continuous reading and writing mode (there is a case of reading data across multiple static fixed graphs). Taking the weight gradient matrix of FIG. 3 (b) as an example, in theory, for each static fixed graph, the upper boundary extension row / lower boundary extension row should be filled with 0 values (i.e. padding operation); but in the embodiment of the present application, the upper boundary extension of each static fixed graph is the effective data of the last row in the previous static fixed graph, and the lower boundary extension is the effective data of the first row in the next static fixed graph, without considering the 0 padding operation corresponding to each static fixed graph in the vertical direction.

[0088] The padding operation is realized by the thread mask (ThreadMask) in the embodiment of the present application. Specifically, because the boundary extension rows of each static fixed graph are theoretically filled with 0 values, and the threads are arranged in rows in the h direction, by disabling the calculation of the corresponding threads, these threads do not participate in the accumulation operation, thereby achieving the effect of filling 0 values. As shown in the left subgraph of FIG. 3 (c), the threads corresponding to the upper boundary extension row of the static fixed graph do not participate in the calculation (similar to T0=0 in FIG. 3 (a)), to generate the thread mask sequence TopMask. As shown in the middle subgraph of FIG. 3 (c), the threads corresponding to the lower boundary extension row of the static fixed graph do not participate in the calculation (similar to T31=31 in FIG. 3 (a)), to generate the thread mask sequence BottomMask. Figure 7 Figure 7 ​​In the middle alignment case, the calculation of the static fixed graph does not involve the boundary expansion row, so there is no need to mask the threads, i.e., all thread mask values in the thread mask sequence MidMask are 1, i.e., all threads are activated (similar to Fig. 3(b)). As shown in the right subgraph of Fig. 3(a), the threads corresponding to the boundary expansion row of the static fixed graph in the lower offset relationship do not participate in the calculation (similar to T7=0 in Fig. 3(c)) to generate the thread mask sequence BottomMask. The thread mask value of 1 indicates that the thread is turned on (participates in the calculation), and the thread mask value of 0 indicates that the thread is turned off (the calculation is shielded). Figure 7

[0089] In steps S14 and S15, the calculation of one or more weight gradient matrices (the number of which is consistent with the number of dynamic sliding graphs that are crossed) related to the mth sliding part and the mth fixed part is performed K h times, each time corresponding to the calculation of a covered area, to output the gradient of the specified row position in the corresponding weight gradient matrix.

[0090] As shown in Figure 6 and Figure 8 , Figure 8 is a schematic diagram of one embodiment of the total thread mask sequence provided by the present application. Taking a 3x3 convolution kernel as an example, Figure 6 The left subgraphs of Part1-Part3 in Fig. 3 correspond to the covered areas (with a size of 32x18) in the upper, middle, and lower vertical offset relationships, respectively; the middle subgraphs of Part1-Part3 are the threads arranged in rows in the h direction to load the corresponding fixed data rows in the covered area, so that each thread is responsible for processing one sliding data row in the sliding part and one fixed data row in the covered area when performing the current calculation; for example, the first thread (T0) loads the first sliding data row of the current sliding part and the first fixed data row of the current covered area; the second thread (T1) loads the second sliding data row of the current sliding part and the second fixed data row of the current covered area; and so on, until the 32nd thread (T31) loads the 32nd sliding data row of the current sliding part and the 32nd fixed data row of the current covered area. Subsequently, the corresponding threads are activated by the thread mask given in Figure 8 ; in each activated thread, the synchronous sliding and dot product operation is performed (as shown in Fig. 3(a) to Fig. 3(c)), and the dot product result generated by each step of sliding is accumulated to the specified row position in the corresponding weight gradient matrix (as shown in the right subgraphs of Part1-Part3 in Fig. 3). Figure 6 .

[0091] ​In step S16, after the traversal and calculation of the fixed data rows and corresponding sliding data rows of all covered areas in all fixed parts corresponding to the overall fixed graph are completed, N1 weight gradient matrices can be obtained; among which, each group of "dynamic sliding graph-static fixed graph" corresponds to a weight gradient matrix.

[0092] From the above, it can be seen that the embodiment of the present invention obtains an overall sliding map and an overall fixed map by continuously splicing multiple input feature maps and output gradient maps in the vertical direction, and extracts the sliding part and the fixed part of the current calculation according to the scale adapted to the thread granularity, which can reduce the number of idle threads. It also introduces a thread mask to skip the invalid calculation of the boundary extension row corresponding to each static fixed map, which can equivalently implement the zero-filling operation, further reducing the computational overhead of existing thread resources.

[0093] It is worth noting that the N1 input feature maps include: input feature maps of all batches of samples in a specified channel, or input feature maps of a specified batch of samples in all channels.

[0094] The composition of N1 input feature maps includes the following two specific situations:

[0095] (1) Input feature maps of all batch samples in a specified channel. When calculating the weight gradient of the convolutional layer during backpropagation, for a specified channel, it is necessary to integrate the input feature maps of all batch samples under that channel to perform gradient calculation.

[0096] For example, if there are 10 batches of samples and the designated channel is channel 3, then the input feature maps of these 10 batches of samples in channel 3 constitute the current N1 input feature maps (in this case, N1 = 10). Since the input feature map of each batch of samples in the designated channel participates in the calculation of the weight gradient, the final weight gradient corresponding to the channel is the average or sum of the gradients calculated for these batches of samples, depending on the optimizer used and the actual implementation method.

[0097] (2) Input feature maps of a specified batch of samples in all channels. When focusing on a specified batch of samples, the input feature maps of the batch samples in all channels participate in the calculation of the weight gradient.

[0098] For example, for a single batch of samples, if there are 15 channels, then the input feature maps of the batch of samples in these 15 channels constitute the current N1 input feature maps (N1=15 at this time).

[0099] See also Figure 9 , is a schematic diagram of an embodiment of obtaining a dot product result provided by the present invention.

[0100] In an alternative embodiment, the performing the synchronous sliding and the point multiplication operation, and accumulating the point multiplication result of each step of sliding to the jth row position in the weight gradient matrix, comprises:

[0101] In each activated thread, the window of the sliding data row and the fixed data row loaded onto the thread is synchronously slid according to the specified sliding step size; wherein the sliding data row is slid in a 1x1 window;

[0102] After each step of sliding, the two window data participating in the current calculation are located, and the point multiplication operation is performed to obtain the corresponding point multiplication result; wherein the size of the point multiplication result is 1xK w ;

[0103] The point multiplication result is accumulated to the jth row position in the weight gradient matrix of the corresponding dynamic sliding map.

[0104] Exemplarily, Figure 9 It is shown that the data (cached in registers) that a certain activated thread can hold, such as R100~R115 storing a sliding data row (size 1x16) in the sliding part, and registers R0~R17 storing the corresponding fixed data row (size 1x18). In Figure 6 The 1x1 yellow area in the sliding data row on the left represents the data points that have been traversed, and the 1x1 pink area represents the new data points that are currently traversed; the gray area in the fixed data row on the right represents the data of the boundary expansion column, and the blue dashed box (size 1x3) represents the sliding window on the fixed data row, and the yellow area represents the data that has been traversed, and the pink area is the new data that is currently traversed; when the specified sliding step size is 1, the sliding window of the fixed data row has an overlapping feature, therefore, the blue dashed box will contain part of the old data that has been traversed and the new data that is currently traversed at the dth sliding; wherein d>1.

[0105] In the process of calculation, the thread traverses along the horizontal direction according to the specified sliding step size, such as synchronously sliding on the sliding data row and the fixed data row with a step size of 1; wherein the window slid on the fixed data row is 1xK w ; and the window slid on the sliding data row is 1x1. Each time the data at the current window is multiplied to obtain the corresponding product result (similar to a 1x3 product result in FIG. 3(a) to FIG. 3(c)), and each element in the product result is accumulated to the corresponding weight gradient position, which can be realized by the FMADW instruction. The FMADW instruction is an instruction for simultaneously performing a multiplication and an addition operation, and the mathematical form is ; wherein, , and is an operand (usually a floating point number); is the result. The accumulated results in the calculation process are sequentially saved in the corresponding accumulation temporary registers (such as R80~R82 of Figure 9 ) of the shared memory / global memory until the accumulated calculation is completed and then written out.

[0106] Specifically, in the calculation process, each thread will cache the partial accumulated results of each gradient in the corresponding weight gradient matrix in the specified register, and after the traversal of the sliding part in the horizontal direction of the whole sliding graph is completed, the partial accumulated results cached in each thread will be merged to obtain the gradient of the specified row position in the corresponding weight gradient matrix. If the split task is assigned to different calculation units (i.e., corresponding thread groups) for parallel execution, that is, each calculation unit is responsible for the corresponding sliding part-fixed part, then after the partial accumulation of the threads in all calculation units is completed, the corresponding summary of the partial accumulated results of all threads is performed across the units, so as to obtain the gradient of the specified row position in the corresponding weight gradient matrix. Therefore, the aforementioned "the dot product result of each step sliding is accumulated to the jth row position in the corresponding weight gradient matrix" can be that the dot product result is first accumulated to the specified register of the corresponding thread, and after the traversal of the sliding part in the horizontal direction of the whole sliding graph is completed, the partial accumulated results in each specified register are merged to the cache position of the corresponding weight gradient matrix in the shared memory; of course, the partial accumulated results can also be temporarily stored without using the specified register, but directly accumulating each element of the dot product result to the cache position of the corresponding weight gradient matrix in the shared memory. The specific scheme can be flexibly determined according to the actual calculation resources and performance requirements.

[0107] It is worth noting that in the embodiment of the present application, the fixed data row held by each thread is reused when it slides in the horizontal direction, and there is no need to repeatedly load the overlapping data area from the shared memory or global memory. Therefore, by using the thread horizontal preloading mechanism and the overlapping sliding window characteristics of the fixed data row, the embodiment of the present application significantly reduces the repeated data reading behavior in the weight gradient calculation process, thereby greatly improving the calculation efficiency and memory access performance.

[0108] In the embodiment of the present application, the acquisition of the thread mask sub-sequence has at least the following two implementation manners:

[0109] The first one is to directly generate according to the reading situation of the current sliding part, and the second one is to locate and cut from the generated thread mask total sequence; but no matter which way is used, as long as the "the jth thread mask sub-sequence corresponds to the thread mask value of the static fixed graph corresponding boundary expansion row in the jth coverage area in the inactive state" is obtained.

[0110] In the first embodiment, the dynamic sliding graph involved in the current sliding part is determined, and according to each vertical offset relationship, the row position of each dynamic sliding graph multiplied by the corresponding boundary extended row of the static fixed graph is further determined, and the thread mask value corresponding to the row position is set to 0 (i.e. not activated).

[0111] Example 1: For the calculation of a 3x3 convolution weight matrix, in the case of the first vertical offset relationship (upward offset relationship), the first row in each dynamic sliding graph involved in the current sliding part is multiplied by the upper boundary extended row of the corresponding static fixed graph, so the thread mask value corresponding to the first row in each dynamic sliding graph is 0, and then the first thread mask sub-sequence required for the calculation of the current sliding part is constructed.

[0112] Example 2: For the calculation of a 5x5 convolution weight matrix, in the case of the first vertical offset relationship (upward offset by 2 rows), the thread mask value corresponding to the first row and the second row in each dynamic sliding graph is 0; in the case of the second vertical offset relationship (upward offset by 1 row), the thread mask value corresponding to the first row in each dynamic sliding graph is 0; in the case of the third vertical offset relationship (middle alignment), all thread mask values corresponding to each dynamic sliding graph are 1; in the case of the fourth vertical offset relationship (downward offset by 1 row), the thread mask value corresponding to the last first row in each dynamic sliding graph is 0; in the case of the fifth vertical offset relationship (downward offset by 2 rows), the thread mask value corresponding to the last first row and the last second row in each dynamic sliding graph is 0.

[0113] In the second embodiment, the K h thread mask sub-sequences are obtained by the following steps:

[0114] S31: based on the kth vertical offset relationship between the overall sliding graph and the overall fixed graph, the row number in the overall sliding graph which needs to be multiplied by the corresponding boundary extended row of the static fixed graph is located, and based on all the row numbers, the kth thread mask total sequence is generated; wherein 1≤k≤K h ;

[0115] S32: repeat step S31 until K h thread mask total sequences are generated; wherein the length of each thread mask total sequence is equal to the number of rows of the overall sliding graph;

[0116] S33: according to the row number interval of the mth sliding part in the overall sliding graph, K h thread mask total sequences are extracted from the K h thread mask sub-sequences.

[0117] It should be noted that the embodiment of the present invention can effectively shield invalid calculations of corresponding boundary extension rows of a static fixed graph by configuring the total thread mask sequence with different vertical offset relationships.

[0118] Since each vertical offset relationship corresponds to a different row of the convolution kernel / weight gradient matrix, there is a difference in the interaction position with the boundary extension row of the static fixed graph, so the generated K h The total sequence of thread masks is different. Figure 8 As shown, H is the height of the static fixed graph / dynamic sliding graph. Taking the 3×3 weight gradient matrix as an example, for the first vertical offset relationship (upper offset relationship), first locate all row numbers in the overall sliding graph that have a multiplication relationship with the upper boundary extension row corresponding to the static fixed graph, and then generate the first thread mask total sequence TopMask based on the row number positioning results. For the second vertical offset relationship (middle alignment), the second thread mask total sequence MidMask is directly set to all 1s, and all threads participate in the calculation by default. For the third vertical offset relationship (lower offset relationship), locate all row numbers in the overall sliding graph that have a multiplication relationship with the lower boundary extension row corresponding to the static fixed graph, and generate the third thread mask total sequence BottomMask based on the row number positioning results.

[0119] By batch generation, we get K h The total thread mask sequence of lengths is equal to the number of rows in the entire sliding graph, providing a complete mask configuration for subsequent block calculations. In actual calculations, the entire sliding graph will be divided into multiple sliding parts according to the thread granularity (such as multiples of 32). For each sliding part, first determine its row number interval in the entire sliding graph (such as the first sliding part corresponds to row numbers 0-31); then, in K h In the total thread mask sequence, the segments of the sliding part in the row number interval are intercepted to obtain the currently calculated K h thread mask subsequences.

[0120] In an optional embodiment, the number of divisions of the sliding portion in the overall sliding diagram is determined by the computing resources of the artificial intelligence processor.

[0121] It should be noted that the AI ​​processor's computing resources include, but are not limited to, the number of available threads, register resources, and shared memory capacity. At the initial stage of a task, the number of sliding parts can be appropriately divided based on the processor's current computing resources to achieve optimal resource allocation and efficient utilization, avoiding idle threads or memory overflows.

[0122] Exemplarily, when the number of threads available to the processor is large and the shared memory capacity is sufficient, the size of the sliding part can be appropriately increased (such as increasing the number of threads in the vertical direction by a multiple, and increasing the number of columns in the horizontal direction), and the number of divisions can be reduced, so as to minimize the number of data loading times of the global memory and improve the data access efficiency. Conversely, when limited by factors such as the shared memory capacity, the number of threads, or the register resources, the size of the sliding part can be correspondingly reduced, and the number of divisions can be increased, so as to avoid resource competition and overflow, thereby ensuring that the calculation of each sliding part can be efficiently completed.

[0123] In an optional embodiment, the order of extracting the sliding parts in the overall sliding graph is vertical direction priority or horizontal direction priority.

[0124] It should be noted that the order of vertical direction priority or horizontal direction priority is helpful for quickly extracting K h thread mask subsequences and quickly positioning the vertical coordinates of the sliding parts and the fixed parts.

[0125] Specifically, when the vertical direction priority traversal order is adopted, the sliding parts are extracted along the vertical direction of the overall sliding graph in sequence first, and after all the sliding parts in the vertical direction are extracted, the next round of extraction is performed in the horizontal direction. Under this order, since the sliding parts are continuously distributed in the vertical direction, the row number interval thereof has continuity, and when the subsequences are extracted from the K h thread mask total sequence, the extraction range can be quickly determined according to the starting and ending row numbers in the current vertical direction. Meanwhile, for the positioning of the vertical coordinates of the sliding parts and the fixed parts, due to the continuity in the vertical direction, the vertical coordinate of the next sliding part and the vertical coordinate range of the corresponding fixed part can be obtained through simple coordinate offset, so as to reduce the complexity of coordinate calculation.

[0126] When the horizontal direction priority traversal order is adopted, the sliding parts are extracted along the horizontal direction of the overall sliding graph in sequence first, and after all the sliding parts in the horizontal direction are extracted, the next round of extraction is performed in the next row in the vertical direction. Under this order, the distribution of the sliding parts in the horizontal direction has continuity, and the column number interval thereof is regular. For the positioning of the vertical coordinates, since the change of the vertical direction in the horizontal direction priority traversal is phased, the vertical coordinate is offset by a simple amount each time a traversal segment in the horizontal direction is completed, so as to quickly position the vertical coordinates of the sliding parts and the fixed parts, and quickly extract the thread mask subsequences.

[0127] In an optional embodiment, the boundary extension reading performed on the corresponding region of the overall fixed graph to obtain the mth fixed part includes:

[0128] According to the first logical coordinate of the left upper corner data point in the mth sliding part in the whole sliding graph and the size of the convolution kernel, the second logical coordinate of the left upper corner data point in the mth fixed part in the whole fixed graph is calculated;

[0129] According to the size of the mth sliding part and the size of the convolution kernel, the size of the mth fixed part is determined, and the mth fixed part is read from the whole fixed graph based on the second logical coordinate.

[0130] It should be noted that the vertical coordinate of the left upper corner data point in the mth fixed part is h total =H total -padH; wherein H total is the vertical coordinate of the left upper corner data point in the mth sliding part, and padH is the number of rows of the upper boundary extension row. Of course, if the order of the sliding part in the whole sliding graph is vertical direction first, then the TMA can directly maintain a continuous global logical index h total , and the global logical index is updated after each reading: new h total =h total + (tileH-2*padH); wherein the initial value of h total is -padH, and tileH is the size of the fixed part in the vertical direction.

[0131] Further, the logical coordinate h total is converted, i.e., the reading is started from the h total =ncoord*proH+hcoord, coord row in the nthcoord static fixed graph. Since h total =h total %proH; when h total coord-padH<0, the address read at this time is illegal, so 1 needs to be borrowed from ncoord to achieve a legal address, i.e., new ncoord=ncoord-1, new h total coord=h total +proH; wherein proH is the height of the static fixed graph; represents the floor of x.

[0132] In an optional embodiment, all the threads load the corresponding sliding data rows and fixed data rows in parallel through a specified batch loading instruction.

[0133] As Figure 6As shown, each Part represents the process of threads loading data in the vertical direction; for example, the left sub-graph of Part 1 represents the covered area (rows 1 to 32) where 32 threads load data in parallel in the fixed part, the middle sub-graph is a partial schematic diagram of the data points loaded into the designated registers of the 32 threads, and the right sub-graph is the data points loaded into these designated registers to calculate the gradient of the corresponding row position in the weight gradient matrix. Figure 6 In the example, the fixed portion is 34×18 in size, and the number of scheduled threads is 32. The corresponding fixed data rows are efficiently loaded into the 32 threads using the ldmatrix32r strider18 instruction. The ldmatrix32r strider18 instruction functions as follows: Each thread is responsible for loading a single piece of data, with all 32 threads executing simultaneously. Adjacent thread address accesses span 18 coordinates (in this embodiment of the present invention, this corresponds to a memory access stride of 18 data points per thread, with the actual physical address spacing being 18 times the number of bytes per data point).

[0134] by Figure 6 Taking Part 1 of the example, we first locate the loading starting point at the first row of the fixed part. Using the ldmatrix32rstrider18 instruction, we load the data from the first to the 32nd position of the first column of the fixed part into 32 threads in parallel (corresponding to the green column in Part 1). We then use the ldmatrix32rstrider18 instruction to load the data from the first to the 32nd position of the second column into 32 threads (corresponding to the yellow column in Part 1). This continues in this manner until all the data from the first to the 32nd position of each column in the fixed part is loaded, completing the loading of all fixed data rows in the corresponding covered area under the upper offset relationship by 32 threads.

[0135] Similarly, when 32 threads load the corresponding sliding data rows from the sliding part (size 32×16), they only need to locate the loading starting point at the first row of the sliding part and use the ldmatrix32r strider16 instruction (adjacent thread address access spans 16 coordinates).

[0136] In an optional embodiment, each of the threads releases cache locations through a register reuse mechanism to scroll and load data in corresponding sliding data rows and fixed data rows.

[0137] It should be noted that when the thread register resources are limited, the cache location can be dynamically managed through the register reuse mechanism to achieve rolling loading of the required data. Figure 9For example, if a thread does not have 18 registers when processing a fixed data row, or if the fixed data row contains a large number of data points, horizontal register rolling reuse can be employed: after completing the calculation of the current sliding window data, the thread promptly releases the register space occupied by the calculated data that is no longer needed for the next sliding window, and uses this freed register space to store data that has not yet been loaded into the thread, ensuring that the fixed data row is fully loaded. This register rolling reuse is particularly applicable when the entire fixed graph and the corresponding entire sliding graph are not split horizontally.

[0138] See also Figure 10 , is a flow chart of another embodiment of the vertically continuously arranged convolution weight gradient calculation method provided by the present invention.

[0139] A second aspect of the present invention provides a method for calculating a vertically continuously arranged convolution weight gradient, which is applied to an artificial intelligence processor. The method includes steps S21 to S22, which are specifically as follows:

[0140] Step S21: Assign roles to the input feature tensors and corresponding output gradient tensors of N2 batches of samples, and divide them into N2 dynamic sliding tensors and corresponding N2 static fixed tensors; wherein each dynamic sliding tensor consists of N C The dynamic sliding graph of channels is composed of N C The static fixed graph of channels N C >1;

[0141] Step S22: Execute the vertically continuously arranged convolution weight gradient calculation method described in any embodiment of the first aspect for the dynamic sliding graph and the corresponding static fixed graph of all the batch samples in each channel to obtain N2 weight gradient matrices for the corresponding channel; where N2=N1.

[0142] It should be noted that 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 10×64×64×192 (i.e., batch×height×width×number of channels, N2=10), and the size of the output gradient tensor is also 10×64×64×192. The input feature maps of these 10 batches of samples in a specified channel constitute the current N2=N1=10 input feature maps, and their output gradient maps constitute the current N2=N1=10 output gradient maps; each channel is independently processed using the vertically consecutive convolution weight gradient calculation method described in any embodiment of the first aspect above to obtain the weight gradient matrix (10×192) of these 10 sample batches in the corresponding channel. The weight gradient matrix of each channel can then be further summed / averaged.

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

[0144] It should be noted that when the thread resources enabled by the processor are sufficient, the weight gradient matrix of multiple channels can be calculated in parallel. Taking the size of the sliding part as 32x16 as an example, when 32 channels are allowed to be calculated in parallel, the data blocks (with a size of 32x16x32) corresponding to the sliding parts of the 32 channels and the data blocks (with a size of 34x18x32) corresponding to the fixed parts can be loaded into the shared memory together to calculate the weight gradient matrix involved in the 32 channels in parallel. The embodiment of the present application can fully exert the parallel computing capability of the thread, and realize efficient synchronous calculation of the weight gradient matrix of multiple channels.

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

[0146] The third aspect embodiment of the present application provides a convolution weight gradient calculation device arranged in a vertical direction, comprising:

[0147] The merging and splicing module 11 is configured to perform role allocation on the N1 input feature maps and the N1 output gradient maps to be processed, divide N1 dynamic sliding maps and corresponding N1 static fixed maps, and perform logical splicing in the vertical direction to obtain corresponding overall sliding maps and overall fixed maps; wherein N1≥1.

[0148] The first loading module 12 is configured to extract the mth sliding part from the overall sliding map according to the thread granularity, and perform boundary extension reading on the corresponding region of the overall fixed map to obtain the mth fixed part; wherein m≥1.

[0149] The mask generation module 13 is configured to determine K h covered regions and K h thread mask subsequences in the mth fixed part according to K h vertical offset relationships; wherein each thread mask subsequence is configured to control the activation state of the thread to shield the calculation of the corresponding boundary extension row of the static fixed map in the corresponding covered region.

[0150] The second loading module 14 is configured to load the ith sliding data row in the mth sliding part and the ith fixed data row in the jth covered region into the ith thread; wherein i≥1; 1≤j≤K h ; K h is the height of the convolution kernel.

[0151] The thread calculation module 15 is configured to activate corresponding threads to perform synchronous sliding and point multiplication operation according to the thread mask subsequence corresponding to the jth covered area, and to accumulate the point multiplication result of each step of sliding to the jth row position in the corresponding weight gradient matrix; wherein the fixed data row is slid with a 1xK w window; K w is the width of the convolution kernel.

[0152] The first result acquisition module 16 is configured to perform traversal calculation on the fixed data row and the corresponding sliding data row in all the covered areas of the whole fixed graph, and to obtain N1 weight gradient matrices.

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

[0154] Referring to Figure 12 , it is a structural schematic diagram of another embodiment of the vertically-continuously-arranged convolution weight gradient calculation device provided in the present application.

[0155] The fourth aspect of the present application provides a vertically-continuously-arranged convolution weight gradient calculation device, comprising:

[0156] The role allocation module 21 is configured to allocate roles to the input feature tensor and the corresponding output gradient tensor of the N2 batches of samples, and to divide N2 dynamic sliding tensors and N2 static fixed tensors; wherein each dynamic sliding tensor is composed of N C channel dynamic sliding graphs; each static fixed tensor is composed of N C channel static fixed graphs; N C >1.

[0157] The second result acquisition module 22 is configured to perform the vertically-continuously-arranged convolution weight gradient calculation method described in any of the embodiments of the first aspect on the dynamic sliding graph and the corresponding static fixed graph of all the batches of samples in each channel respectively, and to obtain N2 weight gradient matrices of the corresponding channel; wherein N2=N1.

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

[0159] The fifth aspect of the present application provides a computer readable storage medium, which includes a stored computer program; wherein the computer program controls the device where the computer readable storage medium is located to execute the vertical direction continuous arrangement convolution weight gradient calculation method described in any embodiment when running.

[0160] The sixth aspect of the present application provides a computer program product, which includes a computer program, and the computer program realizes the vertical direction continuous arrangement convolution weight gradient calculation method described in any embodiment when executed by a processor.

[0161] Referring to Figure 13 is a structural schematic diagram of an embodiment of the electronic device provided by the present application.

[0162] The seventh aspect of the present application provides an electronic device, which includes a processor 31, a memory 32, and a computer program stored in the memory 32 and configured to be executed by the processor 31, and the processor 31 realizes the vertical direction continuous arrangement convolution weight gradient calculation method described in any embodiment when executing the computer program.

[0163] Preferably, the computer program can be divided into one or more modules / units (such as computer program 1, computer program 2, …), 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 that can complete a specific function, which are used to describe the execution process of the computer program in the electronic device.

[0164] 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, connecting various parts of the electronic device through various interfaces and lines.

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

[0166] It should be noted that the above electronic device can include, but is not limited to, a processor and a memory, and those skilled in the art can understand that, Figure 13 The structural block diagram shown is only an example of the structure of the above electronic device, and does not constitute a limitation on the structure of the above electronic device. The above electronic device can include more or fewer components than shown, or combine certain components, or different components.

[0167] Referring to Figure 14 is a structural schematic diagram of an embodiment of the artificial intelligence processor provided by the application.

[0168] The artificial intelligence processor provided by the embodiments of the present application comprises a plurality of computing units (Computing Unit, CU), the CU is a core component for performing a computing task, and a plurality of threads (Threads) can be run in parallel to perform the vertical direction continuous arrangement of the convolution weight gradient calculation method described in any of the above embodiments. The CU further comprises other hardware resources, such as a thread group level shared memory (Group Shared Memory, GSM), a thread level register (Tread Level Register, TLR), etc.; wherein the GSM is used to store the sliding part and the fixed part extracted from the whole sliding graph and the whole fixed graph respectively; the TLR is used to cache the fixed data row and the corresponding sliding data row loaded by the thread in the computing process, and temporarily store the intermediate calculation result. In addition, the artificial intelligence processor further comprises a global memory (Global Memory, GLM) for storing the input feature map and the output gradient map in any of the first aspect embodiments, or the input feature tensor and the output gradient tensor in any of the second aspect embodiments.

[0169] The above only describes the preferred embodiments of the present application, and it should be noted that for those skilled in the art, without departing from the technical principles of the present application, a number of improvements and modifications can be made, and these improvements and modifications should also be considered as the protection scope of the present application.

Claims

1. A method for calculating the gradient of convolution weights arranged continuously in a vertical direction, characterized in that: Applied to an artificial intelligence processor, the method includes: Assign roles to the N1 input feature maps and N1 output gradient maps to be processed, divide them into N1 dynamic sliding maps and corresponding N1 static fixed maps, and perform logical splicing in the vertical direction to obtain the corresponding overall sliding map and overall fixed map; where N1 ≥ 1; Extracting the mth sliding portion from the overall sliding graph according to the thread granularity, and performing boundary extension reading on the corresponding area of ​​the overall fixed graph to obtain the mth fixed portion; wherein m≥1; According to K h The vertical offset relationship determines the K in the mth fixed part. h covered areas, and K h thread mask subsequences; wherein each thread mask subsequence is used to control the activation state of the thread to mask the calculation of the corresponding boundary extension row of the static fixed graph in the corresponding covered area; Load the i-th sliding data row in the m-th sliding part and the i-th fixed data row in the j-th covered area into the i-th thread; where i≥1; 1≤j≤K h ;K h is the height of the convolution kernel; According to the thread mask subsequence corresponding to the j-th covered area, the corresponding thread is activated to perform synchronous sliding and dot multiplication operations, and the dot multiplication results of each sliding step are accumulated to the j-th row position in the corresponding weight gradient matrix; wherein, the fixed data row is 1×K w The window slides; K w is the width of the convolution kernel; The fixed data rows and the corresponding sliding data rows in all the covered areas of the overall fixed graph are traversed and calculated to obtain N1 weight gradient matrices.

2. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: The synchronous sliding and dot multiplication operations are performed, and the dot multiplication results of each sliding step are accumulated to the j-th row position in the corresponding weight gradient matrix, including: In each activated thread, the windows of sliding data rows and fixed data rows loaded onto the thread are synchronously slid according to the specified sliding step size; the sliding data rows are slid in 1×1 windows; After each sliding step, the two window data involved in the current calculation are located and the dot product operation is performed to obtain the corresponding dot product result; wherein the size of the dot product result is 1×K w ; The dot product result is accumulated to the j-th row position in the weight gradient matrix of the corresponding dynamic sliding graph.

3. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: The K h The thread mask subsequence is obtained by the following steps: S31: Based on the kth vertical offset relationship between the overall sliding graph and the overall fixed graph, locate the row number in the overall sliding graph that needs to be multiplied by the boundary extension row corresponding to the static fixed graph, and generate the kth thread mask total sequence based on all the row numbers; where 1≤k≤K h ; S32: Repeat step S31 until K is generated. h thread mask total sequences; wherein the length of each thread mask total sequence is equal to the number of rows of the overall sliding graph; S33: According to the row number interval of the m-th sliding part in the overall sliding graph, h Extract K from the total thread mask sequence h thread mask subsequences.

4. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: The order of traversing and extracting sliding parts in the overall sliding map is vertical direction priority or horizontal direction priority.

5. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: The step of performing boundary expansion reading on the corresponding area of ​​the overall fixed graph to obtain the mth fixed portion includes: Calculate the second logical coordinate of the upper left corner data point in the mth fixed part in the overall fixed map based on the first logical coordinate of the upper left corner data point in the mth sliding part in the overall sliding map and the size of the convolution kernel; The size of the mth fixed part is determined according to the size of the mth sliding part and the size of the convolution kernel, and the mth fixed part is read from the overall fixed map based on the second logical coordinate.

6. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: All the threads load corresponding sliding data rows and fixed data rows in parallel through the specified batch loading instructions.

7. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: Each of the threads releases cache locations through a register reuse mechanism to scroll and load data in corresponding sliding data rows and fixed data rows.

8. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: The number of divisions of the sliding portion in the overall sliding map is determined by the computing resources of the artificial intelligence processor.

9. The method for calculating the vertically continuous convolution weight gradient according to claim 1, wherein: The N1 input feature maps include: input feature maps of all batches of samples in a specified channel, or input feature maps of a specified batch of samples in all channels.

10. A method for calculating the gradient of vertically continuously arranged convolution weights, characterized in that: Applications in artificial intelligence processors include: The input feature tensors and corresponding output gradient tensors of N2 batch samples are assigned roles, and N2 dynamic sliding tensors and corresponding N2 static fixed tensors are divided; wherein each dynamic sliding tensor is composed of N C The dynamic sliding graph of channels is composed of N C The static fixed graph of channels N C >1; For the dynamic sliding graph and the corresponding static fixed graph of all the batch samples in each channel, respectively, the vertically continuously arranged convolution weight gradient calculation method according to any one of claims 1 to 8 is performed to obtain N2 weight gradient matrices of the corresponding channel; wherein N2=N1.

11. The method for calculating the vertically continuously arranged convolution weight gradient according to claim 10, wherein: The number of parallel calculations of the weight gradient matrix is ​​determined by the thread resources enabled by the artificial intelligence processor.

12. A vertically arranged convolution weight gradient calculation device, characterized in that: include: The merging and splicing module is used to assign roles to the N1 input feature maps and N1 output gradient maps to be processed, divide them into N1 dynamic sliding maps and corresponding N1 static fixed maps, and perform logical splicing in the vertical direction to obtain the corresponding overall sliding map and overall fixed map; where N1 ≥ 1; A first loading module is configured to extract an mth sliding portion from the overall sliding graph according to a thread granularity, and perform boundary expansion reading on a corresponding region of the overall fixed graph to obtain an mth fixed portion; wherein m≥1; Mask generation module, used to generate the h The vertical offset relationship determines the K in the mth fixed part. h covered areas, and K h thread mask subsequences; wherein each thread mask subsequence is used to control the activation state of the thread to mask the calculation of the corresponding boundary extension row of the static fixed graph in the corresponding covered area; The second loading module is used to load the i-th sliding data row in the m-th sliding part and the i-th fixed data row in the j-th covered area into the i-th thread; wherein i≥1; 1≤j≤K h ;K h is the height of the convolution kernel; The thread calculation module is used to activate the corresponding thread to perform synchronous sliding and dot multiplication operations according to the thread mask subsequence corresponding to the j-th covered area, and accumulate the dot multiplication results of each sliding step to the j-th row position in the corresponding weight gradient matrix; wherein the fixed data row is 1×K w The window slides; K w is the width of the convolution kernel; The first result acquisition module is used to perform traversal calculations on the fixed data rows and the corresponding sliding data rows in all the covered areas of the overall fixed graph to obtain N1 weight gradient matrices.

13. A vertically arranged convolution weight gradient calculation device, characterized in that: include: The role assignment module is used to assign roles to the input feature tensors and corresponding output gradient tensors of N2 batch samples, dividing them into N2 dynamic sliding tensors and corresponding N2 static fixed tensors; each dynamic sliding tensor consists of N C The dynamic sliding graph of channels; each static fixed tensor consists of N C The static fixed graph of channels N C >1; The second result acquisition module is used to execute the vertically continuously arranged convolution weight gradient calculation method according to any one of claims 1 to 8 on the dynamic sliding graph and the corresponding static fixed graph of all the batch samples in each channel, respectively, to obtain N2 weight gradient matrices of the corresponding channel; wherein N2=N1.

14. A computer-readable storage medium, characterized in that The computer-readable storage medium includes a stored computer program; wherein, when the computer program is running, it controls the device where the computer-readable storage medium is located to execute the vertically continuously arranged convolution weight gradient calculation method as described in any one of claims 1 to 11.

15. A computer program product, characterized in that It includes a computer program, which, when executed by a processor, implements the vertically continuously arranged convolution weight gradient calculation method as described in any one of claims 1 to 11.

16. An electronic device, characterized in that: It includes a processor, a memory, and a computer program stored in the memory and configured to be executed by the processor, and when the processor executes the computer program, it implements the vertically continuously arranged convolution weight gradient calculation method as described in any one of claims 1 to 11.

Images