Image scaling method, processor system, graphics processor and electronic equipment

Abstract

The invention provides an image scaling method, a processor system, a graphics processor and electronic equipment. The method comprises the following steps: acquiring a to-be-scaled image and a target resolution of the to-be-scaled image; the target resolution comprises a height H and a width W; determining the length and the width of a single kernel function block according to the H, the W and the maximum thread number K of the single kernel function block in the GPU; determining the length and the width of a grid in the GPU according to the H, the W and the length and the width of a kernel function block; and starting a CUDA kernel of the GPU according to the length and the width of the grid and the length and the width of the single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm to scale the image to be scaled, and a target image with the target resolution is obtained. According to the scheme provided by the invention, more suitable length and width of the grid and the length and width of the kernel function block can be obtained, so that the calculation efficiency of an image scaling algorithm is further improved.

Description

Technical Field

[0001] The present application relates to the field of graphics processor processing technology, and in particular to an image scaling method, device, electronic device and computer-readable storage medium. Background Art

[0002] Image resizing is a basic operation in computer vision and image processing, which involves changing the size of an image. Image resizing is usually implemented through an image resizing algorithm. Currently, image resizing algorithms are usually run by the CPU (central processing unit), but image resizing algorithms are computationally intensive tasks. The parallel processing capabilities and high bandwidth of GPUs (graphics processing units) can significantly improve algorithm performance and efficiency. To this end, the tasks in the image resizing algorithm can be processed in parallel on the GPU based on the CUDA (Compute Unified Device Architecture, a parallel computing platform and programming model) kernel of the GPU.

[0003] However, in the GPU, the CUDA kernel is executed based on a grid, which is composed of multiple kernel function blocks. In theory, larger kernel function blocks can better utilize the parallel computing capabilities of the GPU, thereby improving the execution efficiency of the program. However, too large a kernel function block may cause a decrease in memory access efficiency, thereby increasing access latency and affecting the execution efficiency of the program. The size of the grid determines the number of parallel tasks. The larger the grid, the higher the degree of parallelism of the computing tasks, thereby increasing the execution speed of the program. However, since each kernel function block requires a certain amount of GPU resources to execute, an overly large grid will lead to insufficient GPU resources, thereby affecting the execution of the program.

[0004] Currently, the kernel function block size (blockDim) and the grid size (gridDim) are directly specified by the developer when writing the CUDA program. As a result, when faced with different image scaling requirements, the kernel function block size (blockDim) and the grid size (gridDim) often cannot match the actual requirements, thus affecting the improvement of the computational efficiency of the image scaling algorithm. Summary of the invention

[0005] The purpose of the embodiments of the present application is to provide an image scaling method, a processor system, a graphics processor and an electronic device to solve the problem in the related art that the size of the kernel function block and the size of the grid often cannot match the actual needs, thereby affecting the improvement of the computational efficiency of the image scaling algorithm.

[0006] An embodiment of the present application provides an image scaling method, which includes: obtaining an image to be scaled and a target resolution of the image to be scaled; the target resolution includes a height H and a width W; determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have; determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block; starting a CUDA kernel of the GPU according to the length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm to scale the image to be scaled, and obtains a target image with the target resolution.

[0007] In the above implementation, when the GPU is used to scale the image, the kernel function length and width and the length and width of the grid that are compatible with the current scaling target are first determined according to the target resolution of the image to be scaled (i.e., the height H and width W of the target image), and then the CUDA kernel of the GPU is started based on the determined length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm based on the determined grid and kernel function block to scale the image to be scaled, and obtains a target image with a target resolution. In this way, each time the image to be scaled is scaled, the length and width of the grid and the length and width of the kernel function block used can be more in line with actual needs, and will not be too large or too small, which can further improve the computational efficiency of the image scaling algorithm.

[0008] Furthermore, when the kernel function block does not have a depth dimension, the length and width of a single kernel function block are determined based on H, W, and the maximum number of threads K that a single kernel function block in the GPU can have, including: calculating the length of a single kernel function block according to the formula sqrt(W*K / H); calculating the width of a single kernel function block according to the formula sqrt(H*K / W); wherein sqrt is the square root.

[0009] It can be understood that, assuming that the length of a single kernel function block is x and the width is y, the length of the grid is W / x (denoted as gridSize.x), and the width of the grid is H / y (denoted as gridSize.y). In the absence of a depth dimension, the total number of threads of the grid is x*y*gridSize.x*gridSize.y=x*(W / x+1)*y*(H / y+1)=(W+x)*(H+y)=H*W+W*y+H*x+x*y. Among them, assuming that one thread can process one pixel, the minimum number of threads required to output the target image of H*W is H*W, and W*y+H*x+x*y is the number of redundant threads. Therefore, the smaller W*y+H*x+x*y is, the smaller the kernel function block and the grid size can be while meeting the processing requirements for the scaled image, thereby avoiding the problem of the kernel function block and the grid size being too large or too small.

[0010] Then, in the absence of a depth dimension, x*y represents the number of threads in the kernel function block. In order to make full use of GPU resources, the number of threads in the kernel function block should be as equal as possible to the maximum number of threads K that a single kernel function block can have. The K value is determined according to the hardware conditions of the GPU. For example, some GPU chip hardware determines that its single thread bundle (warp) is 32 threads (such as NVIDIA's GPU chip), then the maximum number of threads K that its kernel function block can have is 32*32=1024; for another example, some GPU chip hardware determines that its single thread bundle (warp) is 64 threads, then the maximum number of threads K that its kernel function block can have is 64*64=4096.

[0011] Then, the factor that affects the size of W*y+H*x+x*y is W*y+H*x. Since x*y=K, W*y+H*x can be expressed as W*(K / x)+H*x. Since W*(K / x) is greater than 0, and H*x is also greater than 0, W*(K / x)+H*x≥2*sqrt(W*(K / x)*x*H), then W*(K / x)+H*x=2*sqrt(W*(K / x)*x*H) only when W*(K / x)=x*H. That is, when W*(K / x)=x*H, W*y+H*x is the smallest. So in order to calculate x, we can solve W*(K / x)=x*H.

[0012] Based on the above reasoning, it can be proved that the length of a single kernel function block calculated by the formula sqrt(W*K / H) and the width of a single kernel function block calculated by the formula sqrt(H*K / W) can minimize the kernel function block and grid size while meeting the processing requirements of the scaled image, thereby avoiding the problem of the kernel function block and grid size being too large or too small, and further improving the computational efficiency of the image scaling algorithm.

[0013] Furthermore, when the kernel function block has a depth dimension, the length and width of a single kernel function block are determined according to H, W, and the maximum number of threads K that a single kernel function block in the GPU can have, including: calculating the length of a single kernel function block according to the formula sqrt(W*M / H); calculating the width of a single kernel function block according to the formula sqrt(H*M / W); wherein sqrt is the square root, M is equal to the ratio of K to the depth value s of the depth dimension of the kernel function block; and s is a positive integer greater than or equal to 1.

[0014] Referring to the previous analysis, since the kernel function block has a depth dimension, the maximum number of threads K that the kernel function block can have will be equal to x*y*s, which leads to x*y=K / s=M. Then W*y+H*x can be expressed as W*(M / x)+H*x. In addition, W*(M / x) is greater than 0, and H*x is also greater than 0, so W*(M / x)+H*x≥2*sqrt(W*(M / x)*x*H), then and only when W*(M / x)=x*H, W*(M / x)+H*x=2*sqrt(W*(M / x)*x*H). That is, when W*(M / x)=x*H, W*y+H*x is the smallest, so in order to calculate x, it is sufficient to solve W*(M / x)=x*H.

[0015] Therefore, when the kernel function block has a depth dimension, based on the above implementation method, the length of a single kernel function block is calculated according to the formula sqrt(W*M / H), and the width of a single kernel function block is calculated according to the formula sqrt(H*M / W). This can minimize the kernel function block and grid size while meeting the processing requirements for the scaled image, thereby avoiding the problem of the kernel function block and grid size being too large or too small, and can further improve the computational efficiency of the image scaling algorithm.

[0016] Furthermore, the s is a value determined according to a preset number of images to be processed in parallel, or the s is a value determined according to the number of channels of the image to be scaled.

[0017] In the above implementation, the depth dimension value s of the kernel function block is determined according to the number of images to be processed in parallel. After the CUDA kernel of the GPU is started according to the length and width of the grid and the length and width of the kernel function block, the number of images to be processed in parallel can be scaled at the same time to improve the scaling efficiency. The depth dimension value s of the kernel function block is determined according to the number of channels of the image to be scaled. After the CUDA kernel of the GPU is started according to the length and width of the grid and the length and width of the kernel function block, the data of each channel of the image to be scaled can be processed at the same time to meet the processing requirements of the image to be scaled.

[0018] Further, after determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have, and before determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block, the method further includes:

[0019] If one of the length and width of the kernel function block is greater than the maximum number of threads k in the thread warp in the GPU, and the other is less than k, then the one of the length and width of the kernel function block that is greater than k is set to k, and the other is kept unchanged.

[0020] When a kernel function block can contain an integer number of thread bundles in one direction, it means that all thread bundles contained in the kernel function block will be complete, that is, the thread resources in the thread bundle will not be wasted because they are not allocated to the kernel function block. For this reason, in the above implementation, when it is found that one of the length and width of the kernel function block is greater than the maximum number of threads k in the thread bundle in the GPU, and the other is less than k, it means that there may be some thread bundles that cannot be completely contained in the kernel function block. For this reason, the one of the length and width of the kernel function block that is greater than k is set to k, and the other is kept unchanged. This can ensure that the product of the length and width of the kernel function block (the product of the length, width and depth dimensions of the kernel function block when there is a depth dimension) will not be greater than the maximum number of threads K that a single kernel function block can have, while ensuring that all thread bundles contained in the kernel function block will be complete, thereby maximizing the thread utilization and improving the overall performance of the GPU.

[0021] Further, after determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have, and before determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block, the method further includes:

[0022] If both the length and width of the kernel function block are smaller than the maximum number of threads k in the thread bundle in the GPU, the larger one of the length and width of the kernel function block is set to the k; determine whether the product of the other of the length and width of the kernel function block and k is greater than M; if the product of the other of the length and width of the kernel function block and k is greater than M, reduce the value of the other of the length and width of the kernel function block so that the product of the other of the length and width of the kernel function block after reduction and the k is less than or equal to the M; if the product of the other of the length and width of the kernel function block and k is less than or equal to the M, keep the value of the other of the length and width of the kernel function block unchanged.

[0023] It can be understood that when the value of the depth dimension is greater than or equal to 2, the length and width of the kernel function block may be less than the maximum number of threads k in the thread bundle in the GPU. At this time, by setting the larger of the length and width of the kernel function block to k, it can be ensured that all thread bundles contained in the kernel function block will be complete, and the magnitude of the change is relatively small, which can reduce the impact of the modification of the length or width of the kernel function block on the degree of adaptation to the image to be scaled. After that, it is further determined whether the product of the other of the length and width of the kernel function block and k is less than or equal to M. If it is less than or equal to M, it means that the increase in the length or width of the kernel function block will not make the kernel function block unable to be implemented in the GPU. At this time, keeping the other unchanged can reduce the impact on the degree of adaptation to the image to be scaled. If it is greater than M, it means that the increase in the length or width of the kernel function block will make the kernel function block unable to be implemented in the GPU. At this time, the value of the other of the length and width of the kernel function block is reduced so that the product of the other of the length and width of the reduced kernel function block and k is less than or equal to M, thereby ensuring the feasibility of the kernel function block.

[0024] Furthermore, after determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block can have in the GPU, and before determining the length and width of the grid in the GPU according to the H, the W, and the length and width of the single kernel function block, the method also includes: if at least one of the length and width of the kernel function block is equal to the maximum number of threads k of the thread bundle in the GPU, then keeping the values ​​of the length and width of the kernel function block unchanged.

[0025] In the above implementation, if at least one of the length and width of the kernel function block is equal to the maximum number of threads k in the thread bundle in the GPU, it indicates that the size of the kernel function block can ensure that all thread bundles contained in the kernel function block will be complete, so no adjustment is made, thereby avoiding the impact on the degree of adaptation of the scaled image due to modification of the length or width of the kernel function block.

[0026] Furthermore, the length and width of the grid in the GPU are determined based on the H, the W, and the length and width of a single kernel function block, including: determining a first ratio between the W and the length of the kernel function block, and rounding up when the first ratio is not an integer to obtain the length of the grid in the GPU; determining a second ratio between the H and the width of the kernel function block, and rounding up when the second ratio is not an integer to obtain the width of the grid in the GPU.

[0027] In the above implementation, by determining the first ratio between W and the length of the kernel function block, and rounding up when the first ratio is not an integer, and determining the second ratio between H and the width of the kernel function block, and rounding up when the second ratio is not an integer, it can be ensured that the grid can completely accommodate the determined kernel function block.

[0028] The embodiment of the present application also provides an image scaling device, which is applied to a graphics processor GPU, comprising:

[0029] An acquisition module, used to acquire an image to be scaled and a target resolution of the image to be scaled; the target resolution includes a height H and a width W;

[0030] A size determination module, used to: determine the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have; and determine the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block;

[0031] The processing module is used to start the CUDA kernel of the GPU according to the length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm to scale the image to be scaled to obtain a target image with the target resolution.

[0032] An embodiment of the present application also provides a processor system, including: a CPU and a GPU; the CPU is used to execute any of the above-mentioned image scaling methods to determine the length and width of the grid, the length and width of a single kernel function block and notify the GPU, so that the GPU starts the CUDA kernel according to the length and width of the grid, the length and width of a single kernel function block, and executes a preset image scaling algorithm through the CUDA kernel to scale the image to be scaled, so as to obtain a target image with the target resolution.

[0033] An embodiment of the present application also provides a graphics processor, including: a CUDA kernel, a configuration core and a video memory; the configuration core is used to execute one or more programs stored in the video memory to implement any of the above-mentioned image scaling methods.

[0034] An embodiment of the present application also provides an electronic device, including the above-mentioned processor system, or including the above-mentioned graphics processor. BRIEF DESCRIPTION OF THE DRAWINGS

[0035] In order to more clearly illustrate the technical solutions of the embodiments of the present application, the drawings required for use in the embodiments of the present application will be briefly introduced below. It should be understood that the following drawings only show certain embodiments of the present application and therefore should not be regarded as limiting the scope. For ordinary technicians in this field, other related drawings can be obtained based on these drawings without paying creative work.

[0036] Figure 1 A schematic diagram of a flow chart of an image scaling method provided in an embodiment of the present application;

[0037] Figure 2 A schematic diagram of the structure of an image scaling device provided in an embodiment of the present application;

[0038] Figure 3 A schematic diagram of the structure of a processor system provided in an embodiment of the present application;

[0039] Figure 4 A schematic diagram of the structure of a GPU provided in an embodiment of the present application. DETAILED DESCRIPTION

[0040] The technical solutions in the embodiments of the present application will be described below in conjunction with the drawings in the embodiments of the present application.

[0041] Image scaling is a basic operation in computer vision and image processing, which involves changing the size of an image. Image scaling is usually implemented through image scaling algorithms. Currently, the most widely used image processing library is OpenCV, which has CPU implementations of various mainstream image scaling algorithms. Therefore, image scaling algorithms are usually run by the CPU.

[0042] However, image scaling is a computationally intensive task. The parallel processing capability and high bandwidth of GPU can significantly improve the performance and efficiency of the algorithm. Implementing image scaling on GPU has several significant advantages over implementing it on CPU, which are mainly reflected in the following aspects:

[0043] Parallel processing capability: GPU has a large number of computing cores and can process multiple pixels at the same time. This is very beneficial for operations such as image scaling that require processing large amounts of data, while CPUs usually only have a few cores and process them in a more serial manner.

[0044] Data throughput: GPUs typically have higher bandwidth than CPUs, and can read and write large amounts of data more quickly. This is a key advantage for processing large images or performing real-time processing.

[0045] Hardware optimization: Modern GPUs have specialized hardware units such as CUDA cores (NVIDIA GPUs) or ComputeUnits (AMD GPUs) (collectively referred to as CUDA cores in this application) that are dedicated to floating-point operations and other image processing tasks, which makes it more efficient to perform these tasks on the GPU.

[0046] Algorithm parallelization: Many image processing algorithms, such as bilinear interpolation and bicubic interpolation, can be split into independent computing tasks that have no dependencies on each other. This feature is very suitable for parallel processing on GPUs, which can significantly improve performance.

[0047] Energy efficiency: For certain compute-intensive tasks, GPUs can provide a higher performance / watt ratio than CPUs, that is, they can provide higher computing power at the same energy consumption.

[0048] Scalability: For situations where multiple images or higher resolutions need to be processed, GPUs can scale performance by adding more computing cores or using multi-GPU configurations, while CPU scalability is relatively limited.

[0049] However, in the GPU, CUDA kernels are executed on a grid basis, and the grid is composed of multiple kernel function blocks. In theory, larger kernel function blocks can better utilize the parallel computing capabilities of the GPU, thereby improving the execution efficiency of the program. However, too large a kernel function block may lead to a decrease in memory access efficiency, thereby increasing access latency and affecting the execution efficiency of the program. The size of the grid determines the number of parallel tasks. The larger the grid, the higher the degree of parallelism of the computing tasks, thereby increasing the execution speed of the program. However, since each kernel function block requires a certain amount of GPU resources to execute, an overly large grid will lead to insufficient GPU resources, thereby affecting the execution of the program.

[0050] Currently, the size of the kernel function block and the size of the grid are directly specified by the developer when writing the CUDA program. As a result, when faced with different image scaling requirements, the size of the kernel function block and the size of the grid often cannot match the actual requirements, thus affecting the improvement of the computational efficiency of the image scaling algorithm.

[0051] In order to solve the problem that the size of the kernel function block and the size of the grid often cannot match the actual demand, thereby affecting the improvement of the computational efficiency of the image scaling algorithm, an image scaling method applied to a GPU is provided in an embodiment of the present application. Figure 1 As shown, Figure 1 The flowchart of the image scaling method provided in the embodiment of the present application includes:

[0052] S101: Acquire an image to be scaled and a target resolution of the image to be scaled.

[0053] In the embodiment of the present application, the target resolution may be characterized by a height H and a width W, for example, the target resolution may be expressed as H*W.

[0054] In the embodiment of the present application, the size of the target resolution is determined by business requirements. That is, when it is determined that there is an image scaling task, the image to be scaled and the corresponding target resolution of the task are confirmed.

[0055] In the embodiment of the present application, the image to be scaled may be input by a user, may be acquired from other devices, or may be an intermediate round of images generated during multiple rounds of image processing. In the embodiment of the present application, there is no restriction on the source of the image to be scaled.

[0056] In the embodiment of the present application, the image to be scaled may be a YUV image.

[0057] In the embodiment of the present application, the target resolution can be set by the user, or determined by the currently used image scaling software, or pre-set in the used image scaling algorithm. The method for obtaining the target resolution is not limited in the embodiment of the present application.

[0058] S102: Determine the length and width of a single kernel function block according to H, W, and the maximum number of threads K that a single kernel function block can have in the GPU.

[0059] It is understandable that in GPU, the processing of images is often two-dimensional, so there is no need to introduce the depth dimension, and the size of the kernel function block and the grid only involves the length and width. However, in some cases, such as when multiple images to be scaled need to be scaled in parallel, or when the image to be scaled has multiple channels (such as the image output by a convolutional layer after being processed by a neural network model), the processing of the image is often multi-dimensional. At this time, the depth dimension needs to be further introduced in the size representation of the kernel function block and the grid, and the value of the depth dimension can be predetermined according to the needs.

[0060] To this end, in a feasible implementation manner of an embodiment of the present application, when the kernel function block does not have a depth dimension, step S102 may include: calculating the length of a single kernel function block according to the formula sqrt(W*K / H), and calculating the width of a single kernel function block according to the formula sqrt(H*K / W); wherein sqrt is a square root, indicating that a square root operation is performed on the number in the brackets after sqrt.

[0061] It can be understood that the length and width of the kernel function block calculated based on the above feasible implementation method can minimize the kernel function block and grid size while meeting the processing requirements of the scaled image, thereby avoiding the problem of the kernel function block and grid size being too large or too small, and further improving the computational efficiency of the image scaling algorithm. The specific analysis is as follows:

[0062] Assuming that the length of a single kernel function block is x and the width is y, the length of the grid is W / x (denoted as gridSize.x) and the width of the grid is H / y (denoted as gridSize.y). Then, in the absence of a depth dimension, the total number of threads in the grid is x*y*gridSize.x*gridSize.y=x*(W / x+1)*y*(H / y+1)=(W+x)*(H+y)=H*W+W*y+H*x+x*y (where x*y represents the number of threads in a single kernel function block, and gridSize.x*gridSize.y reflects the number of kernel function blocks in the grid).

[0063] Assuming that one thread can process one pixel, the minimum number of threads required to output the target image of H*W is H*W, and W*y+H*x+x*y is the number of redundant threads. Therefore, the smaller W*y+H*x+x*y is, the smaller the kernel function block and grid size can be while meeting the processing requirements of the scaled image, thereby avoiding the problem of the kernel function block and grid size being too large or too small.

[0064] Then, in the absence of a depth dimension, x*y represents the number of threads in the kernel function block. In order to make full use of GPU resources, the number of threads in the kernel function block should be as equal as possible to the maximum number of threads K that a single kernel function block can have. The K value is determined according to the hardware conditions of the GPU. For example, some GPU chip hardware determines that its single thread bundle (warp) is 32 threads (such as NVIDIA's GPU chip), then the maximum number of threads K that its kernel function block can have is 32*32=1024; for another example, some GPU chip hardware determines that its single thread bundle (warp) is 64 threads, then the maximum number of threads K that its kernel function block can have is 64*64=4096.

[0065] Then, the factor that affects the size of W*y+H*x+x*y is W*y+H*x. Since x*y=K, W*y+H*x can be expressed as W*(K / x)+H*x. Since W*(K / x) is greater than 0, and H*x is also greater than 0, W*(K / x)+H*x≥2*sqrt(W*(K / x)*x*H). Then, only when W*(K / x)=x*H, W*(K / x)+H*x=2*sqrt(W*(K / x)*x*H). That is, when W*(K / x)=x*H, W*y+H*x is the smallest. Then, solving W*(K / x)=x*H, we can get x=sqrt(W*K / H). Correspondingly, substituting into y=K / x, we get y=sqrt(H*K / W).

[0066] Therefore, it can be proved that the length of a single kernel function block calculated by the formula sqrt(W*K / H) and the width of a single kernel function block calculated by the formula sqrt(H*K / W) can minimize the kernel function block and grid size while meeting the processing requirements of the scaled image, thereby avoiding the problem of the kernel function block and grid size being too large or too small, and can further improve the computational efficiency of the image scaling algorithm.

[0067] In another feasible implementation of the embodiment of the present application, when the kernel function block has a depth dimension, step S102 may include: calculating the length of a single kernel function block according to the formula sqrt(W*M / H), and calculating the width of a single kernel function block according to the formula sqrt(H*M / W). Wherein, M is equal to the ratio between K and the depth value s of the depth dimension of the kernel function block, and s is a positive integer greater than or equal to 1.

[0068] At this time, similar to the previous article, since the kernel function block has a depth dimension, the maximum number of threads K that the kernel function block can have will be equal to x*y*s, which leads to x*y=K / s=M. Then W*y+H*x can be expressed as W*(M / x)+H*x. In addition, W*(M / x) is greater than 0, and H*x is also greater than 0, so W*(M / x)+H*x≥2*sqrt(W*(M / x)*x*H), then and only when W*(M / x)=x*H, W*(M / x)+H*x=2*sqrt(W*(M / x)*x*H). That is, when W*(M / x)=x*H, W*y+H*x is the smallest, so in order to calculate x, it is sufficient to solve W*(M / x)=x*H.

[0069] Therefore, when the kernel function block has a depth dimension, based on the above implementation method, the length of a single kernel function block is calculated according to the formula sqrt(W*M / H), and the width of a single kernel function block is calculated according to the formula sqrt(H*M / W). This can minimize the kernel function block and grid size while meeting the processing requirements for the scaled image, thereby avoiding the problem of the kernel function block and grid size being too large or too small, and can further improve the computational efficiency of the image scaling algorithm.

[0070] In an embodiment of the present application, the depth value s of the depth dimension of the kernel function block can be a value determined according to a preset number of images to be processed in parallel, for example, s is equal to the number of images to be processed in parallel; or, the depth value s of the depth dimension of the kernel function block can be a value determined according to the number of channels of the image to be scaled, for example, s is equal to the number of channels of the image to be scaled.

[0071] In this way, if the depth dimension value s of the kernel function block is determined according to the number of images to be processed in parallel, then after the CUDA kernel of the GPU is started according to the length and width of the grid and the length and width of the kernel function block, the number of images to be processed in parallel can be scaled at the same time to improve the scaling efficiency. If the depth dimension value s of the kernel function block is determined according to the number of channels of the image to be scaled, then after the CUDA kernel of the GPU is started according to the length and width of the grid and the length and width of the kernel function block, the data of each channel of the image to be scaled can be processed at the same time to meet the processing requirements of the image to be scaled.

[0072] It can be understood that in a GPU, the maximum number of threads k contained in a thread warp is determined based on the hardware structure of the GPU. Therefore, when a kernel function block can completely contain each thread warp, the thread resources in each thread warp will not be wasted due to not being allocated to the kernel function block, so that the overall performance of the GPU will be higher.

[0073] To this end, in a feasible implementation of an embodiment of the present application, after determining the length and width of the kernel function block based on H, W, and the maximum number of threads K that a single kernel function block can have in the GPU, before executing the subsequent step S103, the relationship between the length and width of the kernel function block and the maximum number of threads k in the thread bundle in the GPU can also be determined.

[0074] In an optional implementation, it can be set that: if one of the length and width of the kernel function block is greater than the maximum number of threads k in the thread bundle in the GPU, and the other is less than k, then the one of the length and width of the kernel function block that is greater than k is set to the k, and the other is kept unchanged.

[0075] It can be understood that when it is found that one of the length and width of the kernel function block is greater than the maximum number of threads k in the thread bundle in the GPU, and the other is less than k, this means that there may be some thread bundles that cannot be completely included in the kernel function block. For this reason, the one of the length and width of the kernel function block that is greater than k is set to k, and the other is kept unchanged. This ensures that the product of the length and width of the kernel function block (the product of the values ​​of the length, width and depth dimensions of the kernel function block when there is a depth dimension) will not be greater than the maximum number of threads K that a single kernel function block can have, while ensuring that all thread bundles contained in the kernel function block will be complete, thereby maximizing thread utilization and improving the overall performance of the GPU.

[0076] In another optional implementation, it can be set that if at least one of the length and width of the kernel function block is equal to the maximum number of threads k in the thread warp in the GPU, the length and width of the kernel function block are kept unchanged.

[0077] It can be understood that if at least one of the length and width of the kernel function block is equal to the maximum number of threads k in the thread bundle in the GPU, it indicates that the size of the kernel function block can ensure that all thread bundles contained in the kernel function block will be complete, so no adjustment is made, thereby avoiding the impact on the degree of adaptation of the scaled image due to modification of the length or width of the kernel function block.

[0078] In another optional implementation, it can be set that: if one of the length and width of the kernel function block is greater than the maximum number of threads k in the thread bundle in the GPU, and the other is less than k, and neither value is an integer multiple of k, then the one of the length and width of the kernel function block that is greater than k is set to n*k, and the other is kept unchanged. Wherein, n*k is the integer multiple value of k that is smaller than the one of the length and width of the kernel function block that is greater than k and is closest to the one of the length and width of the kernel function block that is greater than k.

[0079] In this way, it can also ensure that the product of the length and width of the kernel function block (when there is a depth dimension, it is the product of the values ​​of the length, width and depth dimensions of the kernel function block) will not be greater than the maximum number of threads K that a single kernel function block can have. At the same time, it can ensure that all thread bundles contained in the kernel function block will be complete, thereby maximizing the utilization of threads and improving the overall performance of the GPU.

[0080] In another optional implementation, it can be set that if at least one of the length and width of the kernel function block is equal to an integer multiple of the maximum number of threads k in the thread warp in the GPU, the length and width of the kernel function block are kept unchanged.

[0081] It can be understood that if at least one of the length and width of the kernel function block is equal to an integer multiple of the maximum number of threads k in the thread bundle in the GPU, it indicates that the size of the kernel function block can ensure that all thread bundles included in the kernel function block will be complete, so no adjustment is made, thereby avoiding the impact on the degree of adaptation of the scaled image due to modification of the length or width of the kernel function block.

[0082] It can be understood that when there is a depth dimension and the value of the depth dimension is greater than or equal to 2, it is also possible that the length and width of the kernel function block are both smaller than the maximum number of threads k of the thread warp in the GPU.

[0083] Therefore, when there is a depth dimension and the value of the depth dimension is greater than or equal to 2, in an optional implementation, it can be set: if the length and width of the kernel function block are both less than the maximum number of threads k in the thread bundle in the GPU, then the larger one of the length and width of the kernel function block is set to k; determine whether the product of the other of the length and width of the kernel function block and k is greater than M; if the product of the other of the length and width of the kernel function block and k is greater than M, then reduce the value of the other of the length and width of the kernel function block so that the product of the other of the length and width of the reduced kernel function block and k is less than or equal to M; if the product of the other of the length and width of the kernel function block and k is less than or equal to M, then keep the value of the other of the length and width of the kernel function block unchanged.

[0084] It can be understood that at this time, by setting the larger of the length and width of the kernel function block to k, it can be ensured that all thread bundles contained in the kernel function block will be complete, and the magnitude of the change is relatively small, which can reduce the impact of the modification of the length or width of the kernel function block on the degree of adaptation to the image to be scaled. Thereafter, it is further determined whether the product of the other of the length and width of the kernel function block and k is less than or equal to M. If it is less than or equal to M, it means that the increase in the length or width of the kernel function block will not make the kernel function block unable to be implemented in the GPU. At this time, keeping the other unchanged can reduce the impact on the degree of adaptation to the image to be scaled. If it is greater than M, it means that the increase in the length or width of the kernel function block will make the kernel function block unable to be implemented in the GPU. At this time, the value of the other of the length and width of the kernel function block is reduced so that the product of the other of the length and width of the reduced kernel function block and k is less than or equal to M, thereby ensuring the feasibility of the kernel function block.

[0085] S103: Determine the length and width of the grid in the GPU according to H, W, and the length and width of a single kernel function block.

[0086] In an embodiment of the present application, a first ratio between W and the length of the kernel function block can be determined, and when the first ratio is not an integer, it is rounded up to obtain the length of the grid in the GPU; and a second ratio between H and the width of the kernel function block can be determined, and when the second ratio is not an integer, it is rounded up to obtain the width of the grid in the GPU.

[0087] In this way, by determining the first ratio between W and the length of the kernel function block and rounding up when the first ratio is not an integer, and determining the second ratio between H and the width of the kernel function block and rounding up when the second ratio is not an integer, it can be ensured that the grid can completely accommodate the determined kernel function block.

[0088] It can be understood that when there is a depth dimension, the value of the depth dimension of the grid is consistent with the value of the depth dimension of each kernel function block, so it can be directly determined.

[0089] S104: starting the CUDA kernel of the GPU according to the length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm to scale the image to be scaled.

[0090] In a feasible implementation of an embodiment of the present application, the above method can be executed by the CPU. After determining the length and width of the grid and the length and width of a single kernel function block, the determined length and width of the grid and the length and width of the single kernel function block are sent to the GPU, and the GPU starts the CUDA kernel according to the length and width of the grid and the length and width of the single kernel function block, and then executes a preset image scaling algorithm through the CUDA kernel to scale the image to be scaled.

[0091] In another feasible implementation manner of the embodiment of the present application, the above method can be executed by other cores in the GPU other than the CUDA core that executes the current image scaling task, and after determining the length and width of the grid and the length and width of a single kernel function block, the determined length and width of the grid and the length and width of the single kernel function block are written into the startup program of the CUDA core that executes the current image scaling task, so that the CUDA core that executes the current image scaling task is started according to the length and width of the grid and the length and width of the single kernel function block, and then executes the preset image scaling algorithm to scale the image to be scaled.

[0092] In another feasible implementation of the embodiment of the present application, the above method can also be written into the firmware program of the CUDA kernel in the GPU, and executed at the stage of loading the firmware before the CUDA kernel is started, so as to obtain the length and width of the grid and the length and width of a single kernel function block at this startup, and then execute the preset image scaling algorithm after startup to scale the image to be scaled.

[0093] In an embodiment of the present application, the image scaling algorithm may include but is not limited to at least one of nearest neighbor interpolation, bilinear interpolation, bicubic interpolation, Lanczos resampling, anti-aliasing scaling, and deep learning methods (for example, using a convolutional neural network or a generative adversarial network (GAN) for super-resolution image scaling).

[0094] Based on the same inventive concept, an image scaling device 200 is also provided in the embodiment of the present application. Figure 2 As shown, Figure 2 Shows the use of Figure 1 The image scaling device of the method shown. It should be understood that the specific functions of the device 200 can be referred to the description above, and the detailed description is appropriately omitted here to avoid repetition. The device 200 includes at least one software function module that can be stored in the memory in the form of software or firmware or fixed in the operating system of the device 200. The device 200 can be configured in the CPU or applied to each core of the GPU. Specifically:

[0095] See also Figure 2 As shown, the device 200 includes: an acquisition module 201, a size determination module 202 and a processing module 203. Among them:

[0096] An acquisition module 201 is used to acquire an image to be scaled and a target resolution of the image to be scaled; the target resolution includes a height H and a width W;

[0097] The size determination module 202 is used to: determine the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have; and determine the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block;

[0098] The processing module 203 is used to start the CUDA kernel of the GPU according to the length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm to scale the image to be scaled to obtain a target image with the target resolution.

[0099] In a feasible implementation of an embodiment of the present application, the size determination module 202 is specifically used to: when the kernel function block does not have a depth dimension, calculate the length of a single kernel function block according to the formula sqrt(W*K / H), and calculate the width of a single kernel function block according to the formula sqrt(H*K / W); wherein sqrt is the square root.

[0100] In a feasible implementation manner of an embodiment of the present application, the size determination module 202 is specifically used to: when the kernel function block has a depth dimension, calculate the length of a single kernel function block according to the formula sqrt(W*M / H), and calculate the width of a single kernel function block according to the formula sqrt(H*M / W); wherein sqrt is the square root, M is equal to the ratio between K and the depth value s of the depth dimension of the kernel function block; and s is a positive integer greater than or equal to 1.

[0101] In this feasible implementation manner, the s is a value determined according to a preset number of images to be processed in parallel, or the s is a value determined according to the number of channels of the image to be scaled.

[0102] In a feasible implementation manner of the embodiment of the present application, the size determination module 202 is specifically used for: after determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block can have in the GPU, and before determining the length and width of the grid in the GPU according to the H, the W, and the length and width of the single kernel function block, if one of the length and the width of the kernel function block is greater than the maximum number of threads k of the thread bundle in the GPU, and the other is less than the k, then the one of the length and the width of the kernel function block that is greater than the k is set to the k, and the other is kept unchanged.

[0103] In a feasible implementation of the embodiment of the present application, the size determination module 202 is further specifically used for:

[0104] After determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block can have in the GPU, and before determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block, if both the length and width of the kernel function block are smaller than the maximum number of threads k of the thread warp in the GPU, then setting the larger one of the length and width of the kernel function block to the k;

[0105] Determine whether the product of the other of the length and the width of the kernel function block and the k is greater than the M;

[0106] If the product of the other of the length and the width of the kernel function block and the k is greater than the M, the value of the other of the length and the width of the kernel function block is reduced so that the product of the other of the length and the width of the kernel function block after reduction and the k is less than or equal to the M;

[0107] If the product of the other of the length and the width of the kernel function block and k is less than or equal to M, the value of the other of the length and the width of the kernel function block is kept unchanged.

[0108] In a feasible implementation manner of the embodiment of the present application, the size determination module 202 is specifically used for: after determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block can have in the GPU, and before determining the length and width of the grid in the GPU according to the H, the W, and the length and width of the single kernel function block, if at least one of the length and width of the kernel function block is equal to the maximum number of threads k of the thread bundle in the GPU, then keeping the length and width of the kernel function block unchanged.

[0109] In an embodiment of the present application, the size determination module 202 is specifically used to: determine a first ratio between the W and the length of the kernel function block, and round up when the first ratio is not an integer to obtain the length of the grid in the GPU; determine a second ratio between the H and the width of the kernel function block, and round up when the second ratio is not an integer to obtain the width of the grid in the GPU.

[0110] It should be understood that, for the sake of brevity, some of the contents described in the method embodiment section will not be repeated in the device embodiment section.

[0111] Based on the image scaling method and device provided by the embodiment of the present application, when using the GPU to scale the image, first determine the kernel function length and width and the length and width of the grid that are compatible with the current scaling target based on the target resolution of the image to be scaled (i.e., the height H and width W of the target image), and then start the CUDA kernel of the GPU based on the determined length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm based on the determined grid and kernel function block to scale the image to be scaled, and obtain a target image with the target resolution. In this way, each time the image to be scaled is scaled, the length and width of the grid and the length and width of the kernel function block used can be more in line with actual needs, and will not be too large or too small, which can further improve the computational efficiency of the image scaling algorithm.

[0112] To facilitate understanding of the process of determining the length and width of the grid and the length and width of the kernel function block in the embodiment of the present application, the following is an example of a GPU that supports a set of 64 threads in hardware to exemplify the process of determining the length and width of the grid and the length and width of the kernel function block:

[0113] Assume the depth dimension is 1:

[0114] Assume that the height and width of the target resolution are H and W respectively, and set the length of the kernel function block blockSize.x = x, and the width of the kernel function block blockSize.y = y.

[0115] Then the length of the grid gridSize.x = ceil(W / x), and the width of the grid gridSize.y = ceil(H / y).

[0116] Among them, ceil is a ceiling function, ceil(W / x) indicates that when W / x is not divisible, gridSize.x = int(W / x) + 1, ceil(H / y) indicates that when H / y is not divisible, gridSize.y = int(H / y) + 1. int indicates the value of the integer part.

[0117] Then the total number of threads in the grid is: x*y*gridSize.x*gridSize.y=x*(W / x+1)*y*(H / y+1)=(W+x)*(H+y)=H*W+W*y+H*x+x*y.

[0118] Since the target resolution is H*W, that is, only H*W pixels need to be selected during interpolation, the number of threads of interest should be: H*W. Therefore, redundant threads = (H*W+W*y+H*x+x*y)-H*W=W*y+H*x+x*y.

[0119] Since x*y*1=64*64=4096, we need to find the x and y values ​​that make W*y+H*x. Then, after expressing y with x, we have: W*(4096 / x)+H*x>=2*sqrt(W*(4096 / x)*x*H).

[0120] If and only if W*(4096 / x)=x*H, W*(4096 / x)+H*x=2*sqrt(W*(4096 / x)*x*H). Therefore, W*y+H*x is the smallest when W*(4096 / x)=x*H. Then we can calculate that x=sqrt(W*4096 / H) and y=sqrt(H*4096 / W).

[0121] Due to the organizational structure of a warp with 64 threads, blockSize.x and blockSize.y are preferably multiples of 64. At the same time, due to the restriction that the product of the three dimensions of the kernel function block is less than or equal to 4096, there is no situation where both x and y are greater than 64. There are two scenarios at this time:

[0122] The calculated x and y are both equal to 64, so directly take x=y=64.

[0123] If one of the calculated x and y is greater than 64 and the other is less than 64, the larger one is taken as 64 and the other value remains unchanged.

[0124] Calculate block_size (i.e. the size of the kernel function block) = (b_size_x, b_size_y, 1)

[0125] After determining block_size, the corresponding grid_size (grid size) = (grid_size_x, grid_size_y, 1) can be determined according to the formula gridSize.x = ceil(W / x) and gridSize.y = ceil(H / y).

[0126] Based on the solution in the embodiment of the present application, the computational efficiency of the image scaling algorithm can be further improved. Figure 3As shown, Figure 3 The results of the comparative experiments conducted to verify the effect of the scheme in the embodiment of the present application. When the size of the image to be scaled is 1920*1080 and the target resolution is 640*360, the length (blockDim.x) of the kernel function is determined to be 64 and the width (blockDim.y) is 48 based on the method of the embodiment of the present application. At this time, the time consumed by the CUDA kernel to execute the preset image scaling algorithm to scale the image to be scaled (i.e., the kernel function time consumption) is 41.034us. In other comparative examples, the length and width of the kernel function are set according to the engineer's experience. It can be seen that the length and width of the kernel function determined based on the scheme of the embodiment of the present application, when the CUDA kernel executes the preset image scaling algorithm to scale the image to be scaled, the time consumption is the shortest, which can prove that the scheme based on the embodiment of the present application can further improve the computational efficiency of the image scaling algorithm.

[0127] Based on the same inventive concept, a processor system is also provided in the embodiment of the present application, see Figure 3 As shown, it includes: a CPU and a GPU. The CPU is used to execute any of the above-mentioned image scaling methods to determine the length and width of the grid and the length and width of a single kernel function block and notify the GPU, so that the GPU starts the CUDA kernel according to the length and width of the grid and the length and width of a single kernel function block, and scales the image to be scaled by executing a preset image scaling algorithm through the CUDA kernel to obtain a target image with the target resolution.

[0128] It can be understood that a storage unit, such as a memory, may be provided outside the processor system, and a program for implementing the above-mentioned image scaling method may be configured in the storage unit and read and executed by the CPU.

[0129] Based on the same inventive concept, a GPU is also provided in the embodiment of the present application, see Figure 4 As shown, it includes: a CUDA kernel, a configuration core and a video memory; the configuration core is used to execute one or more programs stored in the video memory to implement any of the above-mentioned image scaling methods.

[0130] In a feasible implementation of the embodiment of the present application, the CUDA kernel is the CUDA kernel that performs the current image scaling task. The configuration core may be a kernel other than the CUDA kernel that performs the current image scaling task, which, after determining the length and width of the grid and the length and width of a single kernel function block, writes the determined length and width of the grid and the length and width of the single kernel function block into the startup program of the CUDA kernel that performs the current image scaling task, so that the CUDA kernel that performs the current image scaling task is started according to the length and width of the grid and the length and width of the single kernel function block, and then executes a preset image scaling algorithm to scale the image to be scaled.

[0131] In another feasible implementation of the embodiment of the present application, the configuration core can be a CUDA kernel that performs the image scaling task. In this case, the CUDA kernel executes the above method in the stage of loading the firmware before startup, and executes a preset image scaling algorithm to scale the image to be scaled after startup.

[0132] Based on the same inventive concept, an electronic device is also provided in an embodiment of the present application, which may include the aforementioned processor system, or include the aforementioned GPU.

[0133] It can be understood that the electronic device can be a host, a server, a computer, or other device that can perform image scaling tasks.

[0134] It is also understood that the electronic device may include more components. For example, it may also have an internal communication bus for realizing communication between the GPU and the CPU; for another example, it may also have an external communication interface, such as a USB (Universal Serial Bus) interface, a CAN (Controller Area Network) bus interface, etc.; for another example, it may also have an information display component such as a display screen, but this is not a limitation.

[0135] Based on the same inventive concept, this embodiment further provides a computer-readable storage medium, such as a floppy disk, an optical disk, a hard disk, a flash memory, a USB flash disk, an SD (Secure Digital Memory Card) card, an MMC (Multimedia Card) card, etc., in which one or more programs for implementing the above steps are stored, and the one or more programs can be executed by one or more processors to implement the above image scaling method. No further details will be given here.

[0136] In the embodiments provided in the present application, it should be understood that the disclosed devices and methods can be implemented in other ways. The device embodiments described above are merely schematic. For example, the division of the units is only a logical function division. There may be other division methods in actual implementation. For example, multiple units or components can be combined or integrated into another system, or some features can be ignored or not executed. Another point is that the mutual coupling or direct coupling or communication connection shown or discussed can be through some communication interfaces, and the indirect coupling or communication connection of the devices or units can be electrical, mechanical or other forms.

[0137] In addition, the units described as separate components may or may not be physically separated, and the components shown as units may or may not be physical units, that is, they may be located in one place or distributed on multiple network units. Some or all of the units may be selected according to actual needs to achieve the purpose of the solution of this embodiment.

[0138] Furthermore, the functional modules in the various embodiments of the present application may be integrated together to form an independent part, or each module may exist separately, or two or more modules may be integrated to form an independent part.

[0139] In this document, relational terms such as first and second, etc. are used merely to distinguish one entity or operation from another entity or operation, but do not necessarily require or imply any such actual relationship or order between these entities or operations.

[0140] As used herein, a plurality refers to two or more than two.

[0141] The above description is only an embodiment of the present application and is not intended to limit the protection scope of the present application. For those skilled in the art, the present application may have various modifications and variations. Any modification, equivalent replacement, improvement, etc. made within the spirit and principle of the present application shall be included in the protection scope of the present application.

Claims

1. An image scaling method, characterized in that: include: Acquire an image to be scaled and a target resolution of the image to be scaled; The target resolution includes a height H and a width W; Determine the length and width of a single kernel function block according to H, W, and the maximum number of threads K that a single kernel function block in the GPU can have; Determine the length and width of the grid in the GPU according to the H, the W, and the length and width of a single kernel function block; The CUDA kernel of the GPU is controlled to start according to the length and width of the grid and the length and width of a single kernel function block, so that the CUDA kernel executes a preset image scaling algorithm to scale the image to be scaled to obtain a target image with the target resolution.

2. The image scaling method according to claim 1, wherein: In the case where the kernel function block does not have a depth dimension, determining the length and width of a single kernel function block according to H, W, and the maximum number of threads K that a single kernel function block in the GPU can have includes: Calculate the length of a single kernel function block according to the formula sqrt(W*K / H); The width of a single kernel function block is calculated according to the formula sqrt(H*K / W); Among them, sqrt is the square root.

3. The image scaling method according to claim 1, wherein: When the kernel function block has a depth dimension, the length and width of a single kernel function block are determined according to H, W, and the maximum number of threads K that a single kernel function block in the GPU can have, including: Calculate the length of a single kernel function block according to the formula sqrt(W*M / H); The width of a single kernel function block is calculated according to the formula sqrt(H*M / W); Wherein, sqrt is a square root, M is equal to the ratio between K and the depth value s of the depth dimension of the kernel function block; and s is a positive integer greater than or equal to 1.

4. The image scaling method according to claim 3, wherein: The s is a value determined according to a preset number of images to be processed in parallel, or the s is a value determined according to the number of channels of the image to be scaled.

5. The image scaling method according to any one of claims 1 to 4, characterized in that: After determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have, and before determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block, the method further includes: If one of the length and width of the kernel function block is greater than the maximum number of threads k in the thread warp in the GPU, and the other is less than k, then the one of the length and width of the kernel function block that is greater than k is set to k, and the other is kept unchanged.

6. The image scaling method according to claim 3 or 4, characterized in that: After determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have, and before determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block, the method further includes: If both the length and width of the kernel function block are smaller than the maximum number of threads k in the thread warp of the GPU, the larger one of the length and width of the kernel function block is set as k; Determine whether the product of the other of the length and the width of the kernel function block and the k is greater than the M; If the product of the other of the length and the width of the kernel function block and the k is greater than the M, the value of the other of the length and the width of the kernel function block is reduced so that the product of the other of the length and the width of the kernel function block after reduction and the k is less than or equal to the M; If the product of the other of the length and the width of the kernel function block and k is less than or equal to M, the value of the other of the length and the width of the kernel function block is kept unchanged.

7. The image scaling method according to any one of claims 1 to 4, characterized in that: After determining the length and width of a single kernel function block according to the H, the W, and the maximum number of threads K that a single kernel function block in the GPU can have, and before determining the length and width of a grid in the GPU according to the H, the W, and the length and width of a single kernel function block, the method further includes: If at least one of the length and the width of the kernel function block is equal to the maximum number of threads k of the thread warp in the GPU, the length and the width of the kernel function block are kept unchanged.

8. The image scaling method according to any one of claims 1 to 4, characterized in that: Determining the length and width of the grid in the GPU according to the H, the W, and the length and width of a single kernel function block includes: Determine a first ratio between W and the length of the kernel function block, and round up when the first ratio is not an integer to obtain a length of the grid in the GPU; A second ratio between the H and the width of the kernel function block is determined, and when the second ratio is not an integer, the ratio is rounded up to obtain a width of the grid in the GPU.

9. A processor system, characterized in that: include: A central processing unit (CPU) and a graphics processing unit (GPU); the CPU is used to execute the image scaling method as described in any one of claims 1 to 8 to determine the length and width of the grid and the length and width of a single kernel function block and notify the GPU, so that the GPU starts a CUDA kernel according to the length and width of the grid and the length and width of a single kernel function block, and executes a preset image scaling algorithm through the CUDA kernel to scale the image to be scaled, so as to obtain a target image with the target resolution.

10. A graphics processor, characterized in that: Includes: CUDA kernel, configuration core and video memory; The configuration core is used to execute one or more programs stored in the video memory to implement the image scaling method according to any one of claims 1 to 8.

11. An electronic device, characterized in that: include: The processor system of claim 9, or the graphics processor of claim 10.

Images