Processor operation acceleration method and device based on vectorization and memory alignment, equipment and medium

Abstract

The invention discloses a processor operation acceleration method and device based on vectorization and memory alignment, equipment and a medium, and relates to the technical field of data computation.The method comprises the steps that whether a storage structure of to-be-processed data in a preset memory meets a preset memory alignment condition or not is judged; if the to-be-processed data meets the preset memory alignment condition, determining a vectorization width corresponding to a preset kernel function in a unified computing device architecture; and packaging each data element in the to-be-processed data into a corresponding data vector according to the vectorization width through the preset kernel function, and carrying out vectorization operation on the data vector on a preset processor so as to complete a data operation task corresponding to the to-be-processed data. Based on the unified computing device architecture, the vectorization and memory alignment technology is combined, data can be packaged into vectors to be processed, and the computing performance and the resource utilization rate of programs in the unified computing device architecture are remarkably improved.

Description

Technical Field

[0001] The present invention relates to the field of data computing technology, and in particular to a processor operation acceleration method, device, equipment and medium based on vectorization and memory alignment. Background Art

[0002] With the rapid development of deep learning technology, GPU (Graphics Processing Unit) has become the key hardware for performing large-scale computing tasks. However, traditional computing methods on GPU often face problems such as high memory access latency and low computing efficiency. In order to solve these problems, CUDA (Compute Unified Device Architecture) came into being, allowing developers to use the parallel computing capabilities of GPU for efficient computing. However, relying solely on the parallelism of CUDA is not enough to fully tap the potential of GPU.

[0003] Currently, the computational efficiency can be improved by optimizing the algorithm design and reducing unnecessary computation and memory access. For example, a more efficient algorithm can be used to replace the original algorithm, or the computational complexity can be reduced by algorithm reconstruction. However, this is difficult to develop and requires developers to have an in-depth understanding of the CUDA programming model, GPU hardware characteristics, and data structures, which increases the complexity and difficulty of development and may lead to poor optimization results or introduce new errors. In addition, with the continuous upgrading of GPU hardware and the update of CUDA versions, the optimization algorithm may need to be continuously adjusted and optimized to adapt to the new hardware and software environment, increasing the cost and complexity of maintenance.

[0004] It can be seen that how to further improve the computing efficiency of GPU is a problem that technical personnel in this field need to solve. Summary of the invention

[0005] The purpose of the embodiments of the present invention is to provide a method, device, equipment and medium for accelerating processor operations based on vectorization and memory alignment, which can solve the problem that the computing efficiency is not high enough when only relying on the parallelism of CUDA for data calculation.

[0006] To solve the above technical problems, in a first aspect, an embodiment of the present invention provides a processor operation acceleration method based on vectorization and memory alignment, comprising:

[0007] Determine whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment condition;

[0008] If the data to be processed meets the preset memory alignment condition, determining the vectorization width corresponding to the preset kernel function in the unified computing device architecture;

[0009] Through a preset kernel function, each data element in the data to be processed is packed into a corresponding data vector according to the vectorization width, and the data vector is vectorized on a preset processor to complete the data calculation task corresponding to the data to be processed.

[0010] In some embodiments, determining whether the storage structure of the to-be-processed data in the preset memory meets the preset memory alignment condition includes:

[0011] According to the target pointer in the data to be processed, it is determined whether the starting position of the data to be processed in the preset memory is an integer multiple of the preset number of bytes.

[0012] In some embodiments, each data element in the data to be processed is packed into a corresponding data vector according to the vectorization width by a preset kernel function, including:

[0013] Determine the data type of the data to be processed, and judge whether the data to be processed meets the preset data vectorization conditions according to the data type;

[0014] If the data to be processed meets the preset data vectorization conditions, each data element in the data to be processed is packed into a corresponding data vector according to the vectorization width through a preset kernel function.

[0015] In some embodiments, before packing each data element in the to-be-processed data into a corresponding data vector according to the vectorization width by using a preset kernel function, the method further includes:

[0016] A target structure is constructed based on preset data alignment instructions, data type conversion template classes, operator overloading instructions, and preset code compilation instructions, so that each data element in the data to be processed is packed into a corresponding data vector according to the vectorization width by using the target structure through a preset kernel function.

[0017] In some embodiments, each data element in the data to be processed is packed into a corresponding data vector according to the vectorization width by a preset kernel function, and the data vector is vectorized on a preset processor to complete the data operation task corresponding to the data to be processed, including:

[0018] If the data to be processed meets the preset memory alignment condition and the preset data vectorization condition, compile the preset kernel function;

[0019] Using the compiled preset kernel function to determine the memory area corresponding to the data to be processed through the preset type qualifier;

[0020] Using the preset mandatory type conversion symbol corresponding to the data type conversion template class of the target structure in the memory area, the data type corresponding to the data to be processed is converted into the data type corresponding to the target structure, and the vectorization width corresponding to the data to be processed is determined;

[0021] Packing each data element in the data to be processed into a corresponding data vector based on the vectorization width;

[0022] Determine the global index of the current thread of the unified computing device architecture, and determine whether the data to be processed meets the preset data calculation condition according to the global index and the amount of data to be processed;

[0023] If the data to be processed meets the preset data calculation condition, the operator of the preset operation function is overloaded using the operator overload instruction of the target structure;

[0024] Based on the overloaded operator, the operation function in the preset kernel function is used to perform vectorized operations on the data vector on the preset processor to complete the data operation task corresponding to the data to be processed.

[0025] In some embodiments, before compiling the preset kernel function, the method further includes:

[0026] Add corresponding compilation conditions in the preset kernel function based on several data types and several vectorization widths corresponding to each data to be processed;

[0027] According to the data type and vectorization width corresponding to the current data to be processed, a corresponding target compilation condition is determined from each compilation condition, so as to compile a preset kernel function based on the target compilation condition.

[0028] In some embodiments, the processor operation acceleration method based on vectorization and memory alignment disclosed in the present invention further includes:

[0029] If the data to be processed does not meet the preset memory alignment condition, each data element in the data to be processed is read in sequence so as to directly execute the data calculation task corresponding to the data to be processed based on each data element.

[0030] In a second aspect, the present invention provides a processor operation acceleration device based on vectorization and memory alignment, comprising:

[0031] A condition judgment module, used to judge whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment condition;

[0032] A width determination module, used for determining a vectorization width corresponding to a preset kernel function in a unified computing device architecture if the data to be processed meets a preset memory alignment condition;

[0033] The data computing module is used to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width through a preset kernel function, and perform vectorization operations on the data vector on a preset processor to complete the data calculation task corresponding to the data to be processed.

[0034] In a third aspect, the present invention provides an electronic device, comprising:

[0035] Memory for storing computer programs;

[0036] A processor is used to execute a computer program to implement the steps of the aforementioned processor operation acceleration method based on vectorization and memory alignment.

[0037] In a fourth aspect, the present invention provides a computer-readable storage medium having a computer program stored thereon, which, when executed by a processor, implements the steps of the aforementioned processor operation acceleration method based on vectorization and memory alignment.

[0038] The present invention can determine whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment conditions. If the data to be processed meets the preset memory alignment conditions, the vectorization width corresponding to the preset kernel function in the unified computing device architecture is determined, and each data element in the data to be processed is packaged into a corresponding data vector according to the vectorization width through the preset kernel function, so as to perform corresponding data calculations based on each data vector.

[0039] The beneficial effects of the present invention are as follows: Based on the above technical scheme, the present invention provides a computing method based on the fusion of vectorization and memory alignment of a unified computing device architecture. By combining vectorization and memory alignment, the vectorization technology of the unified computing device architecture can be used to pack multiple data elements into a vector, which helps to significantly reduce the repeated execution of instructions, and through the judgment of memory alignment, it can ensure that the storage method of data in the memory conforms to the access mode of the graphics processor hardware architecture, reducing memory access latency and bandwidth bottlenecks, thereby significantly improving the computing performance and resource utilization of the unified computing device architecture program, meeting the needs of high-performance computing in the fields of big data and artificial intelligence. BRIEF DESCRIPTION OF THE DRAWINGS

[0040] In order to more clearly illustrate the embodiments of the present invention, the following briefly introduces the drawings required for use in the embodiments. Obviously, the drawings described below are only some embodiments of the present invention. For ordinary technicians in this field, other drawings can be obtained based on these drawings without paying creative work.

[0041] Figure 1 A flowchart of a processor operation acceleration method based on vectorization and memory alignment provided by an embodiment of the present invention;

[0042] Figure 2 A compatible processing flow chart of CUDA-based vectorization and memory alignment fusion provided in an embodiment of the present invention;

[0043] Figure 3 A flowchart of a specific processor operation acceleration method based on vectorization and memory alignment provided by an embodiment of the present invention;

[0044] Figure 4 A flow chart of a structure implementation provided by an embodiment of the present invention;

[0045] Figure 5 A schematic diagram of the structure of a processor operation acceleration device based on vectorization and memory alignment provided by an embodiment of the present invention;

[0046] Figure 6 A structural diagram of an electronic device provided in an embodiment of the present invention. DETAILED DESCRIPTION

[0047] The following will be combined with the drawings in the embodiments of the present invention to clearly and completely describe the technical solutions in the embodiments of the present invention. Obviously, the described embodiments are only part of the embodiments of the present invention, not all of the embodiments. Based on the embodiments of the present invention, all other embodiments obtained by ordinary technicians in this field without creative work are within the scope of protection of the present invention.

[0048] The terms "including" and "having" in the specification of the present invention and the above-mentioned drawings, as well as any variations thereof, are intended to cover non-exclusive inclusions. For example, a process, method, system, product or device including a series of steps or units is not limited to the listed steps or units, but may include steps or units that are not listed.

[0049] Current computing methods on GPUs often face problems such as high memory access latency and low computing efficiency. To solve these problems, developers can use the parallel computing capabilities of GPUs for efficient computing. However, relying solely on the parallelism of CUDA is not enough to fully tap the potential of the GPU. The present invention can combine vectorization and memory alignment technology to pack multiple data elements into a vector, and through memory alignment judgment, ensure that the storage method of data in the memory conforms to the access mode of the graphics processor hardware architecture, thereby significantly improving the computing performance and resource utilization of the unified computing device architecture program.

[0050] In order to enable those skilled in the art to better understand the scheme of the present invention, the present invention is further described in detail below in conjunction with the accompanying drawings and specific implementation methods.

[0051] See also Figure 1 As shown, the embodiment of the present application discloses a processor operation acceleration method based on vectorization and memory alignment, including:

[0052] Step S11: determining whether the storage structure of the to-be-processed data in the preset memory meets the preset memory alignment condition.

[0053] This embodiment discloses an accelerated computing method based on the fusion of vectorization and memory alignment of CUDA. By deeply exploring the parallel computing capabilities of the GPU and combining the two optimization methods of vectorization and memory alignment, the CUDA program can be significantly improved in computing performance and resource utilization, thereby meeting the needs of high-performance computing in the fields of big data and artificial intelligence. Therefore, the technical solution in this embodiment can be mainly applied to tasks related to data computing in the field of AI (Artificial Intelligence). At this time, the data to be processed is the data required for the data computing task received by the graphics processor. And it can be understood that this embodiment is not limited to the above-mentioned computing field, and can also be applied to other fields that need to process multidimensional data and complex data formats, such as scientific computing, image processing, signal processing, etc. Through similar implementation methods, the efficiency of data processing and computing in these fields can be expanded and optimized.

[0054] Specifically, Figure 2 As shown, after the operator call preprocessing is performed, this embodiment can determine whether the storage structure of the current data to be processed in the preset memory meets the preset memory alignment condition. Specifically, it can determine whether the starting position of the data to be processed in the preset memory is an integer multiple of the preset number of bytes according to the target pointer in the data to be processed. In a specific embodiment, the above-mentioned preset memory is a graphics processor, and the data to be processed can be the data received when performing a graphics processing task. As shown in the following code:

[0055] auto inp_ptr = reinterpret_cast<std::uintptr_t> (input.data_ptr());

[0056] auto out_ptr = reinterpret_cast<std::uintptr_t> (output.data_ptr());

[0057] bool ptrs_are_aligned =

[0058] inp_ptr % 16 == 0 && out_ptr % 16 == 0 ;

[0059] if (ptrs_are_aligned) {

[0060] kernel<scalar, width=8, op> ;

[0061] } else {

[0062] kernel<scalar, width=0, op> ;

[0063] }.

[0064] The above code can be used to check whether the two pointers (inp_ptr and out_ptr) corresponding to the data to be processed are aligned to 16 bytes, and then the execution of the corresponding kernel function is determined based on the alignment result. It can be understood that memory alignment means that the address of the data in the memory is a multiple of a certain value. In a specific embodiment, the alignment requirement is 16 bytes, which means that the starting address of the data should be a multiple of 16. Memory alignment can optimize memory access performance or meet specific hardware requirements, and aligned data access can more effectively utilize the GPU's memory bandwidth, thereby reducing memory bandwidth consumption.

[0065] Step S12: If the data to be processed meets the preset memory alignment condition, then determine the vectorization width corresponding to the preset kernel function in the unified computing device architecture.

[0066] In this embodiment, if the data to be processed meets the above memory alignment conditions, the vectorization width corresponding to the preset kernel function in the unified computing device architecture is determined. It can be understood that the above vectorization width is the number of elements contained in each vector when the data is packed into a vector. The vectorization width satisfied when the current data to be processed is vectorized can be determined by the template parameter width in the CUDA kernel function kernel, so that the format of the data to be processed is converted into the format of the vector according to the determined vectorization width for subsequent tasks.

[0067] Step S13: Packing each data element in the data to be processed into a corresponding data vector according to the vectorization width through a preset kernel function, and performing vectorization operations on the data vector on a preset processor to complete the data calculation task corresponding to the data to be processed.

[0068] After the vectorization width is determined in this embodiment, each data element in the data to be processed can be packed into a corresponding data vector according to the vectorization width through a preset kernel function, and the data vector can be vectorized on a preset processor to complete the data calculation task corresponding to the data to be processed. And when packing each data element in the data to be processed into a data vector, the data type of the data to be processed can also be determined, and whether the data to be processed meets the preset data vectorization conditions can be judged according to the data type. If the data to be processed meets the preset data vectorization conditions, each data element in the data to be processed can be packed into a corresponding data vector according to the vectorization width through a preset kernel function. In other words, Figure 2 As shown, if the current data to be processed meets the requirements of vectorization and the requirements of the data type for accelerated calculation, the data to be processed can be vectorized through the corresponding specialized version of the kernel function for the specific width and data type to complete the data calculation task corresponding to the data to be processed.

[0069] Accordingly, in another specific embodiment, before executing a computing task, it is possible to check whether the memory address of the input data is aligned to a specified byte boundary (such as 16 bytes). If the data is not aligned, additional steps are taken. That is, if the data to be processed does not meet the preset memory alignment condition, each data element in the data to be processed can be read in sequence, so as to directly execute the data computing task corresponding to the data to be processed based on each data element. Figure 2 As shown, the vectorized general version of the kernel function can be called to perform data computing tasks, as shown in the following code:

[0070] template<typename scalar_t, int width> ;

[0071] __global__ std::enable_if_t<(width == 0) || !_typeConvert<scalar_t> ::exists>;

[0072] kernel(...);

[0073] The above kernel function can handle situations where vectorization is not supported or the vectorization width is 0. The same operation as the specialized version can be performed, but the array is processed element by element instead of vectorized processing. When the data type of the data to be processed does not support vectorized operations or the developer wants to use non-vectorized code, this version of the kernel function can be used for processing. In this way, before executing the calculation task, it is checked whether the memory address of the input data is aligned to the specified byte boundary. If the data is not aligned, the vectorized calculation method is not used, thereby complying with the access mode of the GPU hardware architecture.

[0074] In this embodiment, it can be determined whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment condition. If the data to be processed meets the preset memory alignment condition, the vectorization width corresponding to the preset kernel function in the unified computing device architecture is determined, and the preset kernel function is used to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width, so as to perform corresponding data calculations based on each data vector. In this way, by combining vectorization with memory alignment, the vectorization technology of the unified computing device architecture can be used to pack multiple data elements into a vector, which helps to significantly reduce the repeated execution of instructions, and through the judgment of memory alignment, it can be ensured that the storage method of data in the memory meets the access mode of the graphics processor hardware architecture, reducing the number of memory accesses, delays and bandwidth bottlenecks, so that the unified computing device architecture program can be significantly improved in computing performance and resource utilization, meeting the needs of high-performance computing in the fields of big data and artificial intelligence.

[0075] Based on the previous embodiment, it can be seen that the present application can realize the operation acceleration of the processor by combining vectorization and memory alignment technology. Next, the process of performing related operations will be described in detail in this embodiment. Figure 3 As shown, the embodiment of the present application discloses a specific processor operation acceleration method based on vectorization and memory alignment, including:

[0076] Step S21: constructing a target structure based on preset data alignment instructions, data type conversion template class, operator overloading instructions and preset code compilation instructions.

[0077] In this embodiment, Figure 4 As shown, the target structure can be first constructed based on the preset data alignment instruction (alignas(16) instruction), the data type conversion template class (Convert), the operator overloading instruction, and the preset code compilation instruction (if constexpr and #pragma unroll), so that the target structure can be used by the preset kernel function to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width.

[0078] Specifically, as shown in the following code, you can predefine a template structure for data type conversion:

[0079] template<typename scalar_t>

[0080] struct _typeConvert {

[0081] using hip_type = scalar_t;

[0082] struct packed_hip_type {

[0083] scalar_t x, y;

[0084] __device__ packed_hip_type& operator+=(const packed_hip_type&other) {

[0085] x += other.x;

[0086] y += other.y;

[0087] return *this;

[0088] }

[0089] __device__ packed_hip_type& operator*=(const packed_hip_type&other) {

[0090] x *= other.x;

[0091] y *= other.y;

[0092] return *this;

[0093] }

[0094] };

[0095] };

[0096] The above code defines a template structure _typeConvert, which can accept a type parameter scalar_t, which is usually a numeric type, such as float, double, or other user-defined numeric types. This template structure can be used to implement some conversions and operations related to the scalar_t type. Specifically, an internal structure packed_hip_type can be defined to store two scalar_t type values ​​(x and y), and provide some basic arithmetic operations for these two values.

[0097] Next, you can define a template structure named _f16Vec, which is a vectorized data structure optimized for GPU computing, as shown in the following code:

[0098] template<typename scalar_t, int width>

[0099] struct alignas(16) _f16Vec {

[0100] static_assert(width > 0 && (width & (width - 1)) == 0, "Widthmust be a positive power of 2");

[0101] using Converter = _typeConvert<scalar_t> ;

[0102] using T1 = typename Converter::hip_type;

[0103] using T2 = typename Converter::packed_hip_type;

[0104] T1 data[width];

[0105] / / Use reinterpret_cast to convert part of the data array into an array of type T2 for operation;

[0106] __device__ _f16Vec& operator+=(const _f16Vec& other) {

[0107] if (width % 2 == 0) {

[0108] T2* myData = reinterpret_cast<T2*> (data);

[0109] const T2* otherData = reinterpret_cast<const T2*> (other.data);

[0110] for (int i = 0; i < width / 2; ++i) {

[0111] myData[i] += otherData[i];

[0112] }

[0113] } else {

[0114] for (int i = 0; i < width; ++i) {

[0115] data[i] += other.data[i];

[0116] }

[0117] }

[0118] return *this;

[0119] }

[0120] __device__ _f16Vec& operator*=(const _f16Vec& other) {

[0121] if (width % 2 == 0) {

[0122] T2* myData = reinterpret_cast<T2*>(data);

[0123] const T2* otherData = reinterpret_cast<const T2*>(other.data);

[0124] for (int i = 0; i < width / 2; ++i) {

[0125] myData[i] *= otherData[i];

[0126] }

[0127] } else {

[0128] for (int i = 0; i < width; ++i) {

[0129] data[i] *= other.data[i];

[0130] }

[0131] }

[0132] return *this;

[0133] }

[0134] };

[0135] Through this structure, such as Figure 4As shown. You can improve computing efficiency by packing multiple data elements (scalar_t type) into a vector and taking advantage of the parallel computing capabilities of the GPU. The size of the structure (that is, the number of packed data elements) is specified by the template parameter width, and width is required to be a positive power of 2 to ensure data alignment and efficiency of vectorized operations. This structure is mainly used to implement the following functions:

[0136] Data alignment: The alignas(16) directive forces the structure to be aligned to 16-byte boundaries. This conforms to the memory access patterns of many GPU hardware architectures, helping to reduce memory access latency and bandwidth bottlenecks.

[0137] Type conversion: A template class called Convert is used inside the structure to perform data type conversion, allowing _f16Vec to support multiple basic data types (such as float, half, etc.), and use T1 and T2 type aliases to represent single and packaged data types;

[0138] Vectorized operations: The structure provides operator overloads such as +=, *= (multiply with another _f16Vec vector or scalar), which use vectorization techniques to reduce repeated execution of instructions. In particular, when width is an even number, these operations can try to further optimize performance by processing two data elements at a time (through float2 or similar packed types);

[0139] Compile-time optimization: By using the if constexpr and #pragma unroll directives, the code can be optimized based on the value of width at compile time, thus avoiding the runtime conditional checking and loop unrolling overhead.

[0140] That is to say, this embodiment combines data alignment, type conversion, vectorized operations and CUDA device code and other technologies to provide an efficient, flexible and optimizable vectorized data structure, which can be instantiated according to the specified data type (such as FP16, BF16, etc.) and width (i.e., the number of elements contained in the vector). In addition, the _f16Vec structure provides operator overloading such as addition and multiplication, and vectorized operations can be used as conveniently as basic data types in the GPU kernel, allowing users to perform vector operations conveniently. This design not only protects the core functions of _f16Vec, but also greatly improves its usability and ease of use. By packing multiple data elements into a vector through vectorization, this embodiment significantly reduces the repeated execution of instructions, thereby improving computing efficiency.

[0141] Step S22: using a preset kernel function and a target structure to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width, and performing vectorization operations on the data vector on a preset processor to complete the data calculation task corresponding to the data to be processed.

[0142] Correspondingly, based on the above structure, when a vectorized operation is performed on a data vector on a preset processor to complete a data calculation task corresponding to the data to be processed, specifically, when the data to be processed meets the preset memory alignment condition and the preset data vectorization condition, a preset kernel function can be compiled, and the compiled preset kernel function can be used to determine the memory area corresponding to the data to be processed through the preset type qualifier, so that the preset mandatory type conversion symbol corresponding to the data type conversion template class of the target structure is used in the memory area to convert the data type corresponding to the data to be processed into the data type corresponding to the target structure, and then the vectorization width corresponding to the data to be processed is determined, and then each data element in the data to be processed is packaged into a corresponding data vector based on the above vectorization width, and then the global index of the current thread of the unified computing device architecture is determined, and it is judged whether the data to be processed meets the preset data calculation condition based on the global index and the data amount of the data to be processed. If the data to be processed meets the preset data calculation condition, the operator of the preset operation function is overloaded by using the operator overload instruction of the target structure, and the operation function in the preset kernel function is used to perform a vectorized operation on the data vector on the preset processor based on the overloaded operator to complete the data calculation task corresponding to the data to be processed. That is to say, this embodiment can dynamically select the appropriate computing strategy according to the data type and size through the built-in efficient CUDA kernel optimization module. In particular, this embodiment performs deep optimization for specific data types (such as FP16 / BF16) and takes into account the memory alignment status to ensure the efficiency and performance stability of the kernel implementation, providing users with a simple, easy-to-use and effective CUDA computing experience.

[0143] like Figure 2 As shown, two CUDA kernel functions are implemented in this embodiment to perform an operation. This operation can be written specifically as needed, and this kernel is designed to process different types of numerical data (such as FP16 / BF16, etc.). There are two specialized versions according to different data types. The following code is a specialized version for a specific width and data type:

[0144] template<typename scalar_t, int width, typename Op>

[0145] __global__ void kernel(

[0146] scalar_t* __restrict__ input,

[0147] scalar_t* __restrict__ output,

[0148] Op op / / Add operation function as parameter

[0149] ) {

[0150] auto* __restrict__ input_v = reinterpret_cast<_f16Vec<scalar_t,width> *>(input);

[0151] auto* __restrict__ output_v = reinterpret_cast<_f16Vec<scalar_t,width> *>(output);

[0152] / / Get the global index of the thread (assuming 1D grid and 1D block here)

[0153] int idx = threadIdx.x + blockIdx.x * blockDim.x;

[0154] if (idx < (width * (total number of blocks of input data, which needs to be calculated or passed in as a parameter))) {

[0155] / / Load the input vector

[0156] _f16Vec<scalar_t, width> in_vec = input_v[idx / width];

[0157] / / Apply the operation function (here we assume that op accepts an _f16Vec and returns an _f16Vec, or can be modified to accept a scalar and return a scalar)

[0158] _f16Vec<scalar_t, width> out_vec = op(in_vec);

[0159] / / Store the output vector

[0160] output_v[idx / width] = out_vec; / / Again, you may need to adjust

[0161] }

[0162] };

[0163] The above code mainly includes:

[0164] (1) Template parameters:

[0165] scalar_t: scalar type, indicating the type of input and output data;

[0166] width: an integer indicating the width of the vectorization, that is, the number of elements contained in each vector;

[0167] op: Add the operation function as a parameter.

[0168] By using the template parameters scalar_t, width, and op, the kernel function can handle multiple data types and vectorization widths and be applied to various operations, thereby improving the reusability of the code.

[0169] (2) Conditional compilation:

[0170] Using std::enable_if_t and the conditional expression (width > 0) && _typeConvert<scalar_t> ::exists ensures that this kernel function is only used when width is greater than 0 and the scalar_t type supports some conversion (defined by _typeConvert<scalar_t> ::exists check). In this way, through std::enable_if_t and conditional expressions, it is ensured that this kernel function is compiled only when specific conditions are met, avoiding unnecessary compilation errors and performance losses.

[0171] (3) Input and output:

[0172] scalar_t* __restrict__ input: A pointer to input data. The __restrict__ keyword tells the compiler that the memory area pointed to by this pointer will not be accessed by other pointers, which helps with optimization.

[0173] scalar_t* __restrict__ output: pointer to output data.

[0174] (4) Vectorization:

[0175] Convert the input and output pointers to _f16Vec via reinterpret_cast<scalar_t, width> * type. _f16Vec is the template class used for vectorization, which is used to represent a vector containing width scalar_t type elements. Through the _f16Vec class, the code can use the vectorization instruction set of the GPU to reduce the number of instructions and memory access times, thereby significantly improving the computing efficiency on the GPU.

[0176] (5) Calculation:

[0177] The code omits the specific calculation part (indicated by .......), which will use vectorized operations to perform some calculations on the input data and store the results in the output. In this way, the above code structure is clear and easy to expand. For example, more conditional compilation logic can be added to handle different data types or vectorization widths, or more input / output parameters can be added to support more complex calculations. That is to say, before compiling the preset kernel function, corresponding compilation conditions can be added to the preset kernel function based on several data types and several vectorization widths corresponding to each data to be processed, and the corresponding target compilation condition is determined from each compilation condition according to the corresponding data type and vectorization width of the current data to be processed, so as to compile the preset kernel function based on the target compilation condition. It should also be pointed out that when the data to be processed required by the current image processor does not meet the memory alignment conditions and / or the conditions for data calculation of the CUDA kernel function (the data calculation conditions include vectorization conditions and data type conditions), in this embodiment, it is also possible to choose to pre-process the data to be processed that does not meet the memory alignment so that the current data to be processed meets the corresponding memory alignment conditions, and when the data type of the data to be processed can be converted, the data type of the data to be processed is converted through the above-mentioned template structure _f16Vec to obtain the data type that meets the requirements of the current kernel function for data calculation. In this way, when some of the data to be processed does not meet the requirements of the specialized version of the kernel function for a specific width and data type, this embodiment can also choose to perform corresponding processing on the current data, thereby improving computing efficiency.

[0178] The present invention also provides a system based on the above method, which includes a data preprocessing module, a CUDA kernel execution module and a result post-processing module. The data preprocessing module is responsible for loading and aligning data; the CUDA kernel execution module is responsible for performing vectorized calculations; and the result post-processing module is responsible for transmitting the calculation results from the GPU back to the CPU and performing necessary subsequent processing. Through the above-mentioned accelerated computing method and system based on the fusion of CUDA vectorization and memory alignment, the efficiency of executing complex computing tasks (such as addition and root mean square normalization in deep learning) on ​​the GPU can be improved by combining vectorization technology and memory alignment technology.

[0179] Based on the above technical solution, in this embodiment, a template structure _f16Vec is first defined, which can be instantiated according to the specified data type and width. The structure provides operator overloading such as addition, multiplication and square sum, so that it can be easily used in the GPU kernel, thereby utilizing CUDA's vectorization technology to pack multiple data elements into a vector, so as to significantly reduce the repeated execution of instructions and improve computing efficiency. And after optimizing the CUDA kernel, according to the type and size of the data, the appropriate CUDA kernel can be dynamically selected for calculation. For example, when the data type is FP16 / BF16 and the data is aligned, the optimized kernel will be used to perform packing and vectorization operations to make full use of the hardware characteristics of the GPU (such as SIMD (Single Instruction Multiple Data, single instruction multiple data) instruction set, cache, etc.). And in this embodiment, it is not limited to specific computing tasks or operators. By defining the template structure _f16Vec and providing operator overloading, it can support vectorization operations of multiple data types and widths, thereby adapting to the needs of different algorithms and operators.

[0180] See also Figure 5 As shown, the embodiment of the present application discloses a processor operation acceleration device based on vectorization and memory alignment, the device comprising:

[0181] The condition judgment module 11 is used to judge whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment condition;

[0182] A width determination module 12, configured to determine a vectorization width corresponding to a preset kernel function in a unified computing device architecture if the data to be processed meets a preset memory alignment condition;

[0183] The data calculation module 13 is used to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width through a preset kernel function, and perform vectorization operations on the data vector on a preset processor to complete the data calculation task corresponding to the data to be processed.

[0184] Since the embodiments of the device part correspond to the above embodiments, please refer to the description of the embodiments of the method part for the embodiments of the device part, and will not be repeated here.

[0185] It can be seen that this embodiment can first determine whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment conditions. If the data to be processed meets the preset memory alignment conditions, the vectorization width corresponding to the preset kernel function in the unified computing device architecture is determined, and the preset kernel function is used to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width, so as to perform corresponding data calculations based on each data vector. In this way, by combining vectorization and memory alignment technology, the vectorization technology of the unified computing device architecture can be used to pack multiple data elements into a vector, which helps to significantly reduce the repeated execution of instructions, and through the judgment of memory alignment, it can ensure that the storage method of data in the memory meets the access mode of the graphics processor hardware architecture, reducing memory access delays and bandwidth bottlenecks, so that the unified computing device architecture program can be significantly improved in computing performance and resource utilization, meeting the needs of high-performance computing in the fields of big data and artificial intelligence.

[0186] Furthermore, the present application also discloses an electronic device. Figure 6 It is a structural diagram of an electronic device according to an exemplary embodiment, and the content in the figure cannot be regarded as any limitation on the scope of use of this application. The electronic device may specifically include: at least one processor 21, at least one memory 22, a power supply 23, a communication interface 24, an input and output interface 25 and a communication bus 26. Among them, the memory 22 is used to store a computer program, and the computer program is loaded and executed by the processor 21 to implement the relevant steps in the processor operation acceleration method based on vectorization and memory alignment disclosed in any of the aforementioned embodiments. In addition, the electronic device in this embodiment may specifically be an electronic computer.

[0187] In this embodiment, the power supply 23 is used to provide working voltage for each hardware device on the electronic device; the communication interface 24 can create a data transmission channel between the electronic device and the external device, and the communication protocol it follows is any communication protocol that can be applied to the technical solution of the present application, and is not specifically limited here; the input and output interface 25 is used to obtain external input data or output data to the outside world, and its specific interface type can be selected according to specific application needs and is not specifically limited here.

[0188] In addition, the memory 22, as a carrier for storing resources, can be a read-only memory, a random access memory, a disk or an optical disk, etc. The resources stored thereon can include an operating system 221, a computer program 222, etc., and the storage method can be temporary storage or permanent storage.

[0189] The operating system 221 is used to manage and control various hardware devices on the electronic device and the computer program 222, which can be Windows Server, Netware, Unix, Linux, etc. In addition to including a computer program that can be used to complete the processor operation acceleration method based on vectorization and memory alignment performed by the electronic device disclosed in any of the aforementioned embodiments, the computer program 222 can further include a computer program that can be used to complete other specific tasks.

[0190] Furthermore, the present application also discloses a computer-readable storage medium for storing a computer program; wherein, when the computer program is executed by a processor, the aforementioned processor operation acceleration method based on vectorization and memory alignment is implemented. The specific steps of the method can refer to the corresponding contents disclosed in the aforementioned embodiments, and will not be repeated here.

[0191] Furthermore, an embodiment of the present application also discloses a computer program product, including a computer program / instruction, which, when executed by a processor, implements the steps of the processor operation acceleration method based on vectorization and memory alignment disclosed in any of the aforementioned embodiments.

[0192] In this specification, each embodiment is described in a progressive manner, and each embodiment focuses on the differences from other embodiments. The same or similar parts between the embodiments can be referred to each other. For the device disclosed in the embodiment, since it corresponds to the method disclosed in the embodiment, the description is relatively simple, and the relevant parts can be referred to the method part.

[0193] Professionals may further appreciate that the units and algorithm steps of each example described in conjunction with the embodiments disclosed herein can be implemented in electronic hardware, computer software, or a combination of the two. In order to clearly illustrate the interchangeability of hardware and software, the composition and steps of each example have been generally described in the above description according to function. Whether these functions are performed in hardware or software depends on the specific application and design constraints of the technical solution. Professionals and technicians may use different methods to implement the described functions for each specific application, but such implementation should not be considered to be beyond the scope of this application.

[0194] The steps of the method or algorithm described in conjunction with the embodiments disclosed herein may be implemented directly using hardware, a software module executed by a processor, or a combination of the two. The software module may be placed in a random access memory (RAM), a memory, a read-only memory (ROM), an electrically programmable ROM, an electrically erasable programmable ROM, a register, a hard disk, a removable disk, a CD-ROM, or any other form of storage medium known in the art.

[0195] Finally, it should be noted that, in this article, relational terms such as first and second, etc. are only used to distinguish one entity or operation from another entity or operation, and do not necessarily require or imply any such actual relationship or order between these entities or operations. Moreover, the terms "include", "comprise" or any other variants thereof are intended to cover non-exclusive inclusion, so that a process, method, article or device including a series of elements includes not only those elements, but also other elements not explicitly listed, or also includes elements inherent to such process, method, article or device. In the absence of further restrictions, the elements defined by the sentence "comprise a ..." do not exclude the presence of other identical elements in the process, method, article or device including the elements.

[0196] The technical solution provided by the present application is introduced in detail above. Specific examples are used in this article to illustrate the principles and implementation methods of the present application. The description of the above embodiments is only used to help understand the method of the present application and its core idea. At the same time, for general technical personnel in this field, according to the idea of ​​the present application, there will be changes in the specific implementation method and application scope. In summary, the content of this specification should not be understood as a limitation on the present application.

Claims

1. A processor operation acceleration method based on vectorization and memory alignment, characterized in that: include: Determine whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment condition; If the data to be processed meets the preset memory alignment condition, determining a vectorization width corresponding to a preset kernel function in a unified computing device architecture; The preset kernel function is used to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width, and the data vector is vectorized on a preset processor to complete the data calculation task corresponding to the data to be processed.

2. The processor operation acceleration method based on vectorization and memory alignment according to claim 1, characterized in that: The determining whether the storage structure of the to-be-processed data in the preset memory meets the preset memory alignment condition includes: According to the target pointer in the data to be processed, it is determined whether the starting position of the data to be processed in the preset memory is an integer multiple of a preset number of bytes.

3. The processor operation acceleration method based on vectorization and memory alignment according to claim 1, characterized in that: The step of packing each data element in the to-be-processed data into a corresponding data vector according to the vectorization width by using the preset kernel function includes: Determine the data type of the data to be processed, and determine whether the data to be processed meets the preset data vectorization condition according to the data type; If the data to be processed meets the preset data vectorization condition, each of the data elements in the data to be processed is packed into the corresponding data vector according to the vectorization width through the preset kernel function.

4. The processor operation acceleration method based on vectorization and memory alignment according to claim 3 is characterized in that: Before packing each data element in the to-be-processed data into a corresponding data vector according to the vectorization width by the preset kernel function, the method further includes: A target structure is constructed based on preset data alignment instructions, data type conversion template classes, operator overloading instructions and preset code compilation instructions, so that each data element in the data to be processed is packed into a corresponding data vector according to the vectorization width by using the target structure through the preset kernel function.

5. The processor operation acceleration method based on vectorization and memory alignment according to claim 4 is characterized in that: The step of packing each data element in the data to be processed into a corresponding data vector according to the vectorization width by the preset kernel function, and performing a vectorization operation on the data vector on a preset processor to complete a data operation task corresponding to the data to be processed includes: If the data to be processed meets the preset memory alignment condition and the preset data vectorization condition, compiling the preset kernel function; Determine the memory area corresponding to the data to be processed by using the compiled preset kernel function through a preset type qualifier; Using a preset mandatory type conversion symbol corresponding to the data type conversion template class of the target structure in the memory area, the data type corresponding to the data to be processed is converted into the data type corresponding to the target structure, and the vectorization width corresponding to the data to be processed is determined; Packing each of the data elements in the data to be processed into a corresponding data vector based on the vectorization width; Determine a global index of a current thread of a unified computing device architecture, and determine whether the data to be processed meets a preset data calculation condition according to the global index and the amount of the data to be processed; If the data to be processed meets the preset data calculation condition, overloading the operator of the preset operation function by using the operator overloading instruction of the target structure; Based on the overloaded operator, the operation function in the preset kernel function is used to perform vectorized operations on the data vector on a preset processor to complete the data operation task corresponding to the data to be processed.

6. The processor operation acceleration method based on vectorization and memory alignment according to claim 5, characterized in that: Before compiling the preset kernel function, the method further includes: Adding corresponding compilation conditions in the preset kernel function based on a number of data types and a number of vectorization widths corresponding to each to-be-processed data; According to the data type and the vectorization width corresponding to the current data to be processed, a corresponding target compilation condition is determined from the various compilation conditions, so as to compile the preset kernel function based on the target compilation condition.

7. The processor operation acceleration method based on vectorization and memory alignment according to any one of claims 1 to 6, characterized in that: Also includes: If the data to be processed does not meet the preset memory alignment condition, each of the data elements in the data to be processed is read in sequence, so as to directly execute the data calculation task corresponding to the data to be processed based on each of the data elements.

8. A processor operation acceleration device based on vectorization and memory alignment, characterized in that: include: A condition judgment module, used to judge whether the storage structure of the data to be processed in the preset memory meets the preset memory alignment condition; A width determination module, configured to determine a vectorization width corresponding to a preset kernel function in a unified computing device architecture if the data to be processed meets the preset memory alignment condition; A data computing module is used to pack each data element in the data to be processed into a corresponding data vector according to the vectorization width through the preset kernel function, and perform vectorization operations on the data vector on a preset processor to complete the data calculation task corresponding to the data to be processed.

9. An electronic device, characterized in that: include: Memory for storing computer programs; A processor, used to execute the computer program to implement the steps of the processor operation acceleration method based on vectorization and memory alignment as described in any one of claims 1 to 7.

10. A computer-readable storage medium, characterized in that: The computer-readable storage medium stores a computer program, which, when executed by a processor, implements the steps of the processor operation acceleration method based on vectorization and memory alignment as described in any one of claims 1 to 7.

Images