Method, computing device, medium and program product for memory access of tensor data

Abstract

The invention relates to a method for memory access of tensor data, a computing device, a computer readable storage medium and a computer program product. The method comprises the following steps: converting an input tensor into a one-dimensional array; determining non-tail elements and tail elements of the one-dimensional array based on a vectorization unit and the total element number of the one-dimensional array; processing non-tail elements of the one-dimensional array in a vectorized continuous memory access mode; and processing tail elements of the one-dimensional array in an element-by-element memory access mode. According to the method, the processing efficiency of the tensor data in the non-alignment scene can be remarkably improved, and the code complexity and the maintenance cost are reduced.

Description

TECHNICAL FIELD

[0001] Embodiments of the present application generally relate to the field of artificial intelligence, and more particularly to a method for memory access of tensor data, a computing device, a computer readable storage medium, and a computer program product. BACKGROUND

[0002] In current mainstream deep learning frameworks (such as PyTorch, TensorFlow) and underlying operator libraries (such as cuBLAS, cuDNN, Cutlass), pointwise operators for tensor data are usually implemented through kernel functions.

[0003] The traditional memory access method for tensor data associated with an operator, in a non-aligned scenario, for example, the shape of the input tensor does not meet the alignment state, for example, the total number of elements of the input tensor cannot be divided by the vectorization unit (such as 4 floats). For example, the starting address is not aligned. These non-aligned scenarios will cause the entire kernel function to use scalar read-write for memory access throughout, and cannot use the high-bandwidth burst mode for memory access. In addition, the above-mentioned traditional method also needs to explicitly judge the alignment state on the host (Host) or device (Device) side, and maintain two sets of kernel function logic, which increases the code complexity and maintenance cost. Therefore, the traditional memory access method for tensor data has the defects of low processing efficiency, high code complexity, and high maintenance cost in a non-aligned scenario.

[0004] In summary, the traditional memory access method for tensor data has the defects of low processing efficiency and high maintenance cost in a non-aligned scenario. SUMMARY

[0005] The present application provides a method for memory access of tensor data, a computing device, a computer readable storage medium, and a computer program product, which can significantly improve the processing efficiency of tensor data in a non-aligned scenario, and reduce the code complexity and maintenance cost.

[0006] According to a first aspect of the present application, a method for memory access of tensor data is provided. The method comprises: converting an obtained input tensor associated with a current operator into a one-dimensional array; determining non-tail elements and tail elements of the one-dimensional array based on a vectorization unit and a total number of elements of the one-dimensional array; processing the non-tail elements of the one-dimensional array in a vectorized continuous memory access manner; and processing the tail elements of the one-dimensional array in a pointwise memory access manner.

[0007] According to a second aspect of the present application, there is also provided a computing device. The computing device comprises at least one processor; and a memory communicatively connected with the at least one processor; wherein the memory stores instructions executable by the at least one processor, the instructions being executed by the at least one processor to enable the computing device to perform the method of the second aspect of the present application.

[0008] According to a third aspect of the present application, there is also provided a computer readable storage medium. The computer readable storage medium has stored thereon a computer program which, when executed by a machine, performs the method of the second aspect of the present application.

[0009] According to a fourth aspect of the present application, there is also provided a computer program product comprising a computer program which, when executed by a machine, performs the method of the second aspect of the present application.

[0010] In some embodiments, the method for memory access of tensor data further comprises: determining a vectorization type of the one-dimensional array based on a data type of the input tensor, so as to determine a vectorization unit with respect to the one-dimensional array.

[0011] In some embodiments, converting the input tensor into the one-dimensional array comprises: in response to determining that any of the following conditions is met, converting the input tensor into the one-dimensional array: the current operator does not involve an accumulation or reduction operation across elements, and each output element only depends on an input element at a corresponding position; the current operator is an element-wise operator included in a predetermined operator set.

[0012] In some embodiments, processing the non-tail elements of the one-dimensional array in a vectorized continuous memory access manner comprises: for the non-tail elements, reading or writing a number of elements corresponding to the vectorization unit at a time in a form corresponding to the vectorization type.

[0013] In some embodiments, determining the non-tail elements and the tail elements of the one-dimensional array comprises: dividing a total number of elements of the one-dimensional array by the vectorization unit to obtain a remainder of the division, and determining the tail elements of the one-dimensional array based on the obtained remainder.

[0014] In some embodiments, the method for memory access of tensor data further comprises: based on a total number of elements of the one-dimensional array, allocating threads and thread blocks for processing the one-dimensional array by the current operator, so that each thread and thread block is used to process a continuous index interval of the one-dimensional array; and causing the allocated threads and thread blocks to process the full amount of data of the one-dimensional array in parallel.

[0015] In some embodiments, processing the non-tail elements of the one-dimensional array in a way of vectorized continuous memory access comprises: determining the index of all elements of the one-dimensional array; and in response to determining that the index value of the to-be-processed element in the one-dimensional array is less than the index value of the tail element, processing a corresponding number of to-be-processed elements corresponding to a vectorization unit from the one-dimensional array each time in a form corresponding to the determined vectorization type.

[0016] In some embodiments, based on the data type of the input tensor, determining the vectorization type of the one-dimensional array so as to determine the vectorization unit of the burst transfer of the one-dimensional array comprises one or more of the following: in response to determining that the data type of the input tensor is single-precision floating point number, determining that the vectorization type of the one-dimensional array is four-channel single-precision floating point number, and the vectorization unit is 4; in response to determining that the data type of the input tensor is half-precision floating point number, determining that the vectorization type of the one-dimensional array is two-channel half-precision floating point number, and the vectorization unit is 2; in response to determining that the data type of the input tensor is brain floating point number 16, determining that the vectorization type of the one-dimensional array is two-channel brain floating point number 16, and the vectorization unit is 2; in response to determining that the data type of the input tensor is 32-bit integer, determining that the vectorization type of the one-dimensional array is four-channel integer, and the vectorization unit is 4; in response to determining that the data type of the input tensor is 8-bit integer, determining that the vectorization type of the one-dimensional array is four-channel integer, and the vectorization unit is 16; in response to determining that the data type of the input tensor is 8-bit unsigned integer, determining that the vectorization type of the one-dimensional array is four-channel unsigned integer, and the vectorization unit is 16.

[0017] The present application can unify the processing logic for input tensors, which is conducive to reducing code complexity and maintenance cost, by converting the obtained input tensor associated with the current operator into a one-dimensional array, and determining the non-tail elements and the tail elements of the one-dimensional array based on the vectorization unit and the total number of elements of the one-dimensional array, processing the non-tail elements of the one-dimensional array in a way of vectorized continuous memory access, and processing the tail elements of the one-dimensional array in a way of element-by-element memory access. In addition, even for the non-aligned scenario, most of the data (non-tail elements) of the input tensor can be processed in a way of vectorized continuous memory access, and only a small amount of tail elements are processed in a way of element-by-element memory access. Therefore, the present application can significantly improve the processing efficiency of tensors in a non-aligned scenario, and reduce code complexity and maintenance cost.

[0018] It should be understood that the content described in this part is not intended to identify key or important features of the embodiments of the present application, nor to limit the scope of the present application. Other features of the present application will become apparent from the following description. BRIEF DESCRIPTION OF DRAWINGS

[0019] The above and other features, aspects and advantages of embodiments of the present application will become more apparent from the following detailed description when taken in conjunction with the accompanying drawings. In the drawings similar or common elements of the drawings are denoted by like reference numbers.

[0020] Figure 1 A schematic diagram of a computing device implementing a method for memory access of tensor data according to an embodiment of the present application is shown schematically.

[0021] Figure 2 A flowchart of a method for memory access of tensor data according to an embodiment of the present application is shown.

[0022] Figure 3 A schematic diagram of conversion of a tensor to a one-dimensional array according to some embodiments of the present application is shown.

[0023] Figure 4 A flowchart of a method of processing non-tail elements of a one-dimensional array in a way of vectorized contiguous memory access according to some embodiments of the present application is shown.

[0024] Figure 5 A flowchart of a method of thread and thread block allocation according to an embodiment of the present application is shown.

[0025] In the various drawings, like or common elements and components are denoted by like reference numbers. DETAILED DESCRIPTION

[0026] Preferred embodiments of the present application will be described herein below with reference to the accompanying drawings. While the preferred embodiments of the present application are shown in the drawings, it is understood that the present application can be embodied in various forms and should not be limited by the embodiments set forth herein. Rather, these embodiments are provided so that this disclosure will be thorough and complete, and will fully convey the scope of the present application to those skilled in the art.

[0027] The term "comprising" and variations thereof as used herein are intended to cover a non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements is not necessarily limited to those elements, but can include other elements not expressly listed or inherent to such process, method, article, or apparatus. The term "or" as used herein is intended to mean "and / or" unless otherwise indicated. The term "based on" means "based, at least in part, on" unless otherwise indicated. The terms "one example embodiment" and "an example embodiment" as used herein refer to at least one example embodiment. The term "another embodiment" as used herein refers to at least one additional embodiment. The terms "first," "second," and the like as used herein do not imply importance or significance to one or the other, but merely distinguish one element from another.

[0028] As described previously, the conventional method for memory access of tensor data has the disadvantage that in a non-aligned scenario, the processing efficiency for tensor data is low and the maintenance cost is high.

[0029] To at least partially address one or more of the aforementioned problems and other potential issues, exemplary embodiments of the present invention propose a method for memory access of tensor data. In this method, the acquired input tensor associated with the current operator is converted into a one-dimensional array; the non-tail elements and tail elements of the one-dimensional array are determined based on the vectorization unit and the total number of elements in the one-dimensional array; the non-tail elements of the one-dimensional array are processed using vectorized continuous memory access; and the tail elements of the one-dimensional array are processed using element-by-element memory access. This invention unifies the processing logic for input tensors, reducing code complexity and maintenance costs. Furthermore, even in scenarios involving non-alignment, processing most of the input tensor data (non-tail elements) using vectorized continuous memory access, while only performing element-by-element memory access on a small number of tail elements, significantly improves memory throughput. Therefore, the present invention can significantly improve the processing efficiency of tensors in non-aligned scenarios and reduce code complexity and maintenance costs.

[0030] Figure 1 A schematic diagram of a computing device 100 implementing a method for accessing memory for tensor data according to an embodiment of the present invention is shown. Figure 1 As shown, the computing device 100 may have one or more processing units and a memory communicatively connected to one or more processors; wherein the memory stores instructions executable by at least one processor. The processing units include, for example, dedicated processing units such as Graphics Processing Units (GPUs), Field Programmable Gate Arrays (FPGAs), Application Specific Integrated Circuits (ASICs), General-purpose computing on graphics processing units (GPGPUs), and general-purpose processing units such as CPUs. The computing device 100 also includes at least: a one-dimensional array transformation unit 102, a non-tail element and tail element determination unit 104, a vectorized contiguous memory access unit 106, and an element-by-element memory access unit 108.

[0031] Regarding the one-dimensional array conversion unit 102, it is used to convert the acquired input tensor associated with the current operator into a one-dimensional array.

[0032] Regarding the non-tail element and tail element determination unit 104, it is used to determine the non-tail element and tail element of the one-dimensional array based on the vectorization unit and the total number of elements of the one-dimensional array.

[0033] Regarding the vectorized contiguous memory access unit 106, it is used to process the non-tail elements of the one-dimensional array in a vectorized contiguous memory access manner.

[0034] Regarding the element-by-element memory access unit 108, it is used to process the tail element of the one-dimensional array in an element-by-element memory access manner.

[0035] The following will combine Figure 2 and Figure 3 A method 200 for accessing memory for tensor data, according to an embodiment of the present invention, is described. It should be understood that method 200 can, for example, be implemented in... Figure 1 The described computing device 100 performs the operation. Method 200 may also include additional actions not shown and / or the actions shown may be omitted; the scope of the invention is not limited in this respect.

[0036] At step 202, the computing device 100 converts the acquired input tensor associated with the current operator into a one-dimensional array.

[0037] For example, in some embodiments, the computing device 100 converts the input tensor into a one-dimensional array if it determines that any of the following conditions are met: the current operator does not involve accumulation or reduction operations across elements, and each output element depends only on the input element at the corresponding position; the current operator is an element-wise operator included in a predetermined set of operators.

[0038] Pointwise operators, also known as element-wise operators, do not involve cross-element accumulation or reduction operations, and each output element depends only on the corresponding input element. For example, each output element processed by a pointwise operator is an operator that is independently computed by one or more input elements at the corresponding position, without cross-element dependencies, such as y[i] = f(x[i]) or y[i] = x1[i] + x2[i]. For example, pointwise operators include tensor addition operators (such as the operator named "ADD" in Table 1), tensor subtraction operators (such as the operator named "SUB" in Table 1), and so on. Therefore, regardless of the dimensionality of the input tensor (e.g., an input tensor with shape [B, C, H, W]), it can be losslessly converted into a one-dimensional array. The processing of the input tensor is achieved by uniformly processing the converted one-dimensional array in a one-dimensional index space.

[0039] The predetermined set of operators includes multiple element-wise operators. Table 1 below illustrates the multiple element-wise operators included in the predetermined set of operators. It should be understood that Table 1 only illustrates some element-wise operators, and it should be understood that other operators or other element-wise operators may be applicable to this invention.

[0040] Table 1

[0041] A method for converting an input tensor into a one-dimensional array includes, for example, the computing device 100 flattening the input tensor into a one-dimensional array. For example, an input tensor of arbitrary dimensions (such as shape = [d0, d1, ..., d...)) can be flattened into a one-dimensional array. k-1 Flatten it directly into a logical one-dimensional array. Please refer to [link / reference]. Figure 3 The shape of the input tensor 310 is (2,3), meaning the original shape of the input tensor 310 is 2. A matrix of size 3, with 6 elements, can map the input tensor 310 to a one-dimensional array 312 in row-major order, where the logical order of elements in the one-dimensional array 312 and the input tensor 310 remains unchanged. For example, if the logical order of elements in the input tensor 310 is 0, 1, 2, 3, 4, 5, the logical order of elements in the one-dimensional array 312 is also 0, 1, 2, 3, 4, 5. Furthermore, since the shape of the input tensor 310 is known to be (2,3), the one-dimensional array 312 can also be restored to the input tensor 310, making this a lossless conversion. Therefore, the above scheme can losslessly convert an input tensor into a one-dimensional array, and the resulting one-dimensional array has the same logical order of elements as the input tensor. Thus, efficient processing of the input tensor can be achieved by uniformly processing the converted one-dimensional array in a one-dimensional index space.

[0042] At step 204, the computing device 100 determines the non-tail element and tail element of the one-dimensional array based on the vectorized unit and the total number of elements in the one-dimensional array.

[0043] The calculation of the total number of elements in a one-dimensional array is based on an input tensor of arbitrary dimensions (e.g., shape = [d0, d1, ..., d...)). k-1 For example, the following formula (1) illustrates the algorithm for the total number of elements in a one-dimensional array.

[0044] N=∏ i d i (1) In formula (1) above, N represents the total number of elements in the input tensor. It should be understood that the total number of elements N in the tensor is equal to the number of elements in the transformed one-dimensional array (or the length of the one-dimensional array); i represents the i-th dimension of the input tensor, and d... i ∏ represents the length of the input tensor in the i-th dimension, and ∏ represents the product.

[0045] For example, if the shape of the input tensor is [B, C, H, W], the input tensor can be transformed into a one-dimensional array with a total number of elements N = B. C H W. Please refer to... Figure 3 The shape of input tensor 310 is (2,3), and the total number of elements in input tensor 310 is N=2. 3, meaning N=6, the total number of elements in the one-dimensional array 312 is 6. It should be understood that the shape and size of the input tensor can be very large. Figure 3 The shape and total number of elements of the input tensor 310 are merely examples of the method of the present invention.

[0046] Regarding vectorization, it refers to processing multiple data elements at once. Regarding the vectorization unit, it indicates, for example, how many data elements are read / written at a time when processing data using vectorization. For example, a vectorization unit of 2 indicates reading / writing 2 data elements at a time; a vectorization unit of 4 indicates reading / writing 4 data elements at a time. It should be understood that vectorization units include, for example, 2, 4, 8, 16, etc. The invention does not limit this.

[0047] Regarding the method for determining the vectorization unit, in some embodiments, it is, for example, that the computing device 100 determines the vectorization type of the one-dimensional array based on the data type of the input tensor, in order to determine the vectorization unit for the one-dimensional array.

[0048] Regarding the data type of the input tensor, it includes, for example, at least one or more of the following data types: single-precision floating-point number (Float), half-precision floating-point number (Half), Bfloat16, 32-bit integer (Int32), 8-bit integer (Int8), 8-bit unsigned integer (Unit8), etc.

[0049] Regarding vectorization types, they include at least one or more of the following vectorization types: four-channel single-precision floating-point number (Float), two-channel half-precision floating-point number (Half), two-channel Bfloat16, four-channel integer (Int4), four-channel unsigned integer (Unit4), etc.

[0050] In some embodiments, the determination of the vectorization unit is related to the data type of the input tensor.

[0051] For example, in some embodiments, the computing device 100 can determine the vectorization type of the one-dimensional array based on the data type of the input tensor, so as to determine the vectorization unit. For example, if the computing device 100 determines that the data type of the input tensor is single-precision floating-point number (Float), it determines that the vectorization type of the one-dimensional array is four-channel single-precision floating-point number (Float4) and the vectorization unit is 4.

[0052] For example, if the computing device 100 determines that the data type of the input tensor is half-precision floating-point (Half), and determines that the vectorization type of the one-dimensional array is dual-channel half-precision floating-point (Half2) and the vectorization unit is 2.

[0053] In some embodiments, if the computing device 100 determines that the data type of the input tensor is Bfloat16, it determines that the vectorization type of the one-dimensional array is dual-channel Bfloat162 and the vectorization unit is 2.

[0054] In some embodiments, if the computing device 100 determines that the data type of the input tensor is a 32-bit integer (Int32), it determines that the vectorization type of the one-dimensional array is a four-channel integer (Int4) and the vectorization unit is 4.

[0055] In some embodiments, if the computing device 100 determines that the data type of the input tensor is an 8-bit integer (Int8), it determines that the vectorization type of the one-dimensional array is a four-channel integer (Int4) and the vectorization unit is 16.

[0056] In some embodiments, if the computing device 100 determines that the data type of the input tensor is an 8-bit unsigned integer, it determines that the vectorization type of the one-dimensional array is a four-channel unsigned integer and the vectorization unit is 16.

[0057] For ease of understanding, the following table, Table 2, illustrates the relationship between the data types, vectorization types, and vectorization units of at least some of the input tensors.

[0058] Table 2

[0059] Therefore, this invention can determine the matching vectorization unit for tensor data of different data types used by the operator, which is beneficial to further improve memory access efficiency.

[0060] A method for determining the non-tail elements and tail elements of a one-dimensional array includes, for example, dividing the total number of elements in the one-dimensional array by a vectorized unit to obtain a remainder, and determining the tail elements of the one-dimensional array based on the obtained remainder. Elements other than the tail elements are non-tail elements.

[0061] For example, the total number of elements in the input tensor is N. The vectorization unit is B. Divide the total number of elements N of the one-dimensional array by the vectorization unit B, and the remainder is r, r = N%B. Based on this remainder r, determine the number of tail elements of the one-dimensional array: tail = N%B.

[0062] For example, please refer to Figure 3 The input tensor 310 has a total of 6 elements, its vectorization type is float4, and its vectorization unit is 4. The tail element of the one-dimensional array 312 is tail = 6%4. That is, the tail element of the one-dimensional array 312 is 2. Please refer to [reference needed]. Figure 3 The last two elements of the one-dimensional array 312 are 4 and 5, and the non-last two elements of the one-dimensional array 312 are 0, 1, 2, and 3.

[0063] At step 206, the computing device 100 processes the non-tail elements of the one-dimensional array in a vectorized contiguous memory access manner.

[0064] For example, in some embodiments, processing the non-tail elements of the one-dimensional array using vectorized contiguous memory access includes: for the non-tail elements, reading or writing the number of elements corresponding to the vectorized unit at once in the form corresponding to the vectorization type. For example, if the vectorization type is Float4 and the vectorization unit is 4, for the non-tail elements of the one-dimensional array, reading or writing 4 elements at once in Float4 form. Therefore, the above scheme, by using vectorization and contiguous memory access to process the non-tail elements of the one-dimensional array, can greatly improve processing efficiency and maximize bandwidth utilization for the non-tail data of the one-dimensional array converted from tensor data.

[0065] Regarding vectorized contiguous memory access methods, such as burst transfers, burst transfers can utilize the GPU's vectorized load / store instructions (such as ld.global.v4.f32) to read and write multiple contiguous data items at once. Please refer to [link / reference needed]. Figure 3 In a one-dimensional array 312, the non-tail elements are 0, 1, 2, and 3. If the vectorization unit is 4, sequential memory access allows reading and writing all four elements (0, 1, 2, and 3) at once. Alternatively, if the vectorization unit is 2, sequential memory access allows reading and writing two elements from 0 to 3 at a time, thus performing sequential memory access on the four non-tail elements (0, 1, 2, and 3). The following will combine... Figure 4The method 400 for handling non-tail elements of a one-dimensional array using vectorized contiguous memory access is explained in detail and will not be repeated here.

[0066] At step 208, the computing device 100 processes the tail element of the one-dimensional array in an element-by-element memory access manner.

[0067] For example, please refer to Figure 3 The last two elements of the one-dimensional array 312 are 4 and 5. The last two elements 4 and 5 of the one-dimensional array 312 are processed one by one in scalar mode.

[0068] It should be understood that, based on the above scheme, the input tensor is transformed into a one-dimensional array, and based on the vectorization unit and the total number of elements in the one-dimensional array, the elements of the obtained one-dimensional data are distinguished into non-tail elements and tail elements. The number of tail elements is obviously less than the number of vectorization units, and the number of tail elements is also less than the number of non-tail elements. For example, in the case of vectorization unit B=4, the number of tail elements is at most 3. In the case of vectorization unit B=8, the number of tail elements is at most 7. For example, if N=83, B=4, and tail=3, then in this case, all 80 elements of the one-dimensional array are non-tail elements, which can be processed by vectorized continuous memory access. Only the last 3 elements need to be processed one by one by scalar.

[0069] Test data shows that compared to traditional memory access methods for tensor data, this invention can achieve a significant improvement in memory access throughput. Specific experimental data is shown in Table 3 below.

[0070] Table 3

[0071] In the above method, by converting the acquired input tensor associated with the current operator into a one-dimensional array; and determining the non-tail elements and tail elements of the one-dimensional array based on the vectorization unit and the total number of elements in the one-dimensional array; processing the non-tail elements of the one-dimensional array using vectorized continuous memory access; and processing the tail elements of the one-dimensional array using element-by-element memory access, this invention can unify the processing logic for input tensors, which is beneficial for reducing code complexity and maintenance costs. Furthermore, even in scenarios involving non-alignment, most of the data (non-tail elements) of the input tensor can be processed using vectorized continuous memory access, with only a small number of tail elements requiring element-by-element memory access. Therefore, this invention can significantly improve the processing efficiency of tensors in non-aligned scenarios and reduce code complexity and maintenance costs.

[0072] Furthermore, this invention simplifies kernel function design and is compatible with existing memory allocation mechanisms, requiring no special memory layout. Additionally, this invention is applicable to pointwise operators (which include various operators) and to various data types of tensor data.

[0073] The following will combine Figure 3 and Figure 4 This describes a method 400 for processing the non-tail elements of a one-dimensional array using a vectorized contiguous memory access method, according to embodiments of the present invention. It should be understood that method 400 can, for example, be used in... Figure 1 The described computing device 100 performs the operation. Method 400 may also include additional actions not shown and / or the actions shown may be omitted; the scope of the invention is not limited in this respect.

[0074] At step 402, the computing device 100 determines the indices of all elements of the one-dimensional array.

[0075] At step 404, if the computing device 100 determines that the index value of the element to be processed in the one-dimensional array is less than the index value of the tail element, it processes the number of elements to be processed from the one-dimensional array each time, in the form corresponding to the determined vectorization type.

[0076] The index of the element to be processed is the distance of the current element relative to the starting element in the one-dimensional array. It should be understood that the index of the first element is "0".

[0077] The index of the last element is the index of the first last element. For example, the one-dimensional array [3,2,5,6,7,22,0,4,5,1] has a total of 10 elements. The vectorization unit is 4. The vectorization type is Float4. Its last element is [5,1], and the index of the last element is 8. Regarding the element to be processed [3,2,5,6,7,22,0,4], its index value is less than the index value of the last element "8". Therefore, the computing device 100 processes the number of elements to be processed (e.g., 4) corresponding to the vectorization unit from the one-dimensional array in the form corresponding to the determined vectorization type (e.g., Float4). Regarding the element to be processed [5,1], its index value is greater than or equal to the index value of the last element "8". The computing device 100 processes the element [5,1] in an element-by-element memory access manner.

[0078] It should be understood that, since access is usually continuous starting from index 0, and GPU global memory allocation usually guarantees alignment from the starting address to 256 bytes, the non-tail element region from the starting position (index 0) to the tail element can be processed by vectorized continuous memory access.

[0079] In the above scheme, the input tensor is transformed into a one-dimensional array without changing the logical order of the input tensor data. A unified processing logic is adopted, reducing the maintenance difficulty of the kernel function code. As a result, efficient processing can be achieved for tensor data of arbitrary shapes involved in element-wise operators (including unaligned tensor data). In addition, there is no need to maintain multiple sets of kernel function algorithms, eliminating aligned / unaligned code branches, significantly reducing code complexity and maintenance difficulty.

[0080] The following will combine Figure 3 and Figure 5 This describes a method 500 for allocating threads and thread blocks according to embodiments of the present invention. It should be understood that method 500 can, for example, be implemented in... Figure 1 The described computing device 100 performs the operation. Method 500 may also include additional actions not shown and / or the actions shown may be omitted; the scope of the invention is not limited in this respect.

[0081] At step 502, the computing device 100 allocates threads and thread blocks for processing the one-dimensional array by the current operator based on the total number of elements in the one-dimensional array, such that each thread and thread block is used to process a continuous index range of the one-dimensional array.

[0082] For example, the default number of threads in computing device 100 is 128 or 256. The number of thread blocks allocated is equal to the quotient of the total number of elements in the one-dimensional array divided by the default number of threads. For example, if the default number of threads is thread_num=128, the total number of elements in the one-dimensional array is N=250, the number of thread blocks allocated is block_num=250 / 128=2, and the number of threads allocated is 250.

[0083] Regarding the specific method of making each thread block or thread process a continuous index range of the one-dimensional array, for example, if the total number of elements in the one-dimensional array is N=250, allocate 2 thread blocks, and 250 threads process these 250 elements, with one thread allocated to each element. For example, thread block 0 processes the elements in the index range 0-127 of the one-dimensional array, and thread block 1 processes the elements in the index range 128-249 of the one-dimensional array.

[0084] At step 504, the computing device 100 causes the allocated threads and thread blocks to process the full data of the one-dimensional array in parallel.

[0085] For example, continuing the above example, thread block 0 and thread block 1, along with the 250 threads in the two thread blocks, process their respective elements in a one-dimensional array (250 elements) in parallel. This optimizes the ratio of the number of threads to the number of thread blocks, and all threads and thread blocks process the full data of the one-dimensional array transformed from the input tensor in parallel, maximizing the high-performance advantages of the GPU or GPGPU and greatly improving data processing efficiency.

[0086] Additionally, in some embodiments, the computing device 100 aligns the upper bound of the allocated thread with the starting index of the element of the one-dimensional array; aligns the lower bound of the allocated thread with the ending index of the element of the one-dimensional array; and causes the allocated thread to process the one-dimensional array within the index range of the element of the one-dimensional array. For example, please refer to... Figure 3 The last data indices of the one-dimensional array 312 are 4 and 5. The array has a total of 6 elements, which are allocated to 6 threads. By aligning the upper and lower bounds of the threads, we can prevent a thread from reading the 7th element. This ensures that the indices of the data elements in the one-dimensional array are aligned with the upper and lower bounds of the allocated threads, thus preventing out-of-bounds data readings or thread count errors.

[0087] The various processes and procedures described above, such as methods 200, 400, and 500, can be executed at a computing device. This computing device may include, for example, at least one processor (at least one graphics processor and at least one central processing unit); and a memory communicatively connected to the at least one processor; wherein the memory stores instructions executable by the at least one processor. In some embodiments, methods 200, 400, and 500 may be implemented as a computer software program or program product tangibly contained in a machine-readable medium. In some embodiments, part or all of the computer program may be loaded and / or installed on the computing device via read-only memory (ROM) and / or a communication unit. When the computer program is loaded into random-access memory (RAM) and executed by the GPU and CPU, one or more actions of methods 200, 400, and 500 described above can be performed.

[0088] This invention can be a method, apparatus, system, and / or computer program product. A computer program product may include a computer-readable storage medium having computer-readable program instructions loaded thereon for performing various aspects of the invention. The computer-readable storage medium may be a tangible device capable of holding and storing instructions used by an instruction execution device. The computer-readable storage medium may be, for example, but not limited to, an electrical storage device, a magnetic storage device, an optical storage device, an electromagnetic storage device, a semiconductor storage device, or any suitable combination thereof.

[0089] The computer-readable program instructions described herein can be downloaded from a computer-readable storage medium to various computing / processing devices, or downloaded via a network, such as the Internet, a local area network, a wide area network, and / or a wireless network, to an external computer or external storage device. Various aspects of the invention are described herein with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems), and computer program products according to embodiments of the invention. It should be understood that each block of the flowchart illustrations and / or block diagrams, and combinations of blocks in the flowchart illustrations and / or block diagrams, can be implemented by computer-readable program instructions.

[0090] These computer-readable program instructions can be provided to the central processing unit of a general-purpose computer, a special-purpose computer, or other programmable data processing apparatus to produce a machine such that, when executed by the central processing unit of the computer or other programmable data processing apparatus, they create means for implementing the functions / actions specified in one or more blocks of a flowchart and / or block diagram. These computer-readable program instructions can also be stored in a computer-readable storage medium that causes a computer, programmable data processing apparatus, and / or other device to operate in a particular manner. Thus, the computer-readable medium storing the instructions comprises an article of manufacture that includes instructions for implementing aspects of the functions / actions specified in one or more blocks of a flowchart and / or block diagram.

[0091] The flowcharts and block diagrams in the accompanying drawings illustrate the architecture, functionality, and operation of possible implementations of systems, methods, and computer program products according to various embodiments of the present invention. In this regard, each block in a flowchart or block diagram may represent a module, segment, or portion of an instruction, which contains one or more executable instructions for implementing a specified logical function. In some alternative implementations, the functions marked in the blocks may occur in a different order than those shown in the drawings. For example, two consecutive blocks may actually be executed substantially in parallel, and they may sometimes be executed in reverse order, depending on the functions involved. It should also be noted that each block in the block diagrams and / or flowcharts, and combinations of blocks in the block diagrams and / or flowcharts, may be implemented using a dedicated hardware-based system that performs the specified function or action, or using a combination of dedicated hardware and computer instructions.

[0092] It should be understood that the various forms of processes shown above can be used to reorder, add, or delete steps. For example, the steps loaded in this application can be executed in parallel, sequentially, or in different orders, as long as the desired result of the technical solution disclosed in this application can be achieved, and this is not limited herein.

[0093] The specific embodiments described above do not constitute a limitation on the scope of protection of this application. Those skilled in the art should understand that various modifications, combinations, sub-combinations, and substitutions can be made according to design requirements and other factors.

Claims

1. A method for memory access of tensor data, characterized in that, The method comprises: converting the obtained input tensor associated with the current operator into a one-dimensional array; determining non-tail elements and tail elements of the one-dimensional array based on a vectorization unit, a total number of elements of the one-dimensional array; processing the non-tail elements of the one-dimensional array in a vectorized continuous memory access manner; and processing the tail elements of the one-dimensional array in an element-by-element memory access manner.

2. The method of claim 1, wherein, Further comprising: determining a vectorization type of the one-dimensional array based on a data type of the input tensor, so as to determine the vectorization unit with respect to the one-dimensional array.

3. The method of claim 1, wherein, Converting the input tensor into the one-dimensional array comprises: in response to determining that any one of the following conditions is met, converting the input tensor into the one-dimensional array: the current operator does not involve accumulation or reduction operations across elements, and each output element only depends on an input element at a corresponding position; and the current operator is an element-by-element operator included in a predetermined operator set.

4. The method of claim 2, wherein, Processing the non-tail elements of the one-dimensional array in a vectorized continuous memory access manner comprises: for the non-tail elements, reading or writing elements corresponding to the vectorization unit at a time in a form corresponding to the vectorization type.

5. The method of claim 1, wherein, Determining the non-tail elements and the tail elements of the one-dimensional array comprises: dividing the total number of elements of the one-dimensional array by the vectorization unit to obtain a remainder of the division, and determining the tail elements of the one-dimensional array based on the obtained remainder.

6. The method of claim 3, wherein, Further comprising: allocating threads and thread blocks for processing the one-dimensional array by the current operator based on the total number of elements of the one-dimensional array, so that each thread and thread block is used to process a continuous index interval of the one-dimensional array; and causing the allocated threads and thread blocks to process the full amount of data of the one-dimensional array in parallel. Processing the non-tail elements of the one-dimensional array in a vectorized continuous memory access manner comprises:

7. The method of claim 1, wherein, determining an index of all elements of the one-dimensional array; and in response to determining that an index value of a to-be-processed element in the one-dimensional array is less than an index value of a tail element, processing the to-be-processed element from the one-dimensional array at a time in a form corresponding to the determined vectorization type, the to-be-processed element corresponding to the vectorization unit. Determining the vectorization type of the one-dimensional array based on the data type of the input tensor, so as to determine the vectorization unit with respect to the one-dimensional array comprises one or more of the following:

8. The method of claim 2, wherein, in response to determining that the data type of the input tensor is single-precision floating-point numbers, determining that the vectorization type of the one-dimensional array is four-channel single-precision floating-point numbers, and the vectorization unit is 4; in response to determining that the data type of the input tensor is half-precision floating-point numbers, determining that the vectorization type of the one-dimensional array is two-channel half-precision floating-point numbers, and the vectorization unit is 2; in response to determining that the data type of the input tensor is brain floating-point numbers 16, determining that the vectorization type of the one-dimensional array is two-channel brain floating-point numbers 16, and the vectorization unit is 2; in response to determining that the data type of the input tensor is 32-bit integers, determining that the vectorization type of the one-dimensional array is four-channel integers, and the vectorization unit is 4; in response to determining that the data type of the input tensor is 8-bit integers, determining that the vectorization type of the one-dimensional array is four-channel integers, and the vectorization unit is 16; ​ In response to determining that the data type of the input tensor is 8-bit unsigned integer, determining that the vectorization type of the one-dimensional array is four-channel unsigned integer, and the vectorization unit is 16.

9. A computing device, comprising: Comprising: at least one processor; and a memory connected with the at least one processor in communication; wherein the memory stores instructions executable by the at least one processor, the instructions being executed by the at least one processor to enable the at least one processor to perform the method of any one of claims 1-8.

10. A computer-readable storage medium, characterized in that, A computer program is stored on the computer readable storage medium, the computer program being executed by a machine to perform the method of any one of claims 1-8.

11. A computer program product, characterised in that, A computer program is stored on the computer readable storage medium, the computer program being executed by a machine to perform the method of any one of claims 1-8.

Images