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

Abstract

This application provides an image processing method, apparatus, device, and medium based on element-wise operations, relating to the field of GPU technology. The method includes: determining the number of thread blocks required for the element-wise operations of the multiple input tensors based on the number of elements contained in the multiple input tensors to participate in the operations and the hardware resource information of the GPU device; determining the kernel function of the element-wise operation class operator corresponding to the multiple input tensors to participate in the operations based on the number of thread blocks, the number of elements, and the memory layout of the multiple elements, using a predefined template structure, wherein the parameters of the template structure include element type and 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-wise operation process of the multiple input tensors to obtain the image processing result.

Description

Technical Field

[0001] This application relates to the field of GPU technology, and in particular to an image processing method, apparatus, device and medium based on element-wise operations. Background Technology

[0002] Element-wise operations are widely used in data processing and mathematical computation. They refer to operations that perform calculations on corresponding elements in two or more data structures (such as arrays, vectors, matrices, or higher-dimensional tensors) one by one. This type of operation does not depend on the overall relationship between data structures, but acts independently on each element.

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

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

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

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

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

[0008] In related technologies, separate kernel functions are written for each element-wise operator, extracting elements from the input data one by one for calculation and then writing the result to the output. While this approach can accomplish the computational function, it does not provide a unified execution framework for this type of operator, resulting in low development efficiency. Furthermore, it does not utilize the architectural characteristics of computing devices, failing to fully leverage their performance. Summary of the Invention

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

[0010] The first aspect of this application provides an image processing method based on element-wise operations, the method comprising:

[0011] The image to be processed is determined as the input tensor. The pixel value of each pixel in the image to be processed is assigned to each element in the input tensor. Based on the number of elements in the multiple input tensors that need to participate in the operation, and the hardware resource information of the GPU device, the number of thread blocks required for the multiple input tensors to be used in the element-wise operation is determined.

[0012] Based on the number of thread blocks, the number of elements, and the memory layout of the multiple elements, and using a predefined template structure, determine the kernel function of the element-wise operation operator corresponding to the multiple input tensors that need to participate in the operation. The parameters of the template structure include the element type and the number of elements.

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

[0014] Optionally, determining the number of thread blocks required for element-wise operations of the multiple input tensors based on the number of elements in the multiple elements included in the multiple input tensors participating in the computation, and the hardware resource information of the GPU device, includes:

[0015] Based on the shape of the plurality of input tensors, determine the number of elements in each of the plurality of input tensors;

[0016] The number of first thread blocks is calculated based on the number of elements in each of the multiple input tensors;

[0017] Call the API interface provided by the GPU device at runtime to obtain the hardware resource information of the GPU device;

[0018] Based on the hardware resource information of the GPU device, the number of second thread blocks is calculated, which is the maximum number of thread blocks that the GPU device can provide.

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

[0020] Optionally, the hardware resource information includes: the number of stream processors, the maximum number of threads per stream processor; before determining the number of thread blocks required by the plurality of input tensors in the element-wise operation based on the first thread block number and the second thread block number, it further includes:

[0021] The number of threads in each thread block and the number of thread blocks that can be processed on each stream processor are preset.

[0022] The step of calculating the number of first thread blocks based on the number of elements in each of the plurality of input tensors includes: determining the number of first thread blocks based on the number of threads in each thread block and the number of elements in each of the plurality of input tensors;

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

[0024] Determining the number of thread blocks required for the multiple input tensors during element-wise operations based on the number of the first thread blocks and the number of the second thread blocks includes: taking the minimum value between the number of the first thread blocks and the number of the second thread blocks as the number of thread blocks required for the multiple input tensors during element-wise operations.

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

[0026] The template structure is predefined, and the template structure has type aliases internally, which point to the memory blocks corresponding to the parameters of the template structure.

[0027] A template union is predefined, which is used to store the accessed elements in the memory block in a packaged form;

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

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

[0030] The number of elements to be processed by each thread is determined based on the number of elements contained in the multiple input tensors and the number of thread blocks.

[0031] The acquisition of the plurality of elements includes:

[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] The multiple elements in the obtained memory block are packaged using the template union to obtain the multiple elements;

[0034] Multiple elements in the template union can be accessed by index value through the array.

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

[0036] Based on the shape and stride of the multiple input tensors, and the shape and stride of the output tensor obtained after element-wise operations on the multiple input tensors, determine the index calculation method for the elements in the kernel function;

[0037] When both the input tensor and the output tensor are densely stored and have the same shape, the index calculation method is as follows: obtain the elements in the input tensor using the thread number corresponding to the multiple threads as the index value, and write the elements into the output tensor using the thread number corresponding to the multiple threads as the index value;

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

[0039] When both the input tensor and the output tensor are non-contiguous storage tensors, the index calculation method is as follows: for each input tensor and each output tensor, the index value of the element to be acquired or written in each input tensor and each output tensor is calculated according to the shape and stride.

[0040] Optionally, the method further includes:

[0041] The acquired hardware resource information is stored in the internal memory;

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

[0043] A second aspect of this application provides an image processing apparatus based on element-wise operations, the apparatus comprising:

[0044] The 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 that the multiple input tensors need to use in the element-wise operation process by taking the pixel value of each pixel in the image to be processed as each element in the input tensor, the number of elements in 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.

[0045] The kernel function determination module is used to determine the kernel function of the element-wise 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, and a predefined template structure. The parameters of the template structure include the element type and the number of elements.

[0046] The 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 element-wise operations on the multiple input tensors to obtain the image processing result.

[0047] A third aspect of this application 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-wise operations as described in the first aspect.

[0048] A fourth aspect of this application provides a computer-readable storage medium having a computer program stored thereon, wherein the computer program, when executed by a processor, implements the image processing method based on element-wise operations as described in the first aspect.

[0049] The beneficial effects of this application are:

[0050] This application provides an image processing method, apparatus, device, and medium based on element-wise operations. The method includes: determining an image to be processed as an input tensor; assigning the pixel value of each pixel in the image to each element in the input tensor; determining the number of thread blocks required for the element-wise operations of the multiple input tensors based on the number of elements in the multiple input tensors to participate in the operations, and hardware resource information of the GPU device; determining the kernel function of the element-wise operation class operator corresponding to the multiple input tensors to participate in the operations based on the number of thread blocks, the number of elements, and the memory layout of the multiple elements, using a predefined template structure, wherein the parameters of the template structure include element type and 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-wise operation process of the multiple input tensors to obtain the image processing result.

[0051] The technical solution of this application can determine the optimal number of thread blocks based on the number of elements involved in element-wise operations and the limitations of the hardware itself. This approach ensures the computational requirements of the input tensor while leveraging the architectural features of the computing device to maximize the computing performance of the GPU. Furthermore, after determining the number of thread blocks, a template structure is defined. Each thread can use this template structure to specify how many elements can be retrieved per access, eliminating the need to retrieve each element sequentially and reducing the number of storage device accesses. Moreover, this template structure approach eliminates the need to pre-write a corresponding structure for single-entry element retrieval before each element-wise operation; instead, the variable parameters (i.e., the number of elements and element type) can be directly modified. This unified vector type template approach improves the reusability of the corresponding structure code. In addition, during the use of the template structure, based on the number of threads and the memory layout of the elements contained in the input tensor, the kernel function specifies how multiple threads retrieve elements during computation. Each thread retrieves the elements to be processed in parallel according to the kernel function, improving the computational efficiency of element-wise operations on the GPU. Attached Figure Description

[0052] The accompanying drawings, which form part of this application, are used to provide a further understanding of this application. The illustrative embodiments of this application and their descriptions are used to explain this application and do not constitute an undue limitation of this application.

[0053] To more clearly illustrate the technical solution of this application, the drawings used in the description of this application will be briefly introduced below. Obviously, the drawings described below are only some embodiments of this application. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.

[0054] Figure 1 This is a flowchart illustrating an image processing method based on element-wise operations according to an embodiment of this application;

[0055] Figure 2 This is an architecture diagram of an image processing system based on an element-wise operation method provided in one embodiment of this application;

[0056] Figure 3 This is a flowchart illustrating the calculation process for the number of thread blocks provided in an embodiment of this application;

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

[0058] Figure 5This is a schematic diagram of the framework of an image processing device based on element-wise operations according to an embodiment of this application;

[0059] Figure 6 This is a schematic diagram of an electronic device provided in an embodiment of this application. Detailed Implementation

[0060] It should be noted that, unless otherwise specified, the embodiments and features described in this application can be combined with each other.

[0061] The technical solutions of the embodiments of this application will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some embodiments of this application, not all embodiments. Based on the embodiments of this application, all other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of this application.

[0062] Figure 1 This is a flowchart illustrating an image processing method based on element-wise operations according to an embodiment of this application. Figure 2 This is an architecture diagram of an image processing system based on an element-wise operation method provided in an embodiment of this application.

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

[0064] 1. Intelligent Thread Block Count Calculator: This module has hardware awareness capabilities and can dynamically analyze and calculate the optimal number of thread blocks as a key parameter for kernel function startup, thereby ensuring that the hardware resources of the GPU device are fully utilized and improving parallel computing efficiency.

[0065] 2. Data Vectorization Wrapper Optimizer: By introducing data vectorization wrapping technology, this module significantly reduces the frequency of access to the storage device where the elements are located, effectively alleviating the bandwidth bottleneck problem and further accelerating the data processing speed.

[0066] 3. Adaptive Kernel Scheduler: This module intelligently selects and executes the most suitable kernel function based on the total number of elements in the input data and the memory layout characteristics, thereby achieving precise scheduling and efficient execution of element-wise computational tasks.

[0067] 4. Data Index Calculator with Integrated Fast Division Algorithm: To improve the efficiency of element indexing, this module integrates a high-efficiency division algorithm, enabling rapid completion of complex data indexing calculations, reducing computational latency, and optimizing overall performance.

[0068] like Figure 1 As shown, one embodiment of this application provides an image processing method based on element-wise operations, the method comprising:

[0069] Step S101: Determine the image to be processed as the input tensor, and take the pixel value of each pixel in the image to be processed as each element in the input tensor. Based on the number of elements in 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, determine the number of thread blocks that the multiple input tensors need to use in the element-by-element operation process.

[0070] Step S102: Based on the number of thread blocks, the number of elements, and the memory layout of the multiple elements, determine the kernel function of the element-wise operation class operator corresponding to the multiple input tensors that need to participate in the operation, based on a predefined template structure. The parameters of the template structure include the element type and the number of elements.

[0071] Step S103: Based on the number of thread blocks, initiate multiple threads so that the multiple threads obtain the multiple elements according to the kernel function and perform element-wise operations on the multiple input tensors to obtain the image processing result.

[0072] First, in step S101, before performing element-wise operations on the input tensor, the number of threads to be initiated needs to be determined, and the number of threads is determined based on the number of thread blocks. Tensors are data structures of different dimensions, such as matrices, arrays, and vectors, and each tensor contains multiple elements.

[0073] In GPU programming, to effectively manage parallel computation, threads are organized into multiple layers: Grid: The grid is the top-level structure, representing the entire collection of GPU computing tasks. It can be viewed as a two-dimensional or three-dimensional structure composed of multiple thread blocks. Each grid corresponds to one kernel function call, and all thread operations are performed within this grid. Block: Each thread block within a grid can consist of multiple threads. Each thread block has certain shared memory and synchronization mechanisms, enabling threads within a block to efficiently share data and cooperate. Thread blocks can typically be organized in one-dimensional, two-dimensional, or three-dimensional form to accommodate the processing needs of different types of input tensors. Thread: A thread is the most basic unit of execution, performing specific computational operations within a thread block. Each thread has its own independent execution path (instruction stream), registers, and local memory.

[0074] Therefore, the number of thread blocks needs to be determined to ensure efficient parallel processing of element-wise operations. Taking image processing tasks as an example, the image to be processed is defined as the input tensor, and the pixel value of each pixel in the image is input to each element of the tensor. The number of thread blocks is determined based on the number of elements in the multiple input tensors participating in the element-wise operations and the hardware resource information of the GPU device. The number of elements in the multiple input tensors participating in the element-wise operations represents the total number of elements in all input tensors involved in the operation. The hardware resource information of the GPU device includes its computing power, number of cores, and maximum number of thread blocks. This information is typically obtained through API calls.

[0075] The number of elements in the multiple input tensors involved in the element-wise operation determines the required number of thread blocks during the element-wise operation. The hardware resource information of the GPU device can determine the limit of the number of thread blocks that can be provided during the element-wise operation. Constraining the number of thread blocks with these two factors can determine the most reasonable number of thread blocks suitable for the current element-wise operation process.

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

[0077] Meanwhile, a predefined template structure was used in this process. The template structure contains two variable parameters: element type and number of elements. These two variable parameters are used to determine the number and type of elements accessed in a single access when accessing elements. Using 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. The number of threads contained in each thread block is preset, so the number of threads to be initiated can be determined based on the number of thread blocks.

[0079] By initiating multiple threads and following the kernel function determined in step S102, multiple elements are retrieved from the storage device, and the element-wise operation process of multiple input tensors is executed in parallel. Finally, the image processing result is obtained after performing element-wise operations on each pixel value in the image to be processed.

[0080] Through the above embodiments, the optimal number of thread blocks can be determined based on the number of elements participating in element-wise operations and the limitations of the hardware itself. This approach ensures the computational requirements of the input tensor and leverages the architectural characteristics of the computing device to maximize the computing performance of the GPU. Furthermore, after determining the number of thread blocks, a template structure is defined. Each thread can use this template structure to specify how many elements can be retrieved per access, eliminating the need to retrieve each element sequentially and reducing the number of storage device accesses. Moreover, this template structure approach eliminates the need to pre-write a corresponding structure for single-time element retrieval before each element-wise operation; instead, the variable parameters (i.e., the number of elements and element type) can be directly modified. This unified vector type template approach improves the reusability of the structure's code. In addition, during the use of the template structure, based on the number of threads and the memory layout of the elements, the kernel function specifies how multiple threads retrieve elements during computation. Each thread retrieves the elements to be processed in parallel according to the kernel function, improving the efficiency of element-wise computation tasks.

[0081] In summary, the entire technical solution can improve the computational efficiency of element-wise operations on GPU devices.

[0082] Optionally, step S101 includes:

[0083] Step S101-1: Determine the number of elements in each of the multiple input tensors based on the shape of the multiple input tensors;

[0084] Step S101-2: Calculate the number of first thread blocks based on the number of elements in each of the multiple input tensors;

[0085] Step S101-3: Call the API interface provided by the GPU device at runtime to obtain the hardware resource information of the GPU device;

[0086] Step S101-4: Calculate the second thread block number based on the hardware resource information of the GPU device. The second thread block number is the maximum number of thread blocks that the GPU device can provide.

[0087] Step S101-5: Based on the number of the first thread blocks and the number of the second thread blocks, determine the number of thread blocks that the multiple input tensors need to use in the element-wise operation process.

[0088] Figure 3 This is a flowchart illustrating the calculation process for the number of thread blocks provided in one embodiment of this application.

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

[0090] like Figure 3 As shown, firstly, in step S101-1, it is necessary to determine the total number of elements in the input tensors corresponding to the multiple input tensors. This process involves calculating the number of elements in each tensor based on its shape, and then summing the element counts of each input tensor to obtain the total number of elements in the input tensors corresponding to the multiple input tensors. The shape of the input tensor represents its dimensional information (such as the number of rows, 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 number of first thread blocks is calculated. The number of first 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 element-wise operations can be calculated as the number of first thread blocks based on the total number of elements of the input tensors corresponding to multiple input tensors. The more the total number of elements of the input tensors corresponding to multiple input tensors, the more first thread blocks there are.

[0093] Furthermore, in step S101-3, it is necessary to obtain the hardware resource information of the GPU device. This process involves calling the API interfaces provided by the GPU device at runtime to obtain the specific hardware resource information of the current GPU device. For example, APIs such as CUDA can be used to query the GPU device's hardware resource information, such as the number of cores, maximum number of threads, and maximum number of thread blocks.

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

[0095] Finally, in step S101-5, the final number of thread blocks is determined. Combining the first and second thread block counts, a suitable value is selected by comparing them. The actual number of thread blocks needed during element-wise operations is determined, while ensuring that it does not exceed the second thread block count.

[0096] The above embodiments allow for the detailed determination of the number of thread blocks required for element-wise operations, ensuring the efficient and rational use of thread block computational resources on the GPU. The goal of each step is to systematically evaluate the characteristics of the input data, combine this with the GPU hardware capabilities, and ultimately derive the optimal execution strategy. This approach ensures efficient parallel computing and is suitable for scenarios requiring the processing of large amounts of data.

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

[0098] Step S101-6: Pre-set the number of threads in each thread block and the number of thread blocks that can be processed on each stream processor;

[0099] Step S101-2 includes: Step S101-2-1, determining the number of the first thread blocks based on the number of threads in each thread block and the number of elements in each input tensor among the plurality of input tensors;

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

[0101] Step S101-5 includes: Step S101-5-1, taking the minimum value between the first thread block number and the second thread block number as the number of thread blocks required by the plurality of input tensors in the element-wise operation process.

[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. For 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 can be 32.

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

[0104] Furthermore, step S101-4 specifically includes step S101-4-1, in which the number of second thread blocks is determined based on the number of stream processors, the maximum number of threads per stream processor, and the number of threads per thread block.

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

[0106] Based on this, a formula for calculating the number of thread blocks required during element-wise operations on multiple input tensors is shown below:

[0107]

[0108] The above calculation formula contains 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] In the first part, numel represents the number of elements in each input tensor among multiple input tensors, and threads_per_block represents the number of threads in each thread block. The "threads_per_block-1" in the numerator of the first part is used to round up the calculation result. If there are 1000 elements and each thread block has 256 threads, then the number of first thread blocks required is 4.

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

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

[0114] This calculation formula can guarantee that the minimum number of thread blocks is 1. A reasonable number of thread blocks needs to be calculated based on the number of elements, while taking into account the hardware specifications and not exceeding the resource limitations of the GPU device itself.

[0115] Optionally, the method further includes:

[0116] The acquired hardware resource information is stored in the internal memory;

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

[0118] Specifically, in use, you only need to pass the number of elements to the smart thread block count calculator. During its first initialization, it detects the GPU's hardware resource information and stores it in its internal memory. In subsequent calls, it automatically calculates the number of thread blocks based on the number of elements and the hardware resource information. This method reduces the number of hardware accesses and improves code execution efficiency.

[0119] The above embodiments allow for the reasonable calculation of the number of thread blocks required to perform element-wise operations on a GPU. Each step is based on previous calculations and the hardware characteristics of the GPU, ensuring that the parallel capabilities of the GPU are fully utilized during execution without exceeding its maximum processing power. This systematic design helps optimize performance, especially when processing large-scale data, maximizing the computational efficiency of element-wise operations.

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

[0121] Step S103-1-1: Predefine the template structure. The template structure has a type alias inside, and the type alias points to the memory block corresponding to the parameter of the template structure.

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

[0123] Step S103-1-3: Predefine an array, the array corresponding to the element size and element type in the parameters of the template structure.

[0124] Specifically, in one embodiment, considering that the computing units of modern GPUs operate much faster than storage devices, the reading and writing of elements often becomes a performance bottleneck. Therefore, steps S103-1-1 to S103-1-3 are included before acquiring multiple elements.

[0125] This application proposes an innovative vector type template, VectorType, and its implementation, which combines the features of template metaprogramming and unions to achieve efficient storage and access of fixed-size data sequences.

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

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

[0128] This design aims to provide a flexible and memory-aligned optimized storage solution. Below is a pseudocode example of defining 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 defined in the template structure VectorTypeProto. This allows for easy access via a simple VectorType.<T,number> The operation is used to reference a memory block of a specific size and type, aliased as type. Below is a pseudocode example of defining the type alias VectorType.

[0134] template<typename T,int number>

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

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

[0137] This design allows VectorPack to be processed as a whole, and also allows for easy access to multiple elements within a template union. Below is a pseudocode example of defining a template union VectorPack.

[0138] template<typename T,int number>

[0139] union VectorPack{

[0140] VectorType<T,number> data;

[0141] T elements[number];

[0142] };

[0143] The above embodiments utilize aligned storage, a design that improves access speed by allowing data to be accessed as a whole or as individual elements. Furthermore, type aliases simplify the use of template structs for users. This efficient mechanism for storing and accessing fixed-size memory blocks reduces the number of accesses in GPU computing tasks while providing a flexible interface for easy data access and processing. By defining VectorType and VectorPack, developers can more easily manage and use memory, improving program performance.

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

[0145] Step S103-2-1: Determine the number of elements that each thread needs to process based on the number of elements contained in the multiple input tensors and the number of thread blocks.

[0146] Obtaining the plurality of elements in step S103 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: Using the template union, multiple elements in the obtained memory block are packaged 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 plurality of elements in step S103, step S103-2-1 is further included, which determines the number of elements that each thread needs to process based on the number of elements contained in the plurality of input tensors and the number of thread blocks, and uses the obtained number of elements that each thread needs to process as the integer number in the above template, and uses the type of the element as the type T in the above template.

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

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

[0153] In step S103-3-1, for each thread, based on the number of elements to be processed and the template structure, the memory block pointed to by the type alias of the template structure is accessed. In step S103-3-2, multiple elements in the obtained memory block are packaged using a template union to obtain multiple elements. Using a template union to encapsulate multiple elements allows for more flexible handling of elements of different types. In step S103-3-3, array indexing is used to access specific elements in the template union for subsequent element-wise operations. By placing the elements in the template union in an array, threads can directly access and perform operations using index values.

[0154] The above embodiments first determine the number of elements each thread needs to process based on the input volume and the number of thread blocks. Each thread accesses the corresponding memory block based on its ID, reads the data, and then packages it using a template union. Finally, elements in the union are accessed and manipulated using array indexing, improving the flexibility and efficiency of access. By effectively managing and accessing data, the parallel processing performance of the GPU is improved, ensuring that element-wise computational tasks can be completed more efficiently.

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

[0156] Step S103-4-1: Based on the shape and stride of the multiple input tensors, and the shape and stride of the output tensor obtained after element-wise operation of the multiple input tensors, determine the index calculation method of the elements in the kernel function;

[0157] Step S103-4-1-1: When both the input tensor and the output tensor are densely stored and have the same shape, the index calculation method is as follows: obtain the elements in the input tensor using the thread number corresponding to the multiple threads as the index value, and write the elements into the output tensor using the thread number corresponding to the multiple threads as the index value.

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

[0159] Step S103-4-1-3: When both the input tensor and the output tensor are non-contiguous storage tensors, the index calculation method is as follows: for each input tensor and each output tensor, the index value of the element to be acquired 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 to the element-wise operation operator may contain different numbers of elements and memory layouts. For different situations, the index calculation method of the kernel function is different. This application proposes an adaptive kernel scheduler KernelLauncher that automatically selects efficient kernel functions.

[0161] Figure 4 This is a schematic diagram of a tensor type classification process provided in an embodiment of this application.

[0162] like Figure 4 As shown, firstly, in step S103-4-1, the type of the input tensors needs to be determined based on the shape and stride of multiple input / output tensors, categorized into three types: dense and identical in shape, contiguous storage, and others. The kernel scheduler first classifies the input and output tensors according to their shape and stride, classifying them into these three types. It uses the KernelLauncher's IsDense interface to determine if the tensor is densely stored, the IsSame interface to determine if two tensors have the same shape and stride, and the IsContiguous interface to determine if the tensor is a contiguous storage type. After confirming the type, the corresponding kernel function's index calculation method is determined for element access.

[0163] Among them, "dense and identical shape" means that the input tensor and the output tensor are both densely stored in the storage device and have the same shape. When both the input and output tensors are densely stored and their shape (dimensions and dimensions) is exactly the same as the stride, they are classified as "dense and identical shape".

[0164] Contiguous storage means that at least one tensor (input or output) is stored contiguously, but does not satisfy the condition of being dense and having the same shape.

[0165] Other types: Cases other than dense and identical shapes and contiguous storage.

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

[0167] Furthermore, when there are multiple rounds of iterative calculations, the linear index number of each iteration is the thread number + (number of iterations - 1) * total number of threads, and its code form 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 not densely stored or have different shapes, and at least one input tensor or output tensor is stored contiguously, the kernel function calculates the index as follows: for the input tensor or output tensor stored contiguously, the thread number corresponding to multiple threads is used as the index value to obtain or write the element corresponding to the index value corresponding to the thread number; for tensors stored non-contiguously other than the input tensor or output tensor stored contiguously, the input tensor or output tensor stored contiguously is used as a reference, and the index value of the non-contiguous element is calculated according to the shape and stride of the non-contiguous tensor, and the corresponding element is accessed according to the calculated index value.

[0174] Furthermore, the index calculation process for contiguous storage involves division operations on the stride, which can accumulate and reduce the overall computational performance. Therefore, a fast division algorithm can be introduced, which wraps the stride in the FastDivMod module and accelerates the calculation through the div() and mod() interfaces provided by FastDivMod.

[0175] In step S103-4-1-3, when both the input tensor and the output tensor are non-contiguous storage tensors, the index calculation method of the corresponding kernel function is as follows: for each input tensor and each output tensor, the index value of the element 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 determine how to accurately access elements in the input and output tensors within the GPU kernel. Different cases are categorized based on conditions to ensure correct and efficient data processing, and the corresponding design provides a fundamental guarantee for subsequent computations. This flexibility in indexed computation can adapt to various data structure storage methods to achieve efficient parallel computing.

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

[0178] Specifically, in one embodiment, the image processing method based on element-wise operations can be applied to image processing tasks on GPU devices. 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 in the image to be processed.

[0179] The following example, image brightness adjustment, illustrates the application of element-wise computation-based image processing methods to image processing tasks on GPU devices.

[0180] By adjusting the overall brightness of the image, the image can be made brighter or darker.

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

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

[0183] The formula is: 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 blurring, and edge detection.

[0185] By combining the image processing method based on element-wise operations with specific image processing tasks through the above embodiments, the computational efficiency of element-wise operations in application scenarios that process large-scale image data can be improved.

[0186] This application significantly improves computational efficiency and performance in the implementation of element-wise arithmetic operators on GPU devices, with the following specific benefits:

[0187] 1. Dynamically optimize GPU resource utilization:

[0188] By dynamically analyzing the GPU's hardware resource information (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, avoids resource waste, reduces contention between threads, and improves 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. These elements can be retrieved all at once using multiple predefined templates, significantly reducing the number of accesses to storage devices. Furthermore, the vectorized storage and access method makes data more compact in memory, reducing the probability of cache misses and effectively alleviating bandwidth bottlenecks, thereby accelerating data processing speed.

[0191] 3. Adaptive kernel scheduling:

[0192] Based on the total number of elements in the input data and the characteristics of the memory layout, the system intelligently selects and executes the most suitable index calculation method in the kernel function. This precise scheduling mechanism ensures that the processing of each tensor type can find the optimal computation path, improving overall execution efficiency.

[0193] Furthermore, during index calculation, especially for contiguous storage data, efficient division algorithms (such as the FastDivMod module) are introduced. This algorithm significantly reduces the time consumed by division operations, optimizes the performance of index calculation, and thus improves overall computational efficiency.

[0194] 4. Improve code maintainability and scalability:

[0195] This application employs techniques such as template metaprogramming, template unions, and type aliases to make the code more flexible, easier to maintain, and more extensible. By defining template structures and unions, it achieves unified processing of different data types, reduces code duplication, and improves code reusability.

[0196] 5. Improve user experience:

[0197] This reduces users' reliance on and understanding of GPU devices. Users only need to provide the total number of elements in the input tensor, and the intelligent thread block number calculator will automatically calculate the optimal number of thread blocks, lowering the barrier to entry for users.

[0198] 6. Wide range of application scenarios:

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

[0200] In summary, this application significantly improves the computational efficiency and performance of element-wise arithmetic operators on GPU devices through a series of innovative and optimization measures.

[0201] Figure 5 This is a schematic diagram of the framework of an image processing device based on element-wise operations according to an embodiment of this application, as shown below. Figure 5 As shown.

[0202] Based on the same inventive concept, another embodiment of this application also provides an image processing apparatus based on element-wise operations, the apparatus comprising:

[0203] The thread block number determination module 11 is used to determine the image to be processed as an input tensor, and to determine the number of thread blocks that the multiple input tensors need to use in the element-wise operation process by taking the pixel value of each pixel in the image to be processed as each element in the input tensor, the number of elements in 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.

[0204] Kernel function determination module 12 is used to determine the kernel function of the element-wise 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, and a predefined template structure. 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 element-wise operations on the multiple input tensors to obtain the image processing result.

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

[0207] The element count determination unit is used to determine the number of elements in each of the plurality of input tensors based on the shape of the plurality of input tensors;

[0208] The first thread block number calculation unit is used to calculate the number of the first thread blocks based on the number of elements in each of the plurality of input tensors;

[0209] The hardware resource information acquisition unit is used to call the API interface provided by the GPU device during runtime to obtain the hardware resource information of the GPU device;

[0210] The second thread block number calculation unit is used to calculate the second thread block number based on the hardware resource information of the GPU device. The second thread block number is the maximum number of thread blocks that the GPU device can provide.

[0211] The thread block number determination unit is used to determine the number of thread blocks required for the plurality of input tensors in the element-wise operation process 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 per stream processor; the device further includes:

[0213] The setting unit is used to pre-set the number of threads for each thread block and the number of thread blocks that can be processed on each stream processor before determining the number of thread blocks that the plurality of input tensors need to use in the element-wise operation process based on the number of the first thread blocks and the number of the second thread blocks.

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

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

[0216] The thread block number determination unit includes a thread block number determination subunit, which is used to take the minimum value between the first thread block number and the second thread block number as the number of thread blocks required by the plurality of input tensors in the element-wise operation process.

[0217] Optionally, the device further includes:

[0218] The first definition unit is used to predefine the template structure before obtaining the plurality of elements. The template structure has a type alias inside, and the type alias points to the memory block corresponding to the parameter of the template structure.

[0219] The second definition unit is used to predefine a template union, which is used to store the elements accessed in the memory block in a packaged form;

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

[0221] Optionally, the device further includes:

[0222] The unit for determining the number of elements per thread is used to determine the number of elements each thread needs to process based on the number of elements in the multiple input tensors and the number of thread blocks before acquiring the multiple elements.

[0223] The execution module 13 includes:

[0224] The memory block access unit is used to access the memory block pointed to by the type alias for each thread, based on the number of elements that each thread needs to process and the template structure.

[0225] A packaging unit is used to package multiple elements in the obtained memory block using the template union to obtain the multiple elements;

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

[0227] Optionally, the device further includes:

[0228] The index calculation method determination module is used to determine the index calculation method of the elements in the kernel function before obtaining the plurality of elements, based on the shape and stride of the plurality of input tensors and the shape and stride of the output tensor obtained after element-wise operation of the plurality of input tensors;

[0229] The first index calculation method determination unit is used to determine the index calculation method when the input tensor and the output tensor are both densely stored and have the same shape: obtaining the elements in the input tensor with the thread number corresponding to the multiple threads as the index value, and writing the elements into the output tensor with the thread number corresponding to the multiple threads as the index value.

[0230] The second index calculation method determination unit is used when the input tensor and the output tensor are not densely stored or have different shapes, and at least one input tensor or output tensor is continuously stored. The index calculation method is as follows: for continuously stored input tensors or output tensors, the thread number corresponding to the multiple threads is used as the index value to obtain or write the element corresponding to the index value corresponding to the thread number; for non-contiguous stored tensors other than the continuously stored input tensors or output tensors, the continuously stored input tensors or output tensors are used as a reference, and the index value of the non-contiguous stored elements is calculated according to the shape and stride of the non-contiguous stored tensors.

[0231] The third index calculation method determination unit is used to calculate the index value of the element to be acquired or written in each input tensor and each output tensor according to the shape and stride when both the input tensor and the output tensor are non-contiguous storage tensors.

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

[0233] Optionally, the device further includes:

[0234] A storage unit is used to store 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 it is necessary to obtain the hardware resource information again.

[0236] Based on the same inventive concept, another embodiment of this application 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-wise operations as described in any of the above embodiments.

[0237] Among them, electronic devices refer to Figure 6 , Figure 6 This is a schematic diagram of an electronic device provided in an embodiment of this 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 for communication. The memory 610 stores a computer program that can run on the processor 620 to implement the steps in the image processing method based on element-wise operations disclosed in the above embodiments of this application.

[0238] Based on the same inventive concept, another embodiment of this application also provides a computer program product, including a computer program that is executed by a processor as the image processing method based on element-wise operations as described in any of the above embodiments.

[0239] Based on the same inventive concept, another embodiment of this application provides a computer-readable storage medium having a computer program stored thereon, wherein the program, when executed by a processor, implements the image processing method based on element-wise operations as described in any of the above embodiments.

[0240] As the apparatus is basically similar to the method embodiment, it is described in a relatively simple way. For relevant details, please refer to the description of the method embodiment.

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

[0242] Those skilled in the art will understand that embodiments of this application can be provided as methods, apparatus, or computer program products. Therefore, embodiments of this application can take the form of entirely hardware embodiments, entirely software embodiments, or embodiments combining software and hardware aspects. Furthermore, embodiments of this application can take the form of computer program products 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] This application describes embodiments with reference to flowchart illustrations and / or block diagrams of methods, terminal devices (systems), and computer program products according to embodiments of this application. It should be understood that each block of the flowchart illustrations and / or block diagrams, and combinations of blocks in the flowchart illustrations and / or block diagrams, can be implemented by computer program instructions. These computer program instructions can be provided to a processor of a general-purpose computer, special-purpose computer, embedded processor, or other programmable data processing terminal device to produce a machine, such that the instructions, which execute via the processor of the computer or other programmable data processing terminal device, generate instructions for implementing the flowchart illustrations. Figure 1 One or more processes and / or boxes Figure 1 A device that provides the functions specified in one or more boxes.

[0244] These computer program instructions may also be stored in a computer-readable storage medium that can direct a computer or other programmable data processing terminal device to operate in a particular manner, such that the instructions stored in the computer-readable storage medium produce an article of manufacture including instruction means, which are implemented in a process Figure 1 One or more processes and / or boxes Figure 1 The 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 equipment, causing a series of operational steps to be performed on the computer or other programmable terminal equipment to produce a computer-implemented process, thereby providing instructions that execute on the computer or other programmable terminal equipment for implementing the process. Figure 1 One or more processes and / or boxes Figure 1 The steps of the function specified in one or more boxes.

[0246] Although preferred embodiments of the present application have been described, those skilled in the art, upon learning the basic inventive concept, can make other changes and modifications to these embodiments. Therefore, the appended claims are intended to be interpreted as including the preferred embodiments as well as all changes and modifications falling within the scope of the embodiments of the present application.

[0247] Finally, it should be noted that in this document, relational terms such as "first" and "second" are used only to distinguish one entity or operation from another, and do not necessarily require or imply any such actual relationship or order between these entities or operations. Furthermore, the terms "comprising," "including," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or terminal device that comprises a list of elements includes not only those elements but also other elements not expressly listed, or elements inherent to such a process, method, article, or terminal device. Without further limitations, an element defined by the phrase "comprising one..." does not exclude the presence of other identical elements in the process, method, article, or terminal device that includes the element.

[0248] The foregoing has provided a detailed description of an image processing method, apparatus, device, and medium based on element-wise operations provided in this application. Specific examples have been used to illustrate the principles and implementation methods of this application. The descriptions of the above embodiments are only for the purpose of helping to understand the method and core ideas of this application. At the same time, for those skilled in the art, there will be changes in the specific implementation methods and application scope based on the ideas of this application. Therefore, the content of this specification should not be construed as a limitation of this application.

Claims

1. An image processing method based on element-wise operations, characterized in that, The method includes: The image to be processed is determined as the input tensor. The pixel value of each pixel in the image to be processed is assigned to each element in the input tensor. Based on the number of elements in the multiple input tensors that need to participate in the operation, and the hardware resource information of the GPU device, the number of thread blocks required for the multiple input tensors to be used in the element-wise operation is determined. Based on the number of thread blocks, the number of elements, and the memory layout of the multiple elements, and using a predefined template structure, determine the kernel function of the element-wise operation operator corresponding to the multiple input tensors that need to participate in the operation. The parameters of the template structure include the element type and the number of elements. Based on the number of thread blocks, multiple threads are initiated so that the multiple threads, according to the kernel function, obtain the multiple elements and perform element-wise operations on the multiple input tensors to obtain the image processing results; Before acquiring the multiple elements, the method further includes: determining the index calculation method for the elements in the kernel function based on the shape and stride of the multiple input tensors, and the shape and stride of the output tensor obtained after element-wise operations on the multiple input tensors; when the input tensors and the output tensors are both densely stored and have the same shape, the index calculation method is: acquiring elements in the input tensors using the thread numbers corresponding to the multiple threads as index values, and writing elements to the output tensors using the thread numbers corresponding to the multiple threads as index values; when the input tensors and the output tensors are not densely stored or have different shapes, and at least one input tensor or output tensor is stored contiguously, the index calculation method is... For contiguous input or output tensors, the element corresponding to the index value of the thread number corresponding to the multiple threads is retrieved or written using the thread number as the index value. For non-contiguous tensors other than the contiguous input or output tensors, the index value of the non-contiguous elements is calculated according to the shape and stride of the non-contiguous tensor, with the contiguous input or output tensor as a reference. When both the input and output tensors are non-contiguous tensors, the index calculation method is as follows: for each input tensor and each output tensor, the index value of the element to be retrieved or written in each input tensor and each output tensor is calculated according to the shape and stride.

2. The element-wise operation based image processing method according to claim 1, characterized in that, The determination of the number of thread blocks required for element-wise computation of the multiple input tensors, based on the number of elements contained in the multiple input tensors participating in the computation and the hardware resource information of the GPU device, includes: Based on the shape of the plurality of input tensors, determine the number of elements in each of the plurality of input tensors; The number of first thread blocks is calculated based on the number of elements in each of the multiple input tensors; Call the API interface provided by the GPU device at runtime to obtain the hardware resource information of the GPU device; Based on the hardware resource information of the GPU device, the number of second thread blocks is calculated, which is the maximum number of thread blocks that the GPU device can provide. Based on the number of the first thread block and the number of the second thread block, the number of thread blocks required for the multiple input tensors to be used in the element-wise operation process is determined.

3. The image processing method based on element-wise operations according to claim 2, characterized in that, The hardware resource information includes: the number of stream processors and the maximum number of threads per stream processor; before determining the number of thread blocks required for the multiple input tensors during element-wise operations based on the first thread block number and the second thread block number, it also includes: The number of threads in each thread block and the number of thread blocks that can be processed on each stream processor are preset. The step of calculating the number of first thread blocks based on the number of elements in each of the plurality of input tensors includes: determining the number of first thread blocks based on the number of threads in each thread block and the number of elements in each of the plurality of input tensors; The step of calculating the number of second thread blocks based on the hardware resource information of the GPU device includes: determining the number of second thread blocks based on the number of stream processors, the maximum number of threads per stream processor, and the number of threads per thread block; Determining the number of thread blocks required for the multiple input tensors during element-wise operations based on the number of the first thread blocks and the number of the second thread blocks includes: taking the minimum value between the number of the first thread blocks and the number of the second thread blocks as the number of thread blocks required for the multiple input tensors during element-wise operations.

4. The image processing method based on element-wise operations according to claim 1, characterized in that, Before obtaining the plurality of elements, the method further includes: The template structure is predefined, and the template structure has type aliases internally, which point to the memory blocks corresponding to the parameters of the template structure. A template union is predefined, which is used to store the accessed elements in the memory block in a packaged form; An array is predefined, which corresponds to the element size and element type in the parameters of the template structure.

5. The image processing method based on element-wise operations according to claim 4, characterized in that, Before obtaining the plurality of elements, the process also includes: The number of elements to be processed by each thread is determined based on the number of elements contained in the multiple input tensors and the number of thread blocks. The acquisition of the plurality of elements includes: 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; The multiple elements in the obtained memory block are packaged using the template union to obtain the multiple elements; Multiple elements in the template union can be accessed by index value through the array.

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

7. An image processing apparatus based on element-wise operations, characterized in that, The device includes: The 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 that the multiple input tensors need to use in the element-wise operation process by taking the pixel value of each pixel in the image to be processed as each element in the input tensor, the number of elements in 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. The kernel function determination module is used to determine the kernel function of the element-wise 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, and a predefined template structure. The parameters of the template structure include the element type and the number of elements. The 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 element-wise operations on the multiple input tensors to obtain the image processing result; The device further includes: The index calculation method determination module is used to determine the index calculation method of the elements in the kernel function before obtaining the plurality of elements, based on the shape and stride of the plurality of input tensors and the shape and stride of the output tensor obtained after element-wise operation of the plurality of input tensors; The first index calculation method determination unit is used to determine the index calculation method when the input tensor and the output tensor are both densely stored and have the same shape: obtaining the elements in the input tensor with the thread number corresponding to the multiple threads as the index value, and writing the elements into the output tensor with the thread number corresponding to the multiple threads as the index value. The second index calculation method determination unit is used when the input tensor and the output tensor are not densely stored or have different shapes, and at least one input tensor or output tensor is continuously stored. The index calculation method is as follows: for continuously stored input tensors or output tensors, the thread number corresponding to the multiple threads is used as the index value to obtain or write the element corresponding to the index value corresponding to the thread number; for non-contiguous stored tensors other than the continuously stored input tensors or output tensors, the continuously stored input tensors or output tensors are used as a reference, and the index value of the non-contiguous stored elements is calculated according to the shape and stride of the non-contiguous stored tensors. The third index calculation method determination unit is used to calculate the index value of the element to be acquired or written in each input tensor and each output tensor according to the shape and stride when both the input tensor and the output tensor are non-contiguous storage tensors.

8. An electronic device, characterized in that, It includes 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-wise operations as described in any one of claims 1-6.

9. A computer-readable storage medium, characterized in that, It stores a computer program, wherein the computer program, when executed by a processor, implements the image processing method based on element-wise operations as described in any one of claims 1-6.

Images