Convolution processing method, graphics processor, computer equipment, readable storage medium and program product

Abstract

The invention relates to a convolution processing method and device, computer equipment, a computer readable storage medium and a computer program product. The method comprises the following steps: reading first block data and second block data from a video memory; wherein the first block data is determined by blocking based on a first matrix obtained by expanding the sample graph; the second block data is determined by blocking based on a second matrix obtained by expanding the output feature map; the output feature map is output by the convolutional neural network; determining an initial convolution calculation result based on the first block data and the second block data; repeatedly reading the first block data and the second block data from the video memory; determining an initial convolution calculation result based on the first block data and the second block data; and based on each initial convolution calculation result, determining a convolution calculation result, and storing the convolution calculation result in a video memory. By adopting the method, the computing resources can be saved while the video memory resources are saved.

Description

Technical Field

[0001] This application relates to the field of neural network technologies, and particularly to a convolution processing method, apparatus, computer device, computer-readable storage medium, and computer program product. Background Art

[0002] In a convolutional neural network (CNN), the depthwise separable convolution (DSC) calculation of a feature map and a weight map is often implemented through img2col + GEMM. The depthwise separable convolution mainly realizes that for each channel of the input feature map, a convolution kernel is used respectively, and then the outputs of all convolution kernels are concatenated to obtain the final output result.

[0003] In traditional technologies, common convolution calculation methods include direct convolution, img2col + GEMM (General Matrix Multiplication), etc. Direct convolution multiplies and accumulates the elements inside the convolution kernel (weight) with the corresponding elements in the input feature map to obtain an element in the output feature map. Direct convolution is relatively intuitive but inefficient, and each time only the elements inside the convolution kernel can be multiplied and added to obtain an element in the output feature map. In order to improve efficiency, img2col + GEMM appears. Compared with direct convolution, the operation efficiency is greatly improved.

[0004] However, the current img2col + GEMM calculation method needs to explicitly store the expanded matrix, and for the convolution kernel (weight) data, in order to use GEMM, it is necessary to supplement some zeros to the data and then use the convolution operation, resulting in waste of both video memory resources and computing resources. Summary of the Invention

[0005] Based on this, in view of the above technical problems, it is necessary to provide a convolution processing method, apparatus, computer device, computer-readable storage medium, and computer program product that can save both video memory resources and computing resources.

[0006] In a first aspect, this application provides a convolution processing method, and the method includes:

[0007] Read first block data and second block data from the video memory; wherein, the first block data is determined by partitioning a first matrix obtained by expanding a sample map; the second block data is determined by partitioning a second matrix obtained by expanding an output feature map; the output feature map is output by a convolutional neural network;

[0008] Determine an initial convolution calculation result based on the first chunk of data and the second chunk of data;

[0009] Repeat the steps of reading the first chunk of data and the second chunk of data from the video memory; and determining an initial convolution calculation result based on the first chunk of data and the second chunk of data;

[0010] Based on each of the initial convolution calculation results, determine a convolution calculation result, and save the convolution calculation result to the video memory.

[0011] In one embodiment, the reading of the first chunk of data and the second chunk of data from the video memory includes:

[0012] Read the first chunk of data and the second chunk of data from the video memory through a computing unit, and save the first chunk of data and the second chunk of data to a shared cache;

[0013] Read the first chunk of data and the second chunk of data from the shared cache to a general register through the computing unit;

[0014] The determining of the initial convolution calculation result based on the first chunk of data and the second chunk of data includes:

[0015] Determine an initial convolution calculation result based on the first chunk of data and the second chunk of data in the general register through the computing unit, and save the initial convolution calculation result to the general register.

[0016] In one embodiment, the size of the first chunk of data is a first dimension; the reading of the first chunk of data from the video memory through the computing unit includes:

[0017] Calculate a first video memory read address for reading the first chunk of data from the video memory through the computing unit based on a first horizontal coordinate, a first vertical coordinate, and the width of the output feature map; the first horizontal coordinate is the horizontal coordinate for reading in the video memory; the first vertical coordinate is the vertical coordinate for reading in the video memory;

[0018] Read the first chunk of data of the first dimension from the video memory through the computing unit according to the first video memory read address.

[0019] In one embodiment, the first dimension includes a first preset row; the computing unit includes a warp; the warp includes a plurality of data channels; the method for determining the first horizontal coordinate includes:

[0020] Determine the first horizontal coordinate based on the obtained identification number of the data channel, the identification number of the warp, and the number of warps.

[0021] The method for determining the first vertical coordinate includes:

[0022] Determine the first vertical coordinate based on the obtained identification number of the data channel, the first preset row, the number of data channels included in the warp, the number of current executions, the height of the output feature map, and the width of the sample map.

[0023] In one embodiment, the size of the first block of data is a first dimension; the computing unit includes warps; each warp includes a plurality of data channels; saving the first block of data to the shared cache includes:

[0024] Calculate a first write shared cache address for writing the first block of data to the shared cache based on the obtained second horizontal coordinate, second vertical coordinate, and offset value; each execution operation corresponds to an offset value; the second horizontal coordinate is the read horizontal coordinate in the shared cache; the second vertical coordinate is the read vertical coordinate in the shared cache; and both the second horizontal coordinate and the second vertical coordinate are determined based on the identification number of the data channel and the identification number of the warp.

[0025] Write the first block of data of the first dimension to the shared cache according to the first write shared cache address.

[0026] In one embodiment, the size of the second block of data is a second dimension; reading the second block of data from the video memory by the computing unit includes:

[0027] Calculate a second read video memory address for reading the second block of data from the video memory by the computing unit based on the third horizontal coordinate, third vertical coordinate, and the width of the output feature map; the third horizontal coordinate is the read horizontal coordinate in the video memory; the third vertical coordinate is the read vertical coordinate in the video memory.

[0028] Read the second block of data of the second dimension from the video memory by the computing unit according to the second read video memory address.

[0029] In one embodiment, the computing unit includes warps; each warp includes a plurality of data channels; the method for determining the third horizontal coordinate includes:

[0030] Determine the third horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp reads.

[0031] The method for determining the third longitudinal coordinate includes:

[0032] Based on the obtained identification number of the thread bundle, determine the third longitudinal coordinate.

[0033] In one embodiment, the size of the second block data is the second dimension; saving the second block data to the shared cache includes:

[0034] Based on the obtained fourth horizontal coordinate, fourth longitudinal coordinate, and offset value, calculate the second write shared cache address for writing the second block data into the shared cache; each output feature map corresponds to one offset value; the fourth horizontal coordinate is the write horizontal coordinate in the shared cache; the fourth longitudinal coordinate is the write longitudinal coordinate in the shared cache;

[0035] According to the second write shared cache address, write the second block data with the second dimension into the shared cache.

[0036] In one embodiment, the computing unit includes a thread bundle; the thread bundle includes multiple data channels; the method for determining the fourth horizontal coordinate includes:

[0037] Based on the obtained identification number of the data channel and the number of times the thread bundle is written, determine the fourth horizontal coordinate;

[0038] The method for determining the fourth longitudinal coordinate includes:

[0039] Based on the obtained identification number of the data channel, the identification number of the thread bundle, and the number of data channels included in the thread bundle, determine the fourth longitudinal coordinate.

[0040] In one embodiment, the size of the first block data is the first dimension; the size of the second block data is the second dimension; the method for the computing unit to read the first block data and the second block data from the shared cache into the general register includes:

[0041] The computing unit calculates the first read shared cache address for reading the first block data from the shared cache based on the fifth horizontal coordinate, fifth longitudinal coordinate, and offset value; wherein, each output feature map corresponds to one offset value; the fifth horizontal coordinate is the read horizontal coordinate in the shared cache; the fifth longitudinal coordinate is the read longitudinal coordinate in the shared cache;

[0042] According to the first read shared cache address, read the first block data with the first dimension from the shared cache and store it in the general register;

[0043] Based on the sixth horizontal coordinate, the sixth vertical coordinate, and the offset value, a second read shared cache address for reading the second block data from the shared cache is calculated by a calculation unit; wherein, the sixth horizontal coordinate is the read horizontal coordinate in the shared cache; the sixth vertical coordinate is the read vertical coordinate in the shared cache;

[0044] According to the second read shared cache address, the second block data of the second size is read from the shared cache and stored in the general register.

[0045] In one embodiment, the calculation unit includes a warp; the warp includes a plurality of data channels; the method for determining the fifth horizontal coordinate includes:

[0046] Based on the obtained identification number of the data channel, the number of times the warp reads, and the number of data channels included in the warp, the fifth horizontal coordinate is determined;

[0047] The method for determining the fifth vertical coordinate includes:

[0048] Based on the obtained identification number of the warp, the number of data channels included in the warp, and the number of warps, the fifth vertical coordinate is determined.

[0049] In one embodiment, the calculation unit includes a warp; the warp includes a plurality of data channels; the method for determining the sixth horizontal coordinate includes:

[0050] Based on the obtained identification number of the data channel and the number of times the warp writes, the sixth horizontal coordinate is determined;

[0051] The method for determining the sixth vertical coordinate includes:

[0052] Based on the number of times of the current execution, the sixth vertical coordinate is determined.

[0053] In one embodiment, the size of the convolution calculation result is the third size; the storing the convolution calculation result in the video memory includes:

[0054] Based on the seventh horizontal coordinate, the seventh vertical coordinate, and the number of times of the current execution, a write video memory address for writing the convolution calculation result into the video memory is calculated; the seventh horizontal coordinate is the write horizontal coordinate in the video memory; the seventh vertical coordinate is the write vertical coordinate in the video memory;

[0055] According to the write video memory address, the convolution calculation result of the third size is written into the video memory.

[0056] In one embodiment, the third dimension includes a second preset row and a preset column; the computing unit includes a warp; the warp includes a plurality of data channels; the method for determining the seventh horizontal coordinate includes:

[0057] Based on the obtained identification number of the data channel, the number of times the warp reads, the number of data channels included in the warp, and the preset column, determine the seventh horizontal coordinate;

[0058] The method for determining the seventh vertical coordinate includes:

[0059] Based on the obtained identification number of the warp, the number of warps, and the second preset row, determine the seventh vertical coordinate.

[0060] In one embodiment, the process of the computing unit reading the first block data and the second block data from the video memory and saving the first block data and the second block data to the shared cache includes:

[0061] The computing unit reads the first block data and the second block data from the video memory and saves them to the cache;

[0062] Read the first block data and the second block data from the cache and save them to the shared cache.

[0063] In one embodiment, before the computing unit reads the first block data and the second block data from the shared cache to the general-purpose register, it includes:

[0064] Initialize the general-purpose register so that the initial values in the general-purpose register are all preset values.

[0065] In a second aspect, the present application further provides a graphics processing unit, including:

[0066] An execution unit for reading first block data and second block data from the video memory; wherein, the first block data is determined by partitioning a first matrix obtained by expanding a sample graph; the second block data is determined by partitioning a second matrix obtained by expanding an output feature map; the output feature map is output by a convolutional neural network; based on the first block data and the second block data, determine an initial convolution calculation result; repeatedly execute the steps of reading the first block data and the second block data from the video memory; based on the first block data and the second block data, determining an initial convolution calculation result; based on each initial convolution calculation result, determine a convolution calculation result and save the convolution calculation result to the video memory;

[0067] A video memory for storing the first block data and the second block data; and also for storing the convolution calculation result.

[0068] In a third aspect, the present application further provides a computer device, including a memory and a graphics processor. The memory stores a computer program, and when the graphics processor executes the computer program, the steps of the above method are implemented.

[0069] In a fourth aspect, the present application further provides a computer-readable storage medium, on which a computer program is stored. When the computer program is executed by a graphics processor, the steps of the above method are implemented.

[0070] In a fifth aspect, the present application further provides a computer program product, including a computer program. When the computer program is executed by a graphics processor, the steps of the above method are implemented.

[0071] For the above convolution processing method, device, computer device, computer-readable storage medium, and computer program product, the execution unit first unfolds the sample image through img2col to obtain a first matrix; unfolds the output feature map through img2col to obtain a second matrix; and divides the first matrix into blocks to obtain first block data, and divides the second matrix into blocks to obtain second block data. When the execution unit obtains the first block data and the second block data from the video memory, it is not necessary to unfold the first matrix and the second matrix. Since the data volumes of the first matrix and the second matrix are large, more video memory resources will be consumed during unfolding. Instead, the first block data and the second block data with smaller data volumes are obtained, reducing the occupied video memory resources. Moreover, when reading the first block data and the second block data from the video memory, there is no need to wait for img2col to complete. The already segmented first block data and second block data can be obtained while img2col is performing matrix unfolding and segmentation, reducing the waiting time, thereby improving the calculation efficiency. Then, the first block data and the second block data of each group are multiplied by the point-by-point multiplication and accumulation method to determine the initial convolution calculation result, and based on each initial convolution calculation result, the convolution calculation result is determined. Compared with the traditional technology, since some data needs to be supplemented with zeros in order to use GEMM, and the point-by-point multiplication and accumulation operation is directly performed on the first block data and the second block data, the operation of supplementing 0 to the second matrix is avoided. Since there are no supplementary 0s, when using the point-by-point multiplication and accumulation method, the data transfer and calculation can be greatly reduced, reducing the consumption of computing power resources. Description of the Drawings

[0072] To more clearly illustrate the technical solutions in the embodiments of the present application or related technologies, the following will briefly introduce the drawings required for describing the embodiments of the present application or related technologies. Obviously, the drawings in the following description are only some embodiments of the present application. For those of ordinary skill in the art, without creative efforts, other related drawings can also be obtained based on these drawings.

[0073] Figure 1 It is a schematic diagram of depthwise separable convolution in an embodiment;

[0074] Figure 2 It is a schematic diagram of the expansion from a vector to a matrix in an embodiment;

[0075] Figure 3 It is a schematic diagram of the architecture of a graphics processing unit in an embodiment; (a) is a schematic diagram of the architecture of one kind of graphics processing unit, and (b) is a schematic diagram of the architecture of another kind of graphics processing unit;

[0076] Figure 4 It is a schematic diagram of the flow of a convolution processing method in an embodiment;

[0077] Figure 5 It is a schematic diagram of a pointwise multiplication operation in an embodiment;

[0078] Figure 6 It is a schematic diagram of the flow of reading the first block data and the second block data from the video memory and determining the initial convolution calculation result in an embodiment;

[0079] Figure 7 It is a schematic diagram of the video memory coordinates for reading the first block data in an embodiment;

[0080] Figure 8 It is a schematic diagram of the shared cache coordinates for writing the first block data in an embodiment;

[0081] Figure 9 It is a schematic diagram of the video memory coordinates for reading the second block data in an embodiment;

[0082] Figure 10 It is a schematic diagram of the shared cache coordinates for writing the second block data in an embodiment;

[0083] Figure 11 It is a schematic diagram of the flow of reading the first block data and the second block data from the shared cache to the general-purpose register by a computing unit in an embodiment;

[0084] Figure 12 It is a schematic diagram of the shared cache coordinates for reading the first block data in an embodiment;

[0085] Figure 13 It is a schematic diagram of the shared cache coordinates for reading the second block data in an embodiment;

[0086] Figure 14 Schematic diagram of writing video memory coordinates for the initial convolution calculation result in an embodiment;

[0087] Figure 15 Flow schematic diagram of the convolution processing method in another embodiment;

[0088] Figure 16 Internal structure diagram of a computer device in an embodiment. Detailed implementation manners

[0089] In order to make the objectives, technical solutions and advantages of the present application clearer and more understandable, the present application will be further described in detail below with reference to the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are only used to explain the present application and are not used to limit the present application.

[0090] Depthwise separable convolution mainly realizes that it uses a convolution kernel for each channel of the input feature map respectively, and then splices the outputs of all convolution kernels to obtain its final output result. As Figure 1 shown.

[0091] Common convolution calculation methods include direct convolution, img2col + GEMM (General Matrix Multiplication), etc.; direct convolution multiplies and accumulates the elements inside the convolution kernel (weight) with the corresponding elements in the input feature map to obtain an element in the output feature map, and then moves to the next step according to the stride, repeating the above operations until all elements of the output feature map are obtained. Direct convolution is relatively intuitive but inefficient, and it can only perform multiplication and addition operations on the elements inside the convolution kernel to obtain an element in the output feature map each time.

[0092] img2col + GEMM first unfolds the input feature map and the convolution kernel into matrices A and B through img2col respectively, and then performs A×B to obtain the result matrix C. For the convolution kernel (weight) data, in order to use GEMM, some zero-padding must be done to the data, and then convolution operation is used. As Figure 2 shown, that is, the 1×K vector in each row of the original convolution kernel is converted into a C×K matrix by zero-padding, and in this C×K matrix, only the first row has specific data and the others are all 0. That is, the storage is expanded to C times (C is the number of channels of the input data when calculating the depthwise separable convolution).

[0093] The principle of zero padding is explained here. The size of matrix A is [N, H, W, C], which can be converted to [N*H*W, RSC] through imgcol; the size of matrix B is [R, S, 1, K], which can be converted to [R*S*1, K]. A*B = [N*H*W, R*S*C] * [R*S*1, K]. Since R*S*C is not equal to R*S*1, A*B cannot perform multiplication (in matrix multiplication, the number of columns of matrix A must be equal to the number of rows of matrix B). Therefore, matrix B needs to be padded with zeros to make the number of columns of matrix A equal to the number of rows of matrix B.

[0094] In the subsequent convolution calculation, TileC[i, j] = ∑ k TileA[i, k] * TileB[k, j]. Since there are a large number of zeros in TileB, a large amount of calculation is wasted during the convolution calculation. Finally, C is converted into the output feature map through col2img (the inverse process of img2col). Compared with direct convolution, img2col + GEMM greatly improves the operation efficiency. However, on the one hand, img2col has the disadvantage that it needs to store the expanded large matrix A, and it is necessary to explicitly store the expanded matrix, resulting in a waste of computing resources; on the other hand, since img2col + GEMM cannot be completed in one Kernel, usually one convolution kernel performs img2col expansion, and another Kernel can perform GEMM operation. Moreover, in order to use GEMM, it is necessary to expand and pad the data with zeros and then use convolution operation, all of which result in a waste of computing resources.

[0095] The convolution processing method provided by the embodiments of the present application can be applied to, for example, Figure 3 (a) the graphics processor architecture shown. The graphics processing unit (GPU) includes execution units (EUs) and global memory. The execution units (EUs) read the first block data and the second block data from the global memory; wherein, the first block data is determined by partitioning the first matrix obtained by expanding the sample graph; the second block data is determined by partitioning the second matrix obtained by expanding the output feature map; the output feature map is output by the convolutional neural network; based on the first block data and the second block data, determine the initial convolution calculation result; repeat the steps of reading the first block data and the second block data from the global memory; based on the first block data and the second block data, determine the initial convolution calculation result; based on each initial convolution calculation result, determine the convolution calculation result and save the convolution calculation result to the global memory.

[0096] The convolution processing method provided by the embodiments of the present application can be applied to, for example,Figure 3 (b) In the graphics processor architecture shown. The graphics processing unit (GPU) includes execution units (EUs), caches, and global memory. Among them, Figure 3 (a) and Figure 3 (b) The execution units (EUs) include arithmetic logic units (ALUs), shared memories (SMs), and common register files (CRFs).

[0097] In an exemplary embodiment, as Figure 4 shown, a convolution processing method is provided. Taking the method applied to the Figure 3 (a) or Figure 3 (b) graphics processor as an example, it includes the following steps S402 to S408. Among them:

[0098] Step S402, read the first block data and the second block data from the global memory.

[0099] Among them, the first block data is determined by partitioning the first matrix obtained by expanding the sample graph; the second block data is determined by partitioning the second matrix obtained by expanding the output feature map; the output feature map is output by the convolutional neural network.

[0100] For the convenience of discussion, techniques such as double buffering are not described in this solution. Assume that we agree to perform a depthwise separable convolution on the input data, that is, the sample graph x and the weight data (convolution kernel) w to obtain the output data y. The sample graph x is expanded into the first matrix, which can be matrix A, through img2col, and the convolution kernel w is expanded into the second matrix, which can be matrix B, through img2col. Using the idea of tiling, the first matrix is divided into several first block data, and the second matrix is divided into several second block data. Among them, the sizes of the first block data and the second block data satisfy matrix operations, such as multiplication; the number of the first block data is the same as the number of the second block data.

[0101] x, w, and y are all four-dimensional vector tensors. The dimensions of x are [N, H, W, C], corresponding to the number, height, width, and number of channels respectively. Among them, N is the number of feature maps output by the previous layer of the depth convolutional neural network, that is, the number of sample graphs in a batch; H represents the height of each sample graph (i.e., the number of rows of the image); W represents the width of each sample graph (i.e., the number of columns of the image). C represents the number of channels of each sample graph (for example, the number of channels of an RGB image is 3).

[0102] Each dimension of w is [R, S, C, K], corresponding to height, width, number of channels, and number respectively, where K is the number of output feature maps of depthwise separable convolution; R represents the height of the output feature maps of depthwise separable convolution; S represents the width of the output feature maps of depthwise separable convolution; C represents the number of channels of the output feature maps of depthwise separable convolution.

[0103] The output result data, that is, each dimension of the convolution calculation result y is [N, P, Q, K]. Therefore, depthwise separable convolution is represented by formula (1).

[0104] y[N, P, Q, K] = DSC(x(N, H, W, C], w[R, S, C, K]) Formula (1)

[0105] In depthwise separable convolution, the number of channels C of x and the number K of w are equal, and the number of channels C of w is equal to 1, which is a constant, and formula (2) can be obtained.

[0106] y[N, P, Q, K] = DSC(x(N, H, W, C], w[R, S, C, K]) = DSC(x[N, H, W, K], w[R, S, 1K]) Formula (2)

[0107] In formula (2), the relationships of P, Q, H, W, R, S are as shown in formula (3) and formula (4):

[0108]

[0109] In formula (2), the most common cases for the height R and width S of w are both 3; padw, pad h are the padding parameters for x, corresponding to the width direction and height direction respectively, and are generally all 1; stridew, stride h is the sliding step of the convolution kernel during convolution, corresponding to the width direction and height direction respectively, and is generally all 1; dilation w , dilation h is the dilation size of the convolution kernel, corresponding to the width direction and height direction respectively, and is generally all 1. Therefore, the parameters can be set as: R = S = 3, pad w = pad h = 1, dilation w = dilation h = 1, stride w = stride h = 1, and the size of the input data x is the same as the size of the output data y, that is: Q = W, P = H

[0110] Adopting the idea of tiling, the sample image x is expanded into a first matrix through img2col, which can be matrix A, and the convolution kernel w is expanded into a second matrix through img2col, which can be matrix B. Adopting the idea of tiling, the first matrix is divided into several first block data Tiles A , and the second matrix is divided into several second block data Tiles B .

[0111] Step S404: Determine the initial convolution calculation result based on the first block data and the second block data.

[0112] Optionally, the GPU, through the execution unit, calculates the point-by-point multiplication operation between Tile_A[0, j, i] of the first set of first block data and the corresponding second block data Tile_B[0, 1, i] to determine the initial convolution calculation result Tile_C[0, j, i].

[0113] Optionally, the GPU, through the execution unit, determines the initial convolution calculation result of each group based on the first block data and the second block data through formula (5).

[0114] Tile′ C [k] = Tile C [k, j, i] = Tile A [k, j, i] · Tile B [k, 1, i], 0 ≤ j < Tile_M, 0 ≤ i < Tile_N Formula (5)

[0115] Among them, Tile’c[k] and Tile_C[k, j, i] represent the initial convolution calculation result.

[0116] Step S406: Repeat the steps of reading the first block data and the second block data from the video memory; and determining the initial convolution calculation result based on the first block data and the second block data.

[0117] Among them, the number of executions is determined by the height R of the output feature map of the depthwise separable convolution and the width S of the output feature map of the depthwise separable convolution. For example, R = S = 3, the number of executions is R * S = 3 * 3 = 9. Each execution can read a set of first block data and second block data, so the number of executions is also equal to the number of groups.

[0118] Optionally, the GPU, through the execution unit, repeats the steps of reading the first block data and the second block data in the second group, the second group, the... group from the video memory; and determining the initial convolution calculation result based on the first block data and the second block data according to formula (5).

[0119] Step S408: Based on each initial convolution calculation result, determine the convolution calculation result and save the convolution calculation result to the video memory.

[0120] Optionally, the GPU determines the convolution calculation result through the execution unit based on each initial convolution calculation result using formula (6).

[0121]

[0122] Among them, Tile c [Tile M , Tile N represents the convolution calculation result.

[0123] In an optional embodiment, for the case where the weight kernel size is 3*3 (R = 3, S = 3), the pointwise multiplication operation is as Figure 5 shown. The first block data and the second block data are each divided into 9 groups. The pointwise multiplication operation is performed on the first block data and the second block data in each group to obtain the initial convolution calculation result for each group; the initial convolution calculation results for each group are added together to obtain the convolution calculation result.

[0124] In the above convolution processing method, the execution unit first unfolds the sample image through img2col to obtain the first matrix; unfolds the output feature map through img2col to obtain the second matrix; and divides the first matrix into blocks to obtain the first block data, and divides the second matrix into blocks to obtain the second block data. When the execution unit obtains the first block data and the second block data from the video memory, there is no need to unfold the first matrix and the second matrix. Since the data volumes of the first matrix and the second matrix are large, more video memory resources will be consumed during unfolding. Instead, obtaining the first block data and the second block data with smaller data volumes reduces the occupied video memory resources. Also, when reading the first block data and the second block data from the video memory, there is no need to wait for img2col to complete. The already segmented first block data and second block data can be obtained while img2col is performing matrix unfolding and splitting, reducing the waiting time and thus improving the calculation efficiency. Then, the first block data and the second block data in each group are multiplied through the pointwise multiplication accumulation method to determine the initial convolution calculation result, and based on each initial convolution calculation result, the convolution calculation result is determined. Compared with the traditional technology, since in order to use GEMM, some data needs to be padded with zeros, and by directly performing the pointwise multiplication accumulation operation on the first block data and the second block data, the operation of padding the second matrix with 0 is avoided. Since there are no additional zeros, when using the pointwise multiplication accumulation method, the data transfer and calculation can be greatly reduced, reducing the consumption of computing power resources.

[0125] In an exemplary embodiment, asFigure 6 As shown, reading the first block of data and the second block of data from the video memory includes steps S602 to S604. Among them:

[0126] Step S602, read the first block of data and the second block of data from the video memory through a computing unit, and save the first block of data and the second block of data to the shared cache.

[0127] Among them, the graphics processing unit GPU includes an execution unit (EU) and a video memory (Global Memory). Among them, the execution unit (EU) includes a computing unit (ALU), a shared cache (SM), and a general-purpose register (CRF), as Figure 3 (b) The GPU architecture shown.

[0128] The shared cache is for data sharing among the warps in the computing unit.

[0129] Optionally, the GPU reads the first block of data and the second block of data from the video memory through the computing unit. In order to be able to access the data read by each other in different warps, it is necessary to store the first block of data Tile A and the second block of data Tile B in the shared cache.

[0130] When the execution unit (EU) runs the kernel, it first initializes the relevant general-purpose registers (CRF). Then, it reads the external data from the video memory and temporarily stores it in the shared cache (SM).

[0131] Step S604, read the first block of data and the second block of data from the shared cache to the general-purpose register through the computing unit.

[0132] Among them, the form in which the general-purpose register stores data can be an N-dimensional vector. Generally, N is equal to the number of data channels LaneCount included in the warp.

[0133] Optionally, the GPU uses relevant instructions through the computing unit to implement reading the first block of data Tile A and the second block of data Tile B to the general-purpose register.

[0134] Furthermore, the GPU loads the data into the general-purpose register in multiple times in a block (Tile) manner through each warp in the computing unit.

[0135] Determine the initial convolution calculation result based on the first block data and the second block data, including: Step S606, use the calculation unit to determine the initial convolution calculation result based on the first block data and the second block data in the general register, and save the initial convolution calculation result to the general register.

[0136] Optionally, the GPU uses the calculation unit to determine the initial convolution calculation result based on the first block data and the second block data in the general register, and save the initial convolution calculation result to the general register. The calculation pseudocode is shown in Table 1.

[0137] Table 1 Pseudocode for calculating each initial convolution result

[0138]

[0139] The GPU uses the calculation unit to perform a point-by-point multiplication operation according to the relevant instructions in the (ALU) to obtain the initial convolution calculation result. Loop until each set of initial convolution calculation results is obtained, and save the initial convolution calculation result to the general register. Finally, save and output the respective final results to the video memory.

[0140] In this embodiment, by setting up a shared cache, the data of each thread in the calculation unit can be shared, which can improve the operation speed of the depthwise separable convolution.

[0141] In an exemplary embodiment, the size of the first block data is the first dimension; reading the first block data from the video memory by the calculation unit includes: using the calculation unit to calculate the first video memory read address for reading the first block data from the video memory based on the first horizontal coordinate, the first vertical coordinate, and the width of the output feature map; according to the first video memory read address, use the calculation unit to read the first block data of the first dimension from the video memory.

[0142] Among them, the first horizontal coordinate is the horizontal coordinate for reading in the video memory, which can be denoted as glm_A_L_X; the first vertical coordinate is the vertical coordinate for reading in the video memory, which can be denoted as glm_A_L_Y.

[0143] For the convenience of discussion, in all embodiments of the present application, it is preset in advance that: LaneID is each channel ID. Assuming that there are LaneCount channels in a warp, the range of LaneID is from 0 to LaneCount - 1. WaveID is the warp ID. Assuming that WaveCount is the number of warps, the range of WaveID is from 0 to WaveID - 1. get_local_size() and get_global_id() are built-in functions of OpenCL. It is set that arrB_ID[0] = get_global_id(0) / get_local_size(0) and arrB_ID[1] = get_global_id(1) / get_local_size(1).

[0144] Taking blocks as units, different warps in the computing unit, assuming the number of warps is 4, respectively read the data corresponding to x Tiles from the video memory. A of data. Assuming Tile A has a size of [Tile_M, Tile_N]. The values of Tile_M and Tile_N generally need to be comprehensively considered based on multiple factors such as the performance of the GPU computing unit and the amount of resources such as general-purpose registers. Generally, they can be set to 32, 64 or other values respectively.

[0145] Optionally, the GPU calculates the first video memory read address, denoted as Addr_LTile_A(k, l, global mem), through the computing unit according to the first horizontal coordinate glm_A_L_X and the first vertical coordinate glm_A_L_Y, that is, the address for reading the first block of data from the video memory.

[0146] Further, the GPU substitutes the first horizontal coordinate and the first vertical coordinate into formula (7) through the computing unit to calculate and obtain the first video memory read address.

[0147] Addr_LTile_A(k, I, global mem) = glm_A_L_Y * C + glm_A_L_X; (0 <= I <

[0148] Tile_M * LandCount / 16, 0 <= k < R * S) Formula (7)

[0149] After that, the GPU reads the first block of data with the first size [Tile_M, Tile_N] (taking [32, 64] as an example) from the video memory through the computing unit according to the first video memory read address. As Figure 7 shown, so each warp (wave) reads and writes 32 * 16, that is, 16 columns and 32 rows.

[0150] In this embodiment, by means of the first video memory read address, it is possible to accurately read the first block of data of the first size from the video memory.

[0151] Continuing from the above embodiment, the first size includes a first preset row; the computing unit includes a warp; the warp includes a plurality of data channels; the method for determining the first horizontal coordinate includes: determining the first horizontal coordinate based on the obtained identification number of the data channel, the identification number of the warp, and the number of warps; the method for determining the first vertical coordinate includes: determining the first vertical coordinate based on the obtained identification number of the data channel, the first preset row, the number of data channels included in the warp, the number of current executions, the height of the output feature map, and the width of the sample map.

[0152] Optionally, taking the computing unit including 4 warps, each warp including 64 data channels, and the first preset row being 32 as an example for illustration. The GPU determines the first horizontal coordinate according to the identification number of the data channel LaneID, the identification number of the warp WaveID, the number of warps WaveCount, and the built-in function arrB_ID[0] through formula (8).

[0153] glm_A_L_X = LaneID % 16 + (WaveID + arrB_ID[0] * WaveCount) * 16 Formula (8)

[0154] The GPU determines a vertical coordinate according to the identification number of the data channel LaneID, the first preset row Tile_M, the number of data channels included in the warp LaneCount, the number of current executions k, the height of the output feature map R, and the width of the sample map W through formula (9).

[0155] glm_A_L_Y = LaneID / 16 + arrB_ID[1] * Tile_M + I * LaneCount / 16 + (k % R) +

[0156] (k / R) * W Formula (9)

[0157] In the formula, / is the division operation, * is the multiplication operation, and % is the modulo operation.

[0158] In this embodiment, through parameters such as warps and data channels, the first horizontal coordinate and the first vertical coordinate can be accurately calculated, so that the first video memory read address can be obtained subsequently.

[0159] In an exemplary embodiment, the size of the first block data is a first dimension; the computing unit includes warps; each warp includes a plurality of data channels; saving the first block data to the shared cache includes: calculating a first write shared cache address for writing the first block data into the shared cache based on the obtained second horizontal coordinate, second vertical coordinate, and offset value; each execution operation corresponds to one of the offset values; the second horizontal coordinate is the read horizontal coordinate in the shared cache; the second vertical coordinate is the read vertical coordinate in the shared cache; and both the second horizontal coordinate and the second vertical coordinate are determined based on the identification number of the data channel and the identification number of the warp; writing the first block data of the first dimension into the shared cache according to the first write shared cache address.

[0160] Wherein, the second horizontal coordinate is the read horizontal coordinate in the shared cache, which can be denoted as srm_A_S_X; the first vertical coordinate is the read vertical coordinate in the shared cache, which can be denoted as srm_A_S_Y.

[0161] Optionally, take the example that the computing unit includes 4 warps, each warp includes 64 data channels, the first preset behavior is 32, and the current execution count is the k-th time for illustration. In order to be able to access the data read by each other in different warps, it is necessary to store the TileA data in the shared cache. The GPU calculates a first write shared cache address Addr_STile_A(k, l, share mem) for writing the first block data into the shared cache through the computing unit based on the obtained second horizontal coordinate srm_A_S_X, second vertical coordinate srm_A_S_Y, and offset value Offset_STile_A(k); each output feature map corresponds to one of the offset values; Offset_STile_A is an offset, and for different k, the offset is different. For example, the first execution operation corresponds to the offset value 1, the second execution operation corresponds to the offset value 2, and so on. Among them, the value of the current execution count k is determined by the height R of the output feature map of the depthwise separable convolution and the width S of the output feature map of the depthwise separable convolution.

[0162] Further, the GPU calculates the first write shared cache address by substituting the second horizontal coordinate, second vertical coordinate, and offset value into formula (10) through the computing unit.

[0163] Addr_STile_A(k, I, share mem) = srm_A_S_Y * 64 + srm_A_S_X +

[0164] Offset_STile_A(k); (0 <= k < R * S) Formula (10)

[0165] After that, the GPU writes the first block data of the first size [Tile_M, Tile_N] (taking [32, 64] as an example) into the shared cache through the computing unit according to the first write shared cache address. As Figure 8 shown, each warp reads and writes 32 * 16, that is, 16 columns and 32 rows.

[0166] The GPU determines the second horizontal coordinate srm_A_S_X according to the identification number LaneID of the data channel and the identification number WaveID of the warp through formula (11).

[0167] srm_A_S_X = LaneID % 16 + WaveID * 16 Formula (11)

[0168] The GPU determines the second vertical coordinate srm_A_S_Y according to the identification number LaneID of the data channel and the identification number WaveID of the warp through formula (12).

[0169] srm_A_S_Y = LaneID / 16 + WaveID * 32 Formula (12)

[0170] In this embodiment, by storing the TileA data in the shared cache, the data read by each warp can be accessed mutually.

[0171] In an exemplary embodiment, the size of the second block data is the second size; reading the second block data from the video memory through the computing unit includes: calculating the second read video memory address for reading the second block data from the video memory by the computing unit based on the third horizontal coordinate, the third vertical coordinate, and the width of the output feature map; the third horizontal coordinate is the horizontal coordinate for reading in the video memory; the third vertical coordinate is the vertical coordinate for reading in the video memory; reading the second block data of the second size from the video memory by the computing unit according to the second read video memory address.

[0172] Among them, the third horizontal coordinate is the horizontal coordinate for reading in the video memory, which can be denoted as glm_B_X; the third vertical coordinate is the vertical coordinate for reading in the video memory, which can be denoted as glm_B_Y.

[0173] Taking blocks as units, different warps respectively read the corresponding Tile B data from the video memory. Assuming Tile B has a size of [16, Tile_N]. The value of Tile_N generally needs to be comprehensively considered based on multiple factors such as the performance of the GPU computing unit and the amount of resources such as general registers. Generally, it can be set to 64 or other values respectively.

[0174] Optionally, the GPU calculates the second video memory read address Addr_LTile_B(l, global mem) through the computing unit according to the third horizontal coordinate glm_B_X and the third vertical coordinate glm_B_Y, that is, the address for reading the second block data from the thread.

[0175] Further, the GPU calculates the second video memory read address through the computing unit by substituting the third horizontal coordinate and the fourth vertical coordinate into formula (13).

[0176] Addr_LTile_B(I, global mem) = glm_B_Y * C + glm_B_X; (0 <= I < 2) Formula (13)

[0177] After that, the GPU reads the second block data of the second size [16, Tile_N] (taking [16, 64] as an example) from the video memory through the computing unit according to the second video memory read address. As Figure 9 shown, each warp reads 16 * 64, that is, 64 columns and 16 rows.

[0178] In this embodiment, through the second video memory read address, the second block data of the second size can be accurately read from the video memory.

[0179] Continuing with the above embodiment, the computing unit includes warps; each warp includes multiple data channels; the determination method of the third horizontal coordinate includes: determining the third horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp reads; the determination method of the third vertical coordinate includes: determining the third vertical coordinate based on the obtained identification number of the warp.

[0180] Optionally, taking the computing unit including 4 warps, each warp including 64 data channels, and the second size being [16, 64] as an example for illustration. The GPU determines the third horizontal coordinate according to the identification number LaneID of the data channel, the number of times I the warp reads, and built-in functions through formula (14).

[0181] glm_B_X = LaneID % 32 + arrB_ID[0] * 64 + I * 32 Formula (14)

[0182] In the formula, I * 32 represents reading 32 columns each time in the left - right direction (abscissa), I = 0 or I = 1; * 64 represents each wave reads 64 columns.

[0183] The GPU determines the third horizontal coordinate according to the identification number WaveID of the warp and built - in functions through formula (15).

[0184] glm_B_Y = arrB_ID[1] * 16 + WaveID * 4 Formula (15)

[0185] In the formula, *4 means that each warp (the English term for warp is wave) reads 4 rows.

[0186] In this embodiment, through parameters such as warps and data channels, the third horizontal coordinate and the third vertical coordinate can be accurately calculated, so that the second video memory read address can be obtained subsequently.

[0187] In an exemplary embodiment, the size of the second block data is the second dimension; saving the second block data to the shared cache includes: calculating the second write shared cache address of the second block data written into the shared cache based on the obtained fourth horizontal coordinate, fourth vertical coordinate, and offset value; each output feature map corresponds to an offset value; the fourth horizontal coordinate is the write horizontal coordinate in the shared cache; the fourth vertical coordinate is the write vertical coordinate in the shared cache; writing the second block data of the second dimension into the shared cache according to the second write shared cache address.

[0188] Among them, the fourth horizontal coordinate is the write horizontal coordinate in the shared cache, which can be denoted as srm_B_S_X; the fourth vertical coordinate is the write vertical coordinate in the shared cache, which can be denoted as srm_B_S_Y.

[0189] Optionally, take the computing unit including 4 warps, 64 data channels in the warp, and the first preset behavior of 32 as an example to illustrate. In order to be able to access the data read by each other in different warps, the TileB data needs to be stored in the shared cache. The GPU calculates the second write shared cache address Addr_STile_B(k, I, share mem) of the first block data written into the shared cache through the computing unit based on the obtained fourth horizontal coordinate srm_B_S_X, fourth vertical coordinate srm_B_S_Y, and offset value Offset_STile_B(k); each output feature map corresponds to an offset value; Offset_STile_B is an offset, and for different execution times k, the offset is different.

[0190] Furthermore, the GPU calculates the second write shared cache address by substituting the fourth horizontal coordinate, fourth vertical coordinate, and offset value into Formula (16) through the computing unit.

[0191] Addr_STile_B(k, I, share mem) = srm_B_S_Y * 64 + srm_B_S_X + Offset_STile_B(k);

[0192] (0 <= I < 2) Formula (16)

[0193] After that, the GPU writes the second block data with the second size, taking [16, 64] as an example, into the shared cache through the computing unit according to the second write shared cache address. As Figure 10 shown, each warp reads and writes 4 * 32, that is, 32 columns and 4 rows, and writes twice.

[0194] In this embodiment, by storing the TileB data in the shared cache, the data read by each warp can be accessed mutually.

[0195] Continuing with the above embodiment, the computing unit includes warps; each warp includes a plurality of data channels; the method for determining the fourth horizontal coordinate includes: determining the fourth horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp writes; the method for determining the fourth vertical coordinate includes: determining the fourth vertical coordinate based on the obtained identification number of the data channel, the identification number of the warp, and the number of data channels included in the warp.

[0196] Optionally, taking the computing unit including 4 warps, each warp including 64 data channels, and the second size being [16, 64] as an example for illustration. The GPU determines the fourth horizontal coordinate according to formula (17) based on the identification number LaneID of the data channel and the number of times I the warp writes.

[0197] srm_B_S_X = LaneID % 32 + l * 32 Formula (17)

[0198] In the formula, I * 32 represents reading 32 columns each time in the left - right direction (abscissa), I = 0 or I = 1; % is the modulo operation.

[0199] The GPU determines the fourth vertical coordinate according to formula (18) based on the identification number LaneID of the data channel, the identification number WaveID of the warp, and the number of data channels LandCount included in the warp.

[0200] srm_B_S_Y = LaneID / 32 + WaveID * (LandCount / 32) Formula (18)

[0201] In this embodiment, through parameters such as warps and data channels, the fourth horizontal coordinate and the fourth vertical coordinate can be accurately calculated, so as to obtain the second write shared cache address subsequently.

[0202] In an exemplary embodiment, as Figure 11 shown, the size of the first block data is the first size; the size of the second block data is the second size; reading the first block data and the second block data from the shared cache to the general - purpose register through the computing unit includes steps S1102 to S1108. Among them:

[0203] Step S1102: The calculation unit calculates a first read shared cache address for reading the first block of data from the shared cache based on a fifth horizontal coordinate, a fifth vertical coordinate, and an offset value.

[0204] Among them, each output feature map corresponds to an offset value; the fifth horizontal coordinate is the read horizontal coordinate in the shared cache, which can be denoted as srm_A_L_X; the fifth vertical coordinate is the read vertical coordinate in the shared cache, which can be denoted as srm_A_L_Y.

[0205] Optionally, the GPU calculates a first read shared cache address Addr_LTile_A(l, share mem) for reading the first block of data from the shared cache based on a fifth horizontal coordinate, a fifth vertical coordinate, and an offset value through the calculation unit.

[0206] Further, the GPU substitutes the fifth horizontal coordinate and the fifth vertical coordinate into formula (19) through the calculation unit to calculate and obtain the first read shared cache address.

[0207] Addr_LTile_A(l, share mem) = srm_A_L_Y * 64 + srm_A_L_X + Offset_LTile_A

[0208] (k); (0 <= l < 2) Formula (19)

[0209] Step S1104: According to the first read shared cache address, read the first block of data of the first size from the shared cache and store it in the general register.

[0210] After that, the GPU reads the first block of data of the first size with [32, 64] as an example from the shared cache video memory through the calculation unit according to the first read shared cache address. If there are 4 warps, each warp reads 8 * 64 = 512 data. As Figure 12 shown, so each warp reads and writes 8 * 32, that is, 32 columns and 8 rows, and reads in multiple times. If it is read in 2 times, it will correspond to an offset value; if it is read in 8 times, it will correspond to another offset value. In this example, it needs to be read 2 times to complete. The GPU stores the first block of data of the first size read through the calculation unit in the general register.

[0211] Step S1106: Based on a sixth horizontal coordinate, a sixth vertical coordinate, and an offset value, the calculation unit calculates a second read shared cache address for reading the second block of data from the shared cache.

[0212] Among them, the sixth horizontal coordinate is the read horizontal coordinate in the shared cache, which can be denoted as srm_B_L_X; the sixth vertical coordinate is the read vertical coordinate in the shared cache, which can be denoted as srm_B_L_Y.

[0213] Optionally, the GPU calculates, through a computing unit, a second read shared cache address Addr_LTile_B(k, I, share mem) for reading the second block data from the shared cache based on the sixth horizontal coordinate, the sixth vertical coordinate, and the offset value.

[0214] Furthermore, the GPU calculates, through a computing unit, the second read shared cache address by substituting the sixth horizontal coordinate and the sixth vertical coordinate into formula (20).

[0215] Addr_LTile_B(k, l, share mem) = srm_B_L_Y * 64 + srm_B_S_X + Offset_LTile_B

[0216] (k); (0 <= l < 2,) Formula (20)

[0217] Step S1108: Read the second block data of the second size from the shared cache according to the second read shared cache address, and store it in the general register.

[0218] After that, the GPU reads, through a computing unit, the second block data of the second size with [16, 64] as an example from the shared cache video memory according to the second read shared cache address, and each warp reads 64 data. As Figure 13 shown, so each warp reads and writes 16 * 64, that is, 64 columns and 16 rows, and in this example, it needs to be read 1 time to complete. The GPU stores the second block data of the second size read through the computing unit in the general register.

[0219] In this embodiment, by reading the first block data and the second block data from the shared cache into the general register, the computing unit can process the initial convolution calculation result.

[0220] Continuing from the above embodiment, the computing unit includes warps; each warp includes multiple data channels; the method for determining the fifth horizontal coordinate includes: determining the fifth horizontal coordinate based on the identification number of the obtained data channel, the number of times the warp reads, and the number of data channels included in the warp; the method for determining the fifth vertical coordinate includes: determining the fifth vertical coordinate based on the identification number of the obtained warp, the number of data channels included in the warp, and the number of warps.

[0221] Optionally, taking the computing unit including 4 warps, each warp including 64 data channels, and the first dimension being [32, 64] as an example for illustration. The GPU determines the fifth horizontal coordinate according to the identification number of the data channel LaneID, the number of times the warp reads I, and the number of data channels included in the warp LaneCount through formula (21).

[0222] srm_A_L_X = LaneID + I * LaneCount Formula (21)

[0223] The GPU determines the fifth vertical coordinate according to the identification number of the warp WaveID, the number of data channels included in the warp LaneCount, and the number of warps WaveCount through formula (22).

[0224] srm_A_L_Y = WaveID * LaneCount / WaveCount Formula (22)

[0225] In this embodiment, through parameters such as warps and data channels, the fifth horizontal coordinate and the fifth vertical coordinate can be accurately calculated, so that the first read shared cache address can be obtained subsequently.

[0226] Continuing with the above embodiment, the computing unit includes warps; each warp includes a plurality of data channels; the method for determining the sixth horizontal coordinate includes: determining the sixth horizontal coordinate based on the identification number of the obtained data channel and the number of times the warp writes; the method for determining the sixth vertical coordinate includes: determining the sixth vertical coordinate based on the current number of executions.

[0227] Optionally, taking the computing unit including 4 warps, each warp including 64 data channels, the second dimension being [16, 64], and the current number of executions being the k-th time as an example for illustration. The GPU determines the sixth horizontal coordinate according to the identification number of the data channel LaneID and the number of times the warp writes I through formula (23).

[0228] srm_B_L_X = LaneID + I * 32 Formula (23)

[0229] The GPU determines the sixth vertical coordinate according to the current number of executions k through formula (24).

[0230] srm_B_L_Y = k Formula (24)

[0231] In this embodiment, through parameters such as warps and data channels, the sixth horizontal coordinate and the sixth vertical coordinate can be accurately calculated, so that the second read shared cache address can be obtained subsequently.

[0232] In an exemplary embodiment, the size of the convolution calculation result is the third dimension; storing the convolution calculation result in the video memory includes: calculating the write video memory address for writing the convolution calculation result into the video memory based on the seventh horizontal coordinate, the seventh vertical coordinate, and the current execution count; the seventh horizontal coordinate is the write horizontal coordinate in the video memory; the seventh vertical coordinate is the write vertical coordinate in the video memory; writing the convolution calculation result of the third dimension into the video memory according to the write video memory address.

[0233] Among them, the convolution calculation result Tile C has a size of the third dimension; setting the size of Tile C to be [Tile_M, Tile_N]. The seventh horizontal coordinate is the write horizontal coordinate in the video memory, which can be denoted as glm_C_S_X; the seventh vertical coordinate is the write vertical coordinate in the video memory, which can be denoted as glm_C_S_Y.

[0234] Optionally, the GPU calculates the write video memory address Addr_STile_C(I, global mem) for writing the convolution calculation result into the video memory through the computing unit using the seventh horizontal coordinate glm_C_S_X, the seventh vertical coordinate glm_C_S_Y, and the current execution count k.

[0235] Furthermore, the GPU calculates the write video memory address Addr_STile_C(I, global mem) by substituting the seventh horizontal coordinate glm_C_S_X, the seventh vertical coordinate glm_C_S_Y, and the current execution count k into formula (25) through the computing unit.

[0236] Addr_STile_C(I, global mem) = glm_C_S_Y * K + glm_C_S_X; (0 <= I < Tile_N

[0237] / LandCount) Formula (25)

[0238] After that, the GPU writes the data in the general register corresponding to Tile C into the video memory through the computing unit according to the write video memory address. A total of 32 * 64 = 2048 data are written, that is, the convolution calculation result of the third dimension is written into the video memory. As Figure 14 shown, each warp reads and writes 4 * 64, that is, 32 columns and 4 rows.

[0239] In this embodiment, the convolution calculation result can be stored in the video memory through the write video memory address.

[0240] Continuing from the above embodiment, the third dimension includes a second preset row and a preset column; the computing unit includes a warp; the warp includes a plurality of data channels; the method for determining the seventh horizontal coordinate includes: determining the seventh horizontal coordinate based on the obtained identification number of the data channel, the number of times the warp reads, the number of data channels included in the warp, and the preset column; the method for determining the seventh vertical coordinate includes: determining the seventh vertical coordinate based on the obtained identification number of the warp, the number of warps, and the second preset row.

[0241] Wherein, it is set that the size of Tile C is [Tile_M, Tile_N].

[0242] Optionally, the GPU, through the computing unit, determines the seventh horizontal coordinate according to the identification number of the data channel LaneID, the number of times the warp reads l, the number of data channels included in the warp LaneCount, and the built-in function preset column Tile_N through formula (26).

[0243] glm_C_S_X = LaneID + I * LaneCount + arrB_ID[0] * Tile_N Formula (26)

[0244] The GPU, through the computing unit, determines the seventh vertical coordinate according to the identification number of the warp WaveID, the number of warps WaveCount, and the second preset row Tile_M through formula (27).

[0245] glm_C_S_Y = WaveID * WaveCount + arrB_ID[1] * Tile_M Formula (27)

[0246] In this embodiment, through parameters such as warps and data channels, the seventh horizontal coordinate and the seventh vertical coordinate can be accurately calculated so as to obtain the write video memory address subsequently.

[0247] In an exemplary embodiment, the computing unit reads the first block data and the second block data from the video memory and saves the first block data and the second block data to the shared cache, including: the computing unit reads the first block data and the second block data from the video memory and saves them to the cache; reads the first block data and the second block data from the cache and saves them to the shared cache.

[0248] Wherein, the cache can be a first-level cache or a multi-level cache, such as a second-level cache.

[0249] Optionally, Figure 3Taking the GPU of the (b) architecture as an example for illustration. The GPU reads the first block of data and the second block of data from the video memory through the computing unit, saves them to the cache first, and then the GPU reads the first block of data and the second block of data from the cache through the computing unit, which can accelerate the reading speed. And save them to the shared cache for different warps to access the data read by each other.

[0250] In this embodiment, by setting the cache, the efficiency of reading the first block of data and the second block of data can be accelerated.

[0251] In an exemplary embodiment, before reading the first block of data and the second block of data from the shared cache to the general register through the computing unit, it includes: initializing the general register so that the initial values in the general register are all preset values.

[0252] Optionally, before reading the first block of data and the second block of data from the shared cache to the general register through the computing unit, the GPU initializes the general register, that is, when the execution unit (EU) runs the kernel, it first initializes the relevant general register (CRF), so that the data in the general register meets the preset conditions, such as 0. For the convenience of discussion, it is agreed that the TileC of the calculation result of each warp is stored in the register RC[], then the initialization of TileC is: RC[] = 0.

[0253] In this embodiment, by initializing the general register, the data in the general register meets the preset conditions, such as 0, so as to store each initial convolution calculation result and the final convolution calculation result obtained by the subsequent point-by-point multiplication operation, ensuring the accuracy of the convolution calculation result.

[0254] In an exemplary embodiment, as Figure 15 shown. For the convenience of discussion, in all embodiments of this application, it is preset in advance that: LaneID is each channel ID, it is set that there are LaneCount channels in a warp, then the range of LaneID is from 0 to LaneCount - 1. WaveID is the warp ID, it is set that WaveCount is the number of warps, then the range of WaveID is from 0 to WaveID - 1. get_local_size(), get_global_id() are built-in functions of OpenCL, it is set that arrB_ID[0] = get_global_id(0) / get_local_size(0), arrB_ID[1] = get_global_id(1) / get_local_size(1).

[0255] Initialize TileC = 0. The GPU initializes the general-purpose registers. That is, when the kernel runs on the execution unit (EU), the relevant general-purpose registers (CRF) are initialized first, so that the data in the general-purpose registers meets the preset conditions, such as 0. For the convenience of discussion, it is agreed that the TileC of the calculation result of each warp is stored in the register RC[]. Then, initializing TileC means: RC[] = 0.

[0256] Read the TileA data from the video memory and save the TileA data to the shared cache. Taking the block as a unit, different warps in the computing unit, assuming the number of warps is 4, and each warp includes 64 data channels. Read the corresponding Tile A data from the video memory respectively. Assume that the size of Tile A is [Tile_M, Tile_N]. The values of Tile_M and Tile_N generally need to be comprehensively considered according to multiple factors such as the performance of the GPU computing unit and the amount of resources such as general-purpose registers. Generally, they can be set to 32, 64 or other values respectively. In order to be able to access the data read by each other in different warps, the TileA data needs to be stored in the shared cache. Read the Tile A [32, 64] data into the general-purpose register from the video memory in a block-by-block manner, a total of 2048 numbers are read, and the purpose is to read the x for separable convolution calculation. The first video memory address calculation is as follows:

[0257] The first horizontal coordinate:

[0258] glm_A_L_X = LaneID % 16 + (WaveID + arrB_ID[0] * WaveCount) * 16 Formula (8)

[0259] The first vertical coordinate:

[0260] glm_A_L_Y = LaneID / 16 + arrB_ID[1] * Tile_M + I * LaneCount / 16 + (k % R) +

[0261] (k / R) * W Formula (9)

[0262] The first video memory address:

[0263] Addr_LTile_A(k, I, global mem) = glm_A_L_Y * C + glm_A_L_X; (0 <= I <

[0264] Tile_M * LandCount / 16, 0 <= k < R * S) Formula (7)

[0265] The GPU reads the first block of data of the first size [Tile_M, Tile_N] (taking [32, 64] as an example) from the video memory through the computing unit according to the first video memory read address. As Figure 7 shown, so each warp reads and writes 32 * 16, that is, 16 columns and 32 rows.

[0266] Cache TileA: Save TileA to the shared cache. Each warp writes the data in the general register to the shared cache. The calculation of the first shared cache write address is as follows:

[0267] The second horizontal coordinate:

[0268] srm_A_S_X = LaneID % 16 + WaveID * 16 Formula (11)

[0269] The second vertical coordinate:

[0270] srm_A_S_Y = LaneID / 16 + WaveID * 32 Formula (12)

[0271] The first shared cache write address:

[0272] Addr_STile_A(k, I, share mem) = srm_A_S_Y * 64 + srm_A_S_X +

[0273] Offset_STile_A(k); (0 <= k < R * S) Formula (10)

[0274] The GPU writes the first block of data of the first size [Tile_M, Tile_N] (taking [32, 64] as an example) to the shared cache through the computing unit according to the first shared cache write address. As Figure 8 shown, so each warp reads and writes 32 * 16, that is, 16 columns and 32 rows.

[0275] Read TileB data from the video memory and save the TileB data to the shared cache. In units of blocks, different warps read the corresponding Tile B data from the video memory. Assuming Tile BThe size is [16, Tile_N]. The value of Tile_N generally needs to be comprehensively considered based on multiple factors such as the performance of the GPU computing unit and the amount of resources such as general-purpose registers. It can generally be set to 64 or other values respectively. In order to be able to access the data read by each other in different warps, the TileB data needs to be stored in the shared cache. The w in the video memory is read in a block manner, and the TileB[16, 64] data is read into the general-purpose register from the video memory, a total of 1024 numbers are read, and the purpose is to read the w for separable convolution calculation. The second video memory read address is calculated as follows:

[0276] The third horizontal coordinate:

[0277] glm_B_X = LaneID % 32 + arrB_ID[0] * 64 + I * 32 Formula (14)

[0278] The third horizontal coordinate:

[0279] glm_B_Y = arrB_ID[1] * 16 + WaveID * 4 Formula (15)

[0280] The second video memory read address:

[0281] Addr_LTile_B(I, global mem) = glm_B_Y * C + glm_B_X; (0 <= I < 2) Formula (13)

[0282] The GPU reads the second block data of the second size [16, Tile_N] (taking [16, 64] as an example) from the video memory through the computing unit according to the second video memory read address. As Figure 9 shown, so each warp reads and writes 16 * 64, that is, 64 columns and 16 rows.

[0283] Cache TileB: Save the TileB data to the shared cache. Each warp writes the data in the general-purpose register to the shared cache. The first shared cache write address is calculated as follows:

[0284] The fourth horizontal coordinate:

[0285] srm_B_S_X = LaneID % 32 + l * 32 Formula (17)

[0286] The fourth vertical coordinate:

[0287] srm_B_S_Y = LaneID / 32 + WaveID * (LandCount / 32) Formula (18)

[0288] The second shared cache write address:

[0289] Addr_STile_B(k, I, share mem) = srm_B_S_Y * 64 + srm_B_S_X + Offset_STile_B(k);

[0290] (0 <= I < 2) Equation (16)

[0291] The GPU, through the computing unit, writes the second block data with the second size [16, 64] as an example to the shared cache according to the second write shared cache address. As Figure 10 shown, so each warp reads and writes 4 * 32, that is, 32 columns and 4 rows, and writes twice.

[0292] Read TileA and TileB data from the shared cache.

[0293] Read TileA data. Read TileA data from the shared cache to prepare for the next calculation. Each warp reads the data in the shared cache into the general register, and a total of 8 * 64 = 512 data are read. The first read shared cache address is calculated as follows:

[0294] The fifth horizontal coordinate:

[0295] srm_A_L_X = LaneID + I * LaneCount Equation (21)

[0296] The fifth vertical coordinate:

[0297] srm_A_L_Y = WaveID * LaneCount / WaveCount Equation (22)

[0298] The first read shared cache address:

[0299] Addr_LTile_A(l, share mem) = srm_A_L_Y * 64 + srm_A_L_X + Offset_LTile_A

[0300] (k); (0 <= l < 2) Equation (19)

[0301] The GPU, through the computing unit, reads the first block data with the first size [32, 64] as an example from the shared cache video memory according to the first read shared cache address. If there are 4 warps, each warp reads 8 * 64 = 512 data. As Figure 12As shown, each warp reads and writes 8 * 32, that is, 32 columns and 8 rows. It is read in multiple times. If it is read in 2 times, it will correspond to an offset value; if it is read in 8 times, it will correspond to another offset value. In this example, it needs to be read 2 times to complete. The GPU stores the first block data of the first size read through the computing unit into the general register.

[0302] Read the TileB data. Read the TileB data from the shared cache to prepare for the next calculation. Each warp reads the data in the shared cache into the general register, and a total of 16 * 64 = 1024 data are read. The calculation of the second read shared cache address is as follows:

[0303] The sixth horizontal coordinate:

[0304] srm_B_L_X = LaneID + I * 32 Formula (23)

[0305] The sixth vertical coordinate:

[0306] srm_B_L_Y = k Formula (24)

[0307] The second read shared cache address:

[0308] Addr_LTile_B(k, l, share mem) = srm_B_L_Y * 64 + srm_B_S_X + Offset_LTile_B

[0309] (k); (0 <= l < 2,) Formula (20)

[0310] The GPU reads the second block data of the second size with [16, 64] as an example from the shared cache video memory through the computing unit according to the second read shared cache address, as Figure 13 shown. So each warp reads and writes 16 * 64, that is, 64 columns and 16 rows. In this example, it needs to be read 1 time to complete. The GPU stores the second block data of the second size read into the general register through the computing unit.

[0311] Calculate TileC += TileA · TileB. For the convenience of discussion, the number of repeated executions is taken as 9 times. Each warp takes out the TileA data from the shared cache and stores it in the general register RA[8], takes out the TileB data and stores it in the general register RB, and the calculated result TileC is stored in the register RC[8]. Each general register is similar to an N-dimensional vector. Generally, N is equal to LandCount. The calculation pseudocode is shown in Table 1.

[0312] Table 1 Pseudocode for Calculating Each Initial Convolution Result

[0313]

[0314] Repeat the step of reading the first block data and the second block data from the video memory k times; determine the initial convolution calculation result based on the first block data and the second block data; 0 <= k < R * S.

[0315] Write the data in the general register corresponding to TileC to the video memory, writing a total of 32 * 64 = 2048 data. Set the size of TileC to [Tile_M, Tile_N]. The calculation of the write video memory address for TileC is as follows:

[0316] The seventh horizontal coordinate:

[0317] glm_C_S_X = LaneID + I * LaneCount + arrB_ID[0] * Tile_N Formula (26)

[0318] The seventh vertical coordinate:

[0319] glm_C_S_Y = WaveID * WaveCount + arrB_ID[1] * Tile_M Formula (27)

[0320] The write video memory address:

[0321] Addr_STile_C(I, global mem) = glm_C_S_Y * K + glm_C_S_X; (0 <= I < Tile_N

[0322] / LandCount) Formula (25)

[0323] The GPU writes the data in the general register corresponding to Tile to the video memory according to the write video memory address through the computing unit, writing a total of 32 * 64 = 2048 data, that is, writing the convolution calculation result of the third dimension to the video memory. As C shown, so each warp reads and writes 4 * 64, that is, 32 columns and 4 rows. Figure 14 shown, so each warp reads and writes 4 * 64, that is, 32 columns and 4 rows.

[0324] Depthwise separable convolution refers to the idea of img2col + GEMM, and according to the characteristics of depthwise separable convolution, when calculating the output result, the convolution operation is changed to pointwise multiplication and accumulation operation (img2col + APM), eliminating the extended operation of zero-padding filling for the convolution kernel, and also saving the calculation amount. On the one hand, for the convolution kernel data, there is no need to fill zeros for expansion, saving the storage access of video memory and shared cache; on the other hand, using pointwise multiplication and accumulation operation instead of dot product operation saves the calculation time based on DSC.

[0325] It should be understood that although the steps in the flowcharts involved in the above-described embodiments are shown in sequence according to the arrows, these steps are not necessarily executed in the order indicated by the arrows. Unless there is a clear description in this article, there is no strict order limit for the execution of these steps, and these steps can be executed in other orders. Moreover, at least a part of the steps in the flowcharts involved in the above-described embodiments may include multiple steps or multiple stages. These steps or stages are not necessarily executed at the same time, but can be executed at different times. The execution order of these steps or stages is not necessarily sequential, but can be executed alternately or in turn with at least a part of other steps or steps or stages in other steps.

[0326] Based on the same inventive concept, an embodiment of the present application also provides a graphics processor for implementing the above-mentioned convolution processing method. The solution provided by this graphics processor to solve the problem is similar to the solution described in the above method. Therefore, the specific limitations in one or more of the following graphics processor embodiments can be referred to the limitations on the convolution processing method in the above text, and will not be repeated here.

[0327] In an exemplary embodiment, as Figure 3 (b) shows, a graphics processor is provided, including: an execution unit and a video memory, where:

[0328] The execution unit is configured to read first block data and second block data from the video memory; wherein, the first block data is determined by partitioning a first matrix obtained by expanding a sample graph; the second block data is determined by partitioning a second matrix obtained by expanding an output feature map; the output feature map is output by a convolutional neural network; based on the first block data and the second block data, determine an initial convolution calculation result; repeatedly execute the steps of reading the first block data and the second block data from the video memory; and determining the initial convolution calculation result based on the first block data and the second block data; based on each initial convolution calculation result, determine a convolution calculation result and save the convolution calculation result to the video memory.

[0329] The video memory is configured to store the first block data and the second block data; and is also configured to store the convolution calculation result.

[0330] In an exemplary embodiment, the execution unit includes:

[0331] A calculation unit is configured to read first and second block data from a video memory and save the first and second block data to a shared cache; read the first and second block data from the shared cache to general-purpose registers; determine an initial convolution calculation result based on the first and second block data in the general-purpose registers, and save the initial convolution calculation result to the general-purpose registers.

[0332] The shared cache is used to store the first and second block data.

[0333] The general-purpose registers are used to store the first and second block data; and are also used to save the initial convolution calculation result.

[0334] In an exemplary embodiment, the size of the first block data is a first dimension; the calculation unit is further configured to calculate a first video memory read address for reading the first block data from the video memory based on a first horizontal coordinate, a first vertical coordinate, and the width of an output feature map; the first horizontal coordinate is the horizontal read coordinate in the video memory; the first vertical coordinate is the vertical read coordinate in the video memory; according to the first video memory read address, the calculation unit reads the first block data of the first dimension from the video memory.

[0335] In an exemplary embodiment, the first dimension includes a first preset row; the calculation unit includes a warp; the warp includes a plurality of data channels; the graphics processor further includes a first horizontal coordinate determination module configured to determine the first horizontal coordinate based on the identification number of the acquired data channel, the identification number of the warp, and the number of warps. The graphics processor further includes a first vertical coordinate determination module configured to determine the first vertical coordinate based on the identification number of the acquired data channel, the first preset row, the number of data channels included in the warp, the current execution times, the height of the output feature map, and the width of the sample map.

[0336] In an exemplary embodiment, the size of the first block data is a first dimension; the calculation unit includes a warp; each warp includes a plurality of data channels; the calculation unit is further configured to calculate a first shared cache write address for writing the first block data to the shared cache based on the acquired second horizontal coordinate, second vertical coordinate, and offset value; each output feature map corresponds to an offset value; the second horizontal coordinate is the horizontal read coordinate in the shared cache; the second vertical coordinate is the vertical read coordinate in the shared cache; according to the first shared cache write address, the first block data of the first dimension is written to the shared cache.

[0337] The graphics processor further includes a second horizontal coordinate determination module and a second vertical coordinate determination module configured to determine the second horizontal coordinate and the second vertical coordinate based on the identification number of the data channel and the identification number of the warp.

[0338] In an exemplary embodiment, the size of the second block data is a second dimension; the calculation unit is further configured to calculate a second video memory read address for reading the second block data from the video memory based on a third horizontal coordinate, a third vertical coordinate, and the width of the output feature map; the third horizontal coordinate is the horizontal coordinate for reading in the video memory; the third vertical coordinate is the vertical coordinate for reading in the video memory; according to the second video memory read address, the calculation unit reads the second block data of the second dimension from the video memory.

[0339] In an exemplary embodiment, the calculation unit includes a warp; each warp includes a plurality of data channels; the graphics processor further includes a third horizontal coordinate determination module configured to determine a third horizontal coordinate based on the identification number of the acquired data channel and the number of times the warp reads; a third vertical coordinate determination module configured to determine a third vertical coordinate based on the identification number of the acquired warp.

[0340] In an exemplary embodiment, the size of the second block data is a second dimension; the calculation unit is further configured to calculate a second shared cache write address for writing the second block data into the shared cache based on the acquired fourth horizontal coordinate, fourth vertical coordinate, and offset value; each output feature map corresponds to an offset value; the fourth horizontal coordinate is the horizontal coordinate for writing in the shared cache; the fourth vertical coordinate is the vertical coordinate for writing in the shared cache; according to the second shared cache write address, the second block data of the second dimension is written into the shared cache.

[0341] In an exemplary embodiment, the calculation unit includes a warp; the warp includes a plurality of data channels; the graphics processor further includes a fourth horizontal coordinate determination module configured to determine a fourth horizontal coordinate based on the identification number of the acquired data channel and the number of times the warp writes. A fourth vertical coordinate determination module configured to determine a fourth vertical coordinate based on the identification number of the acquired data channel, the identification number of the warp, and the number of data channels included in the warp.

[0342] In an exemplary embodiment, the size of the first chunk of data is a first dimension; the size of the second chunk of data is a second dimension; the computing unit is further configured to calculate a first read shared cache address for reading the first chunk of data from the shared cache based on a fifth horizontal coordinate, a fifth vertical coordinate, and an offset value; wherein each output feature map corresponds to an offset value; the fifth horizontal coordinate is the read horizontal coordinate in the shared cache; the fifth vertical coordinate is the read vertical coordinate in the shared cache; according to the first read shared cache address, read the first chunk of data of the first dimension from the shared cache and store it in the general register; calculate a second read shared cache address for reading the second chunk of data from the shared cache based on a sixth horizontal coordinate, a sixth vertical coordinate, and the offset value by the computing unit; wherein the sixth horizontal coordinate is the read horizontal coordinate in the shared cache; the sixth vertical coordinate is the read vertical coordinate in the shared cache; according to the second read shared cache address, read the second chunk of data of the second dimension from the shared cache and store it in the general register.

[0343] In an exemplary embodiment, the computing unit includes a warp; the warp includes a plurality of data channels; the graphics processor further includes a fifth horizontal coordinate determination module configured to determine the fifth horizontal coordinate based on the obtained identification number of the data channel, the number of times the warp reads, and the number of data channels included in the warp. A fifth vertical coordinate determination module configured to determine the fifth vertical coordinate based on the obtained identification number of the warp, the number of data channels included in the warp, and the number of warps.

[0344] In an exemplary embodiment, the computing unit includes a warp; the warp includes a plurality of data channels; the graphics processor further includes a sixth horizontal coordinate determination module configured to determine the sixth horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp writes. A sixth vertical coordinate determination module configured to determine the sixth vertical coordinate based on the number of times of the current execution.

[0345] In an exemplary embodiment, the size of the convolution calculation result is a third dimension; the computing unit is further configured to calculate a write video memory address for writing the convolution calculation result to the video memory based on a seventh horizontal coordinate, a seventh vertical coordinate, and the number of times of the current execution; the seventh horizontal coordinate is the write horizontal coordinate in the video memory; the seventh vertical coordinate is the write vertical coordinate in the video memory; according to the write video memory address, write the convolution calculation result of the third dimension to the video memory.

[0346] In an exemplary embodiment, the third dimension includes a second preset row and a preset column; the computing unit includes a warp; the warp includes a plurality of data channels; the graphics processor further includes a seventh horizontal coordinate determination module, configured to determine a seventh horizontal coordinate based on the obtained identification number of the data channel, the number of times the warp reads, the number of data channels included in the warp, and the preset column. A seventh vertical coordinate determination module, configured to determine a seventh vertical coordinate based on the obtained identification number of the warp, the number of warps, and the second preset row.

[0347] In an exemplary embodiment, the computing unit is further configured to read first block data and second block data from the video memory and store them in the cache; read the first block data and the second block data from the cache and store them in the shared cache.

[0348] In an exemplary embodiment, the initialization module is configured to initialize the general-purpose registers so that the initial values in the general-purpose registers are all preset values.

[0349] Each module and unit in the above-mentioned graphics processor can be implemented in whole or in part by software, hardware, and their combination. Each of the above modules can be embedded in the processor in the computer device in hardware form or be independent of the processor, or can be stored in the memory in the computer device in software form, so that the processor can call and execute the operations corresponding to each of the above modules.

[0350] In an exemplary embodiment, a computer device is provided. The computer device can be a server, and its internal structure diagram can be as Figure 16 shown. The computer device includes a processor, a memory, an input / output interface (Input / Output, abbreviated as I / O), and a communication interface. Among them, the processor, the memory, and the input / output interface are connected through a system bus, and the communication interface is connected to the system bus through the input / output interface. Among them, the processor of the computer device is configured to provide computing and control capabilities. The memory of the computer device includes a non-volatile storage medium and an internal memory. The non-volatile storage medium stores an operating system, a computer program, and a database. The internal memory provides an environment for the operation of the operating system and the computer program in the non-volatile storage medium. The database of the computer device is configured to store the first block data and the second block data. The input / output interface of the computer device is configured to exchange information between the processor and external devices. The communication interface of the computer device is configured to communicate with an external terminal through a network connection. When the computer program is executed by the processor, a convolution processing method is implemented.

[0351] Those skilled in the art can understand, Figure 16The structure shown is only a block diagram of some structures related to the solution of this application, and does not constitute a limitation on the computer device to which the solution of this application is applied. The specific computer device may include more or fewer components than those shown in the figure, or combine some components, or have different component arrangements.

[0352] In one embodiment, a computer device is further provided, including a memory and a graphics processor. A computer program is stored in the memory, and when the processor executes the computer program, the steps in the above method embodiments are implemented.

[0353] In one embodiment, a computer-readable storage medium is provided, on which a computer program is stored. When the computer program is executed by the graphics processor, the steps in the above method embodiments are implemented.

[0354] In one embodiment, a computer program product is provided, including a computer program. When the computer program is executed by the graphics processor, the steps in the above method embodiments are implemented.

[0355] Those of ordinary skill in the art can understand that all or part of the processes in the methods of the above embodiments can be completed by instructing relevant hardware through a computer program. The computer program can be stored in a non-volatile computer-readable storage medium. When the computer program is executed, it can include the processes of the embodiments of the above methods. Among them, any reference to a memory, database, or other medium used in the embodiments provided in the present application can include at least one of non-volatile memory and volatile memory. Non-volatile memory can include read-only memory (ROM), magnetic tape, floppy disk, flash memory, optical memory, high-density embedded non-volatile memory, resistive random access memory (ReRAM), magnetoresistive random access memory (MRAM), ferroelectric random access memory (FRAM), phase change memory (PCM), graphene memory, etc. Volatile memory can include random access memory (RAM) or external cache memory, etc. By way of illustration and not limitation, RAM can be in various forms, such as static random access memory (SRAM) or dynamic random access memory (DRAM), etc. The databases involved in the embodiments provided in the present application can include at least one of relational databases and non-relational databases. Non-relational databases can include distributed databases based on blockchain, etc., without limitation.

[0356] The technical features of the above embodiments can be combined arbitrarily. For the sake of brevity of description, not all possible combinations of the technical features in the above embodiments are described. However, as long as there is no contradiction in the combination of these technical features, it should be considered as the scope recorded in the present application.

[0357] The above-described embodiments merely represent several implementation manners of the present application. The description is relatively specific and detailed, but it should not be construed as a limitation on the patent scope of the present application. It should be noted that for those of ordinary skill in the art, without departing from the concept of the present application, several modifications and improvements can still be made, and these all belong to the protection scope of the present application. Therefore, the protection scope of the present application should be subject to the appended claims.

Claims

1. A convolution processing method, characterized in that, The method includes: Reading first block data and second block data from a video memory; wherein, the first block data is determined by partitioning a first matrix obtained by expanding a sample graph; the second block data is determined by partitioning a second matrix obtained by expanding an output feature map; the output feature map is output by a convolutional neural network; Determining an initial convolution calculation result based on the first block data and the second block data; Repeating the steps of reading first block data and second block data from the video memory; and determining an initial convolution calculation result based on the first block data and the second block data; Determining a convolution calculation result based on each of the initial convolution calculation results, and saving the convolution calculation result to the video memory.

2. The method according to claim 1, wherein The reading of the first block data and the second block data from the video memory includes: Reading the first block data and the second block data from the video memory by a computing unit, and saving the first block data and the second block data to a shared cache; Reading the first block data and the second block data from the shared cache to general-purpose registers by the computing unit; The determining of the initial convolution calculation result based on the first block data and the second block data includes: Determining an initial convolution calculation result based on the first block data and the second block data in the general-purpose registers by the computing unit, and saving the initial convolution calculation result to the general-purpose registers.

3. The method according to claim 2, wherein The size of the first block data is a first dimension; the reading of the first block data from the video memory by the computing unit includes: Calculating a first video memory read address for reading the first block data from the video memory by the computing unit based on a first horizontal coordinate, a first vertical coordinate, and the width of the output feature map; the first horizontal coordinate is a horizontal read coordinate in the video memory; the first vertical coordinate is a vertical read coordinate in the video memory; Reading the first block data of the first dimension from the video memory by the computing unit according to the first video memory read address.

4. The method according to claim 3, wherein The first dimension includes a first preset row; the computing unit includes a warp; the warp includes a plurality of data channels; the determining manner of the first horizontal coordinate includes: Determining the first horizontal coordinate based on the obtained identification number of the data channel, the identification number of the warp, and the number of warps; The determining manner of the first vertical coordinate includes: Determining the first vertical coordinate based on the obtained identification number of the data channel, the first preset row, the number of data channels included in the warp, the current execution times, the height of the output feature map, and the width of the sample graph.

5. The method according to claim 2, wherein The size of the first block data is a first dimension; the computing unit includes a warp; each warp includes a plurality of data channels; Saving the first block data to the shared cache includes: Based on the obtained second horizontal coordinate, second vertical coordinate, and offset value, calculate the first write shared cache address for writing the first block of data into the shared cache; each execution of the operation corresponds to one of the offset values; the second horizontal coordinate is the read horizontal coordinate in the shared cache; the second vertical coordinate is the read vertical coordinate in the shared cache; and both the second horizontal coordinate and the second vertical coordinate are determined based on the identification number of the data channel and the identification number of the warp. According to the first write shared cache address, write the first block of data of the first size into the shared cache.

6. The method according to claim 2, wherein The size of the second block of data is the second size; reading the second block of data from the video memory by the computing unit includes: The computing unit calculates a second read video memory address for reading the second block of data from the video memory based on a third horizontal coordinate, a third vertical coordinate, and the width of the output feature map; the third horizontal coordinate is the read horizontal coordinate in the video memory; the third vertical coordinate is the read vertical coordinate in the video memory. According to the second read video memory address, the computing unit reads the second block of data of the second size from the video memory.

7. The method according to claim 6, characterized in that, The computing unit includes warps; each warp includes a plurality of data channels; the method for determining the third horizontal coordinate includes: Determine the third horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp reads. The method for determining the third vertical coordinate includes: Determine the third vertical coordinate based on the obtained identification number of the warp.

8. The method according to claim 2, wherein The size of the second block of data is the second size; saving the second block of data to the shared cache includes: Based on the obtained fourth horizontal coordinate, fourth vertical coordinate, and offset value, calculate the second write shared cache address for writing the second block of data into the shared cache; each output feature map corresponds to one of the offset values; the fourth horizontal coordinate is the write horizontal coordinate in the shared cache; the fourth vertical coordinate is the write vertical coordinate in the shared cache. According to the second write shared cache address, write the second block of data of the second size into the shared cache.

9. The method according to claim 8, wherein The computing unit includes warps; the warp includes a plurality of data channels; the method for determining the fourth horizontal coordinate includes: Determine the fourth horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp writes. The method for determining the fourth vertical coordinate includes: Determine the fourth vertical coordinate based on the obtained identification number of the data channel, the identification number of the warp, and the number of data channels included in the warp.

10. The method according to claim 2, wherein The size of the first block of data is the first size; the size of the second block of data is the second size; reading the first block of data and the second block of data from the shared cache to the general register by the computing unit includes: The computing unit calculates a first read shared cache address for reading the first block of data from the shared cache based on a fifth horizontal coordinate, a fifth vertical coordinate, and an offset value; wherein each of the output feature maps corresponds to an offset value; the fifth horizontal coordinate is the horizontal coordinate for reading in the shared cache; the fifth vertical coordinate is the vertical coordinate for reading in the shared cache; According to the first read shared cache address, the first block of data of the first size is read from the shared cache and stored in the general-purpose register; Based on a sixth horizontal coordinate, a sixth vertical coordinate, and an offset value, the computing unit calculates a second read shared cache address for reading the second block of data from the shared cache; wherein the sixth horizontal coordinate is the horizontal coordinate for reading in the shared cache; the sixth vertical coordinate is the vertical coordinate for reading in the shared cache; According to the second read shared cache address, the second block of data of the second size is read from the shared cache and stored in the general-purpose register.

11. The method according to claim 10, characterized in that The computing unit includes a warp; the warp includes multiple data channels; the method for determining the fifth horizontal coordinate includes: Determining the fifth horizontal coordinate based on the obtained identification number of the data channel, the number of times the warp reads, and the number of data channels included in the warp; The method for determining the fifth vertical coordinate includes: Determining the fifth vertical coordinate based on the obtained identification number of the warp, the number of data channels included in the warp, and the number of warps.

12. The method according to claim 10, wherein The computing unit includes a warp; the warp includes multiple data channels; the method for determining the sixth horizontal coordinate includes: Determining the sixth horizontal coordinate based on the obtained identification number of the data channel and the number of times the warp writes; The method for determining the sixth vertical coordinate includes: Determining the sixth vertical coordinate based on the current execution count.

13. The method according to claim 2, wherein The size of the convolution calculation result is a third size; storing the convolution calculation result in the video memory includes: Calculating a write video memory address for writing the convolution calculation result into the video memory based on a seventh horizontal coordinate, a seventh vertical coordinate, and the current execution count; the seventh horizontal coordinate is the horizontal coordinate for writing in the video memory; the seventh vertical coordinate is the vertical coordinate for writing in the video memory; According to the write video memory address, the convolution calculation result of the third size is written into the video memory.

14. The method according to claim 13, wherein The third size includes a second preset row and a preset column; the computing unit includes a warp; the warp includes multiple data channels; the method for determining the seventh horizontal coordinate includes: Determining the seventh horizontal coordinate based on the obtained identification number of the data channel, the number of times the warp reads, the number of data channels included in the warp, and the preset column; The method for determining the seventh vertical coordinate includes: Determining the seventh vertical coordinate based on the obtained identification number of the warp, the number of warps, and the second preset row.

15. The method according to claim 2, wherein The step of the computing unit reading the first block data and the second block data from the video memory and saving the first block data and the second block data to the shared cache includes: The computing unit reads the first block data and the second block data from the video memory and saves them to the cache; The first block data and the second block data are read from the cache and saved to the shared cache.

16. The method according to claim 2, wherein Before the computing unit reads the first block data and the second block data from the shared cache to the general-purpose register, it includes: Initializing the general-purpose register so that the initial values in the general-purpose register are all preset values.

17. A graphics processor, characterized in that, It includes: An execution unit for reading first block data and second block data from the video memory; wherein, the first block data is determined by partitioning a first matrix obtained by expanding a sample graph; the second block data is determined by partitioning a second matrix obtained by expanding an output feature map; the output feature map is output by a convolutional neural network; based on the first block data and the second block data, an initial convolution calculation result is determined; the steps of repeatedly reading the first block data and the second block data from the video memory; and determining an initial convolution calculation result based on the first block data and the second block data are executed; based on each of the initial convolution calculation results, a convolution calculation result is determined and the convolution calculation result is saved to the video memory; A video memory for storing the first block data and the second block data; and also for storing the convolution calculation result.

18. A computer device, comprising a memory and a graphics processor, the memory storing a computer program, characterized in that, When the graphics processor executes the computer program, the steps of the method according to any one of claims 1 to 16 are implemented.

19. A computer-readable storage medium having a computer program stored thereon, characterized in that, When the computer program is executed by the graphics processor, the steps of the method according to any one of claims 1 to 16 are implemented.

20. A computer program product comprising a computer program, characterized in that, When the computer program is executed by the graphics processor, the steps of the method according to any one of claims 1 to 16 are implemented.

Images