Large-scale 2D convolution operator acceleration method for DSP platform

Abstract

This invention discloses a large-scale two-dimensional convolution operator acceleration method for DSP platforms, belonging to the field of digital signal processing. This method selects the im2col or col2im algorithm to rearrange data based on the convolution type; adopts a three-segment matrix partitioning strategy to adapt to L3 cache capacity; utilizes an EDMA buffer ping-pong architecture to achieve pipeline parallelism for data transmission and computation; implements SIMD instruction-level optimization, using DMPYSP and DADDSP instructions in conjunction with pipeline optimization to achieve four FP32 multiplication and addition operations per cycle; and implements multi-core parallel scheduling, using OpenMP to achieve task-level and data-level parallelism. This method has been tested on a TI TMS320C6678 platform to achieve efficient inference of SAR target detection networks, providing a feasible solution for real-time inference of CNN networks on DSP platforms.

Description

Technical Field

[0001] The present invention relates to the technical field of digital signal processing, and in particular to a large-scale two-dimensional convolution operator acceleration method for a DSP platform. Background Art

[0002] Synthetic Aperture Radar (SAR), with its all-weather, all-day, multi-polarization, and high-resolution capabilities, has made significant contributions to both military and civilian applications. It overcomes the limited azimuth resolution of traditional radar, enabling the acquisition of high-resolution images over large areas at long ranges. It is widely used for high-resolution, real-time imaging on airborne and satellite-based platforms. Furthermore, with the recent advancements in machine learning and deep learning in target detection technology, convolutional neural network (CNN) models, such as ResNet, have become widely used in SAR target detection. The trend is to deploy lightweight SAR target detection networks on low-power, compact, and inexpensive hardware platforms, making real-time SAR target detection on airborne and unmanned aerial vehicles (UAVs) a pressing need.

[0003] DSP (Digital Signal Processor), a high-speed electronic device designed specifically for signal processing, is very suitable for real-time processing of SAR target detection and is also widely used in real-time signal processing on airborne platforms.

[0004] Unlike traditional hardware platforms, implementing real-time object detection algorithms based on convolutional neural networks on DSPs presents several challenges. Convolutional neural networks already have a variety of mature implementation frameworks for CPUs and GPUs. These frameworks typically utilize the NVIDIA cudnn library for GPU acceleration, while CPU implementations utilize multithreading acceleration using OpenMP and other methods. Numerous off-the-shelf acceleration libraries are also available. Therefore, the primary acceleration platforms for deep learning CNNs are mostly GPU or FPGA hardware platforms. Due to the inherent disadvantages of DSPs versus GPUs or FPGAs in parallel computing, traditional deep learning frameworks (such as TensorFlow, Caffe, and PyTorch) lack support for DSP platforms and lack a corresponding ecosystem.

[0005] Therefore, implementing real-time inference of convolutional neural networks on a DSP requires optimization efforts, starting from the underlying DSP hardware and the underlying algorithms of the operators. Common algorithmic implementations of two-dimensional convolution operators include data rearrangement followed by matrix multiplication, traditional two-dimensional convolution combined with a sliding window, and the use of Winograd transforms. Typical data rearrangement methods incur additional memory and performance overhead. Traditional two-dimensional convolution combined with a sliding window implementation incurs additional performance overhead due to repeated movement of the same data. While the use of Winograd transforms can reduce multiplication operations during inference, they also increase addition operations. For DSPs optimized for multiplication and multiply-add instructions, Winograd transforms do not improve efficiency. Furthermore, the relative complexity of traditional two-dimensional convolution operations makes them less suitable for optimization using SIMD (Single Instruction Multiple Data). Therefore, the implementation of data rearrangement followed by matrix multiplication is worth considering. Summary of the Invention

[0006] The present invention provides a large-scale two-dimensional convolution operator acceleration method for DSP platforms. It can realize efficient reasoning of SAR target detection networks on the TITMS320C6678 platform and provide a feasible solution for real-time reasoning of CNN networks on DSP platforms.

[0007] The embodiment of the present invention provides a large-scale two-dimensional convolution operator acceleration method for a DSP platform, comprising the following steps:

[0008] 1) Dynamic data rearrangement: select the im2col algorithm or col2im algorithm to rearrange data according to the convolution type;

[0009] 2) Cache management strategy: Adopt a three-segment matrix partitioning strategy to adapt to the L3 cache capacity;

[0010] 3) Using EDMA (Enhanced Direct Memory Access) buffer ping-pong architecture to parallelize the computation and transmission pipelines;

[0011] 4) SIMD (Single Instruction Multiple Data) instruction-level optimization, based on DSP-based DMPYSP (two sets of single-precision floating-point multiplication instructions) and DADDSP (two sets of single-precision floating-point addition instructions) instructions combined with pipeline optimization to achieve four FP32 multiplication and addition operations in a single cycle;

[0012] 5) Multi-core parallel scheduling, combined with OpenMP dynamic load balancing technology to achieve multi-core parallel acceleration.

[0013] Optionally, in one embodiment of the present invention, selecting the im2col algorithm or the col2im algorithm to rearrange data according to the convolution type includes:

[0014] For standard two-dimensional convolution, the im2col algorithm is used to rearrange the input data from a tensor of shape (IC, IH, IW) to a matrix of shape (IC*KH*KW, OH*OW), and the result is obtained through GEMM (General Matrix Multiplication) operation;

[0015] For 1x1 convolution, skip im2col and directly call GEMM operation to obtain the result;

[0016] For two-dimensional transposed convolution, the weights are pre-converted into a matrix of shape (OC*KH*KW, IC). During inference, GEMM is directly called and the col2im algorithm is used to rearrange the result from a matrix of shape (OC*KH*KW, IH*IW) back to a tensor of shape (OC, OH, OW), where IC, IH, and IW refer to the number of input image channels, input image height, and input image width, KH and KW refer to the convolution kernel height and convolution kernel width, and OC, OH, and OW refer to the number of output image channels, output image height, and output image width.

[0017] Optionally, in one embodiment of the present invention, a three-segment matrix partitioning strategy is used to adapt the L3 cache capacity, including:

[0018] The first block division: divide the M×N output matrix into multiple M1×N1 sub-matrices;

[0019] The second block division: divide the K×N input matrix into multiple K1×N1 sub-matrices;

[0020] The third block division: divide the M×K weight matrix into multiple M1×K1 sub-matrices;

[0021] M1, N1 and K1 meet the following conditions:

[0022] (M1 * N1 + K1 * N1 + M1 * K1) * sizeof(data_type) <L3_Cache_Size *0.8;

[0023] Where M1, N1, and K1 are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively. sizeof(data_type) refers to the size of the data type of the matrix elements, in bytes. L3_Cache_Size refers to the size of the available L3 cache, in bytes.

[0024] Optionally, in one embodiment of the present invention, an EDMA buffer ping-pong architecture is used to execute the calculation and transmission pipelines in parallel, including:

[0025] Cache allocation strategy:

[0026] The three-segment blocks allocated in the L3_Cache by the three-segment matrix block strategy are copied to obtain two three-segment caches with exactly the same size and shape, which are called cache ping and cache pong respectively. The suffix ping means cache ping and the suffix pong means cache pong.

[0027] M_ping, N_ping, and K_ping are required to meet the following conditions:

[0028] (M_ping * N_ping + K_ping * N_ping + M_ping * K_ping) * sizeof(data_type) <L3_Cache_Size * 0.4;

[0029] Where M_ping, N_ping, and K_ping are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively; sizeof(data_type) refers to the size of the data type of the matrix element, in bytes; L3_Cache_Size refers to the size of the available L3 cache, in bytes;

[0030] At the same time, the values ​​of M_pong, N_pong, K_pong and M_ping, N_ping, K_ping are exactly the same;

[0031] EDMA and GEMM combination strategy:

[0032] (1) Select the blocks of output matrix, input matrix and weight matrix according to the three-segment matrix block strategy;

[0033] (2) Select cache ping / pong;

[0034] (3) EDMA transfers the block input sub-matrix and weight sub-matrix from DDR (Double Data Rate) to the buffer ping / pong;

[0035] (4) GEMM calculation to obtain the output sub-matrix;

[0036] (5) EDMA transfers the block output sub-matrix from the buffer ping / pong to DDR;

[0037] (6) If the sub-matrix has not been traversed completely, return to (1) and exchange the cache ping-pong until the sub-matrix traversal is completed and the calculation of the entire matrix is ​​completed.

[0038] Optionally, in one embodiment of the present invention, DSP-based DMPYSP and DADDSP instructions are combined with pipeline optimization to implement four FP32 multiplication and addition operations in a single cycle, including:

[0039] Use DSP SIMD instructions to optimize multiplication, addition, and data transmission in GEMM.

[0040] Optionally, in one embodiment of the present invention, multi-core parallel scheduling is combined with OpenMP dynamic load balancing technology to achieve multi-core parallel acceleration, including:

[0041] Task-level parallelism: The matrix multiplication calculation of the output sub-matrix of each matrix block is regarded as a task unit, and the task is evenly distributed to multiple cores. Each DSP core is responsible for calculating the matrix multiplication task of the assigned output sub-matrix. The corresponding input sub-matrices and output sub-matrices are all placed in the L2_Cache of each core. It is suitable for DSPs without L3 shared cache. In this case, the sub-matrix size meets the following conditions:

[0042] (M1 * K + N1 * K + M1 * N1) * sizeof(data_type) <L2_Cache_Size * 0.8;

[0043] Where M1, N1, and K1 are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively; sizeof(data_type) refers to the size of the data type of the matrix element, in bytes; L2_Cache_Size refers to the size of the available L2 cache, in bytes;

[0044] If cache ping-pong is used, the sub-matrix size must meet the following conditions:

[0045] (M_ping * K + N_ping * K + M_ping * N_ping) * sizeof(data_type) <L2_Cache_Size * 0.4;

[0046] Where M_ping, N_ping, and K_ping are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively; sizeof(data_type) refers to the size of the data type of the matrix element, in bytes; L2_Cache_Size refers to the size of the available L2 cache, in bytes;

[0047] At the same time, the values ​​of M_pong, N_pong, K_pong and M_ping, N_ping, K_ping are exactly the same;

[0048] Data-level parallelism: Directly use the OpenMP precompilation directive #pragma omp parallel for to perform multi-core parallel optimization on the innermost loop of GEMM, which is suitable for DSPs with multi-core shared cache.

[0049] The large-scale two-dimensional convolution operator acceleration method for DSP platforms according to an embodiment of the present invention can achieve efficient inference of SAR target detection networks on the TITMS320C6678 platform, providing a feasible solution for real-time inference of CNN networks on DSP platforms.

[0050] Additional aspects and advantages of the present invention will be set forth in part in the description which follows and, in part, will be obvious from the description which follows, or may be learned through practice of the present invention. BRIEF DESCRIPTION OF THE DRAWINGS

[0051] The above and / or additional aspects and advantages of the present invention will become apparent and readily understood from the following description of the embodiments in conjunction with the accompanying drawings, in which:

[0052] Figure 1 A flowchart of an overall process of optimizing a two-dimensional convolution operator according to an embodiment of the present invention is provided;

[0053] Figure 2 and Figure 3 A schematic diagram of data rearrangement according to an embodiment of the present invention;

[0054] Figure 4 Schematic diagram of matrix multiplication blocks according to an embodiment of the present invention;

[0055] Figure 5 A schematic diagram of cache ping-pong according to an embodiment of the present invention;

[0056] Figure 6 and Figure 7 The embodiment of the present invention is the SAR target detection result of the UAV scene. DETAILED DESCRIPTION

[0057] The following describes embodiments of the present invention in detail, examples of which are shown in the accompanying drawings, wherein the same or similar reference numerals throughout represent the same or similar elements or elements having the same or similar functions. The embodiments described below with reference to the accompanying drawings are exemplary and are intended to be used to explain the present invention, and are not to be construed as limiting the present invention.

[0058] Specifically, Figure 1 A flowchart of a large-scale two-dimensional convolution operator acceleration method for a DSP platform provided by an embodiment of the present invention.

[0059] like Figure 1As shown in FIG, the large-scale two-dimensional convolution operator acceleration method for DSP platform includes the following steps:

[0060] 1) Dynamic data rearrangement: select the im2col algorithm or col2im algorithm to rearrange data according to the convolution type;

[0061] 2) Cache management strategy: Adopt a three-segment matrix partitioning strategy to adapt to the L3 cache capacity;

[0062] 3) Using the EDMA buffer ping-pong architecture, the calculation and transmission pipelines are executed in parallel;

[0063] 4) SIMD instruction-level optimization, based on DSP's DMPYSP and DADDSP instructions combined with pipeline optimization to achieve four FP32 multiplication and addition operations in a single cycle;

[0064] 5) Multi-core parallel scheduling, achieving task-level and data-level parallelism through OpenMP.

[0065] In one embodiment of the present invention, selecting the im2col algorithm or the col2im algorithm to rearrange data according to the convolution type includes:

[0066] For standard two-dimensional convolution, the im2col algorithm is used to rearrange the input data from a tensor of shape (IC, IH, IW) to a matrix of shape (IC*KH*KW, OH*OW), and the result is obtained through GEMM operation;

[0067] For 1x1 convolution, skip im2col and directly call GEMM operation to obtain the result;

[0068] For two-dimensional transposed convolution, the weights are pre-converted into a matrix of shape (OC*KH*KW, IC). During inference, GEMM is directly called and the col2im algorithm is used to rearrange the result from a matrix of shape (OC*KH*KW, IH*IW) back to a tensor of shape (OC, OH, OW), where IC, IH, and IW refer to the number of input image channels, input image height, and input image width, KH and KW refer to the convolution kernel height and convolution kernel width, and OC, OH, and OW refer to the number of output image channels, output image height, and output image width.

[0069] In one embodiment of the present invention, a three-segment block algorithm is used to adapt the GEMM computing unit to the L3 shared cache capacity and reduce the DDR access frequency to the theoretical lower limit, including:

[0070] The first block division: divide the M×N output matrix into multiple M1×N1 sub-matrices;

[0071] The second block division: divide the K×N input matrix into multiple K1×N1 sub-matrices;

[0072] The third block division: divide the M×K weight matrix into multiple M1×K1 sub-matrices;

[0073] M1, N1 and K1 are required to meet the following conditions:

[0074] (M1 * N1 + K1 * N1 + M1 * K1) * sizeof(data_type) <L3_Cache_Size *0.8;

[0075] Where M1, N1, and K1 are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively. sizeof(data_type) refers to the size of the data type of the matrix element, in bytes. For example, the size of a single-precision floating-point data type is 4. L3_Cache_Size refers to the size of the available L3 cache, in bytes.

[0076] M1, N1, and K1 should minimize the number of sub-matrix blocks according to the actual sizes of M, N, and K while meeting the above conditions to reduce the overhead caused by the blocks. The three blocks are allocated in the L3_Cache to accelerate the GEMM operation, and finally 20% of the L3_Cache is reserved for other uses.

[0077] For example, if the data_type of GEMM is FP32 and the L3_Cache_Size is 4MB, M1, N1, and K1 are 128, 512, and 256, respectively. In this case, the left side of the inequality is 0.875MB, which is smaller than the right side, 3.2MB, and satisfies the condition.

[0078] In one embodiment of the present invention, the EDMA buffer ping-pong architecture is used to execute the calculation and transmission pipelines in parallel, including:

[0079] Cache allocation strategy:

[0080] The three-segment blocks allocated in the L3_Cache by the three-segment matrix block strategy are copied to obtain two three-segment caches with exactly the same size and shape, which are called cache ping and cache pong respectively;

[0081] M_ping, N_ping, and K_ping are required to meet the following conditions:

[0082] (M_ping * N_ping + K_ping * N_ping + M_ping * K_ping) * sizeof(data_type) <L3_Cache_Size * 0.4;

[0083] Where M_ping, N_ping, and K_ping are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively. sizeof(data_type) refers to the size of the data type of the matrix elements, in bytes. L3_Cache_Size refers to the size of the available L3 cache, in bytes.

[0084] At the same time, the values ​​of M_pong, N_pong, K_pong and M_ping, N_ping, K_ping are exactly the same.

[0085] For example, if the data_type of GEMM is FP32 and the L3_Cache_Size is 4MB, M_ping, N_ping, and K_ping are 128, 512, and 256, respectively. In this case, the left side of the inequality is 0.875MB, which is smaller than the right side, 1.6MB, and thus satisfies the condition.

[0086] EDMA and GEMM combination strategy:

[0087] (1) Select the blocks of output matrix, input matrix and weight matrix according to the three-segment matrix block strategy;

[0088] (2) Select cache ping / pong;

[0089] (3) EDMA transfers the block input sub-matrix and weight sub-matrix from DDR to the buffer ping / pong;

[0090] (4) GEMM calculation to obtain the output sub-matrix;

[0091] (5) EDMA transfers the block output sub-matrix from the buffer ping / pong to DDR;

[0092] (6) If the sub-matrix has not been traversed completely, return to (1) and exchange the cache ping-pong until the sub-matrix traversal is completed and the calculation of the entire matrix is ​​completed.

[0093] In one embodiment of the present invention, DSP-based DMPYSP and DADDSP instructions are combined with pipeline optimization to implement four FP32 multiplication and addition operations in a single cycle, including:

[0094] Use DSP SIMD instructions to optimize the multiplication, addition and data transmission in GEMM. Take TMS320C6678 as an example:

[0095] The intrinsic function DMPYSP is combined with pipeline optimization to implement four FP32 multiplication operations in a single instruction cycle;

[0096] The intrinsic function DADDSP is combined with pipeline optimization to achieve four FP32 addition operations in a single instruction cycle;

[0097] The inline function AMEMD8 is combined with pipeline optimization to achieve the transmission of 128-bit (4 FP32) data in a single-cycle instruction.

[0098] SIMD instruction-level optimization, based on the DSP's VLIW (Very Long Instruction Word) architecture, implements instruction operations of four FP32 multiplications, additions, and data transfers in a single cycle, increasing the theoretical peak computing power to four times the original.

[0099] In one embodiment of the present invention, multi-core parallel scheduling is combined with OpenMP dynamic load balancing technology to achieve multi-core parallel acceleration (only applicable to multi-core DSPs).

[0100] The following two OpenMP multi-core parallel solutions are proposed:

[0101] Solution 1: Task-level parallelism: The matrix multiplication calculation of the output submatrix of each matrix block is regarded as a task unit, and the task is evenly distributed among multiple cores. Each DSP core is responsible for calculating the matrix multiplication task of the assigned output submatrix. The corresponding input submatrices and output submatrices are all stored in the L2 cache of each core. This solution is suitable for DSPs without L3 shared cache. In this case, the submatrix size must meet the following conditions:

[0102] (M1 * K + N1 * K + M1 * N1) * sizeof(data_type) <L2_Cache_Size * 0.8;

[0103] Where M1, N1, and K1 are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively. sizeof(data_type) refers to the size of the data type of the matrix elements, in bytes. L2_Cache_Size refers to the size of the available L2 cache, in bytes.

[0104] If cache ping-pong is used, the sub-matrix size must meet the following conditions:

[0105] (M_ping * K + N_ping * K + M_ping * N_ping) * sizeof(data_type) <L2_Cache_Size * 0.4;

[0106] Where M_ping, N_ping, and K_ping are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively. sizeof(data_type) refers to the size of the data type of the matrix elements, in bytes. L2_Cache_Size refers to the size of the available L2 cache, in bytes.

[0107] At the same time, the values ​​of M_pong, N_pong, K_pong and M_ping, N_ping, K_ping are exactly the same.

[0108] For example, if GEMM's data_type is FP32, L2_Cache_Size is 512KB, M=256, K=128, and N=196, the possible values ​​of M_ping, N_ping, and K_ping are 128, 128, and 256. In this case, the left side of the inequality is 192KB, which is smaller than the right side, 204.8KB, and thus satisfies the condition.

[0109] Solution 2, data-level parallelism: Use the OpenMP precompilation directive #pragma omp parallel for to optimize the innermost loop of the GEMM for multi-core parallelism, fully leveraging the DSP's multi-core parallel capabilities. This solution is suitable for DSPs with multi-core shared cache, such as the TMS320C6678, which has a 4MB L3 shared cache.

[0110] Compared to Option 2, Option 1 adds a significant amount of repeated sub-matrix transfers. However, these repeated transfers are multi-core parallel operations, and with EDMA fully optimized for data transfer, Option 1 does not incur any additional performance overhead. Furthermore, taking the TMS320C6678 as an example, the data access efficiency of the L3 cache and the L2 cache are the same, both half that of the L1 cache. Therefore, Option 1 and Option 2 theoretically offer comparable performance.

[0111] The large-scale two-dimensional convolution operator acceleration method for DSP platform of the present invention is described in detail below with reference to the accompanying drawings.

[0112] The following is a detailed description of the implementation case on the TI TMS320C6678 DSP platform. Important parameter configuration:

[0113] L3_Cache_Size: 3904KB (full L3 is 4096KB);

[0114] Block parameters: M_ping = M_pong = 256, N_ping = N_pong = 196, K_ping = K_pong = 576;

[0115] EDMA configuration: use EDMA3CC0 to control two DMA channels;

[0116] OpenMP thread number: 8 cores fully activated;

[0117] Test network: CenterNet with 224x224 input and ResNet18 as the backbone.

[0118] The present invention is applicable to general two-dimensional convolution and two-dimensional transposed convolution, and the complete optimization algorithm flow of the two is as follows: Figure 1 As shown in the figure, the left side is a two-dimensional convolution and the right side is a two-dimensional transposed convolution. The algorithm will select the corresponding rearrangement algorithm and process to rearrange the data according to the convolution type. For standard two-dimensional convolution, such as Figure 2 As shown, the im2col algorithm is used to rearrange the feature map of shape (IC, IH, IW) into a matrix of shape (IC*KH*KW, OH*OW), and the two-dimensional convolution result is obtained after the deep optimized GEMM operation. Figure 3 The second row shows an example of rearrangement of the single-channel feature map im2col. The example is a 4x4 feature map with a convolution kernel size of 3x3, a pad of 0, and a stride of 0. After rearrangement, it becomes a matrix of shape (9, 4). Figure 2 The first row also shows the rearrangement of the convolution kernel, but for the convolution weights with input shape (OC, IC, KH, KW), the storage method in memory is equivalent to the rearranged matrix, so generally no additional processing is done on the convolution weights. Figure 2 It is not difficult to find that when the convolution kernel size is 1x1, rearrangement does not make sense. In this case, you can skip im2col and directly execute GEMM. For two-dimensional transposed convolution, the convolution weights are pre-merged into a matrix of (IC, OC*KH*KW) before inference and transposed into (OC*KH*KW,IC), as shown in the following example: Figure 3 The deep optimized GEMM is performed and the col2im algorithm is used to rearrange the result from a matrix of shape (OC*KH*KW, IH*IW) back to a tensor of shape (OC, OH, OW). Figure 3 The single-channel col2im algorithm is shown, which can be regarded as the inverse process of the im2col algorithm.

[0119] In order to be applicable to large-scale matrix multiplication, the cache management strategy of the present invention adopts a three-segment blocking algorithm to adapt the GEMM computing unit to the L3 shared cache capacity, reduce the DDR3 access frequency, and improve the cache hit rate. Figure 4 shown.

[0120] To reduce the overhead of data transmission between L3 and DDR3, this paper employs an EDMA buffer ping-pong architecture, improving cache bandwidth utilization by pipelined and parallel execution of computation and transmission. The space allocated in L3_Cache by the three-segment block algorithm is duplicated. These two copies of the three-segment cache are identical in size and shape, respectively, and are called cache ping and cache pong. The requirements for M_ping, N_ping, and K_ping in this case are as described above.

[0121] In the test, the data_type of GEMM is FP32, the L3_Cache_Size is 3904 KB, and the values ​​of M_ping, N_ping, and K_ping are 256, 196, and 576. In this case, the left side of the inequality is 1213 KB, which is smaller than the right side (1561.6 KB), thus satisfying the condition.

[0122] The EDMA channel controller EDMA3CC0 of the TMS320C6678 is optimized for data transfer between the L3 shared cache and the DDR3 subsystem. It has 16 DMA channels and 8 QDMA channels, with two DMA channels selected for cache ping and cache pong data transfers respectively.

[0123] (1) Selecting blocks of the output matrix, input matrix, and weight matrix according to the method of the present invention;

[0124] (2) Select cache ping (the block input sub-matrix and weight sub-matrix need to be passed into cache ping in the first loop);

[0125] (3) EDMA transfers the block input sub-matrix and weight sub-matrix from DDR3 to buffer pong;

[0126] (4) GEMM calculation to obtain the output sub-matrix;

[0127] (5) EDMA transfers the block output sub-matrix from the cache to DDR3;

[0128] (6) If the sub-matrix has not been traversed completely, return to (1) and swap the caches in ping-pong mode until the sub-matrix traversal is completed.

[0129] Figure 5 The figure shows how cache ping-pong works. DDR3 starts addressing at 80 0000h. A and B represent the GEMM input matrices, and C represents the output matrix. The suffix i represents data transferred from DDR3 to the L3 shared cache, and the suffix o represents data transferred from the L3 shared cache to the DDR3. GEMM starts with cache ping and switches to cache pong after completing the cache ping transfer task. The two transfer tasks are performed alternately until the GEMM completes the multiplication task of the entire matrix.

[0130] The following uses the intrinsics supported by the TMS320C6678 to perform SIMD instruction-level optimization on the GEMM algorithm described above. The intrinsics DMPYSP and DADDSP are used to implement four FP32 multiplications and additions per cycle. The intrinsic AMEMD8, combined with pipeline optimization, enables single-cycle 128-bit (four FP32) data transfers. Theoretically, SIMD instruction-level optimization can quadruple the efficiency of the GEMM algorithm.

[0131] Considering that GEMMs optimized with intrinsic functions require memory alignment of data, there are requirements for the matrix shape in GEMMs: M and K must be multiples of 4, and N must be a multiple of 2. To accommodate common matrix shapes while preserving SIMD instruction optimization as much as possible, separate cases are handled based on the values ​​of M, K, and N, with SIMD and pipeline optimization performed separately for each case.

[0132] At this point, the single-core GEMM optimization task has been completed. Next, we combine OpenMP dynamic load balancing technology to achieve multi-core parallel acceleration. This test adopts a data-level OpenMP parallel solution and directly uses the OpenMP precompilation directive #pragma ompparallel for to perform multi-core parallel optimization on the innermost loop of the GEMM, fully utilizing the multi-core parallel capability of the TMS320C6678.

[0133] So far, the optimization of the two-dimensional convolution operator and the two-dimensional transposed convolution operator has been completed. In order to test the effectiveness of the present invention, CenterNet with ResNet18 as the backbone was selected as the detection network (mainly composed of two-dimensional convolution, and the neck is up-sampled by three layers of two-dimensional transposed convolution). Target detection was performed on the SAR image of the UAV scene. The inference effect on TMS320C6678 is as follows: Figure 6 and Figure 7 As shown, (label description: the first number represents the confidence, the second number represents the category, 0(J6), 1(54TH), 2(63AT), 3(63CAAV), 4(62LT), 5(63APTV), 6(T34-85), 7(59-1TC), 8(59AG)). Figure 6 It is the CenterNet inference effect achieved using the original two-dimensional convolution operator and the two-dimensional transposed convolution operator. Figure 7 The CenterNet inference effect is achieved using the optimized two-dimensional convolution operator and two-dimensional transposed convolution operator of the present invention. The target detection effects of the two are basically the same, and the total inference time is optimized from 178.4s to 0.51s, verifying the effectiveness of the solution of the present invention.

[0134] According to the large-scale two-dimensional convolution operator acceleration method for DSP platforms proposed in an embodiment of the present invention, it can achieve efficient inference of SAR target detection networks on the TITMS320C6678 platform in actual tests, filling the gap in the neural network inference ecosystem on DSP, and providing a feasible solution for real-time inference of CNN networks on DSP platforms. It is suitable for real-time convolutional neural network inference in edge computing scenarios and meets the urgent need for realizing real-time SAR target detection based on DSP on airborne platforms.

[0135] In the description of this specification, the reference terms "one embodiment", "some embodiments", "example", "specific example", or "some examples" mean that the specific features, structures, materials or characteristics described in conjunction with the embodiment or example are included in at least one embodiment or example of the present invention. In this specification, the schematic expressions of the above terms do not necessarily refer to the same embodiment or example. Moreover, the specific features, structures, materials or characteristics described can be combined in any one or N embodiments or examples in a suitable manner. In addition, those skilled in the art can combine and combine different embodiments or examples described in this specification and features of different embodiments or examples without contradiction.

[0136] Furthermore, the terms "first" and "second" are used for descriptive purposes only and should not be understood to indicate or imply relative importance or implicitly specify the number of technical features indicated. Thus, a feature specified as "first" or "second" may explicitly or implicitly include at least one such feature. In the description of the present invention, "N" means at least two, such as two, three, etc., unless otherwise specifically defined.

[0137] Any process or method description in a flowchart or otherwise described herein may be understood to represent a module, segment or portion of code comprising one or N executable instructions for implementing a custom logical function or step of a process, and the scope of the preferred embodiments of the present invention includes alternative implementations in which functions may be performed out of the order shown or discussed, including performing functions in a substantially simultaneous manner or in the reverse order depending on the functions involved, which should be understood by those skilled in the art to which the embodiments of the present invention pertain.

Claims

1. A large-scale two-dimensional convolution operator acceleration method for DSP platform, characterized by: The following steps are involved: 1) Dynamic data rearrangement: select the im2col algorithm or col2im algorithm to rearrange data according to the convolution type; 2) Cache management strategy: Adopt a three-segment matrix partitioning strategy to adapt to the L3 cache capacity; 3) Using the EDMA buffer ping-pong architecture, the calculation and transmission pipelines are executed in parallel; specifically, Cache allocation strategy: The three-segment blocks allocated in the L3_Cache by the three-segment matrix block strategy are copied to obtain two three-segment caches with exactly the same size and shape, which are called cache ping and cache pong respectively. The suffix ping means cache ping and the suffix pong means cache pong. M_ping, N_ping, and K_ping are required to meet the following conditions: (M_ping * N_ping + K_ping * N_ping + M_ping * K_ping) * sizeof(data_type)< L3_Cache_Size * 0.4; Where M_ping, N_ping, and K_ping are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively; sizeof(data_type) refers to the size of the data type of the matrix element, in bytes; L3_Cache_Size refers to the size of the available L3 cache, in bytes; At the same time, the values ​​of M_pong, N_pong, K_pong and M_ping, N_ping, K_ping are exactly the same; EDMA and GEMM combination strategy: (1) Select the blocks of output matrix, input matrix and weight matrix according to the three-segment matrix block strategy; (2) Select cache ping / pong; (3) EDMA transfers the block input sub-matrix and weight sub-matrix from DDR to the buffer ping / pong; (4) GEMM calculation to obtain the output sub-matrix; (5) EDMA transfers the block output sub-matrix from the buffer ping / pong to DDR; (6) If the sub-matrix has not been traversed completely, return to (1) and swap the caches for ping-pong until the sub-matrix traversal is completed and the calculation of the entire matrix is ​​completed; 4) SIMD instruction-level optimization, based on DSP's DMPYSP and DADDSP instructions combined with pipeline optimization to achieve four FP32 multiplication and addition operations in a single cycle; 5) Multi-core parallel scheduling, combined with OpenMP dynamic load balancing technology to achieve multi-core parallel acceleration.

2. The method according to claim 1, characterized in that Select the im2col algorithm or col2im algorithm to rearrange the data according to the convolution type, including: For standard two-dimensional convolution, the im2col algorithm is used to rearrange the input data from a tensor of shape (IC, IH, IW) to a matrix of shape (IC*KH*KW, OH*OW), and the result is obtained through GEMM operation; For 1x1 convolution, skip im2col and directly call GEMM operation to obtain the result; For two-dimensional transposed convolution, the weights are pre-converted into a matrix of shape (OC*KH*KW, IC). During inference, GEMM is directly called and the col2im algorithm is used to rearrange the result from a matrix of shape (OC*KH*KW, IH*IW) back to a tensor of shape (OC, OH, OW), where IC, IH, and IW refer to the number of input image channels, input image height, and input image width, KH and KW refer to the convolution kernel height and convolution kernel width, and OC, OH, and OW refer to the number of output image channels, output image height, and output image width.

3. The method according to claim 1, characterized in that A three-segment matrix partitioning strategy is used to adapt the L3 cache capacity, including: The first block division: divide the M×N output matrix into multiple M1×N1 sub-matrices; The second block division: divide the K×N input matrix into multiple K1×N1 sub-matrices; The third block division: divide the M×K weight matrix into multiple M1×K1 sub-matrices; M1, N1 and K1 meet the following conditions: (M1 * N1 + K1 * N1 + M1 * K1) * sizeof(data_type) < L3_Cache_Size * 0.8; Where M1, N1, and K1 are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively. sizeof(data_type) refers to the size of the data type of the matrix elements, in bytes. L3_Cache_Size refers to the size of the available L3 cache, in bytes.

4. The method according to claim 1, wherein DSP-based DMPYSP and DADDSP instructions, combined with pipeline optimization, implement four FP32 multiplication and addition operations in a single cycle, including: Use DSP SIMD instructions to optimize multiplication, addition, and data transmission in GEMM.

5. The method according to claim 1, wherein Multi-core parallel scheduling, combined with OpenMP dynamic load balancing technology, achieves multi-core parallel acceleration, including: Task-level parallelism: The matrix multiplication calculation of the output sub-matrix of each matrix block is regarded as a task unit, and the task is evenly distributed to multiple cores. Each DSP core is responsible for calculating the matrix multiplication task of the assigned output sub-matrix. The corresponding input sub-matrices and output sub-matrices are all placed in the L2_Cache of each core. It is suitable for DSPs without L3 shared cache. In this case, the sub-matrix size meets the following conditions: (M1 * K + N1 * K + M1 * N1) * sizeof(data_type) < L2_Cache_Size * 0.8; Where M1, N1, and K1 are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively; sizeof(data_type) refers to the size of the data type of the matrix element, in bytes; L2_Cache_Size refers to the size of the available L2 cache, in bytes; If cache ping-pong is used, the sub-matrix size must meet the following conditions: (M_ping * K + N_ping * K + M_ping * N_ping) * sizeof(data_type) < L2_Cache_Size * 0.4; Where M_ping, N_ping, and K_ping are the dimensions of the submatrix, and their values ​​are factors of M, N, and K, respectively; sizeof(data_type) refers to the size of the data type of the matrix element, in bytes; L2_Cache_Size refers to the size of the available L2 cache, in bytes; At the same time, the values ​​of M_pong, N_pong, K_pong and M_ping, N_ping, K_ping are exactly the same; Data-level parallelism: Directly use the OpenMP precompilation directive #pragma omp parallel for to perform multi-core parallel optimization on the innermost loop of GEMM, which is suitable for DSPs with multi-core shared cache.

Images