Data processing method and device, processor and electronic equipment

Abstract

The application relates to the technical field of artificial intelligence chips, and discloses a data processing method and device, a processor and electronic equipment. First, in the case of loading the block data of input data to 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 the first dimension; second, the block data is read by jumping 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 stored in the thread private register in a transposed manner, and the thread direction is mapped on the second dimension, thus obtaining the data basis of convolution-related calculation, and then performing convolution-related calculation based on the transposed data in the thread private register and the weight data in the thread private register, so that repeated data reading in the calculation process is reduced, and the operator is changed from a memory performance bottleneck operation to a calculation 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:

[0005] In a first aspect, an embodiment of the present application provides a data processing method, the method comprising:

[0006] 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;

[0007] 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;

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

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

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

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

[0012] 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;

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

[0014] Optionally, the size of the convolution kernel is denoted as R×S; R and S are positive integers greater than 1;

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

[0016] Optionally, the block data is loaded from a group shared memory into the thread private register; the method further includes:

[0017] The input data is loaded from the chip memory into the group shared memory using a block loading strategy.

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

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

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

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

[0022] Optionally, the method further includes:

[0023] Storing the calculation result of the convolution-related calculation in the thread-private register;

[0024] 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;

[0025] The transposed calculation result in the group shared memory is stored from the group shared memory to the chip memory.

[0026] Optionally, the skipping and reading of the calculation result in the thread-private register includes:

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

[0028] In a second aspect, an embodiment of the present application provides a data processing device, the device comprising:

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

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

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

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

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

[0034] The execution unit executes any of the above data processing methods after the instruction parsing unit parses the data processing instruction.

[0035] In a fourth aspect, an embodiment of the present application provides an electronic device, including:

[0036] a memory that non-transitorily stores computer-executable instructions;

[0037] a processor configured to execute the computer-executable instructions,

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

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

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

[0041] Figure 1 A schematic structural diagram of a general-purpose graphics processor provided in an embodiment of the present application;

[0042] Figure 2 A schematic diagram of the computational logic of the depthwise separable convolution provided in an embodiment of the present application;

[0043] Figure 3(a) is a diagram of data loading in the calculation process of the deep separable convolution provided by the embodiment of the present application;

[0044] Figure 3(b) is a diagram of data storage in the calculation process of the deep separable convolution provided by the embodiment of the present application;

[0045] Figure 4 Figure 4 is a diagram of the calculation process of the deep separable convolution provided by the embodiment of the present application;

[0046] Figure 5 Figure 5 is a diagram of the flow of the data processing method provided by the embodiment of the present application;

[0047] Figure 6(a) is a diagram of the working process of the TMA provided by the embodiment of the present application;

[0048] Figure 6(b) is a diagram of the diagrammatic process of the ldmatrix instruction provided by the embodiment of the present application;

[0049] Figure 7 Figure 7 is a diagram of the convolution sliding process in the convolution-related calculation provided by the embodiment of the present application;

[0050] Figure 8 Figure 8 is a diagram of the calculation process of the convolution-related calculation provided by the embodiment of the present application;

[0051] Figure 9 Figure 9 is the calculation process of the traditional calculation instruction provided by the embodiment of the present application;

[0052] Figure 10 Figure 10 is the calculation process of the new calculation instruction provided by the embodiment of the present application;

[0053] Figure 11 Figure 11 is a diagram of the flow of the data processing method provided by the embodiment of the present application;

[0054] Figure 12 Figure 12 is a diagram of the framework of the data processing apparatus provided by the embodiment of the present application;

[0055] Figure 13 Figure 13 is a diagram of the schematic structure of the processor provided by the embodiment of the present application;

[0056] Figure 14 Figure 14 is a diagram of the schematic structure of the computer device provided by the embodiment of the present application. DETAILED DESCRIPTION

[0057] In order to make the purposes, technical solutions and advantages of the embodiments of the present application clearer, the technical solutions in the embodiments of the present application will be described clearly and completely below with reference to the drawings in the embodiments of the present application. Obviously, the described embodiments are only some of the embodiments of the present application, rather than all the embodiments. Based on the embodiments in the present application, all other embodiments obtained by those skilled in the art without creative work fall within the scope of protection of the present application.

[0058] In order to facilitate the description of the present application, the following introduces the concepts related to the present application.

[0059] The data processing method of the present application is based on artificial intelligence (AI). Artificial intelligence is the use of digital computers or machine controlled by digital computers to simulate, extend and expand human intelligence, perceive the environment, acquire knowledge and use knowledge to obtain the best results. In other words, artificial intelligence is a comprehensive technology of computer science, which tries to understand the essence of intelligence and produce a new intelligent machine that can react in a similar way to human intelligence.

[0060] The data processing method of the present application can be based on deep separable convolution. Deep separable convolution is a widely used convolution operation in the field of computer vision and deep learning. It is a special convolution operation that reduces the number of parameters and computational complexity by decomposing the standard convolution operation, thereby improving the efficiency and performance of the model. Deep separable convolution includes two steps: depth convolution and pointwise convolution. In the depth 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 of the depth convolution is convolved with a 1x1 convolution kernel, thereby integrating and interacting the information between channels. In this way, the number of parameters and computational complexity can be reduced, thereby improving the computational efficiency.

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

[0062] This application exemplifies the parameter expression of depth-wise separable convolution:

[0063] Input data x: [N, D, H, W, C]

[0064] Weight data w: [K, T, R, S, C] → [C, T, R, S, 1]

[0065] Output data y: [N, Z, P, Q, K] → [N, T, P, Q, C]

[0066] Step size stride: [strideD, strideH, strideW]

[0067] Padding: [padD, padH, padW]

[0068] Dilation: [dilationD, dilationH, dilationW]

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

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

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

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

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

[0074] Figure 2is a calculation logic schematic diagram of depth separable convolution. The input data includes two channels of 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 and is 2. Each input feature map 102 is subjected to two-dimensional plane convolution calculation with the corresponding convolution kernel 104. Two output feature maps 106 are generated through channel-by-channel convolution. Please continue to refer to Figure 2 The convolution kernel 104 is represented as a 2x2 weight matrix, and the data layout of each weight matrix in the register group is denoted as a weight data layout 108. The input feature map 102 is represented as a 3x3 input matrix, and the data layout of each input matrix in the register group is denoted as an input data layout 110. The output feature map 106 is represented as a 2x2 output matrix, and the data layout of each output matrix in the register group is denoted as an output data layout 112.

[0075] Through analysis of the calculation logic of depth separable convolution, it is found that for each output point, the amount of data of input data x and weight data w to be read is RxS, and there is a repeated amount of data to be read, and the number of transmitted read instructions needs to be reduced. Further, because of the sliding window type calculation logic characteristics of depth separable convolution, there are many data repeated reading behaviors, and the execution time of the operator is mainly determined by the speed and delay of memory access, and is usually considered as a memory-bound operator. Even if the cache is used to reduce the memory access amount, it may be bound to the speed and delay of cache access. In order to reduce the data memory access amount to improve the performance upper limit of the operator, the present application provides a data processing method to change the operator to a calculation-bound operator.

[0076] It should be noted that the execution time of the calculation-bound operator is mainly determined by the operations required for calculation, rather than limited by memory bandwidth or delay. In this case, the performance of the operator depends more on the speed of the calculation unit, and the memory access becomes a relatively small bottleneck.

[0077] It should be noted that the data processing method provided by the present application is a hardware acceleration scheme applied to depth separable convolution calculation, and the key is to change the operator performance from a memory bottleneck (memory bound) to a calculation bottleneck (compute bound) through data rearrangement and storage optimization. The method is suitable for some hardware architectures, and uses its tensor memory accelerator (TMA) and thread-level parallel mechanism to realize efficient data flow.

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

[0122] The convolution sliding process is exemplarily illustrated by taking the first operation region 702 and the second operation region 704 as examples. The convolution correlation calculation is performed by using the weight element and the target input element of the first operation region 702, and after the convolution correlation calculation is completed, a matrix element of the output data is obtained. When the convolution sliding is performed along the first dimension, the second operation region 704 is obtained by sliding 1 target input element on the first dimension each time, and after the convolution correlation calculation is completed, the next matrix element of the output data is obtained. The convolution correlation calculation is performed by using the weight element and the target input element of the second operation region 704. It can be understood that the convolution correlation calculation includes: 1) incremental reading, only loading the input elements in the non-overlapping part between the second operation region 704 and the first operation region 702; 2) data multiplexing, directly multiplexing the overlapping part data in the thread private register; and 3) continuous calculation, performing point multiplication and accumulation of the newly loaded target input element and the weight element to generate a new matrix element of the output data.

[0123] 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; and the convolution correlation calculation based on the transposed data in the thread private register and the weight data in the thread private register can include: performing two multiplication and accumulation operations by using the first input element, the second input element, the first weight element and the second weight element in a current clock cycle of a current multiplication and accumulation instruction to obtain an accumulation result of the current multiplication and accumulation instruction.

[0124] 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; and the convolution correlation calculation based on the transposed data in the thread private register and the weight data in the thread private register can include: performing two multiplication and accumulation operations by using the first input element, the second input element, the first weight element and the second weight element in a current clock cycle of a current multiplication and accumulation instruction to obtain an accumulation result of the current multiplication and accumulation instruction. Figure 9 Figure 9 The calculation process of the traditional calculation instruction is shown in FIG. 8. Although the data multiplexing rate can be improved based on the transposed data, the number of calculation instructions to be transmitted is to be reduced when the calculation is performed by using the traditional calculation instruction, thereby affecting the upper limit of the calculation instruction transmission efficiency. Therefore, a new calculation instruction is provided in the embodiment, which can be denoted as HDP2.a2b2. Please refer to FIG. 9. Figure 10 Figure 10 The calculation process of the new calculation instruction is shown in FIG. 9. By comparing the traditional calculation instruction and the new calculation instruction, only one multiplication and accumulation operation is completed in a single period for the traditional calculation instruction. For the new calculation instruction, two multiplication and accumulation operations are completed in a single period, which has double efficiency of the traditional multiplication and accumulation instruction.

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

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

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

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

[0129] In some embodiments, see Figure 11 , the method may further include:

[0130] S1210: Store the calculation result of the convolution-related calculation in a thread-private register.

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

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

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

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

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

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

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

[0138] 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;

[0139] 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;

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

[0156] 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, wherein 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