GPU scanning path optimization method

Abstract

The invention relates to the technical field of scientific calculation and GPUs, in particular to a GPU scanning path optimization method. The method comprises the following steps of: constructing a three-dimensional Stencil calculation GPU (Graphics Processing Unit) framework based on an optimal scanning traversal path; a GPU framework is calculated based on three-dimensional Stencil, a scanning traversal path mathematical model is constructed, a scanning traversal path is abstracted into an integer programming problem, and an optimal scanning traversal path of a two-dimensional window is obtained through optimization by taking a total access stock as an optimization target. Compared with traditional direct calculation, the method has the advantages that data handling in a global memory can be remarkably reduced, the utilization rate of accessed data is increased, namely, the data reuse rate is increased, invalid waiting caused by data congestion and high memory access delay of a calculation unit is reduced, and therefore calculation efficiency is improved.

Description

TECHNICAL FIELD

[0001] The present application relates to the field of scientific computing and GPU technology, and particularly relates to a GPU scan path optimization method. BACKGROUND

[0002] In scientific / engineering computing programs on GPU, a very important computing hotspot is Stencil computation, which is a type of computation with memory access as a bottleneck, i.e., a memory access limited type of computation, so the overall computing efficiency of Stencil computation depends largely on the memory access efficiency, and the more grid points in the Stencil template, the higher the requirement for memory access efficiency.

[0003] Taking a GPU program of computational fluid dynamics (CFD) as an example, in the solving process of direct numerical simulation (DNS) of CFD, especially in the spatial discrete format and time advancing solving, Stencil computation accounts for almost the entire computing process, and with the improvement of solving precision, the number of grid points in the Stencil computation template increases significantly. For example, a high-order WENO format in incompressible computation is a typical Stencil computation (star 3D25P), which has 25 grid points in the Stencil template in space, i.e., 25 data are required to initiate a memory access request, and these data in the memory are usually discontinuous, which seriously limits the efficiency of Stencil computation on the hardware platform memory bandwidth, bringing great challenges to the GPU acceleration work of CFD program. SUMMARY

[0004] In view of the GPU optimization problem of Stencil computation in scientific computing, the present application proposes a GPU scan path optimization method, which ensures the memory access efficiency and further improves the data reuse rate in the computing process, thereby accelerating the Stencil computation of GPU.

[0005] The present application adopts the following technical solutions: A GPU scan path optimization method, characterized in that it comprises: constructing a three-dimensional Stencil computation GPU framework based on an optimal scan traversal path; based on the three-dimensional Stencil computation GPU framework, constructing a scan traversal path mathematical model, abstracting the scan traversal path into an integer programming problem, taking the total memory access amount as the optimization target, and optimizing to obtain the optimal scan traversal path of the two-dimensional window.

[0006] Preferably, the three-dimensional Stencil computation GPU framework comprises: First, the memory continuous direction is fixed in space; second, the calculation grid is scanned in the form of window in the two-dimensional plane of the other two directions according to the scanning path.

[0007] Preferably, the three-dimensional Stencil calculation GPU framework comprises: The thread block corresponds to the grid point in the x direction of the memory data continuity, and the grid is divided into a plurality of sub-windows in the yz plane, and the calculation window where the thread block is located scans all the sub-windows according to a specific path, and the Stencil calculation is completed after the scanning is completed.

[0008] Preferably, the plurality of sub-windows divided in the yz plane are assigned unique numbers.

[0009] Preferably, each thread is responsible for calculating a square sub-region of the grid.

[0010] Preferably, the mathematical model of the scanning path comprises: If a square grid sub-region has n sub-windows, then the two-dimensional coordinate coding is performed on each sub-window, and three decision variable queues w[n][2], bf[n][2] and count[n-1][6] are introduced, wherein the variable queue w[n][2] records the row number and column number coordinates of each sub-window in the scanning path, the variable bf[n][2] records the buffer area for saving the sub-window data on the chip, and the row number and column number coordinates of the sub-windows with no dependency relationship are saved in the buffer area, and the variable queue count[n-1][6] records the results of 0 or 1 after comparison between each sub-window saved in the buffer area at the i-th step and each sub-window at the i+1-th step, and the objective function is the sum of all elements in the variable count, and the optimization target of the mathematical model of the scanning path is to minimize the objective function.

[0011] Preferably, the side length row and the theoretical minimum redundancy have a certain mathematical relationship: the theoretical minimum redundancy of the square sub-region with an odd side length row is 6 + 2 × (row × row - 1) + row - 1, and the theoretical minimum redundancy of the square sub-region with an even side length row is 6 + 2 × (row × row - 1) + row - 2.

[0012] Preferably, the GPU scanning path optimization method further comprises: The scanning path of the square region is generated by an algorithm.

[0013] Preferably, for an N × N square region, N is an even number, and the algorithm for generating a scanning path comprises: After encoding each sub-window in the square region with two-dimensional coordinates, a path consisting of a series of window coordinates is generated. The algorithm scans the entire sub-domain one by one in a zigzag pattern by alternating forward and backward scans. Each forward or backward scan generates a series of continuous window coordinates and gradually fills the ROUTE list, i.e. the optimal path list.

[0014] Preferably, for an N×N square region where N is an odd number, the algorithm for generating a scan traversal path includes: After encoding each sub-window in the square region with two-dimensional coordinates, a path consisting of a series of window coordinates is generated. The algorithm generates the bottom path, right path, top path, and left path in a loop, and then fills the core part with a preset pattern. Finally, the entire sub-domain is traversed by a spiral path that scans layer by layer from the outer layer to the center. Each loop scan generates a series of continuous window coordinates and gradually fills the ROUTE list, i.e., the optimal path list.

[0015] Compared with the prior art, the beneficial effects of the present invention include: Compared with traditional direct computation, the method of this invention can significantly reduce data movement in global memory, improve the utilization rate of accessed data (i.e., improve data reuse rate), and reduce the invalid waiting caused by data blocking and high memory access latency in computing units, thereby improving computational efficiency. Attached Figure Description

[0016] The above and other objects, features and advantages of this disclosure will become more apparent from the accompanying drawings, in which like reference numerals generally denote like parts.

[0017] Figure 1 This is a GPU memory access bandwidth test blueprint for a streaming benchmark in one embodiment of the present invention; Figure 2 This is a flowchart of a GPU scanning path optimization method according to one embodiment of the present invention; Figure 3 This is a schematic diagram of a calculation window that progresses alternately in the y and z directions in a polyline manner in one embodiment of the present invention, traversing all grids. Figure 4 This is a schematic diagram of a GPU's required grid size of M × N, and a thread block responsible for a row × col sub-region, in one embodiment of the present invention. Figure 5 In one embodiment of the present invention, a two-dimensional coordinate encoding is performed on each sub-window in the established integer programming model, and a schematic diagram of three decision variables is introduced. Figure 6 Optimized path map for different square sub-regions with different side lengths in one embodiment of the present application; Figure 7 Optimized path map for different rectangular sub-regions in one embodiment of the present application; Figure 8 Optimized path map for the same square sub-region in one embodiment of the present application, where there are multiple optimal scanning traversal paths; Figure 9 Optimized path generation algorithm for square sub-regions with even side length N in one embodiment of the present application; Figure 10 Optimized path generation algorithm for square sub-regions with odd side length N in one embodiment of the present application; Figure 11 Optimized path generation algorithm for square sub-regions with even side length N in one embodiment of the present application; DETAILED DESCRIPTION

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

[0019] The term "comprising" and variations thereof as used herein are intended to cover a non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements is not necessarily limited to those elements, but can include other elements not expressly listed or inherent to such process, method, article, or apparatus. Unless specifically stated, the term "or" as used herein refers to "and / or". The term "based on" is used to mean "based, at least in part, on". The term "one example embodiment" and "an embodiment" mean "at least one example embodiment". The term "another embodiment" means "at least one additional embodiment". The terms "a first", "a second", etc. can refer to different or the same objects. Other explicit or implicit definitions can also be included below.

[0020] For ease of understanding, first, the technical terms that may be involved in the embodiments of the present application are briefly introduced: 1. Stencil calculation Stencil calculation is a calculation mode of iterative update on a regular grid (such as image pixels, spatial discrete points). Its core feature is that the new value of each point in the grid is calculated according to a fixed formula based on the old values of itself and a limited number of adjacent points (referred to as "template" or "neighborhood").

[0021] Stencil computation is highly suitable for parallel architectures such as GPUs due to its highly uniform computation rule and highly localized data access. However, it is also deeply limited by the memory bandwidth. It is a key bridge connecting the physical world model and computer simulation, widely used in solving partial differential equations (such as computational fluid dynamics, heat conduction, multi-physical field), image processing (blur, sharpening), etc. It is one of the cornerstone algorithms in scientific computing and high-performance computing.

[0022] 2. Read-Store-Compute-Throw of data The "Read-Store-Compute-Throw" of data is a typical data flow mode when GPU processes Stencil computation. The core of efficiently executing Stencil computation on GPU is to finely manage the flow of data in multi-level storage structures. A highly optimized technical solution usually follows the four-stage pipeline mode of "Read-Store-Compute-Throw".

[0023] The "Read" stage refers to the cooperative loading of data from the high-latency and large-capacity global memory by threads. Due to the strong locality of Stencil access, the thread blocks on the GPU will adopt a combined access mode, efficiently reading a block of continuous grid data at a time. Threads not only read the data required by the current computation window, but also read the Halo data in the peripheral circle of the template to meet the computing needs of all points within the thread block. The goal of this stage is to maximize the use of GPU memory bandwidth, so the thread blocks will be arranged one-to-one with the grid points in the memory continuous direction to obtain higher memory access efficiency.

[0024] The "Store" stage follows, which is to move the "raw data" read from the global memory to the low-latency and small-capacity registers and shared memory. Registers are the fastest on-chip memory units private to each thread, while shared memory is a fast cache shared by all threads within the same thread block. Its speed can be one or two orders of magnitude faster than global memory. Threads within the block will be divided into different groups, each responsible for copying specific data points from global memory to the corresponding location in registers or shared memory, and ensuring that the Halo region is also correctly filled. After this stage, the data required for the entire thread block computation is completely loaded into the shared memory, and all subsequent computations will directly read data from here, thus avoiding repeated access to slow global memory.

[0025] The "compute" stage is where the core computation happens. At this time, each thread reads all the neighbor data (including the inner point and Halo) it needs for computation from the register or shared memory, and performs the multiply-add operation of the stencil template. Due to the low latency of the register or shared memory, the computation in this stage is more efficient, and the computing unit (CUDA Core / Tensor Core) can operate efficiently, while reducing the blocking caused by memory access. The computation is usually performed in the register, and the result is temporarily stored in the register or shared memory, or written back to the global memory as needed.

[0026] The last "throw" stage refers to the operation of directly discarding or replacing data that has been used and will not be used later. The storage space of the on-chip efficient memory, including the register, shared memory, etc., is very valuable, and replacing useless data with useful data as much as possible will reduce the waiting and blocking of the computing unit due to the absence of data. For example, the data block of the left Halo region of the thread block is directly discarded at a certain stage, because it will be covered by a new data-dependent block after the computation window slides to the next position, and there is no long-term value to save, that is, the storage space vacated will be replaced with a new data block according to the strategy.

[0027] 3. Inner point grid The inner point grid is the core computation area assigned to a certain GPU thread block. In the optimization method proposed in the present application, the inner point grid refers to all grid data at the current position of the computation window.

[0028] 4. Halo grid The Halo grid is one or more layers of additional grid points around the edge of the inner point grid, that is, all computation-dependent data of the non-inner point in the stencil template.

[0029] Next, a GPU scan path optimization method suitable for the embodiments of the present application will be described in detail with reference to the accompanying drawings. It can be understood that the step timing in the following embodiments is only an example, and is not strictly limited.

[0030] The performance bottleneck of Stencil computation is the memory access bandwidth, and its performance is severely limited by the access efficiency of the GPU global memory (Global Memory). Therefore, the first problem to be solved by the present application is to achieve a high enough memory access bandwidth. When evaluating the bandwidth efficiency, both the theoretical peak bandwidth and the actual effective memory bandwidth are considered. Although the theoretical bandwidth provides the maximum transmission capacity of the GPU, the memory access mode of the program in actual operation cannot fully utilize the GPU global memory bandwidth. Therefore, the effective bandwidth that is more consistent with the actual program is also used. The effective bandwidth can better reflect the bandwidth usage of the actual application program on the GPU, thereby improving the accuracy of performance evaluation.

[0031] Taking a domestic GPU processor as an experimental platform, first, the hardware characteristics of the HBM2 of the experimental platform display memory are analyzed, and a HBM2 stack has 4 HBM DRAM Die, and the bit width of each HBM DRAM Die is 1024 bits, that is, 16 double-precision numbers. Assuming that the memory continuous direction of data in the display memory is the x direction, based on this, the x dimension size of the GPU kernel thread block of the program is set to a fixed blockDim.x = 16, and at the same time, the thread block in the x direction is one-to-one corresponding to the CFD physical grid, so that the kernel will simultaneously initiate an access request for the continuous 16 double-precision numbers in the running process, so as to exactly occupy the 1024 bits of the HBM DRAM Die.

[0032] The inventors verify the theoretical analysis of the application through the streaming benchmark. The inventors test the display memory access bandwidth of blockDim.x = 16 and blockDim.x = 32, and it can be observed that the bandwidth represented by the black dots of blockDim.x = 16 is higher than the bandwidth represented by the red dots of blockDim.x = 32.

[0033] Furthermore, the inventors also found through the streaming benchmark that the data size in the x direction has a great influence on the Global memory bandwidth, as shown in Figure 1 With the increase of the data size in the x direction, the bandwidth under two different blockDim.x shows an upward trend, and the bandwidth gradually stabilizes and almost no longer increases when the data size reaches a certain size. It is known that the theoretical display memory peak bandwidth of the experimental platform is 800 GB / s, and the actual bandwidth obtained through the test reaches 50% of the peak value, which ensures a very high memory access efficiency.

[0034] After targeted analysis, it is found that there is a large amount of redundant memory access in the Stencil calculation, that is, the data in each grid point will be repeatedly accessed from the display memory in the calculation process, and all the display memory accesses to the same data from the second time are redundant memory accesses, which needs to be avoided for the Stencil calculation limited by memory access, and therefore, it is necessary to effectively eliminate redundant memory access and improve the on-chip data utilization rate.

[0035] The traditional Stencil calculation flow design is that the data reuse times of each grid point are zero or limited, and the most ideal memory access mode is that the data of each grid point is accessed only once, and the common grid points with data dependency are completed through the data reuse of shared memory.

[0036] However, the on-chip data storage space that can be effectively utilized and manually controlled is very limited (for example, the GPU of the experimental platform only has 64 kB*4 registers and 64 kB shared memory), cannot unlimitedly retain the grid data that has been read from the video memory, and too much on-chip resource occupation will make the occupancy (parallel running warps or wavefronts) too high, seriously affecting the performance of the kernel, and therefore a method combined with the "read-store-calculate-throw" of GPU data is needed to reasonably plan the residence and throwing of the on-chip GPU data.

[0037] In order to better provide calculation acceleration for scientific / engineering calculation programs on the GPU, as shown in the background art, the present application provides a GPU scan path optimization method, comprising: Figure 2 S1, constructing a three-dimensional Stencil calculation GPU framework based on an optimal scan traversal path. S1, constructing a three-dimensional Stencil calculation GPU framework based on an optimal scan traversal path.

[0038] The three-dimensional Stencil calculation GPU framework based on an optimal scan traversal path first fixes the continuous direction of the video memory in space, ensuring the efficiency of video memory access; secondly, in the two-dimensional plane of the other two directions, the calculation grid is scanned and traversed in the form of a window according to the scan path.

[0039] Based on the foregoing experiments and analysis, the present application proposes a Stencil calculation framework for GPU, specifically, a three-dimensional Stencil calculation GPU framework based on an optimal scan traversal path.

[0040] Here, taking the 3D 25P star-shaped Stencil calculation of the high-precision WENO format in the CFD program as an example, as shown in the figure, Figure 3 is an inner point area (the green area is an inner point grid responsible for Stencil calculation of a thread block, and the blue area is a halo grid in the Stencil template) that needs to perform Stencil calculation, y = 8, z = 16, the 8*16 inner point grid is divided into eight sub-windows in the form of a 4*4 grid, and is numbered from 1 to 8, and the thread block is one-to-one corresponding to the grid point in the data continuous x direction, and in the yz plane, the calculation window is in the form of a sub-window, and the eight sub-windows from 1 to 8 are scanned and calculated according to the arbitrary path, and the Stencil calculation is completed after the traversal is completed.

[0041] In the process of scanning promotion, the data is read-stored-calculated- thrown in 4x4 sub-window operation unit. The calculation window needs to decide the number of sub-windows on the chip and the number of sub-windows to be thrown away when switching to a new sub-window. Figure 3 The scanning traversal path of a simple 4x4 grid sub-window is simulated, that is, the broken line ①→②→…→⑧ shown by the red arrow in the figure is scanned to traverse all eight sub-windows, and the Stencil calculation of all 8x16 grids is completed, which replaces the one-to-one correspondence between each thread and each grid point in the xyz space and the calculation strategy in the traditional Stencil calculation.

[0042] In this embodiment, the essence of "read-store-calculate-throw" on the GPU is that the thread block is one-to-one corresponding to the grid point in the continuous x direction of the memory data to obtain high memory access efficiency; and the grid is divided into multiple sub-windows in the yz plane, and the calculation window of the thread block is scanned and calculated all sub-windows in a specific path, and the Stencil calculation is completed after the traversal. Specifically, in the process of scanning promotion, the data is read-stored-calculated- thrown in 4x4 sub-window operation unit. The calculation window needs to decide the number of sub-windows on the chip and the number of sub-windows to be thrown away when switching to a new sub-window; the on-chip high-efficiency memory including registers, shared memory, etc. is used as a high-speed transit station to read the data with spatial locality in batches at a time, and the dense calculation is completed on the chip. This process maximizes the access times of slow global memory, reduces the influence of memory bottleneck on calculation efficiency, and is the key programming paradigm to realize high-performance Stencil calculation.

[0043] Figure 4 The Stencil calculation GPU framework is further described, Figure 4 Each grid represents a sub-window of 4x4 grid size, and a thread block is responsible for calculating the grid sub-area rowxcol. Each thread block scans and traverses a complete sub-area in the form of a sub-window according to the preset red line, and the grid size required to be calculated by a GPU card is MxN. Due to the existence of the halo grid of the Stencil calculation template, the halo grid of the sub-area responsible by each thread block is the inner point of the sub-area grid of the adjacent thread block, that is, these sub-windows will be accessed by the adjacent two thread blocks, resulting in one redundancy. If the redundant memory access between blocks is to be minimized, row = col is required, that is, each thread is responsible for calculating a square sub-area.

[0044] S2, a scanning traversal path mathematical model is constructed, the scanning path is abstracted into an integer programming problem, the total memory access amount is taken as the optimization target, and the optimal scanning traversal path of the two-dimensional window is obtained by optimization, and the data reuse rate is improved.

[0045] The aforementioned Stencil computing GPU framework obviously has multiple completely different computation window scanning traversal paths and different "read-store-calculate-throw" strategies for the same inlier region. Research shows that different path selections and "read-store-calculate-throw" strategies directly affect data reuse rate, and thus the execution efficiency of Stencil computing on the GPU. Therefore, this invention further abstracts the problem and finds the optimal computation window scanning traversal path and "read-store-calculate-throw" strategy by establishing an integer programming mathematical model (scanning traversal path mathematical model). The final output results after optimization of the scanning traversal path mathematical model are consistent with Table 4.

[0046] For 3D25P star-shaped stencil calculations, to complete the stencil calculation of the current window in the yz plane, five 4×4 grid sub-windows are required, including the current window located at the center of the template and four necessary adjacent windows (top, bottom, left, and right). Due to limited on-chip storage resources, an unlimited number of window data cannot be stored; therefore, it is assumed here that a maximum of six sub-windows can be stored on-chip.

[0047] In order to find the optimal computation window scanning traversal path, this embodiment abstracts the problem into a linear integer programming problem. With the help of the mature integer programming solver CPLEX, a mathematical programming model is established, as shown in Table 1. Under the constraints of multiple conditions, the optimal solution of the objective function is solved.

[0048] In the model, a thread block is responsible for calculating a grid sub-region that is a square with row = col, such as... Figure 5 The example shown uses row = col = 5 as an example. There are a total of n = row × col = 25 sub-windows. Each sub-window is encoded with two-dimensional coordinates, and three decision variable queues are introduced: w[n][2], bf[n][2], and count[n-1][6]. The variable queue w[n][2] records the row and column coordinates of each sub-window in the scanning traversal path. The variable bf[n][2] records the buffer area used to store sub-window data on-chip, which stores the row and column coordinates of non-dependent sub-windows. The variable queue count[n-1][6] records the 0 or 1 result of comparing each sub-window in step i with each sub-window in step i+1. The objective function is the sum of all elements in the variable count, which is the total number of redundant memory accesses in the thread block. Minimizing the objective function is the optimization objective of the model.

[0049] Table 1. Constraints, Objective Function, and Optimization Objective of the Integer Programming Model In the optimization process of the integer programming model, the solver uses optimization algorithms including branch bound method to solve, as the size of the sub-region increases, the optimization variables, optimization branches, search space of the model also increase, which increases the difficulty and time-consuming of model solving. By continuously adjusting the integer programming mathematical model, multiple model versions are iterated, the geometric symmetry is used, the rules are summarized, unnecessary branches are cut off, multiple sub-regions of different sizes are optimized, and the optimal solution is obtained, that is, the optimal Stencil calculation scanning traversal path and "read-store-compute-throw" strategy are obtained, and the minimum memory access amount is realized. As shown in Table 2, the sizes of the sub-regions are 3x3, 4x4, 5x5, 6x6, 7x7 and 8x8, respectively. The optimization results give the minimum redundancy, and compared with the traditional Stencil calculation method, there is an obvious advantage of smaller memory access amount. Figure 6 The optimal sub-window scanning traversal paths corresponding to the minimum intra-block redundant memory access obtained after solving the six different sizes are respectively shown.

[0050] Table 2 Optimization result statistics of multiple different size square sub-regions Although it is mentioned in the foregoing that the size of the sub-region is square for optimization, we still try to optimize the rectangular sub-region. After adjusting the mathematical model, the sizes of the sub-regions are 4x3, 5x3, 6x3, 6x4, 8x6 and 10x6, respectively. All the optimization results are better than the simple row scanning path. As shown in Table 3, after analysis, the optimization results of the four sub-regions of 4x3, 5x3, 6x3 and 6x4 are the theoretical optimal (better) path. Figure 7 The optimal sub-window scanning traversal paths corresponding to the minimum intra-block redundant memory access obtained after solving the six different sub-region sizes are respectively shown.

[0051] Table 3 Optimization result statistics of multiple different size rectangular sub-regions In a preferred embodiment, a GPU scanning path optimization method further comprises: S3, generating a path for any row x row size square region by an algorithm.

[0052] It is very time-consuming to solve the model using the integer programming solver CPLEX, especially when the sub-region becomes larger, the search space, variables and branches of the model will quickly rise with the complexity of n factorial (n = row*col), that is, the complexity of the problem is o(n!), the aforementioned square sub-region of row = 8 is the limit of CPLEX optimization using a single CPU, which takes more than a week. If you want to further generalize the window scanning traversal path to larger sub-regions, or even to the general case, you need to analyze and summarize the existing results.

[0053] First of all, through analysis, it is found that there is a certain mathematical relationship between the side length row and the theoretical minimum redundancy: the theoretical minimum redundancy of the square sub-region with an odd side length row is 6 + 2 × (row × row - 1) + row - 1, and the theoretical minimum redundancy of the square sub-region with an even side length row is 6 + 2 × (row × row - 1) + row - 2. Therefore, the six minimum redundancies obtained by optimizing the model in the aforementioned table 2 are completely consistent with the theoretical minimum redundancies, proving the correctness of the aforementioned mathematical model.

[0054] Secondly, for the same sub-region, there are multiple optimal scanning traversal paths, Figure 8 Three different paths of row = 8 square sub-region are shown, all of which have the same memory redundancy 138, and are the theoretical minimum redundancy.

[0055] Finally, the present application summarizes a path generation method for any row × row square region and a "read-store-compute-throw" strategy by observing and analyzing the existing paths obtained by the integer programming optimization method, and the code is implemented as shown in Figure 9 、 Figure 10 The automatically generated path is the theoretical optimal path.

[0056] Figure 9 The algorithm is used to generate a scanning traversal path for an N × N sub-domain (N is an even number). After encoding each window in the sub-domain in two-dimensional coordinates, a path composed of a series of window coordinates is generated. The algorithm alternately performs forward and backward scanning to scan and traverse the entire sub-domain in a zigzag manner. Each forward or backward scanning generates a series of consecutive window coordinates and gradually fills the ROUTE list, which is the optimal path list required by the embodiment. The algorithm is suitable for the case where N is an even number, and ensures that the generated path can completely traverse all windows without boundary overflow. The following is Figure 9 The textual description process of the algorithm is as follows: I. Main process (main function) 1. Initialization: Create an empty list ROUTE for storing the generated path coordinates.

[0057] 2. Loop to generate path: Loop using variable i from 0 to N, and in each loop: ① Call the forward function to generate the forward path, passing i, N, and ROUTE as parameters.

[0058] ② If the row number returned by forward is equal to N, terminate the loop to prevent path out of bounds, and the path generation is complete.

[0059] ③ Call the backward function to generate the backward path, passing i, N, and ROUTE as parameters.

[0060] ④ If the row number returned by backward is equal to N, terminate the loop to prevent path out of bounds, and the path generation is complete.

[0061] II. Forward path generation (forward function) 1. Starting coordinates: row = i*4 + 2, col = 0 2. Loop N times: ① First move right down: row = row - 1, col = col + 1, and add the coordinates to the ROUTE list.

[0062] ② Then move straight up: row = row + 1, col remains unchanged, and add the coordinates to the ROUTE list.

[0063] 3. Return the current row number row.

[0064] III. Backward path generation (backward function) 1. Starting coordinates: row = i*4 + 4, col = N + 1 2. Loop N times: ① First move left down: row = row - 1, col = col - 1, and add the coordinates to the ROUTE list.

[0065] ② Then move straight up: row = row + 1, col remains unchanged, and add the coordinates to the ROUTE list.

[0066] 3. Return the current row number row.

[0067] Figure 10The algorithm is used to generate a scanning path for a sub-domain of N x N (N is an odd number). After encoding the two-dimensional coordinates of each window in the sub-domain, a path composed of a series of window coordinates is generated. The algorithm generates the bottom path, right path, top path, and left path through a loop, and then fills the core part (N = 1 or N = 3) using a preset pattern. Finally, the entire sub-domain is scanned in a spiral path from the outer layer to the center layer by layer. Each loop scan generates a series of continuous window coordinates, and the ROUTE list, which is the optimal path list required by the embodiment, is gradually filled. The algorithm is suitable for the case where N is an odd number, and ensures that the generated path can completely traverse all windows without exceeding the boundary. The following is Figure 10 Text description of the algorithm flow: Step 1: Initialization 1. Initialization: Create an empty list ROUTE to store the generated path coordinates.

[0068] 2. Determine the number of layers layers and the core size core: Determine the number of spiral layers by decreasing core = N - 4*layers until core is 1 or 3.

[0069] Step 2: Generate spiral path layer by layer 1. For each layer, perform the following four edge generation and update the path coordinates: ① Bottom_function(ROUTE): Generate the bottom path of the layer, which scans the bottom grid points of the layer from right to left in a zigzag manner, and adds the path to the ROUTE list.

[0070] ② Right_function(ROUTE): Generate the right path of the layer, which scans the right grid points of the layer from top to bottom in a zigzag manner, and adds the path to the ROUTE list.

[0071] ③ Top_function(ROUTE): Generate the top path of the layer, which scans the top grid points of the layer from left to right in a zigzag manner, and adds the path to the ROUTE list.

[0072] ④ Left_function(ROUTE): Generate the left path of the layer, which scans the left grid points of the layer from bottom to top in a zigzag manner, and adds the path to the ROUTE list.

[0073] After each layer, continue to the next layer.

[0074] Step 3: Core filling 1. Call the preset pattern according to the core value: 1. core == 0: call Function_0, add 3 center points to the end of the ROUTE list.

[0075] 2. core == 1: call Function_1, add center point to the end of the ROUTE list.

[0076] 3. core == 3: call Function_3, add 9 preset path points to the end of the ROUTE list.

[0077] Table 4 shows the "read-store-compute-throw" strategy automatically generated by the code for the 25 sub-windows of a square sub-region with side length 5, in the order of Figure 5 Window 1 is the ith window that needs to be scanned and computed by the stencil, windows 2 to 5 are the data dependent windows determined by the stencil, and window 6 is the prefetch window.

[0078] Table 4 "read-store-compute-throw" strategy automatically generated by the code for the 25 sub-windows of a square sub-region with side length 5 Figure 11 Table 5 shows four minimum-redundancy scanning traversal paths for the square sub-region automatically generated by the code.

[0079] Having described various embodiments of the disclosure, the above description is illustrative and not restrictive, and the description is not exhaustive nor limited to the disclosed embodiments. Many modifications and variations to the described embodiments are possible in light of the above teachings. The selection of terms to be used in the description is not intended to limit the scope of the embodiments described herein, but rather to best explain the principles of the various embodiments, practical application, or technical improvements in the art, or to enable others skilled in the art to understand the various embodiments disclosed herein.

Claims

1. A method for GPU scan path optimization, the method comprising: The application relates to a three-dimensional Stencil calculation GPU framework based on an optimal scanning traversal path. The three-dimensional Stencil calculation GPU framework comprises: First, the continuous direction of the video memory is fixed in space; second, the calculation grid is scanned and traversed in the form of a window in the two-dimensional plane of the other two directions according to the scanning path.

2. The method of claim 1, wherein, The three-dimensional Stencil calculation GPU framework comprises: The thread block is one-to-one corresponding to the grid point in the x direction of the continuous video memory data, and the grid is divided into a plurality of sub-windows in the yz plane; the calculation window where the thread block is located scans all the sub-windows according to a specific path one by one, and the Stencil calculation is completed after the traversal.

3. The method of claim 2, wherein, The plurality of sub-windows in the yz plane are assigned unique numbers. Each thread is responsible for calculating a square sub-region of the grid.

4. The method of claim 3, wherein, The scanning traversal path mathematical model comprises:

5. The method according to any one of claims 1 to 4, characterized in that, If a square grid sub-region has n sub-windows, then each sub-window is two-dimensionally encoded, and three decision variable queues, w[n][2], bf[n][2] and count[n-1][6], are introduced; the variable queue w[n][2] records the row and column coordinates of each sub-window in the scanning traversal path, the variable bf[n][2] records the buffer area for saving the sub-window data on the chip, and the variable queue count[n-1][6] records the 0 or 1 result after comparison between each sub-window of the i-th step and each sub-window of the i+1-th step in the buffer area; the sum of all elements in the variable count is the objective function, and the minimum objective function is the optimization objective of the scanning traversal path mathematical model.

6. The method of claim 5, wherein, There is a certain mathematical relationship between the side length row and the theoretical minimum redundancy: the theoretical minimum redundancy of the square sub-region with an odd-numbered side length row is 6 + 2 * (row * row - 1) + row - 1, and the theoretical minimum redundancy of the square sub-region with an even-numbered side length row is 6 + 2 * (row * row - 1) + row - 2. The GPU scanning path optimization method further comprises:

7. The method of claim 5, wherein, An algorithm is used to generate the scanning traversal path of the square region.

8. The method of claim 5, wherein, For an N * N square region, N is an even number, and an algorithm for generating a scanning traversal path comprises: After two-dimensionally encoding each sub-window in the square region, a path composed of a series of window coordinates is generated; the algorithm generates a series of continuous window coordinates through forward and backward scanning alternately, and fills the ROUTE list, i.e., the optimal path list, gradually in a zigzag manner.

9. The method of claim 8, wherein, The three-dimensional Stencil calculation GPU framework comprises: The thread block is one-to-one corresponding to the grid point in the x direction of the continuous video memory data, and the grid is divided into a plurality of sub-windows in the yz plane; the calculation window where the thread block is located scans all the sub-windows according to a specific path one by one, and the Stencil calculation is completed after the traversal. The plurality of sub-windows in the yz plane are assigned unique numbers. Each thread is responsible for calculating a square sub-region of the grid. The scanning traversal path mathematical model comprises: If a square grid sub-region has n sub-windows, then each sub-window is two-dimensionally encoded, and three decision variable queues, w[n][2], bf[n][2] and count[n-1][6], are introduced; the variable queue w[n][2] records the row and column coordinates of each sub-window in the scanning traversal path, the variable bf[n][2] records the buffer area for saving the sub-window data on the chip, and the variable queue count[n-1][6] records the 0 or 1 result after comparison between each sub-window of the i-th step and each sub-window of the i+1-th step in the buffer area; the sum of all elements in the variable count is the objective function, and the minimum objective function is the optimization objective of the scanning traversal path mathematical model. There is a certain mathematical relationship between the side length row and the theoretical minimum redundancy: the theoretical minimum redundancy of the square sub-region with an odd-numbered side length row is 6 + 2 * (row * row - 1) + row - 1, and the theoretical minimum redundancy of the square sub-region with an even-numbered side length row is 6 + 2 * (row * row - 1) + row - 2. The GPU scanning path optimization method further comprises: An algorithm is used to generate the scanning traversal path of the square region. For an N * N square region, N is an even number, and an algorithm for generating a scanning traversal path comprises: After two-dimensionally encoding each sub-window in the square region, a path composed of a series of window coordinates is generated; the algorithm generates a series of continuous window coordinates through forward and backward scanning alternately, and fills the ROUTE list, i.e., the optimal path list, gradually in a zigzag manner.

10. The method of claim 8, wherein, For N*N square area, N is odd, an algorithm for generating a scanning path is provided, comprising: After encoding each sub-window in the square area in two-dimensional coordinates, a path composed of a series of window coordinates is generated, the algorithm generates a bottom path, a right path, a top path and a left path through a loop, and then the core part is filled using a preset mode, and finally the entire sub-domain is scanned in a spiral path from the outer layer to the center layer by layer; each loop scanning generates a series of continuous window coordinates, and the ROUTE list, i.e. the optimal path list, is gradually filled.

Images