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

Abstract

The invention discloses a convolution weight gradient calculation method and device, a medium, equipment and a product, and the method comprises the steps: carrying out the role distribution of an input feature map and a corresponding output gradient map in a global memory, and determining a dynamic sliding map and a static fixed map; respectively extracting a sliding part and a corresponding fixed part from the dynamic sliding graph and the static fixed graph, and enabling each thread to hold a data row and a corresponding aligned data block in the sliding part; and in each thread, carrying out synchronous sliding and point multiplication operation on the held data row and the aligned data block to obtain a local calculation result of a current sliding position, accumulating the local calculation result to a cache position of the weight gradient matrix, and after synchronous traversal calculation of all sliding parts and corresponding fixed parts is completed, obtaining the weight gradient matrix of the convolution kernel. According to the method, data multiplexing in the sliding direction can be effectively realized, and the data volume read repeatedly is reduced, so that the calculation efficiency of the weight gradient is improved.

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: 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.

[0005] Optionally, the number of divisions of the sliding portion is determined by computing resources of the artificial intelligence processor.

[0006] Optionally, performing boundary extension reading on a corresponding area of ​​the static fixed image to obtain the mth fixed portion 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.

[0007] Optionally, all the threads load corresponding data rows and aligned data blocks in parallel through a specified batch loading instruction.

[0008] Optionally, 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.

[0009] Optionally, 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: 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 .

[0010] A second aspect of the present invention provides a convolution weight gradient calculation method, which is 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 the dynamic sliding graph and static fixed graph corresponding to each channel, execute the convolution weight gradient calculation method described in any embodiment of the first aspect to obtain the weight gradient matrix of the corresponding channel.

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

[0012] A third embodiment of the present invention provides a convolution weight gradient calculation device, comprising: 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.

[0013] A fourth aspect of the present invention provides a convolution weight gradient calculation device, comprising: 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 described in any embodiment of the first aspect on the dynamic sliding graph and static fixed graph corresponding to each channel respectively, to obtain the weight gradient matrix of the corresponding channel.

[0014] A fifth aspect of the present invention provides a computer-readable storage medium, which 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 convolution weight gradient calculation method described in any of the above embodiments.

[0015] A sixth aspect of the present invention provides a computer program product, comprising a computer program, which, when executed by a processor, implements the convolution weight gradient calculation method described in any of the above embodiments.

[0016] An embodiment of the seventh aspect of the present invention 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 implements the convolution weight gradient calculation method described in any of the above embodiments when executing the computer program.

[0017] Compared with the prior art, the embodiments of the present invention provide a convolution weight gradient calculation method, device, computer-readable storage medium, electronic device and computer program product. The embodiments of the present invention first 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; extract the sliding part and the corresponding fixed part from the dynamic sliding map and the static fixed map respectively, and let each thread hold a data row and a corresponding aligned data block in the sliding part; in each thread, the held data row and aligned data block are synchronously slid and dot multiplied to obtain the local calculation result of the current sliding position, and accumulate it to the cache position of the weight gradient matrix. When all sliding parts in the dynamic sliding map and all fixed parts in the static fixed map are synchronously traversed and calculated, the weight gradient matrix of the convolution kernel can be obtained. The present invention can achieve data reuse through the thread-level preloading mechanism and the overlapping sliding window characteristics of the aligned data blocks, without repeatedly loading overlapping data areas from shared memory or global memory, thereby significantly reducing the repeated reading of data during the weight gradient calculation process and greatly improving the computing efficiency and memory access performance. BRIEF DESCRIPTION OF THE DRAWINGS

[0018] Figure 1 It is a schematic diagram of an embodiment of convolution weight gradient calculation provided by the related art; Figure 2 is a schematic diagram of an embodiment of the data association relationship provided by the present invention; Figure 31 is a flow chart of an embodiment of a convolution weight gradient calculation method provided by the present invention; Figure 4 1 is a flow chart of another embodiment of the convolution weight gradient calculation method provided by the present invention; Figure 5 is a schematic diagram of an embodiment of the sliding part and the fixed part provided by the present invention; Figure 6 is a schematic diagram of an embodiment of obtaining local calculation results provided by the present invention; Figure 7 It is a schematic diagram of an embodiment of thread loading aligned data blocks provided by the present invention; Figure 8 1 is a flow chart of another embodiment of the convolution weight gradient calculation method provided by the present invention; Figure 9 1 is a schematic structural diagram of an embodiment of a convolution weight gradient calculation device provided by the present invention; Figure 10 1 is a structural diagram of another embodiment of the convolution weight gradient calculation device provided by the present invention; Figure 11 is a structural diagram of an embodiment of an electronic device provided by the present invention; Figure 12 1 is a schematic structural diagram of an embodiment of an artificial intelligence processor provided by the present invention. DETAILED DESCRIPTION

[0019] The following will clearly and completely describe the technical solutions in the embodiments of the present invention in conjunction with the accompanying drawings. Obviously, the described embodiments are only part of the embodiments of the present invention, not all of the embodiments. All other embodiments obtained by ordinary technicians in this technical field based on the embodiments of the present invention without making any creative efforts shall fall within the scope of protection of the present invention.

[0020] The artificial intelligence processor involved in the present invention 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 (Deep Learning Processing Unit), an APU (Accelerated Processing Unit), and a GPGPU (General-Purpose computing on Graphics Processing Unit), which can be determined when the embodiments of the present invention are applied to specific products or technologies.

[0021] Furthermore, in embodiments of the present invention, a "dynamic sliding image" is defined as an image that slides across a "static fixed image" at a specified sliding step size during a convolution operation, sequentially covering all possible positions. Accordingly, a "static fixed image" is defined as an image that remains stationary during the convolution operation, allowing the "dynamic sliding image" to slide across the entire spatial dimensions (including width and height) of the "static fixed image." Therefore, padding operations are required for the "static fixed image."

[0022] 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: 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.

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

[0024] 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 2 FIG. 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.

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

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

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

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

[0029] 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: 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; 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; 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, ; 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; 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.

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

[0031] For example, Figure 5 FIG. 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 crowding, 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).

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

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

[0034] Thus, each thread holds a row of data in the sliding section and can perform Each thread traverses and calculates the data points in the corresponding data row. After all 32 threads are completed, it means that the data points in the current sliding part have been processed. Repeat the above loading operation and the traversal calculation operation within the thread for the remaining 7 sliding parts. In other words, repeat steps S12 to S14 until the data points of all sliding parts in the dynamic sliding graph are traversed. Figure 5 The weight gradient matrix of the convolution kernel of the example.

[0035] For example, Figure 6 FIG. 1 is a schematic diagram of an embodiment of obtaining local calculation results provided by the present invention; FIG. Figure 6 In the data row on the left, the 1×1 yellow area represents the data points that have been traversed, and the 1×1 pink area represents the data points currently traversed. In the aligned data block on the right, the gray area represents the data with extended boundaries, and the blue dashed box (size 3×3) represents the window sliding in the aligned data block. The yellow area represents the data that has been traversed, and the pink area is the new data currently traversed. When the specified sliding step size is 1, the sliding window of the aligned data block has an overlapping characteristic. Therefore, when the blue dashed box slides for the dth time, it will contain both part of the old data that has been traversed and the new data currently traversed; where d>1.

[0036] Figure 6 The data that a thread can hold (cached in registers) is shown, that is, registers R100~R115 store a data row in the sliding part (size 1×16), and registers R0~R53 store the corresponding aligned data block (size 3×18). During the calculation process, the thread traverses in the horizontal direction, that is, it slides synchronously on the data row and aligned data block with a step size of 1; the aligned data block sliding window is 3×3 (the same as the convolution kernel size); the data row sliding window is 1×1. Each time it slides, a dot multiplication operation is performed on the data at the current window to obtain the corresponding local calculation result (similar to Figure 3 ), and accumulate each element of the local calculation result to the corresponding weight gradient position, which can be achieved through the FMADW instruction. The FMADW instruction is an instruction that performs a multiplication and an addition operation at the same time. The mathematical form is ;in, 、 and is the operand (usually a floating point number); The accumulated results during the calculation process are stored in the accumulated temporary registers in sequence until the accumulation calculation is completed and then written out to the shared memory / global memory.

[0037] 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 6The 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.

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

[0039] 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: 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 .

[0040] It should be noted that if Figure 6As shown, the sliding window of the data row is traversed in units of 1×1, that is, each sliding window only covers one data point in 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 size, the window of the data row in the thread and the window of the aligned data block are slid at the same time. Each time they slide, the two windows will locate the new data area, complete a dot multiplication operation, obtain the local calculation result of the current sliding position, and accumulate the local calculation result to the cache position of the weight gradient matrix. This synchronous sliding mechanism ensures that the data window involved always corresponds to the correct spatial position in each dot multiplication operation (similar to Figure 2 The data point A7 in the figure and the local area surrounded by the pink border ensure the accuracy of the calculation.

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

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

[0043] For example, if the processor has a large number of available threads and ample shared memory capacity, the size of the sliding part can be appropriately increased and the number of partitions can be reduced. If computing resources are extremely abundant, the number of sliding parts can even be 1 (i.e., no partitioning is performed) to minimize the number of data loads from global memory and improve data access efficiency. Conversely, if factors such as shared memory capacity, the number of threads, or register resources are limited, the size of the sliding part can be reduced and the number of partitions increased to avoid resource contention and overflow, thereby ensuring that the computation of each sliding part can be completed efficiently.

[0044] In an optional embodiment, performing boundary extension 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.

[0045] like Figure 2 and 5As shown in the figure, the fixed part consists of the initial corresponding area and the extended area added by the padding operation. The data sources of the extended area include the following two situations: (1) Zero filling value, used to fill in the part outside the static fixed graph; (2) Valid data of adjacent fixed parts is used to complete the internal boundaries of the static fixed graph.

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

[0047] like Figure 7 The figure shows an embodiment of thread loading aligned data blocks provided by the present invention. Each Part in the figure represents the process of threads loading data in the vertical direction. For example, the left sub-figure of Part 1 represents the data area (rows 1 to 32) loaded in parallel by 32 threads in a fixed part. The middle sub-figure is a partial schematic diagram of the data points loaded into the designated registers of the 32 threads, and the right sub-figure shows the positions of the data points loaded into these designated registers in the aligned data block sliding window. Figure 7 In this example, the fixed portion is 34×18 in size, the number of enabled threads is 32, and the ldmatrix32r strider18 instruction is used to efficiently load aligned data blocks into the 32 threads. The ldmatrix32rstrider18 instruction functions as follows: 32 threads are grouped together, each thread is responsible for loading a single piece of data, and all 32 threads execute simultaneously, with adjacent thread address accesses spanning 18 coordinates (in this embodiment of the present invention, this is equivalent 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).

[0048] 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 7In 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.

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

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

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

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

[0053] A second embodiment of the present invention provides a convolution weight gradient calculation method, which is applied to an artificial intelligence processor. The method includes steps S21 to S22, which are specifically as follows: Step S21: 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; Step S22: Execute the convolution weight gradient calculation method described in any embodiment of the first aspect above for the dynamic sliding graph and the static fixed graph corresponding to each channel respectively to obtain the weight gradient matrix of the corresponding channel.

[0054] 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 1×64×64×192 (i.e., batch×height×width×number of channels), and the size of the output gradient tensor is also 1×64×64×192. By using the convolution weight gradient calculation method described in any embodiment of the first aspect above to independently process each channel and obtain the weight gradient matrix of the corresponding channel (192 channels correspond to 192 weight gradient matrices), it can greatly improve parallelism and hardware resource utilization.

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

[0056] It should be noted that when the processor has sufficient thread resources enabled, the weight gradient matrices of multiple channels can be calculated in parallel. Taking the sliding part size of 32×16 as an example, when allowing parallel calculation of 32 channels, the data blocks corresponding to the sliding part of these 32 channels (size 32×16×32) and the data blocks corresponding to the fixed part (size 34×18×32) can be loaded into shared memory together to calculate the weight gradient matrices of the 32 channels in parallel. This embodiment of the present invention can fully utilize the parallel computing capabilities of threads to achieve efficient synchronous calculation of multi-channel weight gradient matrices.

[0057] See also Figure 9 , is a structural diagram of an embodiment of the convolution weight gradient calculation device provided by the present invention.

[0058] A third embodiment of the present invention provides a convolution weight gradient calculation device, comprising: The first role assignment module 11 is used to assign roles to 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; A first loading module 12 is configured to extract an mth sliding portion from the dynamic sliding graph, perform boundary expansion 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 13 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 14 is used to perform synchronous sliding traversal in the loaded data row and the aligned data block in each of the threads with a specified sliding step size to perform a dot multiplication operation to obtain a local calculation result, and 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 15 is used to perform 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.

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

[0060] See also Figure 10 , is a structural diagram of another embodiment of the convolution weight gradient calculation device provided by the present invention.

[0061] A fourth aspect of the present invention provides a convolution weight gradient calculation device, comprising: The second role assignment module 21 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 22 is used to execute the convolution weight gradient calculation method described in any embodiment of the first aspect for the dynamic sliding graph and static fixed graph corresponding to each channel, and obtain the weight gradient matrix of the corresponding channel.

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

[0063] A fifth aspect of the present invention provides a computer-readable storage medium, which 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 convolution weight gradient calculation method described in any of the above embodiments.

[0064] A sixth aspect of the present invention provides a computer program product, comprising a computer program, which, when executed by a processor, implements the convolution weight gradient calculation method described in any of the above embodiments.

[0065] See also Figure 11 , is a structural diagram of an embodiment of an electronic device provided by the present invention.

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

[0067] Preferably, the computer program can be divided into one or more modules / units (e.g., computer program 1, computer program 2, etc.). These one or more modules / units are stored in the memory 32 and executed by the processor 31 to implement the present invention. The one or more modules / units can be a series of computer program instruction segments capable of performing specific functions, and these instruction segments are used to describe the execution process of the computer program in the electronic device.

[0068] The processor 31 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 (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 and connects the various components of the electronic device using various interfaces and lines.

[0069] The memory 32 primarily includes a program storage area and a data storage area. The program storage area can store an operating system, at least one application required for a function, and the data storage area can store related data. Furthermore, the memory 32 can be a high-speed random access memory or a non-volatile memory, such as a plug-in hard drive, a Smart Media Card (SMC), a Secure Digital (SD) card, or a flash card. Alternatively, the memory 32 can be another volatile solid-state memory device.

[0070] It should be noted that the above electronic device may include, but is not limited to, a processor and a memory. Those skilled in the art will understand that Figure 11 The structural block diagram shown is only an example of the structure of the above-mentioned electronic device and does not constitute a structural limitation of the above-mentioned electronic device. The above-mentioned electronic device may include more or fewer components than shown in the figure, or combine certain components, or different components.

[0071] See also Figure 12 , is a structural diagram of an embodiment of the artificial intelligence processor provided by the present invention.

[0072] The artificial intelligence processor provided by an embodiment of the present invention includes multiple computing units (CUs). The CU is a core component for executing computing tasks and can run multiple threads in parallel to execute the convolution weight gradient calculation method described in any of the above embodiments. The CU also includes other hardware resources, such as thread group-level shared memory (GSM) and thread level registers (TLRs). The GSM is used to store the sliding part and the fixed part extracted from the dynamic sliding map and the static fixed map respectively; the TLR is used to cache the data rows and corresponding aligned data blocks loaded by the thread during the calculation process, as well as to temporarily store intermediate calculation results. In addition, the artificial intelligence processor also includes a global memory (GLM) for storing the input feature map and output gradient map in any embodiment of the first aspect, or the input feature tensor and output gradient tensor in any embodiment of the second aspect.

[0073] The above is only a preferred embodiment of the present invention. It should be pointed out that for ordinary technicians in this technical field, several improvements and modifications can be made without departing from the technical principles of the present invention. These improvements and modifications should also be regarded as the scope of protection of the present invention.

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 the dynamic sliding graph and the static fixed graph corresponding to each channel, respectively, the convolution weight gradient calculation method according to any one of claims 1 to 6 is executed to obtain the weight gradient matrix of 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