A gpu scan path optimization method

Abstract

The present application relates to the field of scientific computing and GPU technology, and particularly relates to a GPU scan path optimization method.The method comprises: constructing a three-dimensional Stencil calculation GPU framework based on an optimal scan traversal path; based on the three-dimensional Stencil calculation 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 an optimization target, and optimizing to obtain an optimal two-dimensional window scan traversal path.Compared with the traditional direct calculation, the method of the present application can significantly reduce the data transfer in the global memory, improve the utilization rate of the accessed data, i.e., improve the data reuse rate, reduce the invalid waiting caused by the data blocking and high memory delay of the calculation unit, and thus improve the calculation efficiency.

Description

Technical Field

[0001] This invention relates to the fields of scientific computing and GPU technology, and in particular to a GPU scanning path optimization method. Background Technology

[0002] In scientific / engineering computing programs on GPUs, a very important computational hotspot is Stencil computation, which is a type of computation with memory access as the bottleneck, also known as memory-restricted computation. Therefore, the overall computational efficiency of Stencil computation largely depends on memory access efficiency, and the more dependent grid points in the Stencil template, the higher the requirement for memory access efficiency.

[0003] Take a GPU program for Computational Fluid Dynamics (CFD) as an example. In the Direct Numerical Simulation (DNS) solution process of CFD, especially in spatial discretization schemes and time-progressive solutions, stencil computation occupies almost the entire computational process. Furthermore, as the solution accuracy increases, the number of grid points in the stencil computation template rises significantly. For example, a high-order WENO scheme in inviscid computation is a typical stencil computation (star-shaped 3D25P). Its stencil template has 25 grid points in space, meaning that a single stencil computation requires 25 memory access requests for data points. Moreover, these data points are usually discontinuous in memory, which severely limits the efficiency of stencil computation to the memory bandwidth of the hardware platform, posing a significant challenge to GPU acceleration of CFD programs. Summary of the Invention

[0004] To address the GPU optimization problem of Stencil computation in scientific computing, this invention proposes a GPU scanning path optimization method that ensures memory access efficiency while further improving data reuse during the computation process, thereby accelerating GPU Stencil computation.

[0005] This invention is achieved using the following technical solution:

[0006] A GPU scanning path optimization method, characterized by comprising:

[0007] Construct a 3D Stencil computing GPU framework based on the optimal scan traversal path;

[0008] Based on the 3D Stencil computing GPU framework, a mathematical model of the scan traversal path is constructed. The scan traversal path is abstracted into an integer programming problem. The total memory access is used as the optimization objective to obtain the optimal scan traversal path for the 2D window.

[0009] Preferably, the 3D Stencil computing GPU framework includes:

[0010] First, the continuous direction of the video memory is fixed in space; second, the computational grid is scanned and traversed in the form of windows along the scan path in the two-dimensional planes of the other two directions.

[0011] Preferably, the 3D Stencil computing GPU framework includes:

[0012] The thread block corresponds one-to-one with the grid points in the continuous x-direction of the video memory data, while the grid is divided into multiple sub-windows in the yz plane. The calculation window where the thread block is located scans and calculates all sub-windows one by one according to a specific path. After the traversal is completed, the Stencil calculation is completed.

[0013] Preferably, the multiple sub-windows into which the grid in the yz plane is divided are assigned unique numbers.

[0014] Preferably, the grid sub-region for calculation by each thread is square.

[0015] Preferably, constructing the mathematical model for the scan traversal path includes:

[0016] Suppose a square grid sub-region has n 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 scan 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. Minimizing the objective function is the optimization objective of the mathematical model of the scan traversal path.

[0017] Preferably, there is a definite mathematical relationship between the side length row and the theoretical minimum redundancy: the theoretical minimum redundancy of a square subregion with an odd side length row is 6 + 2 × (row × row - 1) + row - 1, and the theoretical minimum redundancy of a square subregion with an even side length row is 6 + 2 × (row × row - 1) + row - 2.

[0018] Preferably, a GPU scanning path optimization method further includes:

[0019] Use an algorithm to generate the scan traversal path for the square region.

[0020] Preferably, for an N×N square region where N is an even number, the algorithm for generating a scan traversal path includes:

[0021] 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.

[0022] Preferably, for an N×N square region where N is an odd number, the algorithm for generating a scan traversal path includes:

[0023] 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.

[0024] Compared with the prior art, the beneficial effects of the present invention include:

[0025] 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

[0026] 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.

[0027] Figure 1 This is a GPU memory access bandwidth test blueprint for a streaming benchmark in one embodiment of the present invention;

[0028] Figure 2 This is a flowchart of a GPU scanning path optimization method according to one embodiment of the present invention;

[0029] 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.

[0030] Figure 4This 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.

[0031] 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.

[0032] Figure 6 This is an optimized path diagram of square sub-regions with different side lengths in one embodiment of the present invention;

[0033] Figure 7 This is an optimized path diagram for different rectangular sub-regions in one embodiment of the present invention;

[0034] Figure 8 This is a schematic diagram illustrating multiple optimal scanning traversal paths for the same square sub-region in one embodiment of the present invention;

[0035] Figure 9 This is a diagram of an algorithm for generating the optimal path in a square sub-region with an even side length N, according to one embodiment of the present invention.

[0036] Figure 10 This is a diagram of the optimal path generation algorithm for a square sub-region with an odd side length N in one embodiment of the present invention;

[0037] Figure 11 This is a schematic diagram illustrating the automatic generation of the minimum redundancy scan traversal path for a square sub-region of arbitrary side length in one embodiment of the present invention. Detailed Implementation

[0038] Embodiments of the present disclosure will now be described in more detail with reference to the accompanying drawings. While embodiments of the present disclosure are shown in the drawings, it should be understood that the present disclosure may be implemented in various forms and should not be limited to the embodiments set forth herein. Rather, these embodiments are provided so that the present disclosure will be thorough and complete, and will fully convey the scope of the present disclosure to those skilled in the art.

[0039] The term “comprising” and its variations as used herein indicate an open-ended inclusion, i.e., including but not limited to. Unless otherwise stated, the term “or” means “and / or”. The term “based on” means “at least partially based on”. The terms “one example embodiment” and “one embodiment” mean “at least one example embodiment”. The term “another embodiment” means “at least one additional embodiment”. The terms “first,” “second,” etc., may refer to different or the same objects. Other explicit and implicit definitions may also be included below.

[0040] To facilitate understanding, the technical terms that may be involved in the embodiments of this application will be briefly introduced first:

[0041] 1. Stencil calculation

[0042] Stencil computation is a computational paradigm that iteratively updates data on a regular grid (such as image pixels or spatially discrete points). Its core characteristic is that the new value of each point in the grid is calculated from its own value and the old values ​​of a finite number of neighboring points (called "templates" or "neighborhoods") according to a fixed formula.

[0043] Because of its highly unified computational rules and highly localized data access, Stencil computation is well-suited for efficient execution on parallel architectures such as GPUs, but it is also heavily limited by memory bandwidth. It serves as a crucial bridge between physical world models and computer simulations, and is widely used in solving partial differential equations (such as computational fluid dynamics, heat conduction, and multiphysics), image processing (blurring, sharpening), and other fields. It is one of the cornerstone algorithms in scientific computing and high-performance computing.

[0044] 2. Data "Read-Store-Calculate-Discard"

[0045] The typical data flow pattern when GPUs process Stencil computations is "read-store-compute-discard". The core of efficiently executing Stencil computations on GPUs lies in the meticulous management of data flow within a multi-layered storage structure. A highly optimized technical solution typically follows a four-stage pipeline pattern of "read-store-compute-discard".

[0046] The "read" phase refers to the collaborative loading of data from high-latency, large-capacity global memory by threads. Due to the strong locality of Stencil access, thread blocks on the GPU employ a merged access pattern, efficiently reading a contiguous block of grid data at once. A thread reads not only the data required by the inlier of its current computation window but also additionally reads the Halo data surrounding its template to satisfy the computational needs of all points within the entire thread block. The goal of this phase is to maximize the utilization of the GPU's memory bandwidth; therefore, thread blocks are arranged in a one-to-one correspondence with grid points along the data's contiguous memory direction to achieve high memory access efficiency.

[0047] The next stage is the "store" phase, where the raw data read from global memory is moved to low-latency, small-capacity registers and shared memory. Registers are the fastest on-chip memory units private to each thread; while shared memory is an on-chip cache that can be quickly shared by all threads within the same thread block, and its speed can be one or two orders of magnitude faster than global memory. Within the thread block, tasks are divided: each thread is 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 is completed, the data required for the entire thread block's computation is completely loaded into shared memory, and all subsequent computations will directly read data from here, thus avoiding repeated access to slow global memory.

[0048] The "computation" phase is where the core computations occur. Here, each thread reads all the neighbor data (including interior points and Halo points) needed for its computation from registers or shared memory, performing multiply-accumulate operations using the Stencil template. Due to the low latency of registers or shared memory, this phase is more computationally efficient, allowing computational units (CUDA Cores / Tensor Cores) to operate efficiently with less memory access blocking. Computations are typically performed in registers, and results are temporarily stored in registers or shared memory, or written back to global memory as needed.

[0049] The final "discard" phase refers to the operation of directly discarding or replacing data that has been used and will not be used again. On-chip high-efficiency memory, including registers and shared memory, is extremely valuable. Replacing useless data with useful data as much as possible will reduce waiting and blocking of computing units due to data not being ready. For example, a data block in the Halo region on the left side of a thread block might be discarded at a certain stage because it will be overwritten by a new data-dependent block after the computation window slides to the next position, and has no long-term value. The freed-up storage space will be replaced with new data blocks according to a strategy.

[0050] 3. Interior point grid

[0051] An interior grid is the core computational region allocated to a specific GPU thread block. In the optimization method proposed in this invention, the interior grid refers to all grid data at the current location of the computation window.

[0052] 4. Halo grid

[0053] Halo meshes are one or more additional mesh points surrounding the inner-point mesh edges, representing all computational dependencies of non-inner-point meshes in the Stencil template.

[0054] The following is a detailed description of a GPU scanning path optimization method applicable to embodiments of this application, with reference to the accompanying drawings. It should be understood that the timing sequence of the steps in the following embodiments is merely an example and not a strict limitation.

[0055] The performance bottleneck of Stencil computing is memory access bandwidth, which is severely limited by the efficiency of GPU's Global Memory access. Therefore, the primary challenge of this invention is to achieve sufficiently high memory access bandwidth. When evaluating bandwidth efficiency, both theoretical peak bandwidth and actual effective memory bandwidth are considered. Although theoretical bandwidth provides the GPU's maximum transfer capacity, the memory access patterns of programs in actual operation will inevitably prevent them from fully utilizing the GPU's memory bandwidth. Therefore, it is also necessary to adopt effective bandwidth that more closely reflects the actual bandwidth usage of applications on the GPU, thereby improving the accuracy of performance evaluation.

[0056] Using a domestically produced GPU processor as the experimental platform, we first analyzed the hardware characteristics of the HBM2 memory on the platform. An HBM2 stack has four HBM DRAM Dies, each with a bit width of 1024 bits, or 16 double-precision numbers. Assuming the data's contiguous memory direction in the GPU memory is the x-axis, we set the x-dimensional size of the program's GPU kernel thread blocks to a fixed blockDim.x = 16. Simultaneously, we ensured that each thread block corresponded one-to-one with the CFD physical grid in the x-axis direction. Therefore, during kernel execution, each thread block simultaneously initiates access requests to 16 consecutive double-precision numbers, thus perfectly filling the 1024-bit HBM DRAM Die.

[0057] The inventors verified the theoretical analysis of this invention through streaming benchmark tests. They tested the memory access bandwidth for blockDim.x = 16 and blockDim.x = 32, and observed that the bandwidth represented by the black dot for blockDim.x = 16 was higher than that represented by the red dot for blockDim.x = 32.

[0058] Furthermore, the inventors discovered through streaming benchmarks that the data size in the x-direction has a significant impact on global memory access bandwidth, such as... Figure 1As shown, the bandwidth under both different blockDim.x models increases with the increase of data size in the x-direction, and then gradually stabilizes and almost stops increasing after the data size reaches a certain level. The theoretical peak memory bandwidth of the experimental platform is known to be 800 GB / s, and the actual bandwidth obtained in the test reached 50% of the peak, ensuring high memory access efficiency.

[0059] After targeted analysis, it was found that there is a large amount of redundant memory access in Stencil computation. That is, the data in each grid point is repeatedly accessed from the video memory during the computation process, and all video memory accesses to the same data from the second time onwards are redundant memory accesses. This is something that needs to be avoided in memory-constrained Stencil computation. Therefore, it is necessary to effectively eliminate redundant memory accesses and improve on-chip data utilization.

[0060] Traditional Stencil computational pipeline design reuses data at each grid point zero times or a limited number of times. The ideal memory access method is that the data at each grid point is accessed only once, and common grid points that have data dependencies on each other can reuse data through methods such as shared memory.

[0061] However, the on-chip data storage space that can be effectively utilized and manually controlled for each computing unit is very limited (for example, the GPU of the experimental platform only has 64 kB×4 registers and 64 kB of shared memory). It cannot retain grid data read from video memory indefinitely, and excessive on-chip resource usage will result in excessive occupancy (parallel running warps or wavefronts), which will seriously affect the performance of the kernel. Therefore, it is necessary to design a method that combines "read-store-compute-discard" with GPU data to rationally plan the residence and discarding of GPU on-chip data.

[0062] To better provide computational acceleration for scientific / engineering computing programs on GPUs, such as Figure 2 As shown, the present invention provides a GPU scanning path optimization method, comprising:

[0063] S1. Construct a 3D Stencil computing GPU framework based on the optimal scan traversal path.

[0064] A 3D Stencil computing GPU framework based on optimal scan traversal path first fixes the continuous direction of the video memory in space to ensure the efficiency of video memory access; secondly, it scans and traverses the computing grid in the form of windows according to the scan path in two-dimensional planes in the other two directions.

[0065] Based on the aforementioned experiments and analysis, this invention proposes a GPU-oriented Stencil computing framework, specifically a 3D Stencil computing GPU framework based on the optimal scan traversal path.

[0066] This example uses a high-precision WENO format 3D 25P star-shaped Stencil calculation in a CFD program. Figure 3 The diagram shows an interior region with y = 8 and z = 16 that requires Stencil calculation (the green area is the interior mesh for Stencil calculation handled by a thread block, and the blue area is the halo mesh in the Stencil template. The 8×16 interior mesh is divided into eight sub-windows with each sub-window being a 4×4 mesh, and numbered from ① to ⑧). The thread block corresponds one-to-one with the mesh points in the continuous x-direction of the data. In the yz plane, the calculation window is in the form of sub-windows, and the calculation of the eight sub-windows from ① to ⑧ is performed one by one along any path. After the traversal is completed, the Stencil calculation is finished.

[0067] During the scanning process, the data is read, stored, calculated, and discarded in 4×4 sub-windows. Each time the calculation window switches to a new sub-window, the number of the sub-window residing on the chip and the number of the sub-window to be discarded need to be strategically determined. Figure 3 This simulates a simple 4×4 grid sub-window scanning traversal path, that is, scanning and traversing all eight sub-windows according to the broken line ① → ② → … → ⑧ shown by the red arrow in the figure, completing the Stencil calculation of all 8×16 grids. This replaces the one-to-one correspondence and calculation strategy between each thread and each grid point in the xyz space in the traditional Stencil calculation.

[0068] In this embodiment, the essence of the "read-store-compute-discard" process on the GPU lies in the following: Thread blocks correspond one-to-one with grid points along the continuous x-axis of the video memory data to achieve high memory access efficiency. In the yz-plane, the grid is divided into multiple sub-windows, and the computation window containing the thread block scans and computes all sub-windows one by one along a specific path, completing the Stencil computation after traversal. Specifically, during the scanning process, data is processed using 4×4 sub-windows as the operation unit for "read-store-compute-discard." Each time the computation window switches to a new sub-window, it needs to strategically determine the sub-window number residing on the chip and the number of the discarded sub-window. Utilizing efficient on-chip memory, including registers and shared memory, as a high-speed relay station, data with spatial locality is read in batches at once, and intensive computation is performed on-chip. This process minimizes the number of accesses to slow global memory, reducing the impact of memory access bottlenecks on computational efficiency, and is a key programming paradigm for achieving high-performance Stencil computation.

[0069] Figure 4 The Stencil computing GPU framework is further described. Figure 4 Each cell in the algorithm represents a 4×4 grid sub-window. A thread block is responsible for calculating a sub-region of the grid, row×col. Each thread block scans and traverses a complete sub-region in the form of sub-windows along a predefined red line. The grid size required for calculation on a single GPU card is M×N. Due to the existence of the halo grid in the Stencil calculation template, the halo grid of the sub-region handled by each thread block is an interior point of the sub-region grid of the adjacent thread block. That is, these sub-windows will be accessed by the two adjacent thread blocks, resulting in redundancy. To minimize the redundant memory access between blocks, row = col, meaning that the sub-region handled by each thread is a square.

[0070] S2. Construct a mathematical model for the scanning traversal path, abstract the scanning path into an integer programming problem, use the total number of memory accesses as the optimization objective, and optimize to obtain the optimal scanning traversal path for the two-dimensional window, thereby improving data reuse rate.

[0071] 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.

[0072] 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.

[0073] 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.

[0074] 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.

[0075] Table 1. Constraints, Objective Function, and Optimization Objective of the Integer Programming Model

[0076]

[0077] In the optimization process of the integer programming model, the solver uses optimization algorithms, including branch and bound methods, to solve the problem. As the size of the sub-region increases, the number of optimization variables, optimization branches, and search space also increases, increasing the difficulty and time consumption of solving the model. Through continuous adjustment of the integer programming mathematical model, multiple model versions were iterated. By utilizing geometric symmetry, summarizing patterns, and pruning unnecessary branches, optimization was performed on multiple sub-regions of different sizes, obtaining the optimal solution. That is, the optimal Stencil computation scan traversal path and "read-store-calculate-throw" strategy were obtained, achieving the minimum total memory access. As shown in Table 2, optimization solutions were performed for sub-region sizes of 3×3, 4×4, 5×5, 6×6, 7×7, and 8×8, respectively. The optimization results show the minimum redundancy and, compared with the traditional Stencil computation method, have a significant advantage in terms of smaller total memory access. Figure 6 The optimal sub-window scan traversal paths corresponding to redundant memory access within the smallest block are shown after solving these six different sizes.

[0078] Table 2. Statistical analysis of optimization results for multiple square sub-regions of different sizes.

[0079]

[0080] Although it was mentioned earlier that a square subregion is optimal, we still attempted to optimize rectangular subregions. After adjusting the mathematical model, we solved for subregions of 4×3, 5×3, 6×3, 6×4, 8×6, and 10×6 sizes. All optimization results were better than a simple row-by-row scanning path. As shown in Table 3 below, after analysis, the optimization results for the 4×3, 5×3, 6×3, and 6×4 subregions are theoretically optimal (or relatively optimal) paths. Figure 7 The optimal sub-window scan traversal path corresponding to the minimum intra-block redundant memory access obtained after solving for 6 different sub-region sizes is shown.

[0081] Table 3. Statistical analysis of optimization results for multiple rectangular sub-regions of different sizes.

[0082]

[0083] In a preferred embodiment, a GPU scan path optimization method further includes:

[0084] S3. Use an algorithm to generate a path for any row × row-sized square region.

[0085] Solving the problem using the integer programming solver CPLEX is extremely time-consuming, especially as the sub-region becomes larger. The search space, variables, and branches of the model increase rapidly with a complexity of n factorial (n = row * col), meaning the problem complexity becomes O(n!). The aforementioned square sub-region with row = 8 represents the limit of optimization achievable with a single CPU using CPLEX, taking over a week to complete. To further extend the window scanning traversal path to larger sub-regions, or even to the general case, it is necessary to analyze the existing results and summarize the patterns.

[0086] First, analysis reveals a definite mathematical relationship between the side length `row` and the theoretical minimum redundancy: the theoretical minimum redundancy for a square subregion with an odd side length `row` is 6 + 2 × (row × row - 1) + row - 1, and the theoretical minimum redundancy for a square subregion with an even side length `row` is 6 + 2 × (row × row - 1) + row - 2. Therefore, the six minimum redundancies obtained from model optimization in Table 2 perfectly match the theoretical minimum redundancy, proving the correctness of the aforementioned mathematical model.

[0087] Secondly, for the same sub-region, there are multiple optimal scan traversal paths. Figure 8 The diagram shows three different paths for the square subregion with row = 8. They all have the same memory access redundancy of 138 and are theoretically the minimum redundancy.

[0088] Finally, by observing and analyzing existing paths obtained through integer programming optimization methods, this invention summarizes a path generation method for any row × row-sized square region, along with a "read-store-calculate-throw" strategy, and provides code implementation, as follows: Figure 9 , Figure 10 As shown, the automatically generated path is the theoretically optimal path.

[0089] Figure 9 The algorithm generates a traversal path for an N×N subdomain (N is even). It encodes each window in the subdomain using two-dimensional coordinates, generating a path composed of a series of window coordinates. The algorithm traverses the entire subdomain sequentially in a zigzag pattern by alternating forward and backward scans. Each forward or backward scan generates a series of continuous window coordinates, gradually filling the ROUTE list, which is the optimal path list required in this embodiment. This algorithm is suitable for cases where N is even, ensuring that the generated path can completely traverse all windows without exceeding boundaries. The following is... Figure 9 The algorithm's textual description process:

[0090] I. Main Flow (Main Function)

[0091] 1. Initialization: Create an empty list ROUTE to store the generated path coordinates.

[0092] 2. Loop to generate path: Use variable i to loop from 0 to N, and in each loop:

[0093] ① Call the forward function to generate the forward path, passing the parameters i, N, and ROUTE.

[0094] ② If the line number returned by forward is equal to N, then the loop terminates to prevent the path from going out of bounds, and the path generation is complete.

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

[0096] ④ If the row number returned by backward is equal to N, then the loop terminates to prevent the path from going out of bounds, and the path generation is complete.

[0097] II. Forward Path Generation (forward function)

[0098] 1. Starting coordinates: row = i*4 + 2, col = 0

[0099] 2. Repeat N times:

[0100] ① Move to the lower right each time: row = row - 1, col = col + 1, and add the coordinates to the ROUTE list.

[0101] ② Move directly upwards: row = row + 1, keep col unchanged, and add the coordinates to the ROUTE list.

[0102] 3. Return to the current row number.

[0103] III. Backward path generation (backward function)

[0104] 1. Starting coordinates: row = i*4 + 4, col = N + 1

[0105] 2. Repeat N times:

[0106] ① Move to the lower left each time: row = row - 1, col = col - 1, and add the coordinates to the ROUTE list.

[0107] ② Move directly upwards: row = row + 1, keep col unchanged, and add the coordinates to the ROUTE list.

[0108] 3. Return to the current row number.

[0109] Figure 10 The algorithm generates a scanning traversal path for an N×N subdomain (N is odd). It encodes each window in the subdomain using two-dimensional coordinates, generating a path composed of a series of window coordinates. The algorithm iteratively generates bottom, right, top, and left paths, then fills the core (N=1 or N=3) using a preset pattern, ultimately traversing the entire subdomain in a spiral path scanning from the outermost layer towards the center. Each loop scan generates a series of continuous window coordinates, gradually filling the ROUTE list, which is the optimal path list required in this embodiment. This algorithm is suitable for cases where N is odd, ensuring that the generated path can completely traverse all windows without exceeding boundaries. The following is... Figure 10 The algorithm's textual description process:

[0110] Step 1: Initialization

[0111] 1. Initialization: Create an empty list ROUTE to store the generated path coordinates.

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

[0113] Step 2: Generate a spiral path layer by layer

[0114] 1. For each layer, generate the following four sides and update the path coordinates:

[0115] ① Bottom_function(ROUTE): Generates the bottom path of this layer, traversing the bottom grid points of this layer one by one in a zigzag pattern to the right, and adding the path to the ROUTE list.

[0116] ② Right_function(ROUTE): Generates the right-side path of this layer, traversing the right-side grid points of this layer one by one in a zigzag pattern, and adding the path to the ROUTE list.

[0117] ③ Top_function(ROUTE): Generates the top path of the layer by scanning the top grid points of the layer one by one in a zigzag pattern to the left and adding the path to the ROUTE list.

[0118] ④ Left_function(ROUTE): Generates the left path of this layer by scanning down the left grid points of this layer in a zigzag pattern and adding the path to the ROUTE list.

[0119] After each layer is completed, proceed to the next layer.

[0120] Step 3: Core Filling

[0121] 1. Call the preset mode based on the core value:

[0122] ① core == 1: Call Function_1 to add the center point at the end of the ROUTE list.

[0123] ② core == 3: Call Function_3 to add 9 preset path points to the end of the ROUTE list.

[0124] 2. Returns the complete list of routes.

[0125] Table 4 shows the results for a square subregion with a side length of 5, categorized by... Figure 5 After encoding all 25 sub-windows with two-dimensional coordinates one by one, the code automatically generates a "read-store-calculate-throw" strategy. Window 1 is the i-th window that needs to be scanned and processed by Stencil. Windows 2 to 5 are data-dependent windows determined by the template, and window 6 is a prefetch window.

[0126] Table 4 shows the "read-store-calculate-throw" strategy automatically generated by the code for a square subregion with a side length of 5.

[0127]

[0128] Figure 11 This demonstrates four minimum redundancy scan traversal paths for square subregions automatically generated by code.

[0129] The various embodiments of this disclosure have been described above. These descriptions are exemplary and not exhaustive, nor are they limited to the disclosed embodiments. Many modifications and variations will be apparent to those skilled in the art without departing from the scope and spirit of the described embodiments. The terminology used herein is chosen to best explain the principles, practical application, or technical improvements to the embodiments in the market, or to enable others skilled in the art to understand the 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 the following steps: First, the memory continuous direction 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. The thread block is one-to-one corresponding 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; 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 scanning traversal path mathematical model comprises the following steps: A square grid sub-region has n sub-windows, the two-dimensional coordinate coding is carried out on each sub-window, 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 number and column number 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 row number and column number coordinates of the sub-windows in the non-dependent relationship are saved in the buffer area; the variable queue count[n-1][6] records the 0 or 1 result 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 sum of all elements in the variable count is the objective function, and the optimization target of the scanning traversal path mathematical model is to minimize the objective function. The plurality of sub-windows in the yz plane are assigned unique numbers. Each thread is responsible for calculating a square grid sub-region.

2. The method of claim 1, wherein, There is a certain mathematical relationship between the side length row and the theoretical minimum redundancy, that is, the theoretical minimum redundancy of the square sub-region with the odd-numbered side length row is 6+2*(row*row-1)+row-1, and the theoretical minimum redundancy of the square sub-region with the even-numbered side length row is 6+2*(row*row-1)+row-2.

3. The method according to claim 1 or 2, characterized in that, The GPU scanning path optimization method further comprises the following steps:

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

5. The method of claim 3, wherein, For an N*N square region, N is even, and an algorithm for generating a scanning traversal path comprises the following steps: After the two-dimensional coordinate coding of 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 the ROUTE list, that is, the optimal path list, is gradually filled.

6. The method of claim 5, wherein, For an N*N square region, N is odd, and an algorithm for generating a scanning traversal path comprises the following steps: ​ 7. The method of claim 5, wherein, ​ After the two-dimensional coordinate coding of each sub-window in the square region, a path composed of a series of window coordinates is generated, the algorithm generates the bottom path, the right path, the top path and the left path through a loop, and then the core part is filled using a preset pattern, and finally the entire sub-domain is traversed 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