Convolution weight gradient calculation method, device, medium, equipment and product

Abstract

The application discloses a convolution weight gradient calculation method and device, medium, equipment and products, the method comprises: the role distribution is carried out to the input feature map and corresponding output gradient graph in global memory, dynamic sliding graph and static fixed graph are determined; the sliding part and the corresponding fixed part are extracted from the dynamic sliding graph and the static fixed graph respectively, and each thread holds a data row in the sliding part and the corresponding alignment data block; in each thread, the held data row and the alignment data block are synchronously slid and dot multiplication is carried out, the local calculation result of the current sliding position is obtained, and is accumulated to the cache position of the weight gradient matrix, when all the sliding parts and the corresponding fixed parts are synchronously traversed and calculated, the weight gradient matrix of the convolution kernel is obtained. The application can effectively realize data multiplexing in the sliding direction, reduce the amount of repeated data reading, thereby improving the calculation efficiency of the weight gradient.

Description

Technical Field

[0001] The present invention relates to the field of artificial intelligence technology, and in particular to a convolution weight gradient calculation method, device, computer-readable storage medium, electronic device, and computer program product. Background Art

[0002] In traditional convolution weight gradient calculation methods, the weight gradient matrix is ​​obtained by performing convolution calculations on the input feature map during the forward calculation process and the output gradient map during the backward propagation process. Due to the sliding window-type computational logic of convolution calculations, there is a high degree of data duplication, that is, data in overlapping areas needs to be read repeatedly, which leads to memory access bottlenecks. Summary of the Invention

[0003] The purpose of the embodiments of the present invention is to provide a convolution weight gradient calculation method, device, computer-readable storage medium, electronic device and computer program product, which can effectively realize data reuse in the sliding direction by performing synchronous sliding calculation on data rows and aligned data blocks within the thread, and accumulate each element in the calculated local calculation results to the corresponding weight gradient position, thereby reducing the amount of data repeatedly read, thereby improving the calculation efficiency of the weight gradient.

[0004] A first embodiment of the present invention provides a convolution weight gradient calculation method, which is applied to an artificial intelligence processor. The method includes:

[0005] Assign roles to the input feature map and the corresponding output gradient map in the global memory to determine the dynamic sliding map and the static fixed map;

[0006] Extracting an mth sliding portion from the dynamic sliding graph, performing boundary extension reading on a corresponding region of the static fixed graph to obtain an mth fixed portion, and loading the mth sliding portion and the mth fixed portion into a shared memory; wherein m≥1;

[0007] The mth sliding part data rows, and the mth fixed part consists of Go to The row consists of Alignment data blocks are loaded into threads; among them, ;

[0008] In each of the threads, a synchronous sliding traversal is performed in the loaded data row and the aligned data block with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result, and each element in the local calculation result is accumulated to the corresponding weight gradient position; wherein, the aligned data block is based on a window of the data row and the aligned data block is slidingly traversed by a window of the data row and the aligned data block of the sliding part; a height of the convolution kernel; a width of the convolution kernel;

[0009] The data row and the corresponding aligned data block of all the sliding parts in the dynamic sliding graph are traversed and calculated to obtain a weight gradient matrix of the convolution kernel.

[0010] Optionally, the number of divisions of the sliding part is determined by the computing resource of the artificial intelligence processor.

[0011] Optionally, the boundary extension reading of the corresponding region of the static fixed graph to obtain the mth fixed part comprises:

[0012] The corresponding region of the mth sliding part in the static fixed graph is determined, wherein the size of the corresponding region is equal to the size of the mth sliding part;

[0013] The corresponding region is boundary extended based on the size of the convolution kernel to obtain the mth fixed part, wherein the extended region in the mth fixed part comprises zero padding values and / or effective data of adjacent fixed parts.

[0014] Optionally, all the threads load the corresponding data row and the aligned data block in parallel through a specified batch loading instruction.

[0015] Optionally, each thread releases a cache location through a register multiplexing mechanism to rollingly load data in the corresponding data row and the aligned data block.

[0016] Optionally, the synchronous sliding traversal in the loaded data row and the aligned data block at the specified sliding step is performed to execute a dot product operation to obtain a local calculation result, comprising:

[0017] The window of the data row and the aligned data block in the thread is synchronously slid according to the specified sliding step, wherein the data row is slidingly traversed by a window of the data row and the aligned data block;

[0018] After each sliding, two window data participating in the current calculation are located, and a dot product operation is performed to obtain a corresponding local calculation result, wherein the size of the local calculation result is .

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

[0020] ​The input feature tensor and the corresponding output gradient tensor in the global memory are role-assigned to determine a dynamic sliding tensor and a static fixed tensor; wherein the dynamic sliding tensor is composed of dynamic sliding maps of channels; and the static fixed tensor is composed of static fixed maps of channels; >1;

[0021] The convolution weight gradient calculation method of any one of the first aspect embodiments is performed on the dynamic sliding map and the static fixed map corresponding to each channel respectively to obtain a weight gradient matrix of the corresponding channel.

[0022] Optionally, the parallel calculation number of the weight gradient matrix is determined by the thread resource enabled by the artificial intelligence processor.

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

[0024] A first role assignment module is configured to role-assign the input feature map and the corresponding output gradient map in the global memory to determine a dynamic sliding map and a static fixed map.

[0025] A first loading module is configured to extract an mth sliding part from the dynamic sliding map, perform boundary extension reading on the corresponding region of the static fixed map to obtain an mth fixed part, and load the mth sliding part and the mth fixed part to a shared memory; wherein m≥1.

[0026] A second loading module is configured to load the first data row of the mth sliding part and the first aligned data block composed of the first row to the first row in the mth fixed part to the first thread; wherein ;

[0027] A thread calculation module is configured to perform synchronous sliding traversal in the loaded data row and aligned data block with a specified sliding step in each thread to perform dot multiplication operation to obtain a local calculation result, and accumulate each element in the local calculation result to a corresponding weight gradient position respectively; wherein the aligned data block is slidingly traversed with a window of ; is the height of the convolution kernel; is the width of the convolution kernel;

[0028] A first result acquisition module is configured to perform traversal calculation on the data row of all sliding parts in the dynamic sliding map and the corresponding aligned data block to obtain a weight gradient matrix of the convolution kernel.

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

[0030] A second role allocation module is configured to allocate roles to the input feature tensor and the corresponding output gradient tensor in the global memory, and determine a dynamic sliding tensor and a static fixed tensor; wherein the dynamic sliding tensor is composed of dynamic sliding maps of channels; and the static fixed tensor is composed of static fixed maps of channels. >1;

[0031] A second result obtaining module is configured to execute the convolution weight gradient calculation method of any one of the first aspect for each channel to obtain a weight gradient matrix of the corresponding channel.

[0032] The fifth aspect 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.

[0033] The sixth aspect of the present application provides a computer program product, comprising a computer program, which realizes the convolution weight gradient calculation method of any one of the above embodiments when executed by a processor.

[0034] The seventh aspect 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.

[0035] 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. The embodiment of the present application first performs role allocation on the input feature map and the corresponding output gradient map in the global memory, determines a dynamic sliding map and a static fixed map, extracts a sliding part and a corresponding fixed part from the dynamic sliding map and the static fixed map respectively, and makes each thread hold a data row in the sliding part and a corresponding aligned data block. In each thread, the held data row and aligned data block are synchronously slid and multiplied, the local calculation result of the current sliding position is obtained, and is accumulated in the cache position of the weight gradient matrix. When the synchronous traversal calculation of all sliding parts in the dynamic sliding map and all fixed parts in the static fixed map is completed, the weight gradient matrix of the convolution kernel can be obtained. The present application can realize data reuse through the thread horizontal preloading mechanism and the overlapping sliding window characteristics of the aligned data block, without repeatedly loading the overlapping data area from the shared memory or the global memory, thereby significantly reducing the repeated data reading behavior in the weight gradient calculation process, and greatly improving the calculation efficiency and memory access performance. BRIEF DESCRIPTION OF DRAWINGS

[0036] Figure 1 is a schematic diagram of an embodiment of the convolution weight gradient calculation provided by the related art;

[0037] Figure 2 is a schematic diagram of an embodiment of the data association relationship provided by the present application;

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

[0039] Figure 4 is a flowchart of another embodiment of the convolution weight gradient calculation method provided by the present application;

[0040] Figure 5 is a schematic diagram of an embodiment of the sliding part and the fixed part provided by the present application;

[0041] Figure 6 is a schematic diagram of an embodiment of obtaining the local calculation result provided by the present application;

[0042] Figure 7 is a schematic diagram of an embodiment of loading the aligned data block by the thread provided by the present application;

[0043] Figure 8 is a flowchart of another embodiment of the convolution weight gradient calculation method provided by the present application;

[0044] Figure 9is a structural schematic diagram of one embodiment of the convolution weight gradient calculation device provided by the present application;

[0045] Figure 10 is a structural schematic diagram of another embodiment of the convolution weight gradient calculation device provided by the present application;

[0046] Figure 11 is a structural schematic diagram of one embodiment of the electronic device provided by the present application;

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

[0048] 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. Based on the embodiments in the present application, all other embodiments obtained by those skilled in the art without creative labor fall within the scope of protection of the present application.

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

[0050] In addition, in the embodiments of the present application, a “dynamic sliding graph” is defined as a graph that is used to slide on a “static fixed graph” according to a specified sliding step during a convolution operation, and sequentially covers all possible positions. Correspondingly, a “static fixed graph” is defined as a graph that remains in a position during a convolution operation, and is used for a “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 a padding operation.

[0051] The following will expand on the inventor's basic inventive concept, that is, a convolution weight gradient calculation method based on an unsplit input feature map and an output gradient map, and not involving thread resources, to better understand the technical solution in the embodiments of the present invention:

[0052] In back propagation, the calculation formula of the convolution kernel weight gradient matrix is ​​usually: ;in, is the weight gradient matrix; It is the input feature map (Activation) in the forward propagation; It is the output gradient map (ActivationGradient) in back propagation; is the convolution operation; is the loss function. Depend on weight gradients; is the convolution weight The gradient, is the vertical offset between the convolution weight and the center of the convolution kernel, 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.

[0053] Taking a 3×3 weight gradient matrix as an example, the inventors discovered during their research on the calculation of convolution weight gradients that the related technology, Image to Column (Im2Col) method, requires 9 position shifts during calculation. Figure 1 FIG. 1 is a schematic diagram of an embodiment of convolution weight gradient calculation provided by the related art. Figure 1 In the figure, the yellow image represents the input feature map (as a dynamic sliding map), and the blue image represents the output gradient map (as a static fixed map). During the calculation process, the input feature map needs to be offset 9 times on the output gradient map in sequence to obtain 9 weight gradients respectively. Each time the position is offset, it is necessary to re-read the area data covered by the input feature map in the output gradient map (including zero-padding data), which results in a large amount of repeated data reading and significantly affects the calculation efficiency.

[0054] After further research, the inventors found that since the input feature map needs to perform 9 offset operations on the output gradient map, each data point in the input feature map will actually be dot-multiplied with the corresponding local area (size 3×3) in the output gradient map to obtain a local calculation result (also a 3×3 matrix), and each element in this local calculation result is the accumulated component of the corresponding weight gradient. Figure 2FIG. 1 is a schematic diagram of an embodiment of the data association relationship provided by the present invention. Figure 2 In the figure, the data point A0 in the input feature map is multiplied by the local area surrounded by the orange border in the output gradient map; the data point A3 is multiplied by the local area surrounded by the green border; the data point A7 is multiplied by the local area surrounded by the pink border; and so on, until all data points in the input feature map (i.e., A0 to A63) are traversed.

[0055] like Figure 3 FIG. 1 is a flow chart of an embodiment of the convolution weight gradient calculation method provided by the present invention. Figure 3 From the matrix calculation formula of the Im2Col method, we can see that the output gradient map is read 9 times, so there are 9 columns of data to form its expanded matrix; in this expanded matrix, each row corresponds to one of the local regions mentioned above. If the input feature map has 64 data points, then there are 64 local regions corresponding to them. Through the dot multiplication operation, 64 local calculation results (3×3 matrix) are obtained; and the elements at the corresponding positions in the 64 local calculation results are accumulated (that is, all matrices are added together) to obtain the weight gradient matrix of the convolution kernel. .

[0056] The inventors also found that 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 calculation size (such as Figure 2 The red dotted line division method in the middle) only needs to ensure that the divided sliding part and the fixed part meet the coverage constraint relationship of the convolution operation to achieve parallel decomposition of computing tasks (such as point multiplication and accumulation calculations).

[0057] The following uses a GPU as an example to illustrate the convolution weight gradient calculation method, device, computer-readable storage medium, electronic device, and computer program product provided by embodiments of the present invention.

[0058] See also Figure 4 , is a flow chart of another embodiment of the convolution weight gradient calculation method provided by the present invention.

[0059] A first embodiment of the present invention provides a convolution weight gradient calculation method, which is applied to an artificial intelligence processor. The method includes steps S11 to S15, which are specifically as follows:

[0060] Step S11: assign roles to the input feature map and the corresponding output gradient map in the global memory to determine the dynamic sliding map and the static fixed map;

[0061] Step S12: extracting the mth sliding portion from the dynamic sliding graph, performing boundary expansion reading on the corresponding area of ​​the static fixed graph to obtain the mth fixed portion, and loading the mth sliding portion and the mth fixed portion into a shared memory; wherein m≥1;

[0062] Step S13: Set the mth sliding part data rows, and the mth fixed part consists of Go to The row consists of Alignment data blocks are loaded into threads; among them, ;

[0063] Step S14: In each of the threads, a synchronous sliding traversal is performed in the loaded data row and the aligned data block with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result, and each element in the local calculation result is accumulated to the corresponding weight gradient position; wherein, the aligned data block is based on The window is traversed by sliding; is the height of the convolution kernel; is the width of the convolution kernel;

[0064] Step S15: performing traversal calculation on the data rows of all sliding parts and the corresponding aligned data blocks in the dynamic sliding graph to obtain the weight gradient matrix of the convolution kernel.

[0065] It should be noted that step S15 is equivalent to "repeating steps S12 to S14 until all the data rows of the sliding parts and the corresponding aligned data blocks in the dynamic sliding graph are traversed to obtain the weight gradient matrix of the convolution kernel." , we can use the input feature map In the output gradient map Perform sliding calculation to obtain , you can also use the output gradient map In the input feature map Sliding calculation is performed on In other words, if the input feature map is a dynamic sliding map, the output gradient map is the covered static fixed map; if the output gradient map is a dynamic sliding map, the input feature map is the covered static fixed map. Obviously, the static fixed map needs to be padded to ensure the effective calculation of the boundary data. If the dynamic sliding map is divided into multiple sliding parts, the fixed part corresponding to each sliding part needs to be obtained through the padding operation.

[0066] For example, Figure 5FIG. 1 is a schematic diagram of an embodiment of the sliding part and the fixed part provided by the present invention. Figure 5 In the example, the input feature map is a dynamic sliding map (size is 64×64), and the output gradient map is a static fixed map (size is 64×64). The block execution strategy is adopted. The sliding part is extracted by the size of , there are 8 in total. It should be noted that in order to simplify the diagram and avoid visual congestion, Figure 5 Only two sliding parts are selected for display, and the window size is scaled, not drawn strictly according to the actual size (such as Figure 5 The yellow area (sliding portion) in the upper-left subfigure is displayed as 5×5, but its actual size is 32×16. Threads T0 through T31 ​​shown in the figure represent 32 parallel threads, matching all the data rows (32 rows) actually contained in the yellow area (sliding portion).

[0067] exist Figure 5 In the example, since the fixed part must meet the requirements of the sliding part to offset on it, the fixed part needs to supplement the boundary expansion data (i.e., Padding operation) based on its corresponding area. Taking the convolution kernel size of 3×3 as an example, the size of the fixed part should be appropriately expanded based on the corresponding area of ​​the sliding part according to the convolution kernel's requirements for the boundary, that is, the size of the fixed part is The sliding and fixed parts can be asynchronously loaded from global memory to Group Shared Memory (GSM) in advance through the Tensor Memory Access (TMA) accelerator to increase the parallelism of computing and memory access.

[0068] If the shared memory has loaded the first sliding part (size ) and the first fixed part (size is ), 32 threads are enabled. To improve parallel efficiency, the arrangement direction of the threads is consistent with the vertical direction of the original data, that is, each thread is responsible for processing a row of data in the sliding part and the corresponding aligned data block. Figure 5 As shown, the first thread (T0) loads the first data row of the current sliding part and the aligned data block consisting of the first to third rows in the current fixed part; the second thread (T1) loads the second data row of the current sliding part and the aligned data block consisting of the second to fourth rows in the current fixed part; and so on, until the 32nd thread (T31) loads the 32nd data row of the current sliding part and the aligned data block consisting of the 32nd to 34th rows in the current fixed part.

[0069] Thus, each thread holds a row of data in the sliding section and can perform the alignment data block (derived from the fixed part). Each thread traverses the data points in the corresponding data row respectively. After 32 threads complete the traversal, it indicates that the data points in the current sliding part have been processed. The above loading operation and the traversal operation in the thread are repeated for the remaining 7 sliding parts. In other words, steps S12 to S14 are repeatedly executed until the traversal of the data points in all sliding parts in the dynamic sliding graph is completed, and the local calculation result is obtained Figure 5 The weight gradient matrix of the example convolution kernel.

[0070] As shown in the example, the weight gradient matrix of the convolution kernel is obtained by the local calculation result. Figure 6 As shown in the example, the weight gradient matrix of the convolution kernel is obtained by the local calculation result. Figure 6 In the data row on the left side, the 1x1 yellow area represents the data points that have been traversed, and the 1x1 pink area represents the current traversed data points. In the alignment data block on the right side, the gray area represents the boundary expanded data, the blue dashed box (3x3 in size) represents the sliding window in the alignment data block, the yellow area represents the data that has been traversed, and the pink area represents the new data that is currently traversed. When the specified sliding step is 1, the sliding window of the alignment data block has an overlapping characteristic. Therefore, the blue dashed box contains part of the old data that has been traversed and the new data that is currently traversed when it slides for the d-th time. Wherein, d>1.

[0071] Figure 6 It shows the data that a certain thread can hold (cached in registers), that is, registers R100-R115 store a data row (1x16 in size) in the sliding part, and registers R0-R53 store the corresponding alignment data block (3x18 in size). During the calculation, the thread traverses along the horizontal direction, that is, the data row and the alignment data block are synchronously slid with a step of 1. The sliding window of the alignment data block is 3x3 (same as the size of the convolution kernel), and the sliding window of the data row is 1x1. Each time the data in the current window is multiplied to obtain the corresponding local calculation result (similar to a 3x3 matrix in Figure 3 ), and each element in the local calculation result is accumulated to the corresponding weight gradient position. This can be realized by the FMADW instruction. The FMADW instruction is an instruction that simultaneously performs a multiplication and an addition operation, and the mathematical form is ; wherein, , and are operands (usually floating-point numbers); is the result. The accumulation results in the calculation process are sequentially saved in the accumulation temporary register, and then written to the shared memory / global memory after the accumulation calculation is completed.

[0072] When all sliding parts in the dynamic sliding graph and all fixed parts in the static fixed graph are synchronously traversed and calculated (similar to Figure 3 After summing all matrices in , the weight gradient matrix of the convolution kernel can be obtained; specifically, first in each thread, use the specified register (such as Figure 6 The weight gradient matrix is ​​cached in R80-R88 of the thread. After all sliding and fixed parts are traversed, the partial accumulation results cached in each thread are merged to obtain the weight gradient matrix of the convolution kernel. If the split task is assigned to different computation units (i.e., corresponding thread groups) for parallel execution, that is, each computation unit is responsible for its own sliding and fixed parts, the partial accumulation results of each thread must be completed within all computation units, and then the partial accumulation results of all threads across all units are aggregated to obtain the weight gradient matrix of the convolution kernel. Therefore, the aforementioned "accumulating each element of the local calculation result to the corresponding weight gradient location" can be achieved by first accumulating the local calculation result to the designated register of the corresponding thread. After all sliding and fixed parts are traversed, the partial accumulation results in each designated register are merged into the weight gradient matrix cache location in shared memory. Of course, it is also possible to avoid using designated registers for temporary storage of the partial accumulation results and directly accumulate each element of the local calculation result to the weight gradient matrix cache location in shared memory in real time. The specific solution to be adopted can be flexibly determined based on actual computing resources and performance requirements.

[0073] In an embodiment of the present invention, when the aligned data blocks held by each thread slide in the horizontal direction, data reuse is achieved, and there is no need to repeatedly load overlapping data areas from shared memory or global memory. Therefore, the embodiment of the present invention significantly reduces the repeated reading of data during the weight gradient calculation process through the thread horizontal preloading mechanism and the overlapping sliding window characteristics of the aligned data blocks. In particular, when the input feature map and the output gradient map are not blocked in the horizontal direction, the repeated reading of data only occurs in the vertical direction. Compared with the Im2Col method, the amount of repeated reading is reduced to 1 / 3 of the original, thereby greatly improving the computing efficiency and memory access performance.

[0074] In an optional embodiment, performing synchronous sliding traversal in the loaded data rows and aligned data blocks with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result includes:

[0075] According to the specified sliding step, the data rows in the sliding thread are synchronized and the windows of the aligned data blocks are aligned; wherein the data rows are based on The window is traversed by sliding;

[0076] After each sliding, the two window data involved in the current calculation are located, and the dot multiplication operation is performed to obtain the corresponding local calculation result; wherein the size of the local calculation result is .

[0077] It should be noted that, as shown in Figure 6 , the sliding window of the data row is traversed in units of 1x1, that is, each time the sliding window only covers a single data point of a single row, and the sliding window size of the aligned data block is consistent with the convolution kernel size. According to the specified sliding step, the window of the data row and the window of the aligned data block are slid simultaneously. Each time the two windows are located in a new data area, a dot multiplication operation is completed to obtain the local calculation result at the current sliding position, and the local calculation result is accumulated in the cache position of the weight gradient matrix. This synchronous sliding mechanism ensures that the data window involved in each dot multiplication operation always corresponds to the correct spatial position (similar to the data point A7 and the local area surrounded by the pink frame in Figure 2 ), ensuring the accuracy of the calculation.

[0078] In an optional embodiment, the number of divisions of the sliding part is determined by the computing resources of the artificial intelligence processor.

[0079] 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 sliding parts can be reasonably divided according to the current processor computing resources, so as to realize optimal allocation and efficient use of resources, and avoid the situation of idle thread resources or memory overflow.

[0080] For example, if the number of available threads of the processor is large and the shared memory capacity is sufficient, the size of the sliding part can be appropriately increased, and the number of divisions can be reduced. In the case of extremely abundant computing resources, the number of sliding parts can even be 1 (i.e. no division), so as to minimize the number of 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 sliding part can be appropriately 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 sliding part.

[0081] In an optional embodiment, the boundary expansion reading of the corresponding area of the static fixed graph to obtain the mth fixed part comprises:

[0082] determining the corresponding area of the mth sliding part in the static fixed graph; wherein the size of the corresponding area is equal to the size of the mth sliding part;

[0083] Padding the corresponding region according to the size of the convolution kernel to obtain the mth fixed part; wherein the expanded region in the mth fixed part comprises zero padding values and / or valid data of adjacent fixed parts.

[0084] As shown in Figure 2 and 5 , the constituent region of the fixed part comprises the initial corresponding region and the expanded region added by the padding operation. The data source of the expanded region comprises the following two cases:

[0085] (1) zero padding values, used to complete the part outside the static fixed graph;

[0086] (2) valid data of adjacent fixed parts, used to complete the internal boundary divided in the static fixed graph.

[0087] In an optional embodiment, all the threads load the corresponding data row and the aligned data block in parallel through the specified batch loading instruction.

[0088] As shown in Figure 7 , it is a schematic diagram of an embodiment of the thread loading aligned data block provided by the application, wherein each Part represents the process of the thread loading data in the vertical direction; for example, the left subgraph of Part1 represents the data region (1st row to 32nd row) loaded in parallel by 32 threads in the fixed part, the middle subgraph is a partial schematic diagram of the data points loaded in the specified registers of the 32 threads, and the right subgraph is the position of the data points loaded in the specified registers in the sliding window of the aligned data block. In Figure 7 , the size of the fixed part is 34x18, the number of enabled threads is 32, and the ldmatrix32rstrider18 instruction is used to load the aligned data block into the 32 threads efficiently. The function of the ldmatrix32rstrider18 instruction is: taking 32 threads as a group, each thread is responsible for loading 1 data, and 32 threads are executed simultaneously, and the address access of adjacent threads spans 18 coordinates (in the embodiment of the 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).

[0089] Specifically, first locate the loading starting point at the first row of the fixed part, and use the ldmatrix32rstrider18 instruction to load the data from the 1st to the 32nd position points of the first column of the fixed part into 32 threads in parallel (corresponding to the orange number column in Part1); continue to use the ldmatrix32r strider18 instruction to load the data from the 1st to the 32nd position points of the second column into 32 threads (corresponding to the yellow number column in Part1). And so on, until the data from the 1st to the 32nd position points of each column of the fixed part are loaded, that is, the loading of the first row of data of the aligned data block in all threads is completed. If you need to load subsequent rows of the aligned data block, you need to move the loading starting point down row by row. Figure 7 In Part 2, position the load starting point at the second row of the fixed portion and repeat the ldmatrix32r strider18 operation to load data from the second to the 33rd position of all columns into the 32 threads, completing the loading of the second row of the aligned data block in all threads. Similarly, position the load starting point at the third row of the fixed portion and execute the ldmatrix32r strider18 operation again to load data from the third to the 34th position of all columns into the 32 threads, completing the loading of the third row of the aligned data block in all threads, ultimately loading all rows of the aligned data block.

[0090] Similarly, when 32 threads load corresponding 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).

[0091] In an optional embodiment, each of the threads releases cache locations through a register reuse mechanism to scroll and load data in a corresponding data row and aligned data block.

[0092] 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 6 For example, if a thread does not have 54 registers when processing an aligned data block, horizontal register rolling reuse can be employed: after completing calculations on the current sliding window, the thread promptly releases the register space occupied by the calculated data no longer needed for the next sliding window. This freed register space is then used to store data not yet loaded into the thread, ensuring complete loading of the aligned data block. This register rolling reuse is particularly applicable when a static fixed graph and its corresponding dynamic sliding graph are not split horizontally.

[0093] See alsoFigure 8 is a flowchart of another embodiment of the convolution weight gradient calculation method provided by the present application.

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

[0095] Step S21: Role assignment is performed on the input feature tensor and the corresponding output gradient tensor in the global memory, and a dynamic sliding tensor and a static fixed tensor are determined; wherein the dynamic sliding tensor is composed of dynamic sliding maps of channels; and the static fixed tensor is composed of static fixed maps of channels. >1;

[0096] Step S22: The convolution weight gradient calculation method described in any one of the first aspect embodiments is performed on the dynamic sliding map and the static fixed map corresponding to each channel respectively, and a weight gradient matrix of the corresponding channel is obtained.

[0097] 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 the depth separable convolution, the size of the input feature tensor is 1x64x64x192 (i.e. batchxheightxwidthxchannel number), and the size of the output gradient tensor is also 1x64x64x192. By using the convolution weight gradient calculation method described in any one of the first aspect embodiments, each channel is processed independently, and a weight gradient matrix of the corresponding channel (192 weight gradient matrices corresponding to 192 channels) is obtained, which can greatly improve the parallelism and hardware resource utilization.

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

[0099] It should be noted that when the thread resources enabled by the processor are sufficient, the weight gradient matrices 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 calculate in parallel, the data blocks (with a size of 32x16x32) corresponding to the sliding part of the 32 channels and the data blocks (with a size of 34x18x32) corresponding to the fixed part can be loaded into the shared memory together to calculate the weight gradient matrices of the 32 channels in parallel. The embodiments of the present application can fully utilize the parallel computing capability of threads to realize efficient synchronous calculation of the weight gradient matrices of multiple channels.

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

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

[0102] A first role assignment module 11 is configured to assign roles to input feature maps and corresponding output gradient maps in a global memory, and determine dynamic sliding maps and static fixed maps.

[0103] A first loading module 12 is configured to extract an mth sliding part from the dynamic sliding map, and perform boundary extension reading on a corresponding region of the static fixed map to obtain an mth fixed part, and load the mth sliding part and the mth fixed part to a shared memory; wherein m≥1.

[0104] A second loading module 13 is configured to load a first data row of the mth sliding part, and a first aligned data block composed of rows to rows in the mth fixed part, to a first thread. ;

[0105] A thread calculation module 14 is configured to perform synchronous sliding traversal in a specified sliding step in the loaded data row and aligned data block in each thread to perform dot multiplication operation to obtain a local calculation result, and accumulate each element in the local calculation result to a corresponding weight gradient position respectively; wherein the aligned data block is slidingly traversed in a window of ; is the height of the convolution kernel; is the width of the convolution kernel.

[0106] A first result acquisition module 15 is configured to perform traversal calculation on data rows of all sliding parts in the dynamic sliding map and corresponding aligned data blocks to obtain a weight gradient matrix of the convolution kernel.

[0107] It should be noted that the convolution weight gradient calculation device provided by the third aspect of the present application can realize all processes of the convolution weight gradient calculation method described in any one 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 convolution weight gradient calculation method described in any one of the embodiments of the first aspect. Here, the functions and technical effects will not be repeated.

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

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

[0110] The second role assignment module 21 is configured to assign roles to the input feature tensor and the corresponding output gradient tensor in the global memory, and determine a dynamic sliding tensor and a static fixed tensor; wherein the dynamic sliding tensor is composed of dynamic sliding maps of the plurality of channels; and the static fixed tensor is composed of static fixed maps of the plurality of channels. >1;

[0111] The second result obtaining module 22 is configured to execute the convolution weight gradient calculation method of any one of the first aspect for each channel corresponding dynamic sliding map and static fixed map, and obtain the weight gradient matrix of the corresponding channel.

[0112] It should be noted that the convolution weight gradient calculation device provided by the fourth aspect of the present application can realize all processes of the convolution weight gradient calculation method of any one of the second aspect, 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 of any one of the second aspect, which will not be repeated here.

[0113] The fifth aspect of the present application provides a computer readable storage medium, the computer readable storage medium includes stored computer program;Wherein, the computer program when running controls the device where the computer readable storage medium is located executes the convolution weight gradient calculation method of any one of the above embodiments.

[0114] The sixth aspect of the present application provides a computer program product, including computer program, the computer program is executed when the processor realizes the convolution weight gradient calculation method of any one of the above embodiments.

[0115] Referring to Figure 11 , the structural schematic diagram of one embodiment of the electronic device provided by the present application.

[0116] The seventh aspect of the present application provides an electronic device, including processor 31, memory 32 and storage in the memory 32 and configured to be executed by the processor 31 computer program, the processor 31 when executing the computer program realizes the convolution weight gradient calculation method of any one of the above embodiments.

[0117] ​​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.

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

[0119] 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 related 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.

[0120] It should be noted that the above electronic device can include, but is not limited to, a processor, a memory, and those skilled in the art can understand that Figure 11 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.

[0121] Referring to Figure 12FIG. 1 is a structural schematic diagram of an embodiment of the artificial intelligence processor provided by the present application.

[0122] The artificial intelligence processor provided by the embodiment of the present application comprises a plurality of computing units (CUs). The CU is a core component for performing a computing task and can run a plurality of threads in parallel to perform the convolution weight gradient calculation method described in any of the above embodiments. The CU further comprises other hardware resources, such as a group shared memory (GSM) at the thread group level, a thread level register (TLR), and the like. The GSM is used to store the sliding part and the fixed part extracted from the dynamic sliding graph and the static fixed graph respectively. The TLR is used to cache the data row and the corresponding alignment data block 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 (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.

[0123] The above only describes the preferred embodiments 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, 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 convolution weight gradient calculation method, characterized in that: Applied to an artificial intelligence processor, the method includes: Assign roles to the input feature map and the corresponding output gradient map in the global memory to determine the dynamic sliding map and the static fixed map; Extracting an mth sliding portion from the dynamic sliding graph, performing boundary extension reading on a corresponding region of the static fixed graph to obtain an mth fixed portion, and loading the mth sliding portion and the mth fixed portion into a shared memory; wherein m≥1; The mth sliding part data rows, and the mth fixed part consists of Go to The row consists of Alignment data blocks are loaded into threads; among them, ; In each of the threads, a synchronous sliding traversal is performed in the loaded data row and the aligned data block with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result, and each element in the local calculation result is accumulated to the corresponding weight gradient position; wherein, the aligned data block is based on The window is traversed by sliding; is the height of the convolution kernel; is the width of the convolution kernel; The data rows of all sliding parts in the dynamic sliding graph and the corresponding aligned data blocks are traversed and calculated to obtain the weight gradient matrix of the convolution kernel.

2. The convolution weight gradient calculation method according to claim 1, wherein: The number of divisions of the sliding portion is determined by the computing resources of the artificial intelligence processor.

3. The convolution weight gradient calculation method according to claim 1, wherein: The step of performing boundary expansion reading on the corresponding area of ​​the static fixed image to obtain the mth fixed part includes: Determining a corresponding area of ​​the m-th sliding portion in the static fixed map; wherein a size of the corresponding area is equal to a size of the m-th sliding portion; Based on the size of the convolution kernel, the boundary of the corresponding area is expanded to obtain the m-th fixed part; wherein the expanded area in the m-th fixed part includes: zero filling value and / or valid data of adjacent fixed parts.

4. The convolution weight gradient calculation method according to claim 1, wherein: All of the threads load corresponding data rows and aligned data blocks in parallel through specified bulk loading instructions.

5. The convolution weight gradient calculation method according to claim 1, wherein: Each of the threads releases cache locations through a register multiplexing mechanism to scroll and load data in a corresponding data row and aligned data block.

6. The convolution weight gradient calculation method according to claim 1, wherein: The synchronous sliding traversal in the loaded data rows and aligned data blocks with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result includes: According to the specified sliding step, the data rows in the sliding thread are synchronized and the windows of the aligned data blocks are aligned; wherein the data rows are based on The window is traversed by sliding; After each slide, the two window data involved in the current calculation are located and the dot multiplication operation is performed to obtain the corresponding local calculation result; wherein the size of the local calculation result is .

7. A convolution weight gradient calculation method, characterized in that: Applied to an artificial intelligence processor, the method includes: Assign roles to the input feature tensor and the corresponding output gradient tensor in the global memory to determine the dynamic sliding tensor and the static fixed tensor; wherein the dynamic sliding tensor is composed of The dynamic sliding graph of channels is composed of The static fixed graph of the channels; >1; For each channel's corresponding dynamic sliding graph and static fixed graph, respectively, the convolution weight gradient calculation method according to any one of claims 1 to 6 is executed to obtain a weight gradient matrix for the corresponding channel.

8. The convolution weight gradient calculation method according to claim 7, wherein: The number of parallel calculations of the weight gradient matrix is ​​determined by the thread resources enabled by the artificial intelligence processor.

9. A convolution weight gradient calculation device, characterized in that: include: The first role assignment module is used to assign roles to the input feature map and the corresponding output gradient map in the global memory to determine the dynamic sliding map and the static fixed map; a first loading module, configured to extract an mth sliding portion from the dynamic sliding graph, perform boundary extension reading on a corresponding region of the static fixed graph to obtain an mth fixed portion, and load the mth sliding portion and the mth fixed portion into a shared memory; wherein m≥1; The second loading module is used to load the mth sliding part data rows, and the mth fixed part consists of Go to The row consists of Alignment data blocks are loaded into threads; among them, ; The thread calculation module is used to perform synchronous sliding traversal in the loaded data rows and aligned data blocks in each of the threads with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result, and to accumulate each element in the local calculation result to the corresponding weight gradient position; wherein, the aligned data block is based on The window is traversed by sliding; is the height of the convolution kernel; is the width of the convolution kernel; The first result acquisition module is used to traverse and calculate the data rows of all sliding parts in the dynamic sliding graph and the corresponding aligned data blocks to obtain the weight gradient matrix of the convolution kernel.

10. A convolution weight gradient calculation device, characterized in that: include: The second role assignment module is used to assign roles to the input feature tensor and the corresponding output gradient tensor in the global memory, and determine the dynamic sliding tensor and the static fixed tensor; wherein the dynamic sliding tensor is composed of The dynamic sliding graph of channels is composed of The static fixed graph of the channels; >1; The second result acquisition module is used to execute the convolution weight gradient calculation method as described in any one of claims 1 to 6 on the dynamic sliding graph and static fixed graph corresponding to each channel respectively, to obtain the weight gradient matrix of the corresponding channel.

11. A computer-readable storage medium, characterized in that The computer-readable storage medium includes a stored computer program; wherein, when the computer program is run, it controls the device where the computer-readable storage medium is located to execute the convolution weight gradient calculation method according to any one of claims 1 to 8.

12. A computer program product, characterized in that The invention comprises a computer program, which, when executed by a processor, implements the convolution weight gradient calculation method according to any one of claims 1 to 8.

13. An electronic device, characterized in that: The method comprises a processor, a memory, and a computer program stored in the memory and configured to be executed by the processor, wherein the processor implements the convolution weight gradient calculation method according to any one of claims 1 to 8 when executing the computer program.

Images