DSP platform-oriented large-size two-dimensional convolution operator acceleration method

Abstract

The invention discloses a large-size two-dimensional convolution operator acceleration method for a DSP platform, and belongs to the field of digital signal processing. According to the method, an im2col algorithm or a col2im algorithm is selected to rearrange data according to a convolution type; the L3 cache capacity is adapted by adopting a three-section matrix partitioning strategy; an EDMA buffer ping-pong architecture is utilized to realize pipeline parallelism of data transmission and calculation; sIMD instruction level optimization is carried out, and multiplication operation and additive operation of four FP32 in a single cycle are realized based on cooperation of DMPYSP and DADDSP instructions and pipeline optimization; and multi-core parallel scheduling is carried out, and task-level and data-level parallelism is realized through OpenMP. According to the method, efficient reasoning of an SAR target detection network can be realized through actual measurement on a TI TMS320C6678 platform, and a feasible scheme is provided for real-time reasoning of a CNN network on a DSP platform.

Description

Technical Field

[0001] The present invention relates to the technical field of digital signal processing, and particularly relates to a method for accelerating large-size two-dimensional convolution operators for a DSP platform. Background Art

[0002] Synthetic Aperture Radar (SAR) has characteristics such as all-weather, all-day, multi-polarization, and high resolution, and has made outstanding contributions in military and civilian fields. It overcomes the weakness of traditional radars with limited azimuth resolution, can obtain high-resolution images of large areas at long operating distances, and is widely used for high-resolution real-time imaging on airborne and spaceborne platforms. On the other hand, with the development of machine learning and deep learning in target detection technology in recent years, convolutional neural network (CNN) models represented by ResNet have been widely used in the field of SAR target detection. Gradually, it has become a trend to deploy SAR target detection networks in a lightweight manner on low-power, miniaturized, and inexpensive hardware platforms. Implementing real-time SAR target detection on airborne and unmanned aerial vehicle (UAV) platforms has become an urgent practical need.

[0003] DSP (Digital Signal Processor) is a high-speed electronic device designed specifically for signal processing, which 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] Different from traditional hardware platforms, implementing real-time target detection algorithms based on convolutional neural networks on DSPs faces some difficulties. There are already various mature implementation frameworks for convolutional neural networks on CPU and GPU platforms. Most of the running implementations of these implementation frameworks on the GPU side are accelerated using the NVIDIA cudnn library, and the implementations on the CPU use methods such as OpenMP for multi-threaded acceleration. There are also quite a number of ready-made acceleration libraries available. Therefore, most of the main acceleration operation platforms for deep learning CNNs choose GPU or FPGA hardware platforms. Due to the natural disadvantage of DSPs in parallel computing compared to GPUs or FPGAs, traditional deep learning frameworks (such as TensorFlow, caffe, and pytorch) do not support DSP platforms and lack the corresponding ecosystem.

[0005] Therefore, to achieve real-time inference of convolutional neural networks on DSP, optimization work needs to be carried out starting from the hardware bottom layer of DSP and the algorithm bottom layer of operators. Common implementation methods of two-dimensional convolution operators at the algorithm level include rearranging data into matrix multiplication, implementing traditional two-dimensional convolution with a sliding window, using the Winograd transformation method, etc. General data rearrangement methods will bring additional memory expenses and rearrangement performance expenses. The implementation of traditional two-dimensional convolution combined with a sliding window brings additional performance expenses due to the repeated transfer of the same data. As for using the Winograd transformation, although it can reduce the multiplication operations during inference, it also increases the addition operations. For a DSP with optimized multiplication instructions and multiply-accumulate instructions, the Winograd transformation will not bring efficiency improvement. On the other hand, the traditional two-dimensional convolution operation method is relatively complex and not conducive to optimization using SIMD (Single Instruction Multiple Data). At this time, the implementation method of converting data into matrix multiplication after rearrangement is worth considering. Summary of the Invention

[0006] The present invention provides an acceleration method for large-size two-dimensional convolution operators for a DSP platform, which can achieve efficient inference of the SAR target detection network on the TITMS320C6678 platform and provide a feasible solution for real-time inference of CNN networks on the DSP platform.

[0007] An embodiment of the present invention provides an acceleration method for large-size two-dimensional convolution operators for a DSP platform, including the following steps: 1), Dynamic data rearrangement: Select the im2col algorithm or the col2im algorithm to rearrange data according to the convolution type; 2), Cache management strategy: Adopt a three-segment matrix block strategy to adapt to the L3 cache capacity; 3), Utilize the EDMA (Enhanced Direct Memory Access) buffer ping-pong architecture to execute the pipeline parallelism of calculation and transmission; 4), SIMD (Single Instruction Multiple Data) instruction-level optimization: Based on the DMPYSP (two groups of single-precision floating-point multiplication instructions) and DADDSP (two groups of single-precision floating-point addition instructions) instructions of the DSP, cooperate with pipeline optimization to achieve 4 FP32 multiplication operations and addition operations in a single cycle; 5), Multicore parallel scheduling: Combine the OpenMP dynamic load balancing technology to achieve multicore parallel acceleration.

[0008] Optionally, in an embodiment of the present invention, selecting the im2col algorithm or the col2im algorithm to rearrange data according to the convolution type includes: For standard two-dimensional convolution, the im2col algorithm is used to rearrange the input data from a tensor of shape (IC, IH, IW) into a matrix of shape (IC*KH*KW, OH*OW), and the result is obtained through GEMM (General Matrix Multiplication) operation; For 1x1 convolution, skip im2col and directly call the GEMM operation to obtain the result; For two-dimensional transposed convolution, the weight is pre-converted into a matrix of shape (OC*KH*KW, IC). During inference, directly call GEMM and then use the col2im algorithm to rearrange the result from a matrix of shape (OC*KH*KW, IH*IW) back into a tensor of shape (OC, OH, OW), where IC, IH, and IW refer to the number of input image channels, the height of the input image, and the width of the input image, respectively, KH and KW refer to the height and width of the convolution kernel, and OC, OH, and OW refer to the number of output image channels, the height of the output image, and the width of the output image.

[0009] Optionally, in an embodiment of the present invention, adopting a three-segment matrix block strategy to adapt to the L3 cache capacity includes: The first segment of block: Divide the output matrix of M×N into multiple sub-matrices of M1×N1; The second segment of block: Divide the input matrix of K×N into multiple sub-matrices of K1×N1; The third segment of block: Divide the weight matrix of M×K into multiple sub-matrices of M1×K1; M1, N1, and K1 satisfy 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 sub-matrices, and are respectively factors of M, N, and K; 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.

[0010] Optionally, in an embodiment of the present invention, using the EDMA buffer ping-pong architecture to execute the computing and transmission pipelines in parallel includes: Cache allocation strategy: Copy the three - segment block in the L3_Cache using the three - segment matrix block strategy to obtain two three - segment caches with exactly the same size and shape, which are respectively called cache ping and cache pong. The suffix ping represents cache ping, and the suffix pong represents cache pong; It is required that M_ping, N_ping, and K_ping satisfy the following conditions: (M_ping * N_ping + K_ping * N_ping + M_ping * K_ping) * sizeof(data_type) < L3_Cache_Size * 0.4; Among them, M_ping, N_ping, and K_ping are the dimensions of the sub - matrices, which are respectively factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, with the unit of byte; L3_Cache_Size refers to the size of the available L3 cache, with the unit of byte; At the same time, M_pong, N_pong, K_pong have the same values as M_ping, N_ping, K_ping; EDMA and GEMM cooperation strategy: (1) Select the blocks of the output matrix, input matrix, and weight matrix according to the three - segment matrix block strategy; (2) Select cache ping / pong; (3) EDMA transfers the blocked input sub - matrix and weight sub - matrix from DDR (Double Data Rate) to cache ping / pong; (4) GEMM calculates to obtain the output sub - matrix; (5) EDMA transfers the blocked output sub - matrix from cache ping / pong to DDR; (6) If the blocked sub - matrices have not been traversed completely, return to (1), swap cache ping - pong, until the sub - matrices are traversed, and complete the calculation of the entire matrix.

[0011] Optionally, in an embodiment of the present invention, the DMPYSP and DADDSP instructions of the DSP are combined with pipeline optimization to achieve 4 FP32 multiplication operations and addition operations in a single cycle, including: Use the SIMD instructions of the DSP to optimize the multiplication, addition, and data transfer methods in GEMM.

[0012] Optionally, in an embodiment of the present invention, multi - core parallel scheduling is combined with the OpenMP dynamic load - balancing technology to achieve multi - core parallel acceleration, including: Task-level parallelism: Taking the matrix multiplication calculation of each output sub-matrix of the matrix block as a task unit, evenly distributing the tasks to multiple cores for execution. Each DSP core is responsible for calculating the matrix multiplication task of the assigned output sub-matrix. The corresponding several input sub-matrices and output sub-matrices involved are all placed in the L2_Cache of each core. This is applicable to DSPs without L3 shared cache. At this time, the size of the sub-matrix satisfies the following conditions: (M1 * K + N1 * K + M1 * N1) * sizeof(data_type) < L2_Cache_Size * 0.8; Among them, M1, N1, and K1 are the dimensions of the sub-matrix, and they are respectively the factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, with the unit of byte; L2_Cache_Size refers to the size of the available L2 cache, with the unit of byte; If cache ping-pong is used, the size of the sub-matrix satisfies the following conditions: (M_ping * K + N_ping * K + M_ping * N_ping) * sizeof(data_type) < L2_Cache_Size * 0.4; Among them, M_ping, N_ping, and K_ping are the dimensions of the sub-matrix, and they are respectively the factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, with the unit of byte; L2_Cache_Size refers to the size of the available L2 cache, with the unit of byte; At the same time, M_pong, N_pong, K_pong and M_ping, N_ping, K_ping have exactly the same values; Data-level parallelism: Directly use the OpenMP pre-compilation directive #pragma omp parallel for to perform multi-core parallel optimization on the innermost loop of GEMM. This is applicable to DSPs with multi-core shared cache.

[0013] The large-size two-dimensional convolution operator acceleration method for the DSP platform in the embodiments of the present invention can achieve efficient inference of the SAR target detection network through actual measurement on the TITMS320C6678 platform, providing a feasible solution for the real-time inference of the CNN network on the DSP platform.

[0014] The additional aspects and advantages of the present invention will be partially given in the following description, partially become obvious from the following description, or be understood through the practice of the present invention. Description of the Drawings

[0015] The above and / or additional aspects and advantages of the present invention will become apparent and be readily understood from the following description of embodiments in conjunction with the accompanying drawings, in which: Figure 1 It is an overall flowchart for optimizing a two-dimensional convolution operator according to an embodiment of the present invention; Figure 2 And Figure 3 It is a schematic diagram of data rearrangement according to an embodiment of the present invention; Figure 4 It is a schematic diagram of matrix multiplication block division according to an embodiment of the present invention; Figure 5 It is a schematic diagram of cache ping-pong according to an embodiment of the present invention; Figure 6 And Figure 7 It is the SAR target detection result of the UAV scenario according to an embodiment of the present invention. Specific embodiments

[0016] The embodiments of the present invention will be described in detail below. Examples of the embodiments are shown in the accompanying drawings, where the same or similar reference numerals denote the same or similar elements or elements having the same or similar functions throughout. The embodiments described below with reference to the accompanying drawings are exemplary and are intended to explain the present invention and should not be construed as limiting the present invention.

[0017] Specifically, Figure 1 It is a flowchart of a method for accelerating a large-size two-dimensional convolution operator for a DSP platform provided by an embodiment of the present invention.

[0018] As Figure 1 shown, the method for accelerating a large-size two-dimensional convolution operator for a DSP platform includes the following steps: 1), Dynamic data rearrangement: Rearrange data according to the convolution type by selecting the im2col algorithm or the col2im algorithm; 2), Cache management strategy: Adopt a three-segment matrix block division strategy to adapt to the L3 cache capacity; 3), Utilize the EDMA buffer ping-pong architecture to execute the pipeline of calculation and transmission in parallel; 4), SIMD instruction-level optimization, and realize 4 FP32 multiplication operations and addition operations in a single cycle based on the cooperation of the DMPYSP and DADDSP instructions of the DSP with pipeline optimization; 5), Multicore parallel scheduling, and realize task-level and data-level parallelism through OpenMP.

[0019] In an embodiment of the present invention, rearranging data by selecting the im2col algorithm or the col2im algorithm according to the convolution type includes: For the standard two-dimensional convolution, the im2col algorithm is used to rearrange the input data from a tensor of shape (IC, IH, IW) into a matrix of shape (IC*KH*KW, OH*OW), and the result is obtained through GEMM operations; For the 1x1 convolution, the im2col is skipped and the GEMM operation is directly called to obtain the result; For the two-dimensional transposed convolution, the weights are pre-converted into a matrix of shape (OC*KH*KW, IC). During inference, after directly calling GEMM, the col2im algorithm is used to rearrange the result from a matrix of shape (OC*KH*KW, IH*IW) back into a tensor of shape (OC, OH, OW). Here, IC, IH, and IW refer to the number of input image channels, the height of the input image, and the width of the input image, respectively. KH and KW refer to the height and width of the convolution kernel, respectively. OC, OH, and OW refer to the number of output image channels, the height of the output image, and the width of the output image, respectively.

[0020] In an embodiment of the present invention, a three-stage block algorithm is used to adapt the GEMM computing unit to the L3 shared cache capacity, reducing the DDR access frequency to the theoretical lower limit, including: The first-stage block: Divide the output matrix of M×N into multiple sub-matrices of M1×N1; The second-stage block: Divide the input matrix of K×N into multiple sub-matrices of K1×N1; The third-stage block: Divide the weight matrix of M×K into multiple sub-matrices of M1×K1; It is required that M1, N1, and K1 satisfy the following conditions: (M1 * N1 + K1 * N1 + M1 * K1) * sizeof(data_type) < L3_Cache_Size * 0.8; Wherein, M1, N1, and K1 are the dimensions of the sub-matrices, and are respectively factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, in bytes. For example, the size of the single-precision floating-point data type is 4; L3_Cache_Size refers to the size of the available L3 cache, in bytes.

[0021] M1, N1, and K1 should, under the condition of satisfying the above conditions, minimize the number of sub-matrix blocks according to the actual sizes of M, N, and K to reduce the overhead caused by blocking. The three-stage block is allocated in the L3_Cache for accelerating the GEMM operation, and finally 20% of the L3_Cache is reserved for other uses.

[0022] Taking the data_type of GEMM as FP32 and the L3_Cache_Size as 4MB as an example, when M1, N1, and K1 are 128, 512, and 256 respectively, the left side of the inequality is 0.875MB, which is less than 3.2MB on the right side, meeting the condition.

[0023] In an embodiment of the present invention, by using the EDMA buffer ping-pong architecture, the pipelining of calculation and transmission is executed in parallel, including: Cache allocation strategy: Copy the three-segment matrix block strategy to the three segments in the L3_Cache to obtain two three-segment caches with exactly the same size and shape, which are respectively called cache ping and cache pong; It is required that M_ping, N_ping, and K_ping satisfy the following conditions: (M_ping * N_ping + K_ping * N_ping + M_ping * K_ping) * sizeof(data_type) < L3_Cache_Size * 0.4; Wherein, M_ping, N_ping, and K_ping are the dimensions of the sub-matrix, and are respectively the factors of M, N, and K; 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] At the same time, M_pong, N_pong, K_pong have exactly the same values as M_ping, N_ping, K_ping.

[0025] Taking the data_type of GEMM as FP32 and the L3_Cache_Size as 4MB as an example, when M_ping, N_ping, and K_ping are 128, 512, and 256 respectively, the left side of the inequality is 0.875MB, which is less than 1.6MB on the right side, meeting the condition.

[0026] EDMA and GEMM cooperation strategy: (1) Select the blocks of the output matrix, input matrix, and weight matrix according to the three-segment matrix block strategy; (2) Select cache ping / pong; (3) EDMA transfers the blocked input sub-matrix and weight sub-matrix from the DDR to cache ping / pong; [[ID=2**9]] (4) GEMM calculates to obtain the output sub-matrix; (5) EDMA transfers the blocked output sub-matrix from cache ping / pong to the DDR; (6) If the block sub - matrices have not been traversed completely, go back to (1), swap the cache ping - pong until the sub - matrices are traversed, and complete the calculation of the entire matrix.

[0027] In an embodiment of the present invention, based on the DSP's DMPYSP and DADDSP instructions and pipeline optimization, 4 FP32 multiplication operations and addition operations are achieved in a single cycle, including: Use the DSP's SIMD instructions to optimize the multiplication, addition, and data transfer methods in GEMM. Taking TMS320C6678 as an example: The in - line function DMPYSP, combined with pipeline optimization, realizes 4 FP32 multiplication operations in a single instruction cycle; The in - line function DADDSP, combined with pipeline optimization, realizes 4 FP32 addition operations in a single instruction cycle; The in - line function AMEMD8, combined with pipeline optimization, realizes the transfer of 128 - bit (4 FP32) data in a single cycle instruction.

[0028] SIMD instruction - level optimization, based on the DSP's VLIW (Very Long Instruction Word) architecture, realizes the instruction operations of 4 FP32 multiplication operations, addition operations, and data transfer in a single cycle, and the theoretical peak computing power is increased to 4 times the original.

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

[0030] The following two OpenMP multi - core parallel schemes are proposed: Scheme 1, task - level parallelism: Take the matrix multiplication calculation of each output sub - matrix of the matrix block as a task unit, evenly distribute the tasks to multiple cores for execution. Each DSP core is responsible for calculating the matrix multiplication task of the assigned output sub - matrix, and the corresponding several input sub - matrices and output sub - matrices involved are placed in the L2_Cache of each core. This scheme is applicable to DSPs without L3 shared cache. At this time, the size of the sub - matrix needs to meet the following conditions: (M1 * K + N1 * K + M1 * N1) * sizeof(data_type)<L2_Cache_Size * 0.8; Among them, M1, N1, and K1 are the dimensions of the sub - matrix, taking values as 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.

[0031] If cache ping-pong is used, the sub-matrix size needs to meet the following conditions: (M_ping * K + N_ping * K + M_ping * N_ping) * sizeof(data_type) < L2_Cache_Size * 0.4; Among them, M_ping, N_ping, and K_ping are the dimensions of the sub-matrix, and they are respectively the factors of M, N, and K; 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.

[0032] At the same time, M_pong, N_pong, K_pong have the exact same values as M_ping, N_ping, K_ping.

[0033] Taking the data_type of GEMM as FP32, L2_Cache_Size as 512KB, M = 256, K = 128, N = 196 as an example, the possible values of M_ping, N_ping, and K_ping are 128, 128, 256. At this time, the left side of the inequality is 192KB, which is less than 204.8KB on the right side, meeting the conditions.

[0034] Solution 2, data-level parallelism: directly use the pre-compilation directive #pragma omp parallel for of OpenMP to perform multi-core parallel optimization on the innermost loop of GEMM, and give full play to the multi-core parallel capabilities of the DSP. This solution is applicable to DSPs with multi-core shared caches. For example, TMS320C6678 has 4MB of L3 available for shared caching.

[0035] Solution 1 adds a large number of repeated sub-matrix transfers compared to Solution 2. However, these repeated transfers are multi-core parallel operations. In the case of fully optimizing data transfer using EDMA, Solution 1 will not increase additional performance overhead. On the other hand, taking TMS320C6678 as an example, the data access efficiency of L3_Cache is the same as that of L2_Cache, both being half of L1. Therefore, theoretically, Solution 1 and Solution 2 have equivalent performance.

[0036] The following will elaborate in detail on the large-size two-dimensional convolution operator acceleration method of the present invention for the DSP platform in conjunction with the accompanying drawings.

[0037] Taking the case implemented on the TI TMS320C6678 DSP platform as an example for detailed description. Important parameter configuration: L3_Cache_Size: 3904KB (the complete L3 is 4096KB); Chunk parameters: M_ping = M_pong = 256, N_ping = N_pong = 196, K_ping = K_pong = 576; EDMA configuration: Use EDMA3CC0 to control two DMA channels; Number of OpenMP threads: All 8 cores are activated; Test network: CenterNet with ResNet18 as the Backbone and an input of 224x224.

[0038] The present invention is applicable to general two-dimensional convolution and two-dimensional transposed convolution. The complete optimization algorithm processes for both are as Figure 1 shown. The left side is for two-dimensional convolution, and the right side is for 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, as Figure 2 shown, use the im2col algorithm to rearrange the feature map with the shape of (IC, IH, IW) into a matrix with the shape of (IC * KH * KW, OH * OW). After the depth-optimized GEMM operation, the two-dimensional convolution result is obtained. As Figure 3 shown in the second row, a single-channel feature map im2col rearrangement example is given. The example is a 4x4-sized feature map, the convolution kernel size is 3x3, pad is 0, stride is 0, and after rearrangement, it becomes a matrix with the shape of (9, 4). Figure 2 The first row also shows the rearrangement of the convolution kernel. However, for the convolution weights with the input shape of (OC, IC, KH, KW), the storage method in memory is equivalent to the matrix after rearrangement. Therefore, generally no additional processing is done on the convolution weights. Observe Figure 2 and it is not difficult to find that when the convolution kernel size is 1x1, rearrangement is meaningless, and at this time, im2col can be skipped and GEMM can be directly executed. For two-dimensional transposed convolution, before inference, the convolution weights are pre-merged into a matrix of (IC, OC * KH * KW) and transposed into (OC * KH * KW, IC). As Figure 3 shown, after performing the depth-optimized GEMM, use the col2im algorithm to rearrange the result from a matrix with the shape of (OC * KH * KW, IH * IW) back into a tensor with the shape of (OC, OH, OW). <t Figure 3 Shows the single-channel col2im algorithm, which can be regarded as the inverse process of the im2col algorithm.

[0039] To be applicable to large-size matrix multiplication, the cache management strategy applied in the present invention adopts a three-segment chunking algorithm to adapt the GEMM calculation unit to the L3 shared cache capacity, reduce the DDR3 access frequency, and improve the cache hit rate, as Figure 4 shown.

[0040] In order to reduce the cost of data transmission between L3 and DDR3, the present invention adopts an EDMA buffer ping-pong architecture, and improves the cache bandwidth utilization rate by parallel execution of the pipelining of calculation and transmission. Duplicate the space allocated for the above three-segment block algorithm in the L3_Cache. The sizes and shapes of the two three-segment caches are exactly the same, and are respectively called cache ping and cache pong. At this time, it is required that M_ping, N_ping, and K_ping meet the conditions as described in the foregoing records.

[0041] In the test, the data_type of GEMM is FP32, the L3_Cache_Size is 3904KB, and the values of M_ping, N_ping, and K_ping are 256, 196, and 576. At this time, the left side of the inequality is 1213 KB, which is less than the right side of 1561.6 KB, meeting the conditions.

[0042] The EDMA channel controller EDMA3CC0 of TMS320C6678 is optimized for data transmission between the L3 shared cache and the DDR3 subsystem. It has 16 DMA channels and 8 QDMA channels. Select two DMA channels to be used for data transmission of cache ping and cache pong respectively.

[0043] (1) Select the partitioning of the output matrix, input matrix, and weight matrix according to the method of the present invention; (2) Select cache ping (when the first loop is performed, the partitioned input sub-matrix and weight sub-matrix need to be passed into cache ping); (3) EDMA transfers the partitioned input sub-matrix and weight sub-matrix from DDR3 to cache pong; (4) GEMM calculates to obtain the output sub-matrix; (5) EDMA transfers the partitioned output sub-matrix from cache ping to DDR3; (6) If the partitioned sub-matrices have not been traversed completely, return to (1), swap cache ping and pong until the sub-matrices are traversed completely.

[0044] Figure 5 Shows the working mode of cache ping-pong. DDR3 starts addressing from 80 0000h. A and B represent the input matrices of GEMM, C represents the output matrix. The suffix i represents the data transferred from DDR3 to the L3 shared cache, and the suffix o represents the data transferred from the L3 shared cache to DDR3. When GEMM starts, it uses cache ping, and switches to the transmission task of cache pong after completing the transmission task of cache ping. The two alternate until GEMM completes the entire matrix multiplication task.

[0045] The above GEMM algorithm is optimized at the SIMD instruction level using the inline functions supported by TMS320C6678. The inline functions DMPYSP and DADDSP are used to implement 4 FP32 multiplication operations and addition operations in a single cycle, and the inline function AMEMD8 is used in conjunction with pipeline optimization to achieve the transfer of 128-bit (4 FP32) data in a single-cycle instruction. In theory, the efficiency of the GEMM algorithm optimized at the SIMD instruction level can be increased to 4 times the original.

[0046] Considering that the GEMM optimized using inline functions requires memory alignment of data and has requirements for the matrix shape in GEMM, where M and K are required to be multiples of 4 and N is required to be a multiple of 2. To adapt to general matrix shapes and retain SIMD instruction optimization as much as possible, different cases are handled according to the values of M, K, and N, and SIMD and pipeline optimizations are performed separately for each case.

[0047] At this point, the optimization task of the single-core GEMM has been completed. Next, multi-core parallel acceleration is achieved by combining the OpenMP dynamic load balancing technology. In this test, a data-level OpenMP parallel scheme is adopted, and the pre-compilation directive #pragma omp parallel for of OpenMP is directly used to perform multi-core parallel optimization on the innermost loop of GEMM, giving full play to the multi-core parallel capabilities of TMS320C6678.

[0048] So far, the optimization of the two-dimensional convolution operator and the two-dimensional transposed convolution operator has been completed. To test the effectiveness of the present invention, CenterNet with ResNet18 as the Backbone is selected as the detection network (mainly composed of two-dimensional convolutions, and the neck realizes upsampling by three two-dimensional transposed convolutions), and target detection is performed on SAR images of the airborne drone scenario. The inference effect on TMS320C6678 is as Figure 6 and Figure 7 shown. (Label description: The first number represents the confidence level, 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 is the inference effect of CenterNet implemented using the original two-dimensional convolution operator and two-dimensional transposed convolution operator, Figure 7 is the inference effect of CenterNet implemented using the two-dimensional convolution operator and two-dimensional transposed convolution operator optimized by the present invention. The target detection effects of the two are basically the same, and the total inference time is optimized from 178.4 s to 0.51 s, verifying the effectiveness of the proposed solution of the present invention.

[0049] The large-size two-dimensional convolution operator acceleration method for the DSP platform proposed according to the embodiments of the present invention can achieve efficient inference of the SAR target detection network through actual measurement on the TITMS320C6678 platform, filling the gap in the neural network inference ecosystem on the DSP, providing a feasible solution for the real-time inference of CNN networks on the DSP platform, being applicable to the real-time convolutional neural network inference in the edge computing scenario, and meeting the urgent need for realizing DSP-based real-time SAR target detection on the airborne platform.

[0050] In the description of this specification, the description referring to terms such as "one embodiment", "some embodiments", "example", "specific example", or "some examples" means that the specific features, structures, materials, or characteristics described in connection with the embodiment or example are included in at least one embodiment or example of the present invention. In this specification, the schematic representations 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 a suitable manner in any one or N embodiments or examples. In addition, without contradiction, those skilled in the art can combine and combine the different embodiments or examples described in this specification and the features of different embodiments or examples.

[0051] In addition, the terms "first" and "second" are only used for descriptive purposes and cannot be understood as indicating or implying relative importance or implicitly specifying the quantity of the indicated technical features. Thus, the features defined with "first" and "second" can explicitly or implicitly include at least one of the features. In the description of the present invention, the meaning of "N" is at least two, such as two, three, etc., unless otherwise clearly and specifically defined.

[0052] Any process or method description shown in the flowchart or described in other ways herein can be understood as representing a module, segment, or part of code including one or N executable instructions for implementing a customized logic function or process, and the scope of the preferred embodiments of the present invention includes additional implementations, where the functions can be executed in a substantially simultaneous manner or in the reverse order according to the involved functions, rather than in the order shown or discussed, which should be understood by those skilled in the art of the embodiments of the present invention.

Claims

1. An acceleration method for large-size two-dimensional convolution operators for DSP platforms, characterized in that, It includes the following steps: 1), Dynamic data rearrangement: Select the im2col algorithm or the col2im algorithm to rearrange the data according to the convolution type; 2), Cache management strategy: Adopt a three - stage matrix block strategy to adapt to the L3 cache capacity; 3), Utilize the EDMA buffer ping - pong architecture to execute the calculation and transmission pipelines in parallel; 4), SIMD instruction - level optimization: Based on the DSP's DMPYSP and DADDSP instructions, cooperate with pipeline optimization to achieve 4 FP32 multiplication operations and addition operations in a single cycle; 5), Multi - core parallel scheduling: Combine the 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 the col2im algorithm to rearrange the data according to the convolution type, including: For standard 2D convolution, adopt the im2col algorithm to rearrange the input data from a tensor of shape (IC, IH, IW) into a matrix of shape (IC * KH * KW, OH * OW), and obtain the result through GEMM operation; For 1x1 convolution, skip im2col and directly call the GEMM operation to obtain the result; For 2D transposed convolution, pre - convert the weights into a matrix of shape (OC * KH * KW, IC). During inference, directly call GEMM and then use the col2im algorithm to rearrange the result from a matrix of shape (OC * KH * KW, IH * IW) back into a tensor of shape (OC, OH, OW). Here, IC, IH, and IW refer to the number of input image channels, input image height, and input image width respectively, KH and KW refer to the convolution kernel height and convolution kernel width respectively, and OC, OH, and OW refer to the number of output image channels, output image height, and output image width respectively.

3. The method according to claim 1, characterized in that Adopt a three - stage matrix block strategy to adapt to the L3 cache capacity, including: The first - stage block: Divide the output matrix of M×N into multiple sub - matrices of M1×N1; The second - stage block: Divide the input matrix of K×N into multiple sub - matrices of K1×N1; The third - stage block: Divide the weight matrix of M×K into multiple sub - matrices of M1×K1; M1, N1, and K1 satisfy the following conditions: (M1 * N1 + K1 * N1 + M1 * K1) * sizeof(data_type) < L3_Cache_Size * 0.8; Among them, M1, N1, and K1 are the dimensions of the sub - matrices, which are respectively factors of M, N, and K; 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 Utilize the EDMA buffer ping - pong architecture to execute the calculation and transmission pipelines in parallel, including: Cache allocation strategy: Copy the three - stage blocks allocated in the L3_Cache by the three - stage matrix block strategy to obtain two three - stage caches with exactly the same size and shape, which are respectively called cache ping and cache pong. The suffix ping represents cache ping, and the suffix pong represents cache pong; It is required that M_ping, N_ping, and K_ping satisfy the following conditions: (M_ping * N_ping + K_ping * N_ping + M_ping * K_ping) * sizeof(data_type) < L3_Cache_Size * 0.4; Among them, M_ping, N_ping, and K_ping are the dimensions of the sub-matrices, which are respectively factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, with the unit of byte; L3_Cache_Size refers to the size of the available L3 cache, with the unit of byte; At the same time, M_pong, N_pong, K_pong have the exact same values as M_ping, N_ping, K_ping; EDMA and GEMM cooperation strategy: (1) Select the block division of the output matrix, input matrix, and weight matrix according to the three-segment matrix block division strategy; (2) Select cache ping / pong; (3) EDMA transfers the blocked input sub-matrix and weight sub-matrix from DDR to the cache ping / pong; (4) GEMM calculates to obtain the output sub-matrix; (5) EDMA transfers the blocked output sub-matrix from the cache ping / pong to DDR; (6) If the blocked sub-matrices have not been traversed completely, go back to (1), swap the cache ping-pong, until the sub-matrices are traversed, and complete the calculation of the entire matrix.

5. The method according to claim 1, wherein Based on the DSP's DMPYSP and DADDSP instructions, cooperate with pipeline optimization to achieve 4 FP32 multiplication operations and addition operations per cycle, including: Use the DSP's SIMD instructions to optimize the multiplication, addition, and data transfer methods in GEMM.

6. The method according to claim 1, wherein Multi-core parallel scheduling, combined with the OpenMP dynamic load balancing technology to achieve multi-core parallel acceleration, including: Task-level parallelism: Take the matrix multiplication calculation of the output sub-matrix of each matrix block as a task unit, evenly distribute the tasks to multiple cores for execution. Each DSP core is responsible for calculating the matrix multiplication task of the assigned output sub-matrix. The corresponding several input sub-matrices and output sub-matrices involved are all placed in the L2_Cache of each core. This is applicable to DSPs without L3 shared cache. At this time, the sub-matrix size satisfies the following conditions: (M1 * K + N1 * K + M1 * N1) * sizeof(data_type) < L2_Cache_Size * 0.8; Among them, M1, N1, and K1 are the dimensions of the sub-matrices, which are respectively factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, with the unit of byte; L2_Cache_Size refers to the size of the available L2 cache, with the unit of byte; If cache ping-pong is used, the sub-matrix size satisfies the following conditions: (M_ping * K + N_ping * K + M_ping * N_ping) * sizeof(data_type) < L2_Cache_Size * 0.4; Among them, M_ping, N_ping, and K_ping are the dimensions of the sub-matrices, and they respectively take values as factors of M, N, and K; sizeof(data_type) refers to the size of the data type of the matrix elements, with the unit of byte; L2_Cache_Size refers to the size of the available L2 cache, with the unit of byte. At the same time, M_pong, N_pong, K_pong have exactly the same values as M_ping, N_ping, K_ping. Data-level parallelism: Directly use the OpenMP pre-compilation directive #pragma omp parallel for to perform multi-core parallel optimization on the innermost loop of GEMM, which is applicable to DSPs with multi-core shared caches.

Images