DCU architecture-oriented sparse triangle solution dynamic load balancing method and device and medium

Abstract

The invention discloses a DCU architecture-oriented sparse triangle solution dynamic load balancing method and device and a medium, and relates to the technical field of high-performance computing. The method comprises the following steps: rearranging a row sequence of a sparse matrix according to a dependency relationship; wherein the dependency relationship is determined according to an external dependency item and an internal dependency item of the thread block; and based on the sparse matrix of the rearranged row sequence, dividing the sparse matrix into a plurality of sub-matrixes according to the size of a thread scheduling unit, and determining a solving strategy according to the number of average non-zero elements in each sub-matrix. On one hand, the line dependency relationship of the sparse matrix is analyzed and rearranged in the preprocessing stage, so that calculation tasks are more uniformly distributed among DCU threads, and the situation that part of threads enter a waiting state due to the fact that data dependency is not met is reduced. On the other hand, the most suitable solving mode is dynamically selected according to the number of the non-zero elements in each row, and the calculation efficiency of overall solving is greatly improved.

Description

Technical Field

[0001] The present invention relates to the field of high-performance computing technology, and more specifically, to a sparse triangular solve dynamic load balancing method, device and medium for a DCU architecture. Background Art

[0002] Sparse Triangular Solve (abbreviated as SPTRSV) is a basic operator widely used in scientific computing and engineering computing, and plays an important role in fields such as numerical linear algebra, finite element analysis, optimization problems, graph computing, and machine learning. Especially when solving large-scale sparse linear systems (Ax = b), direct methods (such as LU decomposition, Cholesky decomposition) and iterative methods (such as conjugate gradient method, GMRES) usually involve SPTRSV as their core computing steps. For example, in finite element analysis and computational fluid dynamics (CFD), SPTRSV is used as part of matrix decomposition to accelerate the solution process. In the field of graph computing, such as in graph traversal algorithms based on sparse matrices, the SPTRSV operator can efficiently handle data dependencies.

[0003] The core goal of SPTRSV is to solve a linear equation system in the form of, where: L is a sparse upper / lower triangular matrix, most of whose elements are zero. x is the unknown vector to be solved. b is the known vector representing the right-hand side of the equation. Since L is an upper triangular or lower triangular matrix, the elements above or below its diagonal determine the data dependencies in the solution process, and this dependency poses challenges to parallel computing.

[0004] Currently, there are mainly two classic parallelization methods for SPTRSV: the Level-Set method and the Synchronization-Free method. The Level-Set method stratifies the topological structure of the matrix, groups the rows that can be parallelly computed into the same layer and advances layer by layer in order. The Synchronization-Free method relies on dynamic data dependency management, avoids explicit synchronization and directly computes according to the original row order of the matrix. However, the mainstream algorithms of the current SPTRSV operator still have many deficiencies in load balancing. Specifically, when dealing with matrices with uneven distribution of non-zero elements, the unified task partitioning strategy often fails to balance the loads of each thread. Some threads may be overloaded while others are idle, affecting the hardware utilization rate and the utilization rate of computing units is relatively low. The applicability to special matrices is insufficient. For some matrices with special structures or complex dependencies, such as matrices with large differences in depth dependency levels, the algorithm cannot achieve the expected acceleration effect. Summary of the Invention

[0005] To solve the above technical problems, the present invention provides a dynamic load balancing method, device, and medium for sparse triangular solution oriented to the DCU architecture, aiming to solve the pain points of the existing SPTRSV algorithm under the DCU architecture, such as thread load imbalance, waste of computing resources, and insufficient adaptability to special matrices. Through relying on the deep reconstruction and dynamic dual-mode scheduling mechanism, efficient parallel solution is achieved. Based on the Dawning DCU (Deep Computing Unit) architecture, the present invention provides an optimized dynamic scheduling scheme, which integrates fine-grained thread scheduling and cooperative thread bundle Wavefront scheduling strategies. In addition, the present invention also adds a simple and efficient preprocessing mechanism. By constructing reasonable row dependency relationships and optimizing row order in advance, unnecessary waiting can be avoided in the subsequent sparse triangular solution process, improving the parallel execution efficiency and load balancing.

[0006] In the first aspect, the present invention provides a dynamic load balancing method for sparse triangular solution oriented to the DCU architecture, and the method includes:

[0007] Rearranging the row order of the sparse matrix according to the dependency relationship; wherein, the dependency relationship is determined according to the external dependencies and internal dependencies of the thread block.

[0008] Based on the sparse matrix with the rearranged row order, dividing it into several sub-matrices according to the size of the thread scheduling unit, and determining the solution strategy according to the average number of non-zero elements in each sub-matrix.

[0009] Further, rearranging the row order of the sparse matrix according to the dependency relationship includes:

[0010] Initializing the completion flags of each row within the thread block using a shared memory array; wherein, the thread block includes multiple thread scheduling units, each thread scheduling unit includes multiple threads, and each thread can access the shared memory within the same computing unit. The shared memory array is a flag data for marking the calculation status of each row of the sparse matrix.

[0011] Processing the external dependencies of the thread block to determine the first local maximum dependency depth of each row.

[0012] Processing the internal dependencies of the thread block to determine the second local maximum dependency depth of each row.

[0013] Calculating the final dependency depth of each row according to the first local maximum dependency depth and the second local maximum dependency depth.

[0014] Reordering the sparse matrix from small to large according to the size of the final dependency depth of each row of the sparse matrix.

[0015] Further, processing the external dependencies of the thread block to determine the first local maximum dependency depth of each row includes:

[0016] Each thread traverses the non-zero elements in the corresponding row. For the dependent rows outside the current thread block, atomic operations and polling are used to ensure that the dependent data is ready, and the first local maximum dependency depth is updated.

[0017] Furthermore, the dependencies within the thread block are processed to determine the second local maximum dependency depth for each row, including:

[0018] For the dependencies within the same thread block, the threads are synchronized through flags in shared memory. Using the polling waiting method, it is ensured that the calculation of the dependent rows is completed, and the second local maximum dependency depth is updated.

[0019] Furthermore, according to the first local maximum dependency depth and the second local maximum dependency depth, the final dependency depth for each row is calculated through an in-reduction operation.

[0020] Furthermore, the solution strategy is determined according to the average number of non-zero elements in each sub-matrix, including:

[0021] Based on a preset threshold, when the average number of non-zero elements exceeds the preset threshold, a warp-level solution strategy is adopted;

[0022] When the average number of non-zero elements does not exceed the preset threshold, a thread-level solution strategy is adopted.

[0023] Furthermore, the warp-level solution strategy includes:

[0024] Calculate the forward accumulation term of the current row, that is, traverse all non-zero elements in the i-th row of the sparse matrix;

[0025] For each non-zero element, query the status array get_value[col] corresponding to its column index to determine whether the solution vector x[col] has been calculated;

[0026] If x[col] has been calculated, directly solve x[i] and store it. At the same time, update get_value[i]=1 to notify subsequent calculation tasks that depend on x[i];

[0027] If x[col] has not been calculated, put the thread into the sleep state.

[0028] Furthermore, the thread-level solution strategy includes:

[0029] Let the threads within a thread scheduling unit jointly calculate a row of the sparse matrix;

[0030] The different threads of each thread scheduling unit are responsible for different non-zero elements and calculate the partial sums of the corresponding rows respectively;

[0031] Use Wavefront-level reduction to merge the partial sums calculated by all threads, and finally calculate x[i] and update the status array get_value.

[0032] In a second aspect, the present invention provides a sparse triangular solution dynamic load balancing device for a DCU architecture. The device includes:

[0033] A preprocessing unit configured to rearrange the row order of the sparse matrix according to the dependency relationship; wherein, the dependency relationship is determined according to the external dependencies and internal dependencies of the thread blocks;

[0034] A solving unit configured to divide the sparse matrix with the rearranged row order into several sub-matrices according to the size of the thread scheduling unit, and determine the solving strategy according to the average number of non-zero elements in each sub-matrix.

[0035] In a third aspect, the present invention provides a readable storage medium storing one or more programs, and the one or more programs can be executed by one or more processors to implement the method as described above.

[0036] The present invention has at least the following beneficial effects:

[0037] (1) Load balancing optimization: By analyzing and rearranging the row dependency relationship of the sparse matrix in the preprocessing stage, the present invention makes the computing tasks more evenly distributed among the DCU threads, thereby reducing the situation where some threads enter the waiting state due to unmet data dependencies. For matrices with complex dependency relationships, the thread idle waiting time of the present invention is reduced by about 40% compared with the traditional non-synchronous method. The utilization rate of the computing unit is improved, and higher parallel execution efficiency is achieved.

[0038] (2) Computation efficiency improvement: The present invention integrates a dual-mode strategy of fine-grained thread scheduling and cooperative thread warp Wavefront scheduling, and dynamically selects the most suitable solving method according to the number of non-zero elements in each row. It effectively overcomes the problems of thread idleness and uneven load caused by a single solving strategy in traditional non-synchronous algorithms, and greatly improves the overall solving computation efficiency. Under the same hardware and example conditions, as the matrix scale increases, the peak floating-point operation performance of the present invention also increases; compared with the traditional method, the average floating-point operation performance is increased by about 2.2 times, and the overall solving time is shortened by about 30%.

[0039] (3) Dependence reduction: In the preprocessing stage, through the initialization of shared memory and the fast processing of local dependencies, combined with atomic operations and local synchronization mechanisms, the reduction of local dependency depth and flag update are achieved, so that the dependence on global synchronization in the subsequent solving process is greatly reduced. Description of the Drawings

[0040] Figure 1 The flowchart of a dynamic load balancing method for sparse triangular solution oriented to the DCU architecture according to an embodiment of the present invention is shown.

[0041] Figure 2 The flowchart of preprocessing analysis in a dynamic load balancing method for sparse triangular solution oriented to the DCU architecture according to an embodiment of the present invention is shown.

[0042] Figure 3 The flowchart of the solution step in a dynamic load balancing method for sparse triangular solution oriented to the DCU architecture according to an embodiment of the present invention is shown.

[0043] Figure 4 The solution example diagram of the SPTRSV operator according to an embodiment of the present invention is shown.

[0044] Figure 5 The sparse lower triangular matrix L stored in CSR format according to an embodiment of the present invention is shown, as well as the schematic diagram of the dependency level information between each row of the sparse matrix. Among them, (a) is the lower triangular matrix L; (b) is the matrix in CSR storage format; (c) is the dependency level information.

[0045] Figure 6 The flowchart of the dynamic load balancing method for sparse triangular solution oriented to the DCU architecture in the preprocessing stage according to an embodiment of the present invention is shown.

[0046] Figure 7 The flowchart of the dynamic load balancing method for sparse triangular solution oriented to the DCU architecture in the solution stage according to an embodiment of the present invention is shown.

[0047] Figure 8 The structure diagram of a dynamic load balancing device for sparse triangular solution oriented to the DCU architecture according to an embodiment of the present invention is shown. Detailed implementation manners

[0048] To enable those skilled in the art to better understand the technical solutions of the present invention, the present invention will be described in detail below in conjunction with the accompanying drawings and specific implementation manners. The embodiments of the present invention will be further described in detail below in conjunction with the accompanying drawings and specific embodiments, but it is not a limitation to the present invention. For the various steps described herein, if there is no necessity for the front-back relationship between them, the order in which they are described as examples herein should not be regarded as a limitation. Those skilled in the art should know that they can be adjusted in order as long as the logic between them is not destroyed and the entire process cannot be realized.

[0049] Embodiment 1:

[0050] An embodiment of the present invention provides a dynamic load balancing method for sparse triangular solution oriented to the DCU architecture, as follows Figure 1 As shown, the dynamic load balancing method for sparse triangular solution oriented to the DCU architecture includes steps S100 and S200, where steps S100 and S200 are a preprocessing analysis step and a solution step respectively.

[0051] S100: Rearrange the row order of the sparse matrix according to the dependency relationship; wherein, the dependency relationship is determined according to the external dependencies and internal dependencies of the thread block.

[0052] In this embodiment, step S100 is a preprocessing step. Through dependency analysis and sorting, the row order of the sparse matrix is rearranged according to the dependency relationship. Compared with directly processing in the original order of the matrix, the unsorted order may cause some rows to be scheduled for calculation before the dependencies are completed, forcing the DCU threads to enter the waiting state. After preprocessing and sorting, the computing tasks can be more evenly distributed, and the subsequent solution phase can make more full use of the parallel capabilities of the DCU, avoiding partial thread idleness caused by uneven data dependencies, and reducing the frequency of global synchronization and polling waiting.

[0053] In some embodiments, as Figure 2 shown, it is a flowchart of preprocessing analysis, which can be specifically implemented through the following steps S101 - S105.

[0054] S101: Initialize the completion flags of each row within the thread block using a shared memory array.

[0055] In this embodiment, the purpose of step S101 is to achieve shared memory initialization. The done_array array in the shared memory is used to initialize the "completion flags" of each row within the block, preparing for subsequent local dependency processing. In the DCU architecture, a block is a thread block, which is the basic unit of thread scheduling and resource management. A block consists of multiple Wavefronts (thread scheduling units), and each Wavefront contains 64 threads, and these 64 threads can access the shared memory within the same Compute Unit (CU). In addition, done_array is a flag array used to mark the calculation status of each row. Through this array, it can be quickly determined whether each row within the block has completed the calculation, thus providing support for the quick determination and processing of subsequent local dependency relationships.

[0056] S102: Process the external dependencies of the thread block and determine the first local maximum dependency depth of each row.

[0057] Exemplarily, the way to process external dependencies of a block is as follows: Each thread traverses the non-zero elements in the corresponding row. For the dependent rows outside the current block, atomic operations and polling are used to ensure that the dependent data is ready, and the local maximum dependency depth is updated.

[0058] S103: Process the internal dependencies of the thread block to determine the second local maximum dependency depth of each row.

[0059] Exemplarily, the way to process internal dependencies of a block is as follows: For the dependencies within the same block, the threads are synchronized through flags in the shared memory. Similarly, the polling waiting method is adopted to ensure that the calculation of the dependent rows is completed, and the second local maximum dependency depth is updated.

[0060] S104: Calculate the final dependency depth of each row based on the first local maximum dependency depth and the second local maximum dependency depth.

[0061] Exemplarily, based on the first local maximum dependency depth and the second local maximum dependency depth, the final dependency depth of this row is calculated through the reduction operation within the Wavefront, thereby realizing the reduction of the dependency depth and the update of the flag.

[0062] S105: Reorder the sparse matrix from small to large according to the size of the final dependency depth of each row of the sparse matrix.

[0063] S200: Based on the reordered sparse matrix, divide it into several sub-matrices according to the size of the thread scheduling unit, and determine the solution strategy according to the average number of non-zero elements in each sub-matrix.

[0064] Step S200 is a solution step, in which an adaptive fine-grained thread scheduling and cooperative thread bundle Wavefront scheduling dual-mode solution strategy is implemented. Specifically, the input sparse matrix is divided into several sub-matrices according to the size of the Wavefront of the DCU hardware, and then the most suitable solution strategy is selected according to the average number of non-zero elements in each sub-matrix. When the average number of non-zero elements exceeds the preset threshold, the thread bundle-level solution strategy is adopted, that is, one Wavefront processes one row in the sparse matrix; otherwise, the thread-level solution strategy is adopted, that is, a single thread processes one row, and one Wavefront is responsible for processing 64 rows. Considering the characteristic that each Wavefront in the DCU architecture contains 64 threads, through testing a large number of matrices, it is found that when the preset threshold is 16, the overall effect obtained by selecting the solution strategy is the best. This strategy can avoid the problems of thread idleness and load imbalance faced by a single solution strategy, enable each thread to participate in the calculation as much as possible, maximize the DCU computing utilization rate, and improve the throughput.

[0065] In some embodiments, as Figure 3 shown, it is a flowchart of the solution steps. Step S200 specifically includes the following steps:

[0066] Step S201, strategy selection: For the rearranged sparse matrix, dynamically select a suitable solution strategy according to the number of non-zero elements in each row.

[0067] Step S202, calculate thread index: Determine the calculation task responsible for the current thread by calculating the global thread index (global_id) and the Wavefront (wf_id) to which the thread belongs, so as to ensure reasonable task allocation.

[0068] Step S203, perform thread-level calculation based on the warp-level solution strategy: Under this strategy, a single thread is responsible for solving the unknowns of a certain row. First, calculate the forward accumulation term of the current row, that is, traverse all non-zero elements in the i-th row of the sparse matrix. For each non-zero element, query the status array get_value[col] corresponding to its column index to determine whether the solution vector x[col] has been calculated. If x[col] has been calculated, directly solve x[i] and store it, and at the same time update get_value[i]=1 to notify subsequent calculation tasks that depend on x[i]. If x[col] has not been calculated, due to the special nature of the sleep mechanism supported by the DCU hardware, the thread enters the sleep state to release computing resources.

[0069] Step S204, perform Wavefront-level calculation based on the thread-level solution strategy: In this strategy, the threads within a Wavefront jointly calculate a row of the sparse matrix. Different threads in each Wavefront are responsible for different non-zero elements and calculate the partial sums of this row respectively. Finally, use Wavefront-level reduction to merge the sums calculated by all threads. Finally, calculate x[i] and update the get_value array. The thread dependency waiting mechanism in this process is the same as the thread-level strategy.

[0070] Embodiment 2:

[0071] Based on the sparse triangular solution dynamic load balancing method for the DCU architecture provided in Embodiment 1, this embodiment provides a specific application example of the above method. As Figure 4 shown, it is a solution example diagram of the SPTRSV operator, Figure 3 showing the relationship between the sparse matrix L and the solution vector X, Figure 4 in which, L mn represents the input sparse upper / lower triangular matrix, X m represents the output dense vector, B mDenote the input dense vector, where m = 1, 2, 3... 9; n = 1, 2, 3... 9.

[0072] As Figure 5 shown, Figure 5 shows the sparse lower triangular matrix L stored in CSR format with L = b as an example, and the dependency level information between each row of this sparse matrix. Figure 4 In (a) is the lower triangular matrix L; (b) is the matrix in CSR storage format; (c) is the dependency level information. Among them, the Row_Ptr row pointer array stores the starting position of the non-zero elements of each row in the values array. Row_ptr[i] points to the starting index of the i-th row in the values array. The Col_Ind array stores the column index position of each non-zero element, corresponding one-to-one with the Values array. The Value array only stores the non-zero values of the matrix, arranged in row-major order. The dependency level information of the sparse matrix is as Figure 4 shown in (c), which represents the first layer including the 0th and 1st rows of the sparse matrix. These rows do not depend on other rows and can be calculated immediately. The second layer includes the 2nd and 4th rows, which depend on row 1 of Level 0. And so on for the following several levels. This indicates that in the sparse triangular solution process, Level 0 must be calculated first before calculating Level 1, and so on, to ensure the correct data dependency relationship. This level division can reduce the synchronization waiting and improve the parallel computing efficiency.

[0073] As Figure 6As shown, it is the flowchart of the dynamic load balancing method for sparse triangular solution in the preprocessing stage for the DCU architecture. First, in the initialization stage, the dependency_Depth array is established, which is used to store the dependency depths of each row of the sparse matrix. At the same time, the local ID (lid) of each thread in the Wavefront, the warp ID (wid) where the Wavefront is located in the block, and the starting row and the row to be processed (row) of the current thread block are initialized. Subsequently, the non-zero elements of the current row are traversed, and the column index of each non-zero element is obtained according to the csr_col_ind array. If the column index is less than the starting row of the current calculation block, it indicates that this row depends on other rows outside the block; otherwise, this row depends on the rows within the block. For the processing of dependencies outside the block, the method of polling the global shared memory done_array is adopted - when the completion flag of the dependent row has not been set to 1, the thread needs to wait; once the flag is set, the local dependency depth size of this row is updated and stored in the variable local_max. For dependencies within the block, the local shared memory local_done_array within the block is polled for processing, and local_max is also updated to ensure synchronization. Next, at the last thread of each warp (Wavefront), a reduction is performed to perform a maximum reduction on the local_max calculated by each thread within the same warp (Wavefront) to obtain the global maximum dependency depth of the current row and write it back to the dependency_Depth array. Finally, according to the calculated matrix dependency depth array, the matrix rows are sorted so that they are rearranged in ascending order of dependency depth.

[0074] In the polling waiting mechanism of the present invention, a Sleep mechanism and an exponential increment strategy are introduced to optimize the occupation of computing resources by threads during the waiting process. The Sleep mechanism is applicable to DCU hardware. When this mechanism is enabled (i.e., SLEEP == true), the statement __builtin_amdgcn_s_sleep(1) allows the thread to sleep briefly for 1 clock cycle, thereby reducing the continuous occupation of computing units and improving resource utilization. In addition, the present invention also adopts an exponential increment strategy to dynamically adjust the sleep time to avoid unnecessary computing overhead caused by the thread frequently polling the done_array. Specifically, initially times_through = 1 and it gradually increases as the loop is executed. When times_through < 32, a linear growth (+1) is adopted to ensure that the thread can quickly respond to status changes in the early stage; when times_through ≥ 32, an exponential growth (×2) is adopted to further reduce excessive polling. To prevent times_through from being too large and causing the thread to sleep for a long time and affecting the overall computing efficiency, a maximum value of 1024 is set to ensure that the waiting process can adapt to different execution situations and effectively reduce the waste of DCU computing resources.

[0075] As Figure 7 shown, it is a flowchart of the dynamic load balancing method for sparse triangular solution oriented to the DCU architecture during the solution stage. First, a done_array array is established to store the calculation status of each row of the sparse matrix, where done_array[row] = 1 indicates that the calculation of this row has been completed, while done_array[row] = 0 indicates that the calculation is still in progress. Subsequently, the average number of non-zero elements in each row of each sub-matrix is calculated and obtained for subsequent determination of the processing strategy for this row. According to different calculation methods, the solution process can be divided into thread-level solution and Wavefront-level solution.

[0076] In the thread-level solution method, each thread uses local_id to calculate its corresponding row number. For the current row, the thread starts traversing all non-zero elements of this row from Row_ptr[row] and uses atomic operations to check whether its dependent rows have completed the calculation. If the dependency is satisfied, an accumulative calculation is performed to update the partial sum local_sum = Value[j] * x[col]. When encountering the diagonal element of this row, solve the variable: xi = (d_b[row] - local_sum) / diagonal element value. And write the calculation result into x[row]. To ensure global memory consistency, __threadfence_system() is called for synchronization, and finally done_array[row] is set to 1 to mark the completion of the calculation of this row.

[0077] In the Wavefront-level solution method, each thread only processes a part of the non-zero elements in the current row and calculates the corresponding partial sum separately. Subsequently, the sum of the row is calculated through in-Wavefront reduction and finally written into the x array. Similar to the thread-level solution, this calculation process also calls __threadfence_system() to ensure global synchronization, and sets done_array[row] to 1 after the calculation is completed, indicating that the calculation of this row is completed.

[0078] Finally, through the above process, the solution of the SPTRSV operator is achieved.

[0079] Embodiment 3:

[0080] The embodiment of the present invention also provides a sparse triangular solution dynamic load balancing device for a DCU architecture, as Figure 8 shown, the device includes:

[0081] A preprocessing unit 801, configured to rearrange the row order of the sparse matrix according to the dependency relationship; wherein, the dependency relationship is determined according to the external dependencies and internal dependencies of the thread blocks;

[0082] A solution unit 802, configured to divide the sparse matrix with the rearranged row order into several sub-matrices according to the size of the thread scheduling unit, and determine the solution strategy according to the average number of non-zero elements in each sub-matrix.

[0083] In some embodiments, the preprocessing unit is further configured to:

[0084] Initialize the completion flags of each row within the thread block using a shared memory array; wherein, the thread block includes multiple thread scheduling units, each thread scheduling unit includes multiple threads, and each thread can access the shared memory within the same computing unit, and the shared memory array is a flag data for marking the calculation status of each row of the sparse matrix;

[0085] Process the external dependencies of the thread block to determine the first local maximum dependency depth of each row;

[0086] Process the internal dependencies of the thread block to determine the second local maximum dependency depth of each row;

[0087] Calculate the final dependency depth of each row according to the first local maximum dependency depth and the second local maximum dependency depth;

[0088] Reorder the sparse matrix from small to large according to the size of the final dependency depth of each row of the sparse matrix.

[0089] In some embodiments, the preprocessing unit is further configured to:

[0090] Each thread traverses the non-zero elements in the corresponding row. For the dependent rows outside the current thread block, atomic operations and polling are used to ensure that the dependent data is ready, and the first local maximum dependence depth is updated.

[0091] In some embodiments, the preprocessing unit is further configured to:

[0092] For the dependencies within the same thread block, the threads are synchronized through flags in shared memory, and polling is used to ensure that the calculation of the dependent rows is completed, and the second local maximum dependence depth is updated.

[0093] In some embodiments, the preprocessing unit is further configured to:

[0094] According to the first local maximum dependence depth and the second local maximum dependence depth, the final dependence depth of each row is calculated through an in-reduction operation.

[0095] In some embodiments, the solving unit is further configured to:

[0096] Based on a preset threshold, when the average number of non-zero elements exceeds the preset threshold, a warp-level solving strategy is adopted;

[0097] When the average number of non-zero elements does not exceed the preset threshold, a thread-level solving strategy is adopted.

[0098] In some embodiments, the warp-level solving strategy includes:

[0099] Calculate the forward accumulation term of the current row, that is, traverse all non-zero elements in the i-th row of the sparse matrix;

[0100] For each non-zero element, query the status array get_value[col] corresponding to its column index to determine whether the solution vector x[col] has been calculated;

[0101] If x[col] has been calculated, directly solve x[i] and store it, and at the same time update get_value[i]=1 to notify subsequent calculation tasks that depend on x[i];

[0102] If x[col] has not been calculated, put the thread into the sleep state.

[0103] In some embodiments, the thread-level solving strategy includes:

[0104] Let the threads within a thread scheduling unit jointly calculate a row of the sparse matrix;

[0105] The different threads of each thread scheduling unit are responsible for different non-zero elements and calculate the partial sums of the corresponding rows respectively;

[0106] Use Wavefront-level reduction to merge the partial sums calculated by all threads, and finally calculate x[i] and update the status array get_value.

[0107] It should be noted that the structures of the various sparse triangular solution dynamic load balancing devices for the DCU architecture described in this embodiment belong to the same technical concept as the sparse triangular solution dynamic load balancing method for the DCU architecture described previously, and achieve the same beneficial effects through the same principle, which will not be elaborated here.

[0108] The embodiment of the present invention also provides a readable storage medium, which stores one or more programs, and the one or more programs can be executed by one or more processors to implement the method described in any of the above embodiments.

[0109] Finally, it should be noted that the above embodiments are only used to illustrate the technical solutions of the present invention, rather than to limit them; although the present invention has been described in detail with reference to the foregoing embodiments, those of ordinary skill in the art should understand that they can still modify the technical solutions recorded in the foregoing embodiments, or perform equivalent replacements on some or all of the technical features; and these modifications or replacements do not cause the essence of the corresponding technical solutions to deviate from the scope of the technical solutions of the embodiments of the present invention.

Claims

1. A dynamic load balancing method for sparse triangular solution oriented to the DCU architecture, characterized in that The method includes: Rearranging the row order of the sparse matrix according to the dependency relationship; wherein, the dependency relationship is determined according to the external dependencies and internal dependencies of the thread block; Based on the sparse matrix with rearranged row order, dividing it into several sub-matrices according to the size of the thread scheduling unit, and determining the solution strategy according to the average number of non-zero elements in each sub-matrix.

2. The dynamic load balancing method for sparse triangular solution for DCU architecture according to claim 1, wherein Rearranging the row order of the sparse matrix according to the dependency relationship includes: Initializing the completion flags of each row within the thread block using a shared memory array; wherein, the thread block includes multiple thread scheduling units, each thread scheduling unit includes multiple threads, and each thread can access the shared memory within the same computing unit, and the shared memory array is a flag data for marking the calculation status of each row of the sparse matrix; Processing the external dependencies of the thread block to determine the first local maximum dependency depth of each row; Processing the internal dependencies of the thread block to determine the second local maximum dependency depth of each row; Calculating the final dependency depth of each row according to the first local maximum dependency depth and the second local maximum dependency depth; Reordering the sparse matrix from small to large according to the size of the final dependency depth of each row of the sparse matrix.

3. The dynamic load balancing method for sparse triangular solution oriented to the DCU architecture according to claim 2, characterized in that, Processing the external dependencies of the thread block to determine the first local maximum dependency depth of each row, including: Each thread traverses the non-zero elements in the corresponding row, and for the dependent rows outside the current thread block, ensures that the dependent data is ready through atomic operations and polling waits, and updates the first local maximum dependency depth.

4. The dynamic load balancing method for sparse triangular solution oriented to the DCU architecture according to claim 2, wherein Processing the internal dependencies of the thread block to determine the second local maximum dependency depth of each row, including: For the dependencies within the same thread block, the threads are synchronized through the flags in the shared memory, and in a polling wait manner, ensures that the calculation of the dependent rows is completed, and updates the second local maximum dependency depth.

5. The sparse triangular solution dynamic load balancing method for a DCU architecture according to claim 2, wherein Calculating the final dependency depth of each row through an in-reduction operation according to the first local maximum dependency depth and the second local maximum dependency depth.

6. The sparse triangular solution dynamic load balancing method for the DCU architecture according to claim 1, wherein Determining the solution strategy according to the average number of non-zero elements in each sub-matrix, including; Based on a preset threshold, when the average number of non-zero elements exceeds the preset threshold, adopting a warp-level solution strategy; When the average number of non-zero elements does not exceed the preset threshold, adopting a thread-level solution strategy.

7. The sparse triangular solution dynamic load balancing method for the DCU architecture according to claim 6, characterized in that The warp-level solution strategy includes: Calculating the forward accumulation term of the current row, that is, traversing all non-zero elements in the i-th row of the sparse matrix; For each non-zero element, query the status array get_value[col] corresponding to its column index to determine whether the solution vector x[col] has been calculated; If x[col] has been calculated, directly solve x[i] and store it, and at the same time update get_value[i]=1 to notify subsequent calculation tasks that depend on x[i]; If x[col] has not been calculated, put the thread into a sleep state.

8. The dynamic load balancing method for sparse triangular solution oriented to the DCU architecture according to claim 6, wherein The thread-level solution strategy includes: Making the threads within a thread scheduling unit jointly calculate a row of the sparse matrix; The different threads of each thread scheduling unit are responsible for different non-zero elements and calculate the partial sums of the corresponding rows respectively. Use Wavefront-level reduction to merge the partial sums calculated by all threads, and finally calculate x[i] and update the status array get_value.

9. A dynamic load balancing device for sparse triangular solution oriented to the DCU architecture, characterized in that The device includes: A preprocessing unit configured to rearrange the row order of the sparse matrix according to the dependency relationship; wherein the dependency relationship is determined according to the external dependencies and internal dependencies of the thread block; A solving unit configured to divide the sparse matrix with the rearranged row order into several sub-matrices according to the size of the thread scheduling unit, and determine the solving strategy according to the average number of non-zero elements in each sub-matrix.

10. A non-transitory computer-readable storage medium storing instructions that, when executed by a processor, perform the method according to any one of claims 1 to 8.

Images