Data processing method and device, processor and electronic equipment

Abstract

The invention relates to the technical field of artificial intelligence chips, and discloses a data processing method and device, a processor and electronic equipment. Determining a first jump step length required for rearranging the block data based on the block size of the input data in the first dimension; secondly, performing jump reading on the block data according to a first jump step length, rearranging and loading read input elements into a thread private register so as to realize transposition storage of the block data in the thread private register, and mapping a thread direction on a second dimension so as to obtain a data basis of convolution correlation calculation; convolution correlation calculation is executed based on the transposed data in the thread private register and the weight data in the thread private register, repeated data reading in the calculation process is reduced during sliding window calculation, and an operator is converted into calculation performance bottleneck operation from memory access performance bottleneck operation.

Description

Technical Field

[0001] The present application relates to the field of artificial intelligence chip technology, and in particular to a data processing method, device, processor, and electronic device. Background Art

[0002] The depthwise separable convolution operator decomposes the standard convolution operation into two steps: depthwise convolution and pointwise convolution. This reduces the number of parameters and computational complexity, thereby improving computational efficiency. However, in practical applications, this operator suffers from frequent data duplication due to its sliding window-type computational logic. Therefore, an effective data processing method is needed to optimize the depthwise separable convolution operator. Summary of the Invention

[0003] The present application provides a data processing method, device, processor and electronic device, which solves the technical problem of repeated data reading caused by the computational logic characteristics of the sliding window type in the related art, and achieves the technical effect of transforming the depthwise separable convolution operator from a memory access performance bottleneck operation to a computational performance bottleneck operation.

[0004] In order to achieve the above objectives, the main technical solutions adopted in this application include: In a first aspect, an embodiment of the present application provides a data processing method, the method comprising: When block data of input data needs to be loaded into a thread-private register, determining a first jump step required for rearranging the block data based on a block size of the input data in a first dimension; Performing jump reading on the block data according to the first jump step size, reordering and loading the read input elements into the thread-private register, so that the block data is stored in a transposed manner in the thread-private register, and mapping the thread direction onto a second dimension, where the first dimension is an adjacent lower dimension of the second dimension; A plurality of convolution-related calculations are performed based on the transposed data in the thread-private register and the weight data in the thread-private register; and duplicate data exists between corresponding portions of data in the transposed data between adjacent convolution-related calculations.

[0005] Optionally, performing a plurality of convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register includes: The transposed data in the thread-private register is traversed along a first dimension of the input data, and a plurality of convolution-related calculations are performed using the traversed input elements and weight elements included in the weight data.

[0006] Optionally, the transposed data is divided into a plurality of operation regions distributed along the first dimension, the operation regions corresponding to target input elements in the transposed data; and traversing the transposed data in the thread-private register along the first dimension of the input data, and performing a plurality of convolution-related calculations using the traversed input elements and weight elements included in the weight data, includes: performing a first convolution-related calculation using the weight element and a target input element corresponding to a first operating area among the multiple operating areas; After completing the first convolution-related calculation, convolution sliding is performed along the first dimension, and different operation areas overlap with each other. Subsequent convolution-related calculations are performed using the weight elements and the target input elements corresponding to each operation area after the first operation area.

[0007] Optionally, the size of the convolution kernel is denoted as R×S; R and S are positive integers greater than 1; When performing convolution sliding along the first dimension, N target input elements are slid along the first dimension each time; where N is a positive integer greater than or equal to 1 and less than S.

[0008] Optionally, the block data is loaded from a group shared memory into the thread private register; the method further includes: The input data is loaded from the chip memory into the group shared memory using a block loading strategy.

[0009] Optionally, determining a first jump step required for rearranging the block data based on a block size of the input data in a first dimension includes: The first jump step size is determined based on a block size of the input data in a first dimension and a padding operation range.

[0010] Optionally, the transposed data includes a first input element and a second input element, and the weight data includes a first weight element and a second weight element; and performing a plurality of convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register includes: In a current clock cycle of a current multiply-add instruction, two multiply-add operations are performed using the first input element, the second input element, the first weight element, and the second weight element to obtain an accumulation result of the current multiply-add instruction.

[0011] Optionally, the method further includes: Storing the calculation result of the convolution-related calculation in the thread-private register; Performing jump reading on the calculation result in the thread private register, and reordering and storing the read output elements in the group shared memory, so that the calculation result is transposed and restored in the group shared memory; The transposed calculation result in the group shared memory is stored from the group shared memory to the chip memory.

[0012] Optionally, the skipping and reading of the calculation result in the thread-private register includes: The block data is skipped and read according to a second skip step size, wherein the second skip step size is determined based on the block size of the input data in the first dimension.

[0013] In a second aspect, an embodiment of the present application provides a data processing device, the device comprising: a step size determining module, configured to determine, when block data of input data needs to be loaded into a thread-private register, a first jump step size required for rearranging the block data based on a block size of the input data in a first dimension; a data rearrangement module, configured to perform jump reading of the block data according to the first jump step size, rearrange and load the read input elements into the thread-private register, so that the block data is stored in a transposed manner in the thread-private register, and map the thread direction onto a second dimension, where the first dimension is an adjacent lower dimension of the second dimension; A calculation execution module is used to perform multiple convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register.

[0014] In a third aspect, an embodiment of the present application provides a processor, including an instruction parsing unit, an execution unit, a memory, and a thread-private register, wherein: The instruction parsing unit is used to receive and parse data processing instructions, wherein the data processing instructions are used to load input data from the memory into the thread-private register and perform convolution-related operations; The execution unit executes any of the above data processing methods after the instruction parsing unit parses the data processing instruction.

[0015] In a fourth aspect, an embodiment of the present application provides an electronic device, including: a memory that non-transitorily stores computer-executable instructions; a processor configured to execute the computer-executable instructions, Wherein, when the computer executable instructions are executed by the processor, the data processing method according to any one of the above items is implemented.

[0016] In an embodiment of the present application, first, when it is necessary to load block data of input data into a thread-private register, a first jump step required for rearranging the block data is determined based on the block size of the input data in a first dimension; secondly, the block data is jump-read according to the first jump step, and the read input elements are rearranged and loaded into the thread-private register, so that the block data is transposed and stored in the thread-private register, and the thread direction is mapped onto the second dimension, thereby obtaining a data basis for convolution-related calculations, and thus performing convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register, reducing repeated data reading in the calculation process during sliding window calculation, and transforming the operator from a memory access performance bottleneck operation to a computing performance bottleneck operation. BRIEF DESCRIPTION OF THE DRAWINGS

[0017] In order to more clearly illustrate the specific implementation methods of the present application or the technical solutions in the prior art, the following is a brief introduction to the drawings required for use in the specific implementation methods or the description of the prior art. Obviously, the drawings described below are some implementation methods of the present application. For ordinary technicians in this field, other drawings can be obtained based on these drawings without any creative work.

[0018] Figure 1 A schematic structural diagram of a general-purpose graphics processor provided in an embodiment of the present application; Figure 2 A schematic diagram of the computational logic of the depthwise separable convolution provided in an embodiment of the present application; FIG3( a ) is a schematic diagram of data loading during a depthwise separable convolution calculation process according to an embodiment of the present application; FIG3( b ) is a schematic diagram of data storage during a depthwise separable convolution calculation according to an embodiment of the present application; Figure 4 Schematic diagram of the calculation process of depth-wise separable convolution provided in an embodiment of the present application; Figure 5 A flowchart of a data processing method provided in an embodiment of the present application; FIG6 (a) is a schematic diagram of the TMA working process provided in an embodiment of the present application; FIG6( b ) is a schematic diagram of the illustrated process of the ldmatrix instruction provided in an embodiment of the present application; Figure 7 A schematic diagram of the convolution sliding process in the convolution-related calculation provided in an embodiment of the present application; Figure 8 A schematic diagram of the convolution-related calculation process provided in an embodiment of the present application; Figure 9 The calculation process of the traditional calculation instructions provided in the embodiments of the present application; Figure 10 The calculation process of the new calculation instruction provided in the embodiment of the present application; Figure 11 A flowchart of a data processing method provided in yet another embodiment of the present application; Figure 12 A schematic diagram of the framework of a data processing device provided in an embodiment of the present application; Figure 13 A schematic structural diagram of a processor provided in an embodiment of the present application; Figure 14 A schematic structural diagram of a computer device provided in an embodiment of the present application. DETAILED DESCRIPTION

[0019] To make the purpose, technical solutions, and advantages of the embodiments of the present application more clear, the technical solutions in the embodiments of the present application will be clearly and completely described below in conjunction with the drawings in the embodiments of the present application. Obviously, the described embodiments are part of the embodiments of the present application, not all of the embodiments. Based on the embodiments in the present application, all other embodiments obtained by those skilled in the art without making creative efforts shall fall within the scope of protection of this application.

[0020] To facilitate the description of the present application, concepts related to the present application are introduced below.

[0021] The data processing method of this application is based on artificial intelligence (AI). Artificial intelligence is the theory, methods, techniques, and application systems that use digital computers or machines controlled by digital computers to simulate, extend, and expand human intelligence, perceive the environment, acquire knowledge, and use that knowledge to achieve optimal results. In other words, artificial intelligence is a comprehensive technology in computer science that attempts to understand the essence of intelligence and produce new intelligent machines that can respond in a manner similar to human intelligence.

[0022] The data processing method of the present application can be based on depthwise separable convolution. Depthwise separable convolution is a convolution operation widely used in the fields of computer vision and deep learning. It is a special convolution operation that reduces the number of parameters and the amount of computation by decomposing the standard convolution operation, thereby improving the efficiency and performance of the model. Depthwise separable convolution includes two steps: depthwise convolution and pointwise convolution. In the depthwise convolution stage, each input channel is convolved with an independent convolution kernel, which helps to extract the features of the input data. In the pointwise convolution stage, the output result of the depthwise convolution is convolved with a 1x1 convolution kernel to integrate and interact the information between channels. In this way, the number of parameters and the amount of computation can be reduced, thereby improving computational efficiency.

[0023] During operator development on a general-purpose graphics processing unit (GPGPU), three memory spaces are typically used: global shared memory (GSM), thread local registers (TLRs), and high-bandwidth memory (HBM). The TLRs are closest to the compute units and have the highest speed, typically used to store temporary data for operator calculations. The GSM is the second fastest and is typically used to temporarily store data that needs to be used repeatedly. The HBM, the GPGPU's external memory, is farther from the compute units and has the lowest speed. It is typically used to store input and output data for operator calculations, as well as all other data involved. It should be noted that thread local registers, as registers unique to each thread, can be used to store thread-specific data and temporary variables. In multithreaded processing, each thread has its own TLR, which speeds up data access and reduces contention, thereby improving computing performance and concurrency.

[0024] This application exemplifies the parameter expression of depth-wise separable convolution: Input data x: [N, D, H, W, C] Weight data w: [K, T, R, S, C] → [C, T, R, S, 1] Output data y: [N, Z, P, Q, K] → [N, T, P, Q, C] Step size stride: [strideD, strideH, strideW] Padding: [padD, padH, padW] Dilation: [dilationD, dilationH, dilationW] The dimensions of the input data x are [N, D, H, W, C], where N represents batch, D represents depth, H represents height, W represents width, and C represents input channel. The dimensions of the weight data w are [K, T, R, S, C], where K represents output channel, T represents filter depth, R represents filter height, and S represents filter width. The dimensions of the output data y are [N, Z, P, Q, K], where Z represents output depth, P represents output height, and Q represents output width. It can be understood that depthwise separable convolution includes depthwise convolution and pointwise convolution. In the depthwise convolution stage, [K, T, R, S, C] → [C, T, R, S, 1] represents the dimensionality change during the depthwise convolution stage, while in the pointwise convolution stage, [N, Z, P, Q, K] → [N, T, P, Q, C] represents the dimensionality change during the pointwise convolution stage. [strideD, strideH, strideW] represents the sliding stride of the convolution kernel in depth, height, and width. [padD, padH, padW] represents the amount of border padding in depth, height, and width, and [dilationD, dilationH, dilationW] represents the amount of dilation in depth, height, and width.

[0025] The data processing method of this application can be applied to an artificial intelligence processor, which can be any of a GPU (Graphics Processing Unit), a TPU (Tensor Processing Unit), or a GPGPU (General-Purpose Graphics Processing Unit). A general-purpose graphics processor (GPGPU) is used as an example for illustration.

[0026] Figure 1 The following is a schematic diagram of the structure of a general-purpose graphics processing unit (GPGPU). Figure 1 ,A general-purpose graphics processor is actually an array of ,programmable multi-processors, such as Figure 1 As shown, a general purpose graphics processor includes multiple computing units, such as Figure 1Computational unit 1, computational unit 2, ..., computational unit N, where N is a positive integer. Each Compute Unit (CU) is used to perform arithmetic and logical operations, such as accumulation, reduction, conventional addition, subtraction, multiplication, and division. A Computational Unit includes multiple Processing Units (EUs, also known as Execution Units). Each Processing Unit includes an Arithmetic Logic Unit (ALU), a Floating Point Unit, etc. The Processing Unit is used to perform specific computing tasks. In addition, the Processing Unit includes registers (e.g. Figure 1 The register group in the computing unit is used to store source data and destination data related to the computing task. The shared memory in a computing unit is used to share data between the execution units of the computing unit.

[0027] In parallel computing, computing tasks are typically executed by multiple threads. These threads are divided into multiple thread blocks before execution on a general-purpose graphics processing unit (GPU), also known as a parallel computing processor. These thread blocks are then distributed to various compute units (CUs). All threads in a thread block must be assigned to the same CU for execution. Furthermore, thread blocks are split into minimum execution warps (warps), each containing a fixed number (or less) of threads, for example, 32 threads. Multiple thread blocks can execute in the same CU or in different CUs. It should be noted that multiple threads in each warp execute the same instructions. A warp can contain 32 threads or any other number of threads. ldmatrix is ​​a warp-level thread coordination instruction that loads 32 elements sequentially into an intermediate register based on the starting address for all threads within the warp. ldmatrix also provides the ability to load data at a specific step size and can work with the Tensor Memory Accelerator (TMA) to implement data swizzling in the group shared memory (GSM) to avoid GSM bank conflicts.

[0028] When using a general-purpose graphics processor to perform depth-separable convolution, the general implementation method is to load the required input data (activation) into shared memory or registers, load the weight parameters of the convolution layer into registers, and then perform multiplication and addition operations one by one to perform a single-channel two-dimensional planar convolution operation, thereby realizing the convolution operation.

[0029] Figure 2This is a schematic diagram of the computational logic of depthwise separable convolution. The input data includes two-channel input feature maps 102. Each channel has a corresponding convolution kernel 104. The number of convolution kernels is the same as the number of channels, which is 2. Each input feature map 102 is convolved with the corresponding convolution kernel 104 in a two-dimensional plane. After channel-by-channel convolution, two output feature maps 106 are generated. Please continue to refer to Figure 2 The convolution kernel 104 is represented as a 2×2 weight matrix, and the data layout of each weight matrix in the register group is recorded as weight data layout 108. The input feature map 102 is represented as a 3×3 input matrix, and the data layout of each input matrix in the register group is recorded as input data layout 110. The output feature map 106 is represented as a 2×2 output matrix, and the data layout of each output matrix in the register group is recorded as output data layout 112.

[0030] By analyzing the computational logic of depthwise separable convolution, it is found that for each output point, the amount of input data x and weight data w that needs to be read is R×S. There is a lot of data that needs to be read repeatedly, and the number of read instructions issued needs to be reduced. Furthermore, due to the sliding window type computational logic characteristics of depthwise separable convolution, there are many data repeated read behaviors, so the execution time of the operator is mainly determined by the speed and latency of memory access, and is generally considered to be a memory-bound operator. Even if a cache is used to reduce memory access, it may still be bounded by the speed and latency of cache access. In order to reduce the amount of data access to improve the performance ceiling of the operator, this application provides a data processing method to convert the operator into a computationally bound operator.

[0031] It should be noted that the execution time of a computationally bound operator is primarily determined by the number of operations required for the computation, rather than by memory bandwidth or latency. In this case, the performance of the operator depends more on the speed of the computation unit, and memory access becomes a relatively less significant bottleneck.

[0032] It should be noted that the data processing method provided in this application is a hardware acceleration solution for depthwise separable convolution computations. The key lies in transforming operator performance from a memory-bound bottleneck to a compute-bound bottleneck through data rearrangement and storage optimization. This method is compatible with some hardware architectures, leveraging their tensor memory accelerators (TMAs) and thread-level parallelism to achieve efficient data flow.

[0033] Furthermore, embodiments of the present application provide a scenario example of this data processing method. In this scenario example, as shown in Figure 3(a), chip memory 302 stores input data and weight parameters. The input data is loaded from chip memory 302 into group shared memory 304 using the Tensor Memory Access (TMA) instruction cpglm2gsm. The input data is then loaded and transposed from group shared memory 304 into thread-private registers 306 using the stride method in the ldmatrix instruction.

[0034] The weight parameters are loaded from chip memory 302 into group shared memory 304 using the TMA (Tensor Memory Access) instruction cpglm2gsm. The weight parameters are then loaded from group shared memory 304 into thread-private registers 306 using the ldmatrix instruction. A depthwise separable convolution calculation is performed in thread-private registers 306 based on the weight parameters and the transposed input data. This depthwise separable convolution calculation result is then stored in the thread-private registers.

[0035] Referring to Figure 3(b), after the depthwise separable convolution calculation is completed, the depthwise separable convolution calculation result in the thread-private register 306 needs to be stored in the chip memory 302. The stride method in the stmatrix instruction is used to read and transpose the depthwise separable convolution calculation result from the thread-private register and restore it to the group shared memory 304. The instruction cpgsm2glm is used to store the transposed depthwise separable convolution calculation result in the chip memory 302.

[0036] It should be noted that in this scenario example, the parallelism of the calculation process and memory access process is increased by asynchronously loading the input data and weight parameters into the group shared memory in advance.

[0037] As you can understand, the instructions cpglm2gsm, ldmatrix, cpgsm2glm, and stmatrix are part of the general-purpose graphics processor architecture. The cpglm2gsm instruction copies data from on-chip memory (Global Memory) to group shared memory (Shared Memory). The ldmatrix instruction loads data from group shared memory into thread-private registers at a specified stride, automatically transposing the data as needed to match computational requirements. The cpgsm2glm instruction copies data from group shared memory back to on-chip memory (Global Memory). The stmatrix instruction reads data from thread-private registers, transposes it according to the stride, and stores it into group shared memory. The stride method organizes data using a specified stride to optimize memory access patterns.

[0038] For example, see Figure 4 , assuming that the input data 402 is a 4×4 matrix; the dimension of the input data 402 is recorded as W×H; the weight parameter 404 is a 3×3 matrix. For the input data 402, the stride method in the instruction ldmatrix is ​​used to load and transpose the input data 402 from the group shared memory GSM to the thread private register TLR of the register group. The transposed input data stored in the thread private register TLR of the register group is recorded as 406. The thread bundle warp includes the first thread th0, the second thread th1, the third thread th2, and the fourth thread th3. It should be noted that the first thread th0, the second thread th1, the third thread th2, and the fourth thread th3 are independent and parallel. Please continue to refer to Figure 4 When the thread direction is placed in the height direction (H) of the input data, adjacent threads process data at adjacent height positions on the input data, which helps to increase the locality of data access because the data along the height direction is usually in the same memory area, reducing the need for cross-region reading.

[0039] Let's take the first thread th0 as an example. First, the corresponding matrix is ​​retrieved from the transposed input data, denoted as the first input matrix. The matrix elements in the first input matrix include a11, a12, a13, a21, a22, a23, a31, a32, and a33. First thread th0 uses the first input matrix and weight parameter 404 to perform calculations, obtaining the first result element (0×a11+1×a12+2×a13+3×a21+4×a22+5×a23+6×a31+7×a32+8×a33). Next, during the calculation, the multiplication-addition window corresponding to weight parameter 404 is shifted along the W direction, and another corresponding matrix is ​​retrieved from the transposed input data, denoted as the second input matrix. There are duplicate matrix elements between the first and second input matrices. The matrix elements in the second input matrix include a12, a13, a14, a22, a23, a24, a32, a33, and a34. The first thread th0 uses the second input matrix and weight parameter 404 to perform a calculation, obtaining the second result element (0×a12+1×a13+2×a14+3×a22+4×a23+5×a24+6×a32+7×a33+8×a34). It is understood that the first and second result elements are stored in thread-private registers. It should be noted that the matrix elements a12, a13, a22, a23, a32, and a33 are reused during the call to the first and second input matrices, improving data reuse, reducing the number of instruction issues, and transforming the depthwise separable convolution operator into a bounded computation instruction.

[0040] Let's take the second thread th1 as an example. First, the multiplication-add window corresponding to weight parameter 404 is moved to a portion of the matrix in the transposed input data, denoted as the third input matrix. The matrix elements in the third input matrix include a21, a22, a23, a31, a32, a33, a41, a42, and a43. The second thread th1 uses the third input matrix and weight parameter 404 to perform calculations, obtaining the third result element (0×a21+1×a22+2×a23+3×a31+4×a32+5×a33+6×a41+7×a42+8×a43). Next, during the calculation process, the multiplication-add window corresponding to weight parameter 404 is moved along the W direction, and the corresponding portion of the matrix is ​​retrieved from the transposed input data, denoted as the fourth input matrix. There are duplicate matrix elements between the third and fourth input matrices. The matrix elements in the fourth input matrix include a22, a23, a24, a32, a33, a34, a42, a43, and a44. The second thread th1 uses the fourth input matrix and weight parameter 404 to perform calculations, obtaining the fourth result element (0×a22+1×a23+2×a24+3×a32+4×a33+5×a34+6×a42+7×a43+8×a44). It is understood that the third and fourth result elements are stored in thread-private registers. It should be noted that the matrix elements a22, a23, a32, a33, a42, and a43 are reused when calling the third and fourth input matrices, improving data reuse, reducing instruction issuance, and transforming the depthwise separable convolution operator into a bounded computation instruction.

[0041] It is understandable that the calculation process of any of the above threads is similar, and no further description will be given for other threads (such as the third thread th2, the fourth thread th3, etc.).

[0042] According to an embodiment of the present application, a data processing method embodiment is provided. It should be noted that the steps shown in the flowchart of the accompanying drawings can be executed in a computer system such as a set of computer executable instructions, and although a logical order is shown in the flowchart, in some cases, the steps shown or described can be executed in an order different from that shown here. Figure 5 , Figure 5 The data processing method in the embodiment of the present application is shown as a flow chart. The data processing method includes the following steps: S110 : When it is necessary to load block data of input data into a thread-private register, determine a first jump step required for rearranging the block data based on a block size of the input data in a first dimension.

[0043] The input data can be the input feature map (activation) to be processed in the depthwise separable convolution, such as an image pixel matrix or a multi-dimensional tensor collected by a sensor. The block data of the input data can be a fixed-size data block obtained by dividing the input data. The block size can be determined by the hardware storage level, such as cutting the input data into multiple tiles based on the capacity limit requirements of GSM and the pipeline arrangement requirements of the calculation. For example, the input data can be an input tensor of size [H, W, C], which is loaded from the chip memory to the GSM using a block specification of 32 (height) × 16 (width) × 32 (channels).

[0044] In some embodiments, the block data is loaded from the group shared memory into the thread private register. The data processing method may further include: loading the input data from the chip memory into the group shared memory using a block loading strategy.

[0045] Warp shared memory refers to the on-chip cache (Global Shared Memory, GSM) shared by warps in some architectures. For example, this GSM is shared by 32 threads and is used to temporarily store data blocks loaded from on-chip memory.

[0046] The chip memory can be an off-chip high-capacity storage device (such as HBM) that stores the complete input data. For example, the original input data resides in the chip memory and can be loaded into the group shared memory through asynchronous transfer.

[0047] The block loading strategy can refer to a method of transferring input data in batches from chip memory to group shared memory through TMA, which includes block partitioning and boundary processing.

[0048] For example, referring to Figure 6(a), the input data is divided into grid-like blocks of a fixed size (32×16×32), and the block boundaries (shown as dashed boxes in Figure 6(a)) are expanded based on the convolution pad value. For example, when pad = 1, the actual loading size of the block is expanded to 34 (H) × 18 (W) × 32 (C), ensuring boundary computation integrity. It can be understood that the load size is (34, 18), and the number of load channels is 32.

[0049] In this embodiment, the instruction cpglm2gsm is first used to load the input data from the chip memory into the group shared memory using the block loading strategy; secondly, the block data is loaded from the group shared memory into the thread private register through the stride method in the quality ldmatrix, thereby moving the data to the private storage space with lower latency access.

[0050] The Thread Local Register (TLR) is a high-speed memory unit unique to a hardware thread, with significantly lower access latency than off-chip memory. Loading block data into the TLR is particularly important during the convolution phase, where data from the GSM must be moved to the TLR for multiplication and accumulation.

[0051] The first dimension can be a specific dimension of the input data, such as the width (w) of the feature map. The block size can be the length of the block data in the first dimension, such as 16 columns or 18 columns. The first jump step is a step value calculated based on the block size and is used to control the interval between data jumps. For example, when the block size is 32×16, the step size can be equal to 16.

[0052] Here, rearranging the block data may refer to arranging the read input elements in a new order, for example, converting a matrix stored in row-major order to a matrix stored in column-major order through a transposition operation.

[0053] In order to ensure that the non-continuous layout of data in memory can be efficiently loaded by hardware instructions, it is necessary to calculate the interval (step value) used to control the data jump reading. Furthermore, when it is necessary to load the block data of the input data into the thread-private register, in order to optimize the data layout so that the subsequent convolution calculation can be efficiently performed along the thread direction and reduce the data transfer overhead between registers, it is necessary to use non-continuous reading and transposition operations to form a continuous transposed matrix in the register. Therefore, in this embodiment, the first jump step required for rearranging the block data is determined based on the block size of the input data in the first dimension. For example, using the LDMATRIX32r stride18 instruction, 32 rows and 18 columns of block data are read from GSM with a step size of 18 to rearrange the block data.

[0054] In some embodiments, determining a first jump step required for rearranging the block data based on the block size of the input data in the first dimension may include: determining the first jump step based on the block size of the input data in the first dimension and a padding operation range. For example, for block data with a block size of 32×16, to ensure that all data required for sequential calculations is obtained, the block data may be padded in the w direction by inserting corresponding columns before the first column and after the last column in the w direction, and inserting corresponding rows before the first row and after the last row in the h direction, and padding with zero values, to obtain data-padded block data of 34×18. The step size may be equal to 18.

[0055] S120 , performing jump reading on the block data according to the first jump step length, rearranging and loading the read input elements into the thread-private register, so that the block data is transposed and stored in the thread-private register, and mapping the thread direction onto the second dimension.

[0056] The first dimension is the adjacent lower dimension of the second dimension. Jump reading can be extracting data from block data in a non-contiguous manner, for example, reading 18 elements apart from each row of data. Thread direction mapping in the second dimension can be aligning the execution direction of the thread with the second dimension, for example, binding the thread index (such as the thread serial number thx) to the height direction (h direction) so that each thread processes continuous height position data. In other words, adjacent threads process data at adjacent height positions on the same feature map, which helps to increase the locality of data access because the data along the height direction is usually in the same memory area, reducing the need for cross-region reads.

[0057] Specifically, jump reading can be implemented through the hardware instruction ldmatrix. This instruction is a warp-level collaborative instruction that coordinates 32 threads to read data from GSM non-contiguous addresses in a stride manner. The stride parameter of the ldmatrix instruction must be consistent with the first jump step size to ensure the correct arrangement of the data in the register. Furthermore, in order to achieve data transposition storage, through jump reading, the data stored in the original order is reorganized in the TLR into a transposed layout arranged by thread index. After reorganization, the thread direction is aligned with the second dimension (H dimension), so that the same thread continues to process data at the same height position, thereby reducing the redundant data copied in the sliding window calculation and reducing the amount of memory access.

[0058] For example, please refer to Figure 6 (b), which shows the diagrammatic process of the ldmatrix instruction. Figure 6 (b) shows that a warp thread group contains 32 threads (th0, th1, th2...th31). Note that the method of this example can also be used for other numbers of threads. ldmatrix is ​​a warp-level thread collaboration instruction, which means that all threads in the warp can load data corresponding to 32 coordinates into registers (r1, r2, r3...r52) based on the starting address. It can be understood that the 32 coordinates here can be an independent address that is respectively responsible for the 32 threads in the warp, and there is a fixed interval between each coordinate.

[0059] The input data is three-dimensional, but each channel is independent. Therefore, the input data is processed as two-dimensional data, such as the input data 610 being transferred to the GSM shown in Figure 6(b). The input data is transferred to the GSM, and the input data size in the GSM is 34×18, that is, H=34 and W=18. The input data transferred to the GSM includes a first data block 612 (H=0 to H=31), a second data block 614 (H=1 to H=32), and a third data block 616 (H=2 to H=33). It can be understood that the first data block 612, the second data block 614, and the third data block 616 represent data in different row ranges of the input data. When using a 3×3 sliding window 620, the first data block 612 corresponds to the first row 622 of the sliding window, the second data block 614 corresponds to the second row 624 of the sliding window, and the third data block 616 corresponds to the third row 626 of the sliding window.

[0060] To reduce repeated data reading during sliding window calculations, the input data in the GSM needs to be rearranged when entering the thread-private register TLR, adjusting the H and W directions of the input data. When using a thread to read data in the W direction, a jump read with stridew=18 is performed. Thread th0 reads data block A in the upper left corner, thread th1 reads data block B with 1 added to the H direction, thread th2 reads data block C with 2 added to the H direction, and so on. 32 threads (th0, th1, th2...th31) are used to read a column of data in the GSM. For example, the data read by register r1 is the leftmost first column of data 630 in the first data block 612. Similarly, the data read by the corresponding register is the leftmost first column of data in the second data block 614 and the leftmost first column of data in the third data block 616. No further details will be given.

[0061] In summary, first, the first jump stride is determined based on the width of the input data in the GSM. Specifically, the first jump stride in the width direction (W) is equal to 18, with stridew = 18. Second, by reading the input data in the GSM through jumps, the input data is rearranged, with the thread direction aligned with the h-direction of the input data. This satisfies the requirements of sliding window calculations and ensures efficient data reading.

[0062] S130 , performing multiple convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register.

[0063] The weight data can be a parameter matrix used in a convolution operation, such as the filter weights in a depthwise separable convolution. Convolution-related calculations involve performing multiplication and addition operations on the transposed input data and the weight data, such as performing convolution operations on feature maps through matrix multiplication. Adjacent convolution-related calculations may contain duplicate data between corresponding portions of the transposed data.

[0064] Specifically, partial data that matches the weight data structure can be read from the transposed data, and convolution-related calculations can be completed for the read partial data and weight data through warp-level parallel calculations. The layout of the transposed data in the thread-private register provides a data basis for the subsequent multiple convolution-related calculations, and there is duplicate data between the partial data corresponding to each of the adjacent convolution-related calculations in the transposed data. Exemplarily, multiple convolution-related calculations are performed based on the transposed data in the thread-private register and the weight data in the thread-private register. Among the multiple convolution-related calculations, the adjacent convolution-related calculations include a first convolution-related calculation and a second convolution-related calculation. The partial data corresponding to the first convolution-related calculation in the transposed data is duplicated with the partial data corresponding to the second convolution-related calculation in the transposed data.

[0065] In the above embodiment, first, when it is necessary to load the block data of the input data into the thread-private register, the first jump step required for rearranging the block data is determined based on the block size of the input data in the first dimension; secondly, the block data is jump-read according to the first jump step, and the read input elements are rearranged and loaded into the thread-private register, so that the block data is transposed and stored in the thread-private register, and the thread direction is mapped on the second dimension, thereby obtaining the data basis for convolution-related calculations, and then performing convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register, reducing repeated data reading in the calculation process during sliding window calculation, and transforming the operator from a memory-bound operation to a compute-bound operation, thereby significantly improving the calculation efficiency.

[0066] In some embodiments, performing multiple convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register may include: traversing the transposed data in the thread-private register along a first dimension of the input data, and performing multiple convolution-related calculations using the traversed input elements and the weight elements included in the weight data.

[0067] The weight elements in the weight data can be convolution kernel parameters loaded into thread-private registers. For example, weight data of a fixed size (such as R×S×C) can be loaded from global memory to thread-private registers in advance through asynchronous TMA transmission to ensure zero-latency access during calculation. The traversed input elements can refer to the input elements covered by the current position of the sliding window. For example, when the convolution kernel size is R×S, when traversing to the current position (W=k, H=m), in the W direction, the input elements in the sliding window are from the kth column to the k+S-1th column; in the H direction, the input elements in the sliding window are from the mth column to the m+R-1th column.

[0068] Specifically, the transposed data in the thread-private register is traversed along the first dimension of the input data, and each time the transposed data is traversed, convolution-related calculations are performed using the input elements traversed each time and the weight elements included in the weight data. To exemplify the traversal operation in this embodiment, the first dimension can be the width direction (W) of the block data. In this way, the traversal operation slides the convolution window along the W direction at a fixed step size, for example, starting from the W=0 position, and sliding 1 input element each time until the block boundary. It should be noted that by sliding sequentially along the first dimension (W direction), the spatial continuity of the data in the thread-private register is utilized, so that each time the convolution window slides, 1 column of new data needs to be incrementally loaded, thereby reducing the memory access amount of the sliding window operation. During the traversal process, the weight data and the transposed data reside in the thread-private register throughout the process, reducing the repeated reading of data during the calculation process.

[0069] In some embodiments, see Figure 7 , Figure 7 The figure shows a schematic diagram of the convolution sliding process in convolution-related calculations. The transposed data is divided into multiple operation areas distributed along the first dimension, and the multiple operation areas include a first operation area 702, a second operation area 704, ... until the final N-th operation area 706. Each operation area corresponds to a target input element in the transposed data. Two adjacent operation areas have repeated input elements. The operation area can be the input data calculation range defined for a single convolution output point in the thread private register. For example, Figure 7 As shown, when the convolution kernel width S=3, each operation area contains 3 columns of data in the W direction (such as the first operation area 702 and the second operation area 704), and its size is determined by the convolution kernel parameters. The target input element can be the subset of input data participating in the convolution calculation in the operation area. It can be understood that Figure 7 R0, R17, R18, R35, R36, and R53 in can be zero values ​​obtained by a padding operation.

[0070] See also Figure 8 , Figure 8The figure shows a schematic diagram of the convolution-related calculation process. The transposed data in the thread-private register is traversed along the first dimension of the input data, and multiple convolution-related calculations are performed using the traversed input elements and the weight elements included in the weight data, including: S810: Perform a first convolution-related calculation using the weight element and a target input element corresponding to a first operation area among the multiple operation areas.

[0071] S820. After completing the first convolution-related calculation, convolution sliding is performed along the first dimension, and different operation areas overlap with each other. Subsequent convolution-related calculations are performed using weight elements and target input elements corresponding to each operation area after the first operation area.

[0072] The first operation region is the operation region starting along the first dimension (W direction), such as first operation region 702. The first convolution-related calculation may be the first convolution operation based on first operation region 702. The target input element of the first operation region is dot-accumulated with the weight element in the thread-private register to generate the first matrix element of the output data. Each operation region after the first operation region includes a second operation region 704 and an Nth operation region 706.

[0073] In this embodiment, sliding the convolution along the first dimension may be an operation of moving the convolution window at a fixed step size (usually 1). After completing the first convolution-related calculation, the convolution window slides along the first dimension by one input element, forming a second operation area following the first operation area. The second operation area overlaps with the previous area (the first operation area). It should be noted that the overlap between different operation areas is an inevitable result of the convolution kernel size S being larger than the step size. The data of the overlapping part already resides in the thread-private register and does not need to be loaded repeatedly.

[0074] In some embodiments, the size of the convolution kernel is denoted as R×S, where R and S are positive integers greater than 1. When performing convolution sliding along the first dimension, N target input elements are slid along the first dimension each time, where N is a positive integer greater than or equal to 1 and less than S.

[0075] The convolution sliding process is illustrated using the adjacent first operating area 702 and second operating area 704 as examples. Convolution-related calculations are performed using weight elements and the target input elements of the first operating area 702. After the convolution-related calculations are completed, a matrix element of the output data is obtained. When performing convolution sliding along the first dimension, one target input element is slid along the first dimension each time to obtain the second operating area 704. After the convolution-related calculations are completed, the next matrix element of the output data is obtained. Convolution-related calculations are performed using weight elements and the target input elements of the second operating area 704. It can be understood that convolution-related calculations include: 1) incremental reading, which only loads the non-overlapping input elements between the second operating area 704 and the first operating area 702; 2) data reuse, which directly reuses the overlapping data in the thread-private register; and 3) continuous calculation, which performs dot multiplication and accumulation on the newly loaded target input elements and the weight elements to generate new matrix elements of the output data.

[0076] In some embodiments, the transposed data includes a first input element and a second input element, and the weight data includes a first weight element and a second weight element; performing convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register may include: within the current clock cycle of the current multiplication-addition instruction, using the first input element, the second input element, the first weight element and the second weight element to perform two multiplication-addition operations to obtain the accumulated result of the current multiplication-addition instruction.

[0077] The transposed data refers to the input data in the thread-private register that has undergone the rearrangement operation. In some cases, traditional computing instructions can be used to perform multiplication and addition operations based on the transposed data and weight data. Figure 9 , Figure 9 The figure shows the calculation process of the traditional calculation instruction. Although the data reuse rate can be improved based on the transposed data, the number of issuances of the calculation instruction needs to be reduced when the traditional calculation instruction is used for calculation, which affects the upper limit of the efficiency of the issuance of the calculation instruction. Therefore, a new calculation instruction is provided in this embodiment, which can be recorded as HDP2.a2b2. Figure 10 , Figure 10 The figure shows the calculation process of the new calculation instruction. Comparing the traditional calculation instruction with the new calculation instruction, the traditional calculation instruction completes only one multiplication and addition operation per cycle. The new calculation instruction completes two multiplication and addition operations per cycle, achieving double the efficiency of the traditional multiplication and addition instruction.

[0078] Executing two multiply-add operations in the current clock cycle involves the following hardware behaviors: parallel multiplication and accumulation fusion. For parallel multiplication, the first input element multiplied by the first weight element and the second input element multiplied by the second weight element are calculated simultaneously. For accumulation fusion, the two multiplication results are added together to form the accumulated result of the current multiply-add instruction. For example, when the input elements are [0.2, 0.5] and the weight elements are [1.1, 0.8], the accumulated result = 0.2 × 1.1 + 0.5 × 0.8 = 0.22 + 0.40 = 0.62.

[0079] Understandably, please continue to refer to Figure 9 , Figure 9 src30 and src31 can be considered as initial values. src1, src2, and src30 are multiplied and added to get dst0. src1, src2, and src31 are multiplied and added to get dst1.

[0080] Understandably, please continue to refer to Figure 10 , Figure 10 src30 and src3 in the equation can be considered initial values. src10 and src11 are the first and second weight elements, respectively. src20 and src21 are the first and second input elements, respectively. Performing two multiplication-add operations on src10, src20, src11, src21, and src30 yields dst0. Performing two multiplication-add operations on src10, src20, src11, src21, and src3 yields dst1.

[0081] In the above embodiment, based on the fact that rearranging the input data can improve the data reuse rate, the use of the proposed new multiply-add instruction can reduce the number of instruction issuances. Figure 7 , which not only reduces the number of computational instructions by 1 / 3, but also transforms the depthwise separable convolution operator into a computationally bound instruction.

[0082] In some embodiments, see Figure 11 , the method may further include: S1210: Store the calculation result of the convolution-related calculation in a thread-private register.

[0083] S1220 , performing jump reading on the calculation result in the thread private register, and reordering and storing the read output elements into the group shared memory, so that the calculation result can be transposed and restored in the group shared memory.

[0084] S1230 , storing the transposed calculation result in the group shared memory from the group shared memory to the chip memory.

[0085] The results of convolution-related computations refer to the output tensor data generated by the depthwise separable convolution operation. For example, after the sliding window computation is complete, each output point is stored in a thread-private register with the data layout in transposed format. Storing the computation results in thread-private registers avoids frequent writes back to memory.

[0086] In this embodiment, computation results are stored in three stages: 1) thread-private registers; 2) group shared memory; and 3) on-chip memory. It should be noted that the computation result storage process is symmetrical with the aforementioned data rearrangement process. Because the input data is rearranged during the input phase prior to computation, the input transposition operation is reversed during the output phase through a jump read, restoring the original data layout. From transposed storage during the input phase to transposed restoration during the output phase, a closed-loop data layout transformation is formed.

[0087] Jump reads may refer to extracting data from a thread-private register using a non-contiguous address pattern. A stride access strategy symmetrical to the input data rearrangement may be employed, with a second jump stride length determined based on the block size of the input data along a first dimension. Jump reads of calculation results in the thread-private register may include: jump reading the block data according to the second jump stride length.

[0088] In this embodiment, the transposed layout in the thread-private register is restored to the original dimensional order through layout conversion. The output elements of each thread are written to the group shared memory according to the second jump step through the Warp collaboration instruction (such as stmatrix) to obtain the transposed calculation result in the group shared memory. The transposed calculation result can be understood as the data with the original layout restored in the group shared memory. Storage from the group shared memory to the chip memory can be achieved through TMA asynchronous transmission. For example, data can be transferred in batches at a block granularity (such as 32×16×32), and the transmission process is executed in parallel with the calculation of the next block data.

[0089] The present application provides a data processing device. Figure 12 13 is a schematic diagram of a data processing device according to an embodiment of the present invention. The data processing device 1300 may include a step size determination module 1310 , a data rearrangement module 1320 , and a calculation execution module 1330 .

[0090] a step size determining module 1310 for determining, when block data of input data needs to be loaded into a thread-private register, a first jump step size required for rearranging the block data based on a block size of the input data in a first dimension; a data rearrangement module 1320 configured to perform jump reading on the block data according to the first jump step size, rearrange and load the read input elements into the thread-private register, so that the block data is stored in a transposed manner in the thread-private register, and map the thread direction onto the first dimension; The calculation execution module 1330 is configured to perform convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register.

[0091] For the convenience of description, the above devices are described as being divided into various modules according to their functions. Of course, when implementing this application, the functions of each module can be implemented in the same or multiple software and / or hardware.

[0092] The embodiment of the present application provides a processor, Figure 13 : is a schematic structural diagram of the processor provided in the embodiment of the present application. Figure 13 As shown, the processor includes an instruction parsing unit, an execution unit, a memory, and thread-private registers.

[0093] The instruction parsing unit is used to receive and parse data processing instructions, wherein the data processing instructions are used to load input data from the memory into the thread-private register and perform convolution-related operations.

[0094] The execution unit executes the data processing method according to any of the above embodiments after the instruction parsing unit parses the data processing instruction.

[0095] The data processing instructions include, for example, data loading instructions and calculation instructions. The data loading instructions are used to load input data from the memory into the thread private register, and the calculation instructions are used to perform convolution-related operations on the input data and weight data. Regarding the specific process of using the execution unit to execute the data processing instructions, reference can be made to the various steps in the data processing method described above, and the repeated parts will not be repeated. The processor can achieve technical effects similar to the aforementioned data processing method, and the repeated parts will not be repeated.

[0096] The processor involved in the embodiments of the present application can be any one of a CPU (Central Processing Unit), a GPU (Graphics Processing Unit), a TPU (Tensor Processing Unit), an NPU (Neural Network Processing Unit), a DPU (Deep Learning Processing Unit), an APU (Accelerated Processing Unit), and a GPGPU (General-Purpose computing on Graphics Processing Unit), and is determined when the embodiments of the present application are applied to a specific product or technology.

[0097] Embodiments of the present application provide a non-transitory computer-readable storage medium. The storage medium may be a non-transitory computer-readable storage medium that can non-transitorily store one or more computer-readable instructions. For example, when the computer-readable instructions are executed by a processor, one or more steps of the aforementioned data processing method may be performed. The storage medium may be applied to an electronic device, for example, the storage medium may include a storage device in the electronic device.

[0098] The storage device may include any combination of one or more computer program products, which may include various forms of computer-readable storage media, such as volatile memory and / or non-volatile memory. Volatile memory may include, for example, random access memory (RAM) and / or cache memory. Non-volatile memory may include, for example, read-only memory (ROM), a hard disk, an erasable programmable read-only memory (EPROM), a portable compact disc read-only memory (CD-ROM), a USB memory, a flash memory, etc. One or more computer-readable instructions may be stored on the computer-readable storage medium, and the processor may execute the computer-readable instructions to implement various functions of the processor. The storage medium may also store various application programs and various data.

[0099] The storage medium may include a memory card of a smart phone, a cache component of a tablet computer, a hard disk of a personal computer, a random access memory (RAM), a read-only memory (ROM), an erasable programmable read-only memory (EPROM), a portable compact disk read-only memory (CD-ROM), a flash memory, or any combination of the above storage media, or other applicable storage media.

[0100] The present application embodiment provides a computer device, such as Figure 14 As shown, it includes at least one artificial intelligence chip 1510 and a memory 1520 connected to the at least one artificial intelligence chip 1510. The specific connection medium between the artificial intelligence chip 1510 and the memory 1520 is not limited in the embodiment of the present application. Figure 14 For example, the artificial intelligence chip 1510 and the memory 1520 are connected via a bus. Buses can be divided into address buses, data buses, control buses, etc. In the embodiment of the present application, the memory 1520 stores instructions that can be executed by at least one artificial intelligence chip 1510. By executing the instructions stored in the memory 1520, the at least one artificial intelligence chip 1510 can perform the steps of the above-mentioned data processing method.

[0101] Those skilled in the art will appreciate that the embodiments of the present application may be provided as methods, systems, or computer program products. Therefore, the present application may take the form of an entirely hardware embodiment, an entirely software embodiment, or an embodiment combining software and hardware. Furthermore, the present application may take the form of a computer program product implemented on one or more computer-usable storage media (including but not limited to magnetic disk storage, CD-ROM, optical storage, etc.) containing computer-usable program code.

[0102] The present application is described with reference to the flowcharts and / or block diagrams of the methods, devices (systems), and computer program products according to the embodiments of the present application. It should be understood that each process and / or block in the flowchart and / or block diagram, as well as the combination of processes and / or blocks in the flowchart and / or block diagram, can be implemented by computer program instructions. These computer program instructions can be provided to a processor of a general-purpose computer, a special-purpose computer, an embedded processor, or other programmable data processing device to produce a machine, so that the instructions executed by the processor of the computer or other programmable data processing device generate instructions for implementing the processes in the flowchart and / or block diagram. Figure 1 a process or multiple processes and / or boxes Figure 1 A device that provides the functions specified in a block or multiple blocks.

[0103] It should also be noted that the terms "comprises," "includes," or any other variations thereof are intended to encompass non-exclusive inclusion, such that a process, method, commodity, or apparatus that includes a series of elements includes not only those elements but also other elements not explicitly listed, or includes elements inherent to such process, method, commodity, or apparatus. In the absence of further limitations, an element defined by the phrase "comprises a ..." does not exclude the presence of other identical elements in the process, method, commodity, or apparatus that includes the element.

[0104] The various embodiments in this specification are described in a progressive manner. Similar parts between the various embodiments can be referred to in conjunction with each other. Each embodiment focuses on the differences between the other embodiments. In particular, the system embodiments are generally similar to the method embodiments, so the description is relatively simple. For relevant parts, refer to the description of the method embodiments.

[0105] The foregoing is merely an embodiment of the present application and is not intended to limit the present application. For those skilled in the art, the present application may have various changes and variations. Any modifications, equivalent replacements, improvements, etc. made within the spirit and principles of the present application should all be included within the scope of the claims of the present application.

[0106] Although the embodiments of the present application have been described with reference to the accompanying drawings, those skilled in the art may make various modifications and variations without departing from the spirit and scope of the present application, and such modifications and variations shall fall within the scope defined by the appended claims.

Claims

1. A data processing method, characterized in that: The method comprises: When block data of input data needs to be loaded into a thread-private register, determining a first jump step required for rearranging the block data based on a block size of the input data in a first dimension; Performing jump reading on the block data according to the first jump step size, reordering and loading the read input elements into the thread-private register, so that the block data is stored in a transposed manner in the thread-private register, and mapping the thread direction onto a second dimension, where the first dimension is an adjacent lower dimension of the second dimension; A plurality of convolution-related calculations are performed based on the transposed data in the thread-private register and the weight data in the thread-private register; and duplicate data exists between corresponding portions of data in the transposed data between adjacent convolution-related calculations.

2. The method according to claim 1, characterized in that The performing of multiple convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register includes: The transposed data in the thread-private register is traversed along a first dimension of the input data, and a plurality of convolution-related calculations are performed using the traversed input elements and weight elements included in the weight data.

3. The method according to claim 2, characterized in that The transposed data is divided into a plurality of operation regions distributed along the first dimension, the operation regions corresponding to target input elements in the transposed data; The traversing the transposed data in the thread-private register along the first dimension of the input data, and performing a plurality of convolution-related calculations using the traversed input elements and the weight elements included in the weight data, includes: performing a first convolution-related calculation using the weight element and a target input element corresponding to a first operating area among the multiple operating areas; After completing the first convolution-related calculation, convolution sliding is performed along the first dimension, and different operation areas overlap with each other. Subsequent convolution-related calculations are performed using the weight elements and the target input elements corresponding to each operation area after the first operation area.

4. The method according to claim 3, characterized in that The size of the convolution kernel is recorded as R×S; R and S are positive integers greater than 1; When performing convolution sliding along the first dimension, N target input elements are slid along the first dimension each time; where N is a positive integer greater than or equal to 1 and less than S.

5. The method according to claim 1, characterized in that The block data is loaded from the group shared memory into the thread private register; the method further comprises: The input data is loaded from the chip memory into the group shared memory using a block loading strategy.

6. The method according to claim 1, characterized in that The determining, based on the block size of the input data in the first dimension, a first jump step required for rearranging the block data, comprises: The first jump step size is determined based on a block size of the input data in a first dimension and a padding operation range.

7. The method according to claim 1, characterized in that The transposed data includes a first input element and a second input element, and the weight data includes a first weight element and a second weight element; performing a plurality of convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register, comprising: In a current clock cycle of a current multiply-add instruction, two multiply-add operations are performed using the first input element, the second input element, the first weight element, and the second weight element to obtain an accumulation result of the current multiply-add instruction.

8. The method according to any one of claims 1 to 7, characterized in that The method further comprises: Storing the calculation result of the convolution-related calculation in the thread-private register; Performing jump reading on the calculation result in the thread private register, and reordering and storing the read output elements in the group shared memory, so that the calculation result is transposed and restored in the group shared memory; The transposed calculation result in the group shared memory is stored from the group shared memory to the chip memory.

9. The method according to claim 8, characterized in that The skipping and reading of the calculation result in the thread-private register includes: The block data is skipped and read according to a second skip step size, wherein the second skip step size is determined based on the block size of the input data in the first dimension.

10. A data processing device, characterized in that: The device comprises: a step size determining module, configured to determine, when block data of input data needs to be loaded into a thread-private register, a first jump step size required for rearranging the block data based on a block size of the input data in a first dimension; a data rearrangement module, configured to perform jump reading of the block data according to the first jump step size, rearrange and load the read input elements into the thread-private register, so that the block data is stored in a transposed manner in the thread-private register, and map the thread direction onto a second dimension, where the first dimension is an adjacent lower dimension of the second dimension; A calculation execution module is used to perform multiple convolution-related calculations based on the transposed data in the thread-private register and the weight data in the thread-private register.

11. A processor, characterized in that: It includes instruction parsing unit, execution unit, memory and thread private registers, among which, The instruction parsing unit is used to receive and parse data processing instructions, wherein the data processing instructions are used to load input data from the memory into the thread-private register and perform convolution-related operations; The execution unit executes the data processing method according to any one of claims 1 to 9 after the instruction parsing unit parses the data processing instruction.

12. An electronic device, characterized in that: include: a memory that non-transitorily stores computer-executable instructions; a processor configured to execute the computer-executable instructions, Wherein, when the computer executable instructions are executed by the processor, the data processing method according to any one of claims 1 to 9 is implemented.

Images