Image processing method and device based on element-by-element operation, equipment and medium

Abstract

The invention provides an image processing method and device based on element-by-element operation, equipment and a medium, and relates to the technical field of GPUs (Graphics Processing Unit), and the method comprises the following steps: according to the number of a plurality of elements contained in a plurality of input tensors needing to participate in operation and hardware resource information of GPU equipment, determining the number of the plurality of elements contained in the plurality of input tensors needing to participate in operation; determining the number of thread blocks required to be used by the plurality of input tensors in the element-by-element operation process; according to the number of the thread blocks, the number of the elements and the memory layout of the multiple elements, kernel functions of element-by-element operation class operators corresponding to the multiple input tensors needing to participate in operation are determined on the basis of a predefined template structure body, and parameters of the template structure body comprise element types and the number of the elements; and based on the number of the thread blocks, initiating a plurality of threads to enable the plurality of threads to obtain the plurality of elements according to the kernel function, and executing an element-by-element operation process of the plurality of input tensors to obtain an image processing result.

Description

Technical Field

[0001] The present application relates to the field of GPU technology, and in particular to an image processing method, device, equipment and medium based on element-by-element operation. Background Art

[0002] Element-wise operation is an operation widely used in data processing and mathematical calculations. It refers to the operation of calculating the elements in corresponding positions in two or more data structures (such as arrays, vectors, matrices or higher-dimensional tensors) one by one. This operation does not rely on the overall relationship between the data structures, but acts on each element independently.

[0003] Element-wise operators have the following characteristics:

[0004] Independence: Operations are applied independently to each element in the data structure, without regard to the positional relationship between the elements or the overall structure.

[0005] Parallelism: Due to the independence of operations, element-by-element operations are very suitable for large-scale parallel computing, which can significantly improve computing efficiency.

[0006] Versatility: Element-wise operations can be applied to various mathematical and data processing operations such as addition, subtraction, multiplication, division, exponential, logarithm, etc.

[0007] In machine learning fields such as neural networks and deep learning, element-by-element operations are the basis for implementing various activation functions, loss functions, and optimization algorithms.

[0008] In the related art, a separate kernel function is written for each element-by-element operator, and elements are taken out from the input data one by one for calculation, and then written to the output. Although this method can complete the calculation function, it does not provide a unified execution framework for this type of operator, and the development efficiency is low. In addition, the architectural characteristics of the computing device are not utilized, and the performance cannot be fully utilized. Summary of the invention

[0009] In view of this, embodiments of the present application provide an image processing method, apparatus, device and medium based on element-by-element operation to overcome the above-mentioned problems or at least partially solve the above-mentioned problems.

[0010] A first aspect of an embodiment of the present application provides an image processing method based on element-by-element operation, the method comprising:

[0011] Determine an image to be processed as an input tensor, and determine the number of thread blocks required to be used in an element-by-element operation of the multiple input tensors according to the number of elements of the multiple elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device for each element in the input tensor of the pixel value of each pixel point in the image to be processed;

[0012] Determine, according to the number of thread blocks, the number of elements, and the memory layout of the multiple elements, a kernel function of an element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on a predefined template structure, wherein the parameters of the template structure include an element type and an element number;

[0013] Based on the number of thread blocks, multiple threads are initiated so that the multiple threads obtain the multiple elements according to the kernel function and perform an element-by-element operation process of the multiple input tensors to obtain an image processing result.

[0014] Optionally, determining the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device includes:

[0015] Determining the number of elements of each input tensor in the multiple input tensors according to the shapes of the multiple input tensors;

[0016] Calculating the number of first thread blocks according to the number of elements of each input tensor in the multiple input tensors;

[0017] Calling an API interface provided by the GPU device during operation to obtain hardware resource information of the GPU device;

[0018] Calculate the second number of thread blocks according to the hardware resource information of the GPU device, where the second number of thread blocks is the maximum number of thread blocks that the GPU device can provide;

[0019] Based on the first number of thread blocks and the second number of thread blocks, the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors is determined.

[0020] Optionally, the hardware resource information includes: the number of stream processors and the maximum number of threads of each stream processor; before determining the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors based on the first number of thread blocks and the second number of thread blocks, further comprising:

[0021] Preset the number of threads per thread block and the number of thread blocks that can be processed on each stream processor;

[0022] The calculating the first number of thread blocks according to the number of elements of each input tensor in the multiple input tensors comprises: determining the first number of thread blocks based on the number of threads in each thread block and the number of elements of each input tensor in the multiple input tensors;

[0023] The calculating the second number of thread blocks according to the hardware resource information of the GPU device includes: determining the second number of thread blocks based on the number of stream processors, the maximum number of threads of each stream processor, and the number of threads of each thread block;

[0024] The determining, based on the first number of thread blocks and the second number of thread blocks, the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors comprises: taking the minimum value of the first number of thread blocks and the second number of thread blocks as the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors.

[0025] Optionally, before obtaining the multiple elements, the method further includes:

[0026] The template structure is predefined, a type alias is set inside the template structure, and the type alias points to a memory block corresponding to a parameter of the template structure;

[0027] pre-define a template union, the template union is used to store the elements accessed in the memory block in a packaged form;

[0028] An array is predefined, and the array corresponds to the element size and element type in the parameters of the template structure.

[0029] Optionally, before obtaining the multiple elements, the method further includes:

[0030] Determining the number of elements that each thread needs to process according to the number of elements of the multiple elements included in the multiple input tensors and the number of thread blocks;

[0031] The obtaining of the plurality of elements comprises:

[0032] For each thread, based on the number of elements that each thread needs to process and the template structure, access the memory block pointed to by the type alias;

[0033] Packing the multiple elements in the acquired memory block through the template union to obtain the multiple elements;

[0034] Through the array, multiple elements in the template union are accessed in the form of index values.

[0035] Optionally, before obtaining the multiple elements, the method further includes:

[0036] Determine an index calculation method for elements in the kernel function based on shapes and strides of the multiple input tensors and shapes and strides of output tensors obtained after element-by-element operations on the multiple input tensors;

[0037] When the input tensor and the output tensor are both densely stored and have the same shape, the index calculation method is: obtaining elements in the input tensor by using thread numbers corresponding to the multiple threads as index values, and writing elements into the output tensor by using thread numbers corresponding to the multiple threads as index values;

[0038] When the input tensor and the output tensor are non-densely stored or have different shapes, and at least one of the input tensor or the output tensor is continuously stored, the index calculation method is: for the continuously stored input tensor or output tensor, the thread numbers corresponding to the multiple threads are used as index values ​​to obtain or write elements corresponding to the index values ​​corresponding to the thread numbers; for the non-continuously stored tensors other than the continuously stored input tensor or output tensor, the continuously stored input tensor or output tensor is used as a reference, and the index values ​​of the non-continuously stored elements are calculated according to the shape and stride of the non-continuously stored tensor;

[0039] When the input tensor and the output tensor are both non-continuously stored tensors, the index calculation method is: for each input tensor and each output tensor, the index value of the element that needs to be obtained or written in each input tensor and each output tensor is calculated according to the shape and stride.

[0040] Optionally, the method further comprises:

[0041] The obtained hardware resource information is stored in an internal memory;

[0042] When the hardware resource information needs to be obtained again, the internal memory is directly accessed to obtain the hardware resource information.

[0043] A second aspect of an embodiment of the present application provides an image processing device based on element-by-element operation, the device comprising:

[0044] A thread block number determination module is used to determine the image to be processed as an input tensor, and to determine the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements of the multiple elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device for each element in the input tensor of the pixel value of each pixel point in the image to be processed;

[0045] A kernel function determination module, configured to determine, according to the number of thread blocks, the number of elements, and the memory layout of the multiple elements, a kernel function of an element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on a predefined template structure, wherein the parameters of the template structure include an element type and an element number;

[0046] An execution module is used to initiate multiple threads based on the number of thread blocks, so that the multiple threads obtain the multiple elements according to the kernel function and perform an element-by-element operation process of the multiple input tensors to obtain an image processing result.

[0047] A third aspect of an embodiment of the present application provides an electronic device, comprising a memory, a processor, and a computer program stored in the memory, wherein the processor executes the computer program to implement the image processing method based on element-by-element operation as described in the first aspect.

[0048] A fourth aspect of an embodiment of the present application provides a computer-readable storage medium having a computer program stored thereon, wherein when the computer program is executed by a processor, the image processing method based on element-by-element operation as described in the first aspect is implemented.

[0049] Beneficial effects of this application:

[0050] An embodiment of the present application provides an image processing method, apparatus, device and medium based on element-by-element operation, the method comprising: determining an image to be processed as an input tensor, converting the pixel value of each pixel point in the image to be processed into each element in the input tensor, and determining the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements of the multiple elements contained in the multiple input tensors that need to participate in the operation, and the hardware resource information of the GPU device; determining the kernel function of the element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on the number of thread blocks, the number of elements and the memory layout of the multiple elements, based on a predefined template structure, wherein the parameters of the template structure include the element type and the number of elements; and initiating multiple threads based on the number of thread blocks, so that the multiple threads, according to the kernel function, obtain the multiple elements and execute the element-by-element operation process of the multiple input tensors to obtain the image processing result.

[0051] Through the technical solution of the present application, the most suitable number of thread blocks can be determined according to the number of elements involved in the element-by-element operation and the limitations of the hardware itself. This method ensures the computing requirements of the input tensor and utilizes the architectural characteristics of the computing device to give full play to the best computing performance of the GPU device. In addition, after determining the number of thread blocks, a template structure is defined, and each thread can use this template structure to specify how many elements can be obtained each time it is accessed, without having to obtain each element in turn, thereby reducing the number of accesses to the storage device. In addition, by adopting this template structure, there is no need to pre-write the corresponding structure for obtaining elements once before each element-by-element operation, but the variable parameters (i.e., the number of elements and the type of elements) can be directly modified. This unified vector type template method improves the reuse rate of the code corresponding to the structure. In addition, in the process of using the template structure, according to the number of threads and the memory layout of the elements contained in the input tensor, a method for multiple threads to obtain elements during the calculation process is formulated in the kernel function, and each thread obtains the elements that need to be calculated in parallel according to the kernel function, thereby improving the calculation efficiency of the element-by-element operation process on the GPU device. BRIEF DESCRIPTION OF THE DRAWINGS

[0052] The drawings constituting a part of the present application are used to provide a further understanding of the present application. The illustrative embodiments of the present application and their descriptions are used to explain the present application and do not constitute an improper limitation on the present application.

[0053] In order to more clearly illustrate the technical solution of the present application, the drawings required for use in the description of the present application will be briefly introduced below. Obviously, the drawings described below are only some embodiments of the present application. For ordinary technicians in this field, other drawings can be obtained based on these drawings without paying any creative labor.

[0054] Figure 1 It is a flowchart of an image processing method based on element-by-element operation provided by an embodiment of the present application;

[0055] Figure 2 is an architecture diagram of an image processing system based on an element-by-element operation method provided by an embodiment of the present application;

[0056] Figure 3 is a flowchart of a process for calculating the number of thread blocks provided by an embodiment of the present application;

[0057] Figure 4 is a schematic diagram of a classification process of a tensor type provided in an embodiment of the present application;

[0058] Figure 5is a schematic diagram of a framework of an image processing device based on element-by-element operation provided by an embodiment of the present application;

[0059] Figure 6 It is a schematic diagram of an electronic device provided by an embodiment of the present application. DETAILED DESCRIPTION

[0060] It should be noted that, in the absence of conflict, the embodiments and features in the embodiments of the present application may be combined with each other.

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

[0062] Figure 1 is a flowchart of an image processing method based on element-by-element operation provided by an embodiment of the present application. Figure 2 This is an architecture diagram of an image processing system based on an element-by-element operation method provided by one embodiment of the present application.

[0063] like Figure 2 As shown, the present application proposes an image processing method, apparatus, device and medium based on element-by-element operation, which is an efficient implementation scheme of element-by-element operation operators on GPU (Graphics Processing Unit) devices, aiming to significantly improve computing efficiency. The scheme is applied to an image processing system based on element-by-element operation method, which includes the following modules:

[0064] 1. Intelligent thread block number calculator: This module has hardware perception capabilities and can dynamically analyze and calculate the optimal number of thread blocks as a key parameter for starting kernel functions, thereby ensuring that the hardware resources of the GPU device are fully utilized and improving parallel computing efficiency.

[0065] 2. Data vectorization packaging optimizer: By introducing data vectorization packaging technology, this module significantly reduces the frequency of access to the storage device where the elements are located, effectively alleviates the bandwidth bottleneck problem, and further accelerates the data processing speed.

[0066] 3. Adaptive kernel scheduler: This module intelligently selects and executes the most matching kernel function based on the total number of elements and memory layout characteristics of the input data, achieving accurate scheduling and efficient execution of element-by-element computing tasks.

[0067] 4. Data index calculator with integrated fast division algorithm: In order to improve the efficiency of element index processing, this module integrates an efficient division algorithm, so that complex data index calculations can be completed quickly, reducing calculation delays and optimizing overall performance.

[0068] like Figure 1 As shown, an embodiment of the present application provides an image processing method based on element-by-element operation, the method comprising:

[0069] Step S101, determining an image to be processed as an input tensor, taking the pixel value of each pixel point in the image to be processed as each element in the input tensor, and determining the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements included in the multiple input tensors required to participate in the operation and the hardware resource information of the GPU device;

[0070] Step S102, determining the kernel function of the element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on a predefined template structure according to the number of thread blocks, the number of elements, and the memory layout of the multiple elements, wherein the parameters of the template structure include the element type and the number of elements;

[0071] Step S103: Initiate multiple threads based on the number of thread blocks, so that the multiple threads obtain the multiple elements according to the kernel function and perform an element-by-element operation process of the multiple input tensors to obtain an image processing result.

[0072] First, in step S101, before performing element-by-element operations on the input tensor, it is necessary to determine the number of threads to be initiated, and the number of threads is determined based on the number of thread blocks. A tensor is a data structure of different dimensions such as a matrix, array, or vector, and each tensor contains multiple elements.

[0073] Among them, in GPU programming, in order to effectively manage parallel operations, threads are organized into multiple levels: Grid: The grid is the top-level structure, representing the collection of the entire GPU computing task. It can be regarded as a two-dimensional or three-dimensional structure composed of multiple thread blocks. Each grid corresponds to a kernel function call, and all thread operations are performed in this grid. Thread block: Each thread block in a grid can be composed of multiple threads. Each thread block has a certain shared memory and synchronization mechanism, which enables the threads within the thread block to efficiently share data and work together. Thread blocks can usually be organized into one-dimensional, two-dimensional or three-dimensional forms to adapt to the processing requirements of different types of input tensors. Thread: A thread is the most basic execution unit and performs specific computing operations within a thread block. Each thread has an independent execution path (instruction stream), registers, and local memory.

[0074] Therefore, it is necessary to determine the number of thread blocks to ensure the efficiency of parallel processing that can efficiently execute element-by-element operations. Taking image processing tasks as an example, the image to be processed is determined as an input tensor, and the pixel value of each pixel in the image to be processed is input into each element in the tensor, where the number of thread blocks of the thread block is determined based on the number of elements contained in the multiple input tensors participating in the element-by-element operation and the hardware resource information of the GPU device. The number of elements contained in the multiple input tensors participating in the element-by-element operation represents the total number of elements in all the input tensors participating in the operation, and the hardware resource information of the GPU device is the computing power, number of cores, maximum number of thread blocks, and other information of the current GPU device. This information is usually obtained through API interface calls.

[0075] The number of elements contained in the multiple input tensors participating in the element-by-element operation can determine the requirement for the number of thread blocks in the element-by-element operation process. The hardware resource information of the GPU device can determine the limit of the number of thread blocks that can be provided in the element-by-element operation process. These two constraints on the number of thread blocks can determine the most reasonable number of thread blocks suitable for the current element-by-element operation process.

[0076] Further, in step S102, it is necessary to determine the kernel function suitable for different element-by-element operation processes, according to the number of thread blocks calculated in step S101, the number of elements contained in the multiple input tensors participating in the element-by-element operation and the memory layout of the multiple elements contained in the multiple input tensors of the element-by-element operation in the storage device, based on the predefined template structure, determine the kernel function required to perform the element-by-element operation of the input tensor, thereby ensuring that the selected kernel function can efficiently process the element-by-element operation process of the input tensor.

[0077] At the same time, in this process, a predefined template structure is used, which contains two variable parameters: element type and number of elements. The two variable parameters, element type and number of elements, are used to determine the number and type of elements accessed in a single access when accessing elements. The use of a template structure can improve the reusability of this structure.

[0078] Finally, in step S103, based on the number of thread blocks determined in step S101, multiple threads are initiated accordingly. The number of threads contained in each thread block is preset. Therefore, the number of threads that need to be initiated can be known according to the number of thread blocks.

[0079] And through the initiated multiple threads, according to the kernel function determined in step S102, multiple elements are obtained from the storage device, and the element-by-element operation process of multiple input tensors is executed in parallel, and finally an image processing result is obtained after the element-by-element operation is performed on each pixel value in the image to be processed.

[0080] Through the above embodiment, the most suitable number of thread blocks can be determined according to the number of elements involved in the element-by-element operation and the limitations of the hardware itself. This method ensures the computing requirements of the input tensor and utilizes the architectural characteristics of the computing device to give full play to the best computing performance of the GPU device. In addition, after determining the number of thread blocks, a template structure is defined, and each thread can use this template structure to specify how many elements can be obtained each time it is accessed, without having to obtain each element in turn, thereby reducing the number of accesses to the storage device. In addition, by adopting this template structure, there is no need to pre-write the corresponding structure for obtaining elements once before each element-by-element operation, but the variable parameters (i.e., the number of elements and the type of elements) can be directly modified. This unified vector type template method improves the reuse rate of the code corresponding to the structure. In addition, in the process of using the template structure, according to the number of threads and the memory layout of the elements, a method for multiple threads to obtain elements during the calculation process is formulated in the kernel function, and each thread obtains the elements that need to be calculated in parallel according to the kernel function, thereby improving the efficient execution of the calculation task of the element-by-element operation.

[0081] In summary, the entire technical solution can improve the computational efficiency of the element-by-element operation process on the GPU device.

[0082] Optionally, the step S101 includes:

[0083] Step S101-1, determining the number of elements of each input tensor in the multiple input tensors according to the shapes of the multiple input tensors;

[0084] Step S101-2, calculating the number of first thread blocks according to the number of elements of each input tensor in the multiple input tensors;

[0085] Step S101-3, calling an API interface provided by the GPU device during operation to obtain hardware resource information of the GPU device;

[0086] Step S101-4, calculating the second number of thread blocks according to the hardware resource information of the GPU device, where the second number of thread blocks is the maximum number of thread blocks that the GPU device can provide;

[0087] Step S101 - 5 , based on the first number of thread blocks and the second number of thread blocks, determining the number of thread blocks required to be used in the element-by-element operation process of the multiple input tensors.

[0088] Figure 3 It is a flowchart of a process for calculating the number of thread blocks provided in an embodiment of the present application.

[0089] Specifically, in one embodiment, the step S101 further specifically includes steps S101 - 1 to S101 - 5, and through steps S101 - 1 to S101 - 5, specific operations for determining the number of thread blocks required in the element-by-element operation process are described in detail.

[0090] like Figure 3 As shown, first, in step S101-1, it is necessary to determine the total number of elements of the input tensors corresponding to the multiple input tensors. In this process, the number of elements in each tensor is calculated according to the shape of each input tensor in the multiple input tensors, and then the number of elements of each input tensor is accumulated to obtain the total number of elements of the input tensors corresponding to the multiple input tensors. Among them, the shape of the input tensor represents the dimensional information of the input tensor (such as the number of rows, the number of columns, etc.).

[0091] For example, a tensor of shape 3×4 contains 12 elements, and a tensor of shape 2×5 contains 10 elements.

[0092] Further, in step S101-2, the first number of thread blocks is calculated. The first number of thread blocks represents the number of thread blocks required for the input tensor in the calculation process. In this process, the number of thread blocks participating in the element-by-element operation can be calculated as the first number of thread blocks based on the total number of elements of the input tensors corresponding to the multiple input tensors. The more the total number of elements of the input tensors corresponding to the multiple input tensors is, the more the first number of thread blocks is.

[0093] Further, in step S101-3, it is necessary to obtain the hardware resource information of the GPU device. In this process, the API interface provided by the GPU device at runtime is called to obtain the specific hardware resource information of the current GPU device. For example, the hardware resource information of the GPU device, such as the number of cores, the maximum number of threads, the maximum number of thread blocks, etc., is queried using an API such as CUDA.

[0094] Further, in step S101-4, the second number of thread blocks needs to be calculated. In this process, the maximum number of thread blocks that the GPU device can provide is calculated as the second number of thread blocks based on the hardware resource information of the GPU device obtained in step S101-3.

[0095] Finally, in step S101-5, the final number of thread blocks is determined. The first number of thread blocks and the second number of thread blocks are combined, and a suitable value is selected by comparing the first number of thread blocks and the second number of thread blocks. The number of thread blocks actually required to be used in the element-by-element operation process is determined. At the same time, it is necessary to ensure that the number of thread blocks does not exceed the second number of thread blocks.

[0096] Through the above embodiments, the number of thread blocks required for element-by-element operations can be determined in detail to ensure that the computing resources of the thread blocks are used reasonably and efficiently on the GPU device. The goal of each step is to systematically evaluate the characteristics of the input data, combined with the GPU hardware capabilities, and ultimately derive the best execution strategy. The design of this method ensures efficient parallel computing and is suitable for scenarios that require processing large amounts of data.

[0097] Optionally, the hardware resource information includes: the number of stream processors and the maximum number of threads of each stream processor; before step S101-5, it also includes:

[0098] Step S101-6, presetting the number of threads in each thread block and the number of thread blocks that can be processed by each stream processor;

[0099] The step S101-2 comprises: step S101-2-1, determining the number of the first thread blocks based on the number of threads of each thread block and the number of elements of each input tensor in the multiple input tensors;

[0100] The step S101-4 comprises: step S101-4-1, determining the number of the second thread blocks based on the number of stream processors, the maximum number of threads of each stream processor, and the number of threads of each thread block;

[0101] The step S101-5 comprises: step S101-5-1, taking the minimum value between the first number of thread blocks and the second number of thread blocks as the number of thread blocks required to be used in the element-by-element operation process of the multiple input tensors.

[0102] Specifically, before step S101-5 is executed, step S101-6 is also included. In step S101-6, the number of threads in each thread block and the number of thread blocks that can be processed on each stream processor are preset. As an example, the number of threads in each thread block can be 256, and the number of thread blocks that can be processed on each stream processor is 32.

[0103] Further, step S101-2 specifically includes step S101-2-1, in which the first number of thread blocks is determined according to the number of threads in each thread block and the total number of elements of each input tensor in the multiple input tensors.

[0104] Further, step S101-4 specifically includes step S101-4-1, in which the second number of thread blocks is determined according to the number of stream processors, the maximum number of threads of each stream processor, and the number of threads of each thread block.

[0105] Finally, step S101-5 specifically includes step S101-5-1. In step S101-5-1, after determining the first number of thread blocks and the second number of thread blocks, it is necessary to ensure that the number of thread blocks required to be used in the element-by-element operation of multiple input tensors does not exceed the limit of the second number of thread blocks. Therefore, the minimum value of the first number of thread blocks and the second number of thread blocks is used as the number of thread blocks required to be used in the element-by-element operation of multiple input tensors. In addition, it is also necessary to ensure that the number of thread blocks required to be used in the element-by-element operation of multiple input tensors is not less than 1.

[0106] Based on this, a calculation formula for the number of thread blocks required to be used in the element-by-element operation of multiple input tensors is shown, and the calculation formula is as follows:

[0107]

[0108] The above calculation formula includes two parts. The first part is used to calculate the number of first thread blocks, and the second part is used to calculate the number of second thread blocks.

[0109] The first part is:

[0110] The second part is:

[0111] Among them, numel in the first part represents the number of elements of each input tensor in multiple input tensors, threads_per_block represents the number of threads in each thread block, and the numerator "threads_per_block-1" in the first part is used to round up the calculation result. If there are 1000 elements and 256 threads in each thread block, the number of first thread blocks required is 4.

[0112] The second part, sm_count, indicates the number of stream processors, threads_per_sm indicates the maximum number of threads per stream processor, and num_waves indicates the number of thread blocks that can be processed on the stream processor.

[0113] “min” means taking the minimum value between the first number of thread blocks and the second number of thread blocks, and “max” means that the obtained number of thread blocks is not less than 1.

[0114] This calculation formula can ensure that the minimum number of thread blocks is 1. It is necessary to calculate a reasonable number of thread blocks based on the number of elements, while taking into account the hardware specifications and not exceeding the resource limitations of the GPU device hardware itself.

[0115] Optionally, the method further comprises:

[0116] The obtained hardware resource information is stored in an internal memory;

[0117] When the hardware resource information needs to be obtained again, the internal memory is directly accessed to obtain the hardware resource information.

[0118] Specifically, when using it, you only need to pass the number of elements to the intelligent thread block number calculator, which senses the hardware resource information of the GPU device and saves it in the internal memory during the first initialization. In each subsequent call, it automatically completes the calculation of the number of thread blocks based on the number of elements and hardware resource information. This method reduces the number of hardware accesses and can improve the running efficiency of the code.

[0119] Through the above embodiments, the number of thread blocks required for element-by-element operations on a GPU device can be reasonably calculated. Each step is based on the previous calculation and the hardware characteristics of the GPU device, thereby ensuring that the parallel capabilities of the GPU can be fully utilized during the execution process without exceeding the maximum processing capacity that its hardware can provide. This systematic design helps to optimize performance, especially when processing large-scale data, and can maximize the computational efficiency of element-by-element operations.

[0120] Optionally, before obtaining the multiple elements in step S103, the method further includes:

[0121] Step S103-1-1, pre-define the template structure, wherein a type alias is set inside the template structure, and the type alias points to a memory block corresponding to a parameter of the template structure;

[0122] Step S103-1-2, pre-defining a template union, wherein the template union is used to store the elements accessed in the memory block in a packaged form;

[0123] Step S103-1-3, pre-define an array, wherein the array corresponds to the element size and element type in the parameters of the template structure.

[0124] Specifically, in one embodiment, considering that the computing unit of a modern GPU runs much faster than a storage device, the reading and writing of elements often becomes a performance bottleneck. Therefore, before obtaining multiple elements, steps S103-1-1 to S103-1-3 are also included.

[0125] This application proposes an innovative vector type template VectorType and its implementation method, which realizes efficient storage and access to fixed-size data sequences by combining the characteristics of template metaprogramming and union.

[0126] In step S103-1-1, the present application defines a template structure VectorTypeProto, which accepts two variable parameters: type T and integer number, type T represents the type of elements in the input tensor, such as integer, floating point, etc., and integer number represents the number of elements.

[0127] The template structure VectorTypeProto uses std::aligned_storage to define a type alias type, which represents an uninitialized memory block large enough to store number elements of type T.

[0128] This design aims to provide a flexible and memory-aligned storage solution. The following is a pseudocode example that defines the template structure VectorTypeProto.

[0129] template<typename T,int number>

[0130] struct VectorTypeProto{

[0131] using type=typename std::aligned_storage<number*sizeof(T),number*sizeof(T)> ::type;

[0132] };

[0133] In step S103-1-2, for ease of use, this application also defines a type alias VectorType, which directly references the type alias type defined in the template structure VectorTypeProto.<T,number> Operations are used to reference a memory block of a specific size and type of type alias type. The following is a pseudocode example that defines a type alias VectorType.

[0134] template<typename T,int number>

[0135] using VectorType=typename VectorTypeProto<T,number> ::type;

[0136] In step S103-1-3, in order to further improve the flexibility of element access, this application also defines a template union VectorPack. The template union VectorPack contains a VectorType<T,number> The data member of the type can store multiple elements in the accessed memory block in a packaged form; at the same time, an array elements of type T is also defined, whose size is number. The array elements allows direct access to multiple elements obtained through indexes.

[0137] This design allows VectorPack to be processed as a whole while also making it easy to access multiple elements in the template union. The following is a pseudocode example that defines the template union VectorPack.

[0138] template<typename T,int number>

[0139] union VectorPack{

[0140] VectorType<T,number> data;

[0141] T elements[number];

[0142] };

[0143] Through the above embodiment, aligned storage is used. This design can improve access speed. Data can be treated as a whole (data) or elements can be accessed one by one (elements). Moreover, through type aliases, it becomes simple and clear for users to use template structures. This mechanism of efficiently storing and accessing fixed-size memory blocks can reduce the number of accesses in GPU computing tasks, while providing a flexible interface to facilitate access and processing of data. By defining VectorType and VectorPack, developers can more easily manage and use memory and improve program performance.

[0144] Optionally, before obtaining the multiple elements in step S103, the method further includes:

[0145] Step S103-2-1, determining the number of elements that each thread needs to process according to the number of elements of the multiple elements included in the multiple input tensors and the number of thread blocks;

[0146] The step S103 of acquiring the multiple elements includes:

[0147] Step S103-3-1, for each thread, based on the number of elements that each thread needs to process and the template structure, access the memory block pointed to by the type alias;

[0148] Step S103-3-2, packing the multiple elements in the acquired memory block through the template union to obtain the multiple elements;

[0149] Step S103-3-3: access multiple elements in the template union by index value through the array.

[0150] Specifically, in one embodiment, before obtaining the multiple elements in step S103, it also includes step S103-2-1, determining the number of elements that each thread needs to process based on the number of elements of the multiple elements contained in the multiple input tensors and the number of thread blocks, and using the obtained number of elements that each thread needs to process as the integer number in the above template, and using the type of the element as the type T in the above template.

[0151] The number of elements that each thread needs to process = the number of elements contained in the multiple input tensors / the number of thread blocks / the number of threads contained in each thread block.

[0152] Furthermore, in the process of acquiring the multiple elements in step S103, the multiple elements are acquired using the above-mentioned predefined template through steps S103-3-1 to S103-3-3.

[0153] In step S103-3-1, for each thread, the memory block pointed to by the type alias of the template structure is accessed according to the number of elements that each thread needs to process and the template structure, and in step S103-3-2, multiple elements in the acquired memory block are packaged through the template union to obtain multiple elements. The use of the template union to encapsulate multiple elements can more flexibly process different types of elements. In step SS103-3-3, the specific elements in the template union are accessed by array indexing to perform subsequent element-by-element operations. By placing the elements in the template union in an array, the thread can directly access and operate through the index value.

[0154] Through the above embodiment, firstly, the number of elements that each thread needs to process is determined according to the input amount and the number of thread blocks. Each thread accesses the corresponding memory block according to its ID, and uses the template union to package the data after reading it. Finally, the elements in the union are accessed and operated by array indexing, which improves the flexibility and efficiency of access. By effectively managing and accessing data, the efficiency of parallel processing of GPU devices is improved, ensuring that the computing tasks of element-by-element operations can be completed more efficiently.

[0155] Optionally, before obtaining the multiple elements in step S103, the method further includes:

[0156] Step S103-4-1, determining an index calculation method for elements in the kernel function according to shapes and strides of the multiple input tensors, and shapes and strides of output tensors obtained after element-by-element operations on the multiple input tensors;

[0157] Step S103-4-1-1, when the input tensor and the output tensor are both densely stored and have the same shape, the index calculation method is: obtaining elements in the input tensor using thread numbers corresponding to the multiple threads as index values, and writing elements into the output tensor using thread numbers corresponding to the multiple threads as index values;

[0158] Step S103-4-1-2, when the input tensor and the output tensor are non-densely stored or have different shapes, and at least one input tensor or output tensor is continuously stored, the index calculation method is: for the continuously stored input tensor or output tensor, use the thread numbers corresponding to the multiple threads as the index value to obtain or write the element corresponding to the index value corresponding to the thread number; for the non-continuously stored tensors other than the continuously stored input tensor or output tensor, use the continuously stored input tensor or output tensor as a reference, and calculate the index value of the non-continuously stored element according to the shape and stride of the non-continuously stored tensor;

[0159] Step S103-4-1-3, when the input tensor and the output tensor are both non-continuously stored tensors, the index calculation method is: for each input tensor and each output tensor, the index value of the element that needs to be obtained or written in each input tensor and each output tensor is calculated according to the shape and stride.

[0160] Specifically, in one embodiment, the input tensor passed into the element-by-element operator may contain different numbers of elements and memory layouts. For different situations, the index calculation method of the kernel function is different. The present application proposes an adaptive kernel scheduler KernelLauncher that automatically selects efficient kernel functions.

[0161] Figure 4 It is a schematic diagram of a tensor type classification process provided in one embodiment of the present application.

[0162] like Figure 4 As shown, first, in step S103-4-1, it is necessary to first determine the type of the input tensor according to the shapes and strides of multiple input tensors / output tensors, and divide it into three types: dense and same shape, continuous storage, and other. The kernel scheduler first classifies the input and output tensors according to the shapes and strides of the input and output tensors into three types. Through the IsDense interface of KernelLauncher, it is determined whether the tensor is densely stored, through the IsSame interface to determine whether the two tensors have the same shape and stride, and through IsContiguous to determine whether the tensor is a continuous storage type. After confirming the type, determine the index calculation method of the corresponding kernel function to access the element.

[0163] Among them, dense and same shape means: the input tensor and output tensor are densely stored in the storage device and have the same shape. When both the input and output tensors are densely stored and their shapes (dimensions and sizes) and strides are exactly the same, they are classified as dense and same shape types.

[0164] Continuous storage means that at least one tensor (input or output) is stored contiguously, but does not meet the conditions of dense and identical shapes.

[0165] Other types: Other cases except dense, same shape and continuous storage.

[0166] In step S103-4-1-1, when the input tensor and the output tensor are both densely stored and have the same shape, the corresponding index calculation method is: obtaining elements in the input tensor using the thread numbers corresponding to multiple threads as index values, and writing the elements obtained after the operation into the output tensor using the thread numbers corresponding to multiple threads as index values.

[0167] Moreover, when there are multiple rounds of iterative operations, the linear index number of each iteration is thread number + (number of iterations - 1) * total number of threads. The code is as follows: idx is the linear index number, tid is the thread number, and step is the total number of threads.

[0168] int64_t tid=blockDim.x*blockIdx.x+threadIdx.x;

[0169] int64_t step=blockDim.x*gridDim.x;

[0170] for(int64_t idx = tid; idx <numel;idx+=step){ ...

[0172] }

[0173] In step S103-4-1-2, when the input tensor and output tensor are non-densely stored or have different shapes, and at least one input tensor or output tensor is continuously stored, the index calculation method of the kernel function is as follows: for the continuously stored input tensor or output tensor, the thread numbers corresponding to multiple threads are used as index values ​​to obtain or write the elements corresponding to the index values ​​corresponding to the thread numbers; for the non-continuously stored tensors other than the continuously stored input tensors or output tensors, the continuously stored input tensor or output tensor is used as a reference, and according to the shape and stride of the non-continuously stored tensor, the index values ​​of the non-continuously stored elements are calculated, and the corresponding elements are accessed according to the calculated index values.

[0174] In addition, in the index calculation process for continuous storage, the division operation of the stride is involved, which will cumulatively reduce the performance of the overall calculation. Therefore, a fast division algorithm can be introduced to package the stride to the FastDivMod module, and the div() and mod() interfaces provided by FastDivMod can be used to accelerate the calculation.

[0175] In step S103-4-1-3, when both the input tensor and the output tensor are non-continuously stored tensors, the index calculation method of the corresponding kernel function is: for each input tensor and each output tensor, the index value of the element that needs to be obtained or written in each input tensor and each output tensor is calculated according to its shape and stride, and the corresponding element is accessed according to the calculated index value.

[0176] Through the above embodiments, different storage methods and shapes of tensors are used to determine how to accurately access elements in input tensors and output tensors in the GPU kernel. Different situations are divided according to conditions to ensure that data can be processed correctly and efficiently, and the corresponding design provides a basic guarantee for subsequent calculations. This flexibility of index calculation can adapt to the storage methods of various data structures to achieve efficient parallel computing.

[0177] Optionally, the image processing method based on element-by-element operation is applied to an image processing task of the GPU device, the input tensor is an image, and each element in the input tensor corresponds to a pixel value of a pixel point in the image.

[0178] Specifically, in one embodiment, the image processing method based on element-by-element operation can be specifically applied to image processing tasks of a GPU device, where the input vector is a specific image to be processed, and each element in the input tensor corresponds to the pixel value of a pixel point in the image to be processed.

[0179] The following takes image brightness adjustment as an example to illustrate the application of the image processing method based on element-by-element operation to the image processing task of the GPU device.

[0180] Makes an image lighter or darker by adjusting the overall brightness of the image.

[0181] The value (brightness) of each pixel is adjusted by adding a constant value.

[0182] The input tensor I represents the image, which contains the brightness value [i][j] of each pixel, and sets a constant brightness adjustment value C to adjust the brightness of each pixel.

[0183] Calculation formula: Out[i][j]=C+input[i][j].

[0184] In addition, it can also be used for image processing tasks such as image contrast adjustment, image blur, and edge detection.

[0185] Through the above embodiments, the image processing method based on element-by-element operation is combined with a specific image processing task, which can improve the computational efficiency of the element-by-element operation in the application scenario of processing large-scale image data.

[0186] This application significantly improves the computing efficiency and performance in the implementation of element-by-element operators on GPU devices, and the specific beneficial effects include:

[0187] 1. Dynamically optimize GPU resource utilization:

[0188] By dynamically analyzing the hardware resource information of the GPU device (such as the number of stream processors and the maximum number of threads per stream processor), the optimal number of thread blocks is automatically calculated. This ensures full utilization of GPU resources and avoids resource waste, while also reducing contention between threads and improving parallel computing efficiency.

[0189] 2. Reduce memory access latency:

[0190] By introducing data vectorization packaging technology, it is possible to determine how many elements each thread needs to process, and directly obtain these elements at one time through multiple pre-defined templates, significantly reducing the number of accesses to storage devices. In addition, vectorized storage and access methods make data more compact in memory, reducing the probability of cache misses, effectively alleviating bandwidth bottlenecks, and thus accelerating data processing speed.

[0191] 3. Adaptive kernel scheduling:

[0192] According to the total number of elements of the input data and the characteristics of the memory layout, the index calculation method in the most matching kernel function is intelligently selected and executed. This precise scheduling mechanism enables the processing of each tensor type to find the optimal calculation path, improving the overall execution efficiency.

[0193] In addition, in the index calculation process, especially for continuous storage type data, an efficient division algorithm (such as FastDivMod module) is introduced. This algorithm significantly reduces the time consumption of division operations, optimizes the performance of index calculation, and thus improves the overall calculation efficiency.

[0194] 4. Improve code maintainability and scalability:

[0195] The implementation of this application uses technologies such as template metaprogramming, template unions and type aliases to make the code more flexible, easy to maintain and expand. By defining template structures and unions, unified processing of different types of data is achieved, duplication of code is reduced, and code reusability is improved.

[0196] 5. Improve user experience:

[0197] It reduces the user's reliance on and need to understand GPU devices. Users only need to provide the total number of elements of the input tensor, and the intelligent thread block number calculator can automatically calculate the optimal number of thread blocks, which lowers the user's usage threshold.

[0198] 6. Wide range of application scenarios:

[0199] The solution proposed in this application is not only applicable to element-by-element operators, but can also be extended to other types of GPU computing tasks, providing an effective performance optimization solution for various complex GPU calculations.

[0200] In summary, this application has significantly improved the computational efficiency and performance of element-by-element operators on GPU devices through a series of innovations and optimization measures.

[0201] Figure 5 is a schematic diagram of a framework of an image processing device based on element-by-element operation provided by an embodiment of the present application, such as Figure 5 shown.

[0202] Based on the same inventive concept, another embodiment of the present application further provides an image processing device based on element-by-element operation, the device comprising:

[0203] The thread block number determination module 11 is used to determine the image to be processed as an input tensor, and determine the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements of the multiple elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device for each element in the input tensor of the pixel value of each pixel point in the image to be processed;

[0204] A kernel function determination module 12 is used to determine the kernel function of the element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on the number of thread blocks, the number of elements and the memory layout of the multiple elements, based on a predefined template structure, wherein the parameters of the template structure include the element type and the number of elements;

[0205] The execution module 13 is used to initiate multiple threads based on the number of thread blocks, so that the multiple threads obtain the multiple elements according to the kernel function and perform the element-by-element operation process of the multiple input tensors to obtain the image processing result.

[0206] Optionally, the thread block quantity determination module 11 includes:

[0207] an element number determining unit, configured to determine the number of elements of each input tensor in the plurality of input tensors according to shapes of the plurality of input tensors;

[0208] A first thread block number calculation unit, configured to calculate the first thread block number according to the number of elements of each input tensor in the multiple input tensors;

[0209] A hardware resource information acquisition unit, used to call an API interface provided by the GPU device during operation to acquire hardware resource information of the GPU device;

[0210] A second thread block quantity calculation unit, configured to calculate a second thread block quantity according to hardware resource information of the GPU device, wherein the second thread block quantity is a maximum number of thread blocks that can be provided by the GPU device;

[0211] A thread block number determination unit is used to determine the number of thread blocks required to be used in the element-by-element operation process of the multiple input tensors based on the first thread block number and the second thread block number.

[0212] Optionally, the hardware resource information includes: the number of stream processors and the maximum number of threads of each stream processor, and the device further includes:

[0213] A setting unit, configured to pre-set the number of threads in each thread block and the number of thread blocks that can be processed on each stream processor before determining the number of thread blocks to be used in the element-by-element operation process of the multiple input tensors based on the first number of thread blocks and the second number of thread blocks;

[0214] The first thread block number calculation unit includes: a first thread block number calculation subunit, configured to determine the first thread block number based on the number of threads in each thread block and the number of elements in each input tensor in the multiple input tensors;

[0215] The second thread block number calculation unit includes: a second thread block number calculation subunit, configured to determine the second thread block number based on the number of stream processors, the maximum number of threads of each stream processor, and the number of threads of each thread block;

[0216] The thread block number determination unit includes: a thread block number determination subunit, which is used to use the minimum value of the first thread block number and the second thread block number as the thread block number required to be used in the element-by-element operation process of the multiple input tensors.

[0217] Optionally, the device further comprises:

[0218] A first definition unit is used to predefine the template structure before acquiring the multiple elements, wherein a type alias is provided inside the template structure, and the type alias points to a memory block corresponding to a parameter of the template structure;

[0219] A second definition unit, used for predefining a template union, wherein the template union is used for storing elements accessed in the memory block in a packaged form;

[0220] The third definition unit is used to predefine an array, where the array corresponds to the element size and element type in the parameters of the template structure.

[0221] Optionally, the device further comprises:

[0222] a unit element number determination unit, used for determining the number of elements that each thread needs to process according to the number of elements of the multiple elements included in the multiple input tensors and the number of thread blocks before obtaining the multiple elements;

[0223] The execution module 13 includes:

[0224] A memory block access unit, configured to access the memory block pointed to by the type alias for each thread based on the number of elements to be processed by each thread and the template structure;

[0225] A packing unit, used for packing the multiple elements in the acquired memory block through the template union to obtain the multiple elements;

[0226] The element access unit is used to access multiple elements in the template union by index value through the array.

[0227] Optionally, the device further comprises:

[0228] an index calculation mode determination module, used for determining, before obtaining the plurality of elements, the index calculation mode of the elements in the kernel function according to the shapes and strides of the plurality of input tensors and the shapes and strides of the output tensors obtained after the element-by-element operation of the plurality of input tensors;

[0229] A first index calculation method determining unit, configured to, when both the input tensor and the output tensor are densely stored and have the same shape, determine the index calculation method as follows: obtaining elements in the input tensor by using thread numbers corresponding to the multiple threads as index values, and writing elements into the output tensor by using thread numbers corresponding to the multiple threads as index values;

[0230] A second index calculation method determination unit is used for, when the input tensor and the output tensor are non-densely stored or have different shapes, and at least one of the input tensor or the output tensor is continuously stored, the index calculation method is: for the continuously stored input tensor or output tensor, using the thread numbers corresponding to the multiple threads as index values ​​to obtain or write elements corresponding to the index values ​​corresponding to the thread numbers; for non-continuously stored tensors other than the continuously stored input tensor or output tensor, using the continuously stored input tensor or output tensor as a reference, according to the shape and stride of the non-continuously stored tensor, calculating the index values ​​of the non-continuously stored elements;

[0231] The third index calculation method determination unit is used for, when the input tensor and the output tensor are both non-continuously stored tensors, the index calculation method is: for each input tensor and each output tensor, the index value of the element that needs to be obtained or written in each input tensor and each output tensor is calculated according to the shape and stride.

[0232] Optionally, the image processing method based on element-by-element operation is applied to an image processing task of the GPU device, the input tensor is an image, and each element in the input tensor corresponds to a pixel value of a pixel point in the image.

[0233] Optionally, the device further comprises:

[0234] A storage unit, used for storing the acquired hardware resource information in an internal memory;

[0235] The access unit is used to directly access the internal memory to obtain the hardware resource information when the hardware resource information needs to be obtained again.

[0236] Based on the same inventive concept, another embodiment of the present application also provides an electronic device, including a memory, a processor, and a computer program stored in the memory, wherein the processor executes the computer program to implement the image processing method based on element-by-element operation as described in any of the above embodiments.

[0237] Among them, electronic equipment refers to Figure 6 , Figure 6 is a schematic diagram of an electronic device provided by an embodiment of the present application. Figure 6 As shown, the electronic device 600 includes: a memory 610 and a processor 620. The memory 610 and the processor 620 are connected via a bus communication. A computer program is stored in the memory 510. The computer program can be run on the processor 620 to implement the steps in the image processing method based on element-by-element operation disclosed in the above embodiment of the present application.

[0238] Based on the same inventive concept, another embodiment of the present application further provides a computer program product, including a computer program, which is executed by a processor to perform the image processing method based on element-by-element operation as described in any of the above embodiments.

[0239] Based on the same inventive concept, another embodiment of the present application further provides a computer-readable storage medium on which a computer program is stored, wherein when the program is executed by a processor, the image processing method based on element-by-element operation as described in any of the above embodiments is implemented.

[0240] As for the device, since it is basically similar to the method embodiment, the description is relatively simple, and the relevant parts can be referred to the partial description of the method embodiment.

[0241] The various embodiments in this specification are described in a progressive manner, and each embodiment focuses on the differences from other embodiments. The same or similar parts between the various embodiments can be referenced to each other.

[0242] Those skilled in the art will appreciate that the embodiments of the present application can be provided as methods, devices, or computer program products. Therefore, the present application can adopt the form of a complete hardware embodiment, a complete software embodiment, or an embodiment in combination with software and hardware. Moreover, the present application can adopt the form of a computer program product implemented on one or more computer-usable storage media (including but not limited to disk storage, CD-ROM, optical storage, etc.) containing computer-usable program code.

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

[0244] These computer program instructions may also be stored in a computer-readable memory capable of directing a computer or other programmable data processing terminal device to operate in a specific manner, so that the instructions stored in the computer-readable memory produce a manufactured product including an instruction device, which implements the process Figure 1 A process or multiple processes and / or boxes Figure 1 A function specified in one or more boxes.

[0245] These computer program instructions can also be loaded onto a computer or other programmable data processing terminal device so that a series of operating steps are executed on the computer or other programmable terminal device to produce a computer-implemented process, thereby providing instructions for implementing the process in the computer or other programmable terminal device. Figure 1 A process or multiple processes and / or boxes Figure 1 The steps for the functions specified in one or more boxes.

[0246] Although the preferred embodiments of the present application have been described, those skilled in the art may make additional changes and modifications to these embodiments once they have learned the basic creative concept. Therefore, the appended claims are intended to be interpreted as including the preferred embodiments and all changes and modifications that fall within the scope of the embodiments of the present application.

[0247] 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 terminal 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 terminal device. In the absence of further restrictions, the elements defined by the sentence "comprise a ..." do not exclude the existence of other identical elements in the process, method, article or terminal device including the elements.

[0248] The above is a detailed introduction to the image processing method, device, equipment and medium based on element-by-element operation provided by the present application. 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. An image processing method based on element-by-element operation, characterized in that: The method comprises: Determine an image to be processed as an input tensor, and determine the number of thread blocks required to be used in an element-by-element operation of the multiple input tensors according to the number of elements of the multiple elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device for each element in the input tensor of the pixel value of each pixel point in the image to be processed; Determine, according to the number of thread blocks, the number of elements, and the memory layout of the multiple elements, a kernel function of an element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on a predefined template structure, wherein the parameters of the template structure include an element type and an element number; Based on the number of thread blocks, multiple threads are initiated so that the multiple threads obtain the multiple elements according to the kernel function and perform an element-by-element operation process of the multiple input tensors to obtain an image processing result.

2. The image processing method based on element-by-element operation according to claim 1, characterized in that: The determining the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device includes: Determining the number of elements of each input tensor in the multiple input tensors according to the shapes of the multiple input tensors; Calculating the number of first thread blocks according to the number of elements of each input tensor in the multiple input tensors; Calling an API interface provided by the GPU device during operation to obtain hardware resource information of the GPU device; Calculate the second number of thread blocks according to the hardware resource information of the GPU device, where the second number of thread blocks is the maximum number of thread blocks that the GPU device can provide; Based on the first number of thread blocks and the second number of thread blocks, the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors is determined.

3. The image processing method based on element-by-element operation according to claim 2, characterized in that: The hardware resource information includes: the number of stream processors and the maximum number of threads of each stream processor; before determining the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors based on the first number of thread blocks and the second number of thread blocks, further including: Preset the number of threads per thread block and the number of thread blocks that can be processed on each stream processor; The calculating the first number of thread blocks according to the number of elements of each input tensor in the multiple input tensors comprises: determining the first number of thread blocks based on the number of threads in each thread block and the number of elements of each input tensor in the multiple input tensors; The calculating the second number of thread blocks according to the hardware resource information of the GPU device includes: determining the second number of thread blocks based on the number of stream processors, the maximum number of threads of each stream processor, and the number of threads of each thread block; The determining, based on the first number of thread blocks and the second number of thread blocks, the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors comprises: taking the minimum value of the first number of thread blocks and the second number of thread blocks as the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors.

4. The image processing method based on element-by-element operation according to claim 1, characterized in that: Before obtaining the plurality of elements, the method further includes: The template structure is predefined, a type alias is set inside the template structure, and the type alias points to a memory block corresponding to a parameter of the template structure; pre-define a template union, the template union is used to store the elements accessed in the memory block in a packaged form; An array is predefined, and the array corresponds to the element size and element type in the parameters of the template structure.

5. The image processing method based on element-by-element operation according to claim 4, characterized in that: Before obtaining the multiple elements, the method further includes: Determining the number of elements that each thread needs to process according to the number of elements of the multiple elements included in the multiple input tensors and the number of thread blocks; The obtaining of the plurality of elements comprises: For each thread, based on the number of elements that each thread needs to process and the template structure, access the memory block pointed to by the type alias; Packing the multiple elements in the acquired memory block through the template union to obtain the multiple elements; Through the array, multiple elements in the template union are accessed in the form of index values.

6. The image processing method based on element-by-element operation according to claim 5, characterized in that: Before obtaining the multiple elements, the method further includes: Determine an index calculation method for elements in the kernel function based on shapes and strides of the multiple input tensors and shapes and strides of output tensors obtained after element-by-element operations on the multiple input tensors; When the input tensor and the output tensor are both densely stored and have the same shape, the index calculation method is: obtaining elements in the input tensor by using thread numbers corresponding to the multiple threads as index values, and writing elements into the output tensor by using thread numbers corresponding to the multiple threads as index values; When the input tensor and the output tensor are non-densely stored or have different shapes, and at least one of the input tensor or the output tensor is continuously stored, the index calculation method is: for the continuously stored input tensor or output tensor, the thread numbers corresponding to the multiple threads are used as index values ​​to obtain or write elements corresponding to the index values ​​corresponding to the thread numbers; for the non-continuously stored tensors other than the continuously stored input tensor or output tensor, the continuously stored input tensor or output tensor is used as a reference, and the index values ​​of the non-continuously stored elements are calculated according to the shape and stride of the non-continuously stored tensor; When the input tensor and the output tensor are both non-continuously stored tensors, the index calculation method is: for each input tensor and each output tensor, the index value of the element that needs to be obtained or written in each input tensor and each output tensor is calculated according to the shape and stride.

7. The image processing method based on element-by-element operation according to any one of claims 1 to 6, characterized in that: The image processing method based on element-by-element operation is applied to the image processing task of the GPU device, the input tensor is an image, and each element in the input tensor corresponds to the pixel value of a pixel point in the image.

8. An image processing device based on element-by-element operation, characterized in that: The device comprises: A thread block number determination module is used to determine the image to be processed as an input tensor, and to determine the number of thread blocks required to be used in the element-by-element operation of the multiple input tensors according to the number of elements of the multiple elements included in the multiple input tensors that need to participate in the operation and the hardware resource information of the GPU device for each element in the input tensor of the pixel value of each pixel point in the image to be processed; A kernel function determination module, configured to determine, according to the number of thread blocks, the number of elements, and the memory layout of the multiple elements, a kernel function of an element-by-element operation class operator corresponding to the multiple input tensors that need to participate in the operation based on a predefined template structure, wherein the parameters of the template structure include an element type and an element number; An execution module is used to initiate multiple threads based on the number of thread blocks, so that the multiple threads obtain the multiple elements according to the kernel function and perform an element-by-element operation process of the multiple input tensors to obtain an image processing result.

9. An electronic device, characterized in that: The invention comprises a memory, a processor and a computer program stored in the memory, wherein the processor executes the computer program to implement the image processing method based on element-by-element operation as claimed in any one of claims 1 to 7.

10. A computer-readable storage medium, characterized in that: A computer program is stored thereon, wherein when the computer program is executed by a processor, the image processing method based on element-by-element operation as described in any one of claims 1 to 7 is implemented.

Images