An Efficient GPU-Accelerated Algorithm for Solving Dynamic Response of Fluid-Saturated Porous Media

Abstract

The traditional finite element program is executed on the CPU; however, it is challenging for the CPU to compute the ultra-large scale finite element model. In this paper, we present a set of efficient algorithms based on GPU acceleration technology for the dynamic response of fluid-saturated porous media, named PNAM, encompassing the assembly of the global matrix and the iterative solution of equations. In the assembly part, the CSR storage format of the global matrix is directly obtained from the element matrix. For data with two million degrees of freedom, it merely takes approximately 1 s to generate all the data of global matrices, which is significantly superior to the CPU version. Regarding the iterative solution of equations, a novel algorithm based on the CUDA kernel function is proposed. For a data set with two million degrees of freedom, it takes only about 0.05 s to compute an iterative step and transfer the data to the CPU. The program is designed to calculate either in single or double precision. The change in precision has little impact on the assembly of the global matrix, but the calculation time of double precision is generally 1.5 to 2 times that of single precision in the iterative solution part for a model with 2 million degrees of freedom. PNAM has high computational efficiency and great compatibility, which can be used to solve not only saturated fluid problems but also a variety of other problems.

1. Introduction

Finite element calculations are usually performed on the central processing unit (CPU), but this sequential execution makes it difficult to handle large amounts of data. At present, the amount of data needed for large-scale scientific computing is often very large, and the traditional CPU computing resources are gradually becoming difficult to meet with modern computing needs. The CPU achieves lower computational latency and higher computational efficiency by increasing the clock rate, increasing the number of transistors, increasing the cache size, and increasing the manufacturing power, but this improvement is limited by the physical properties. The graphics processing unit (GPU) is a good solution to the shortcomings of the CPU in processing large-scale data. As a new technical solution, parallelization technology based on GPU plays an important role in today’s scientific computing field, which can improve the performance of computing [1,2]. CPUs sacrifice the computational throughput of processors to improve the performance of single cores, while GPUs are basically multicore processors with thousands of cores designed to deliver high computational throughput. This different design concept requires the new development of specialized algorithms to take advantage of the potential performance of GPU hardware [3].

At present, many researchers have focused on GPUs in various scientific and engineering computing fields and have made significant breakthroughs [4,5,6]. The arrival of NVIDIA’s CUDA also provides various developers with the convenience of general-purpose computing on the GPU [7], allowing it to be applied to a variety of scenarios [8,9]. However, the complexity of GPU programming, the special requirements of memory access patterns, and the high demand for algorithmic parallelization make it challenging to port applications from traditional CPU programming to GPU programming. The finite element method can be divided into four parts: the discretization of the model, the generation of the element matrix, the assembly of the global matrix, and the solution of the equations. The most time-consuming part is the assembly of the global matrix and the solution of the equations, which generally consumes 70–90% of the entire finite element simulation time [10]. Furthermore, in some cases, the generation of the element matrices dominates the performance for higher-order elements [11].

Although the discretization of the model does not take a lot of time in most instances, it still takes a certain amount of time to complete for the very large model. GPU has been widely used in the discrete domain by applying a grid of elements [12,13,14], and the acceleration of some algorithms can reach 100–200 times compared to general commercial software [15].

While the assembly process seems well-suited for parallelization on GPU, it is subject to race conditions when updating the global matrix. These race conditions arise due to parallel memory location updates caused by common nodes between elements. There is no proper global synchronization primitive on the GPU to handle race conditions. To solve this problem, three different techniques have been proposed: atomic manipulation, graph coloring [16,17,18], and task dependency graph [19]. The global matrix is a sparse matrix with many zero elements that can be compressed in some way. To save storage and memory space, the coordinate (COO) format was initially proposed and then replaced by the compressed sparse row (CSR) format with higher compression efficiency. CSR format can save more than 99% of storage and memory space, allowing the program to compute very large models. In solving equations, conjugate gradient (CG) is a widely used method [20], but this method has the problem of poor convergence [21]. In order to solve the problem of poor convergence, preconditioned CG (PCG) was proposed and achieved very good effects [22,23].

In the previous literature, the finite element calculation of the dynamic response of saturated fluid porous media was mostly addressed by implicit, semi-explicit, or explicit methods [24,25,26], which entailed issues such as low solution efficiency and low calculation accuracy. Based on a fully explicit time-domain integration method with second-order accuracy [27], this paper employs GPU acceleration technology to expedite calculation and solve the problem of traditional CPU programs. One of the challenges in parallel computing within finite element applications is to handle the thread conflict problem. In this paper, a novel approach is proposed to resolve the thread collision issue, namely the preprocessed non-assembly method (PNAM). The algorithm directly acquires the CSR storage format data of the global matrix from the element matrix and omits the intermediate process of generating the global matrix. This saves a considerable amount of GPU memory space, as the global matrix is a sparse matrix with numerous zero elements that are not necessary for the computation. Thus, it can be utilized to calculate the model with high mesh precision and further enhance the accuracy of the calculation results. In the iterative solution section, a new scheme is devised. It computes one formula at a time and is fast but demands a large GPU memory space. Nevertheless, due to the utilization of the CSR storage format for data storage, to a certain extent, the problem of the iterative solution requiring a large GPU memory space is mitigated. For a GPU with 4 GB of memory, it can be employed to compute the dynamic response problem of a two-dimensional saturated fluid porous media with a size close to 10 million elements (single precision). When transferring the resulting data from the GPU to the CPU, the data are encapsulated, which significantly reduces the data transfer time.

The GPU of the test computer is the RTX-3050, the CPU is the i5-12450H, and the Matlab version is R2021b. The RTX-3050 has 4 GB of RAM and 2560 CUDA cores. The i5-12450H has 16 GB of memory and has eight cores. The CUDA version is 12.3, the driver version is 546.17, and the nvcc release is 12.2, V12.2.91.

2. Theory and Methods

In the study of near-field wave propagation of the fluid-saturated porous media (FSPM), the u-w-p model proposed by Zienkiewicz is the most extensive [28]. The model adopts the displacement of solid soil skeleton, the relative displacement between pore fluid and soil skeleton, and the pore pressure. Through different simplifying assumptions, the model can be further transformed into other simplified forms, such as the u-w model, u-U model, and u-p model [29,30,31]. Since the u-p form equation has fewer unknowns and the key variable p can be directly obtained when solving compared with other formal equations, the u-p form is adopted when solving the problem of near-field wave propagation of FSPM. When solving the u-p formal equation, a complete explicit interlacing algorithm for FSPM near-field wave propagation in the time domain presented in the literature [27] is adopted. This algorithm can greatly reduce the computational workload and significantly improve computational efficiency.

2.1. Wave Equation of Saturated Two-Phase Media in u-p Form

It is assumed that the solid phase in a saturated two-phase medium is isotropic linear elastic medium, and its stress–strain relation obeys Hooke’s law. The liquid phase (pore fluid) is an ideal compressible fluid, and the elastic wave equation of saturated two-phase media in the form of u-p with solid phase displacement u and pore fluid pressure p as the basic unknowns can be expressed in the following form when the volume force is ignored:

$${\mathit{L}}^{\mathit{T}}\mathit{D}\mathit{L}\mathit{u}+{\mathit{L}}^{\mathit{T}}\mathit{m}p-\rho \ddot{\mathit{u}}=0$$

(1)

$${\mathit{m}}^{\mathit{T}}\mathit{L}\dot{\mathit{u}}+{\displaystyle \frac{k_f}{{\mu }_f}}{\nabla }^T\nabla p-p_f{\displaystyle \frac{k_f}{{\mu }_f}}{\mathit{m}}^{\mathit{T}}\mathit{L}\ddot{\mathit{u}}-{\displaystyle \frac{\mathit{n}}{{\mathit{E}}_{\mathit{w}}}}\dot{p}=0$$

(2)

Equation (1) represents the overall momentum balance of saturated two-phase media, and Equation (2) is the conservation equation of pore fluid flow. Where: $\mathit{u}$, $\dot{\mathit{u}}$, and $\ddot{\mathit{u}}$ are the displacement, velocity, and acceleration vectors of the solid phase, respectively.

Using the standard Galerkin finite element method, the weak Galerkin formula of the dynamic equation of a saturated two-phase medium in u-p format after finite element discretization is obtained:

$$\mathit{M}\ddot{\mathit{u}}+\mathit{K}\mathit{u}-\mathit{Q}\mathit{p}={\mathit{f}}_{\mathit{u}}$$

(3)

$$\mathit{J}\mathit{p}+\mathit{S}\dot{\mathit{p}}-{\mathit{Q}}^{\mathit{T}}\dot{\mathit{u}}={\mathit{f}}_{\mathit{p}}$$

(4)

where $\mathit{M}$ is the mass matrix; $\mathit{K}$ is the stiffness matrix; $\mathit{Q}$ is the coupling matrix of solid phase and pore hydrodynamic response; $\mathit{J}$ is the percolation matrix of pore fluid; $\mathit{S}$ is the compression matrix of pore fluid; and ${\mathit{f}}_{\mathit{u}}$ and ${\mathit{f}}_{\mathit{p}}$ are boundary load vectors acting on solid phase and pore fluid, respectively.

The expressions of each matrix and vector in Equations (3) and (4) are as follows:

$$\mathit{M}={\int }_{\Omega }\rho N_u^T{N}_{u}d\Omega$$

(5)

$$\mathit{K}={\int }_{\Omega }{\left(L{N}_u\right)}^{T}D(L{N}_u)d\Omega$$

(6)

$$\mathit{Q}={\int }_{\Omega }{\left(L{N}_u\right)}^{T}m(L{N}_u)d\Omega$$

(7)

$$\mathit{S}={\displaystyle \frac{E_w}{n}}{\int }_{\Omega }{N}_p^T{N}_{p}d\Omega$$

(8)

$$\mathit{J}={\displaystyle \frac{k_f}{{\mu }_f}}{\int }_{\Omega }{\left(\nabla N_p\right)}^T(\nabla N_p)d\Omega$$

(9)

$${\mathit{f}}_{\mathit{u}}={\oint }_{\mathsf{\Gamma }}{N}_u^T{N}_u{f}_{ue}d\mathsf{\Gamma }$$

(10)

$${\mathit{f}}_{\mathit{p}}={\oint }_{\mathsf{\Gamma }}{N}_p^T{N}_p{f}_{pe}d\mathsf{\Gamma }$$

(11)

where ${\mathit{N}}_{\mathit{u}}$ and ${\mathit{N}}_{\mathit{p}}$ are interpolation form functions for solid phase displacement $\mathit{u}$ and pore fluid pressure $\mathit{p}$, respectively; ${\mathit{f}}_{\mathit{u}\mathit{e}}$ and ${\mathit{f}}_{\mathit{p}\mathit{e}}$ are load vectors corresponding to solid phase and pore fluid acting on boundary elements respectively; $\mathit{\rho }$ is the density of solid and liquid media; ${\mathit{E}}_{\mathit{w}}$ is the volume compression modulus of pore fluid; n is the porosity, ${\mathit{k}}_{\mathit{f}}$ is the permeability coefficient; ${\mathit{\mu }}_{\mathit{f}}$ is the viscosity coefficient of pore fluid; and $\mathit{D}$ is the coefficient matrix of the wave equation. For two-dimensional wave problems, its expression is as follows:

$$\mathit{D}=\left[\begin{array}{ccc}\lambda +2G& \lambda & 0\\ \lambda & \lambda +2G& 0\\ 0& 0& G\end{array}\right]$$

(12)

where $\mathit{\lambda }$ and $\mathit{G}$ are the Lame constant and shear modulus of the solid phase, respectively.

The expressions for the gradient operator vector $\mathbf{\nabla }$ and the differential operator matrix $\mathit{L}$ are:

$${\mathbf{\nabla }}^{\mathit{T}}=\left[\begin{array}{cc}{\displaystyle \frac{\partial }{\partial x }}& {\displaystyle \frac{\partial }{\partial y }}\end{array}\right]$$

(13)

$${\mathbf{L}}^{\mathit{T}}=\left[\begin{array}{ccc}{\displaystyle \frac{\partial }{\partial x }}& 0& {\displaystyle \frac{\partial }{\partial y }}\\ 0& {\displaystyle \frac{\partial }{\partial y }}& {\displaystyle \frac{\partial }{\partial x }}\end{array}\right]$$

(14)

The wave Equations (3) and (4) of saturated two-phase media obtained by finite element space discretization are coupled equations with dynamic coupling and coupling between each node. In order to realize the display and calculation of the u-p equation in the time domain, the literature [27] makes two assumptions: (1) Since the inertial coupling between nodes in the finite element calculation system is a high-order small quantity relative to the inertia of the nodes themselves, the inertial coupling between nodes can be disregarded, and the inertia of each node is considered to be independent of one another, with the element mass concentrated on each node; (2) The coupling of the time rate of change of fluid pressure at each node in the finite element calculation system is ignored, and it is assumed that the time rate of change of fluid pressure at each node is independent of each other, and the non-diagonal element of the matrix related to the time rate of change of fluid pressure is zero. The mass matrix $\mathit{M}$ and the fluid compression matrix $\mathit{S}$ in the equations are diagonalized. The inverse operation of the coefficient matrix after diagonalization is only for the principal diagonal elements, which avoids the whole inverse of the matrix in the process of solving and solves one of the main factors restricting the computational efficiency. The diagonalized element mass matrix ${\mathit{M}}^{\mathit{e}}$ and the element fluid compression matrix ${\mathit{S}}^{\mathit{e}}$ are expressed as follows:

$${\left({\mathit{M}}_{\mathit{l}}^{\mathit{e}}\right)}_{ij}=\left\{\begin{array}{ll}a{\left(M^e\right)}_{ii}=a{\int }_{V_e}\rho N_i^T{N}_{i}dV,& i=j\\ 0,& i\ne j\end{array}\right.$$

(15)

$${\left({\mathit{S}}_{\mathit{l}}^{\mathit{e}}\right)}_{ij}=\left\{\begin{array}{ll}b{\left(S^e\right)}_{ii}=B{\int }_{V_e}{\displaystyle \frac{E_w}{n}}{N}_i^T{N}_{i}dV,& i=j\\ 0,& i\ne j\end{array}\right.$$

(16)

The scale coefficient a is determined according to the principle of conservation of mass, that is, the sum of the elements in $M_l^e$ corresponding to the concentrated mass at each node of the same cell should be equal to the mass of the whole cell, that is written as follows:

$${\sum }_{i=1}^{ne}{\left({\mathit{M}}_{\mathit{l}}^{\mathit{e}}\right)}_{ii}=a{\sum }_{i=1}^{ne}{\left(M^e\right)}_{ii}=W{I}_d=\rho V_e{I}_d$$

(17)

The proportionality factor $b$ is determined according to the following relationship:

$${\sum }_{i=1}^{ne}{\left({\mathit{S}}_{\mathit{l}}^{\mathit{e}}\right)}_{ii}=b{\sum }_{i=1}^{ne}{\left(S^e\right)}_{ii}={\displaystyle \frac{E_w}{n}}{V}_e{I}_d$$

(18)

where $V_e$ is the volume of the unit; $I_d$ is the unit vector; and $d$ is the dimension of the unit in geometric space.

2.2. Explicit Time-Domain Solution

The action time t of the dynamic load is discretized according to a fixed step size $\Delta t$, and a series of discrete time points are obtained. The corresponding matrices and vectors are also discretized. By diagonalizing the matrix, we can rewrite Equations (3) and (4) as follows:

$${\mathit{M}}_{\mathit{l}}{\ddot{\mathit{u}}}_{\mathit{k}}+\mathit{K}{\mathit{u}}_{\mathit{k}}-\mathit{Q}{\mathit{p}}_{\mathit{k}}={\mathit{f}}_{\mathit{u}\mathit{k}}$$

(19)

$$\mathit{J}{\mathit{p}}_{\mathit{k}}+{\mathit{S}}_{\mathit{l}}{\dot{\mathit{p}}}_{\mathit{k}}-{\mathit{Q}}^{\mathit{T}}{\dot{\mathit{u}}}_{\mathit{k}}={\mathit{f}}_{\mathit{p}\mathit{k}}$$

(20)

where ${\mathit{M}}_{\mathit{l}}$ and ${\mathit{S}}_{\mathit{l}}$ are the diagonalized global mass matrix and fluid compression matrix, respectively. The angular mark $k$ of each matrix and vector indicates that the number of time steps is $k$ ($k$ = 1,2,3 …).

Based on the central difference method, the solid skeleton displacement of step $k$ is calculated as follows:

$${\ddot{\mathit{u}}}_{\mathit{k}}={\displaystyle \frac{2}{\Delta t^2}}\left({\mathit{u}}_{\mathit{k}+1}-{\mathit{u}}_{\mathit{k}}\right)-{\displaystyle \frac{2}{\Delta t}}{\dot{\mathit{u}}}_{\mathit{k}}$$

(21)

Substitute Equation (21) into Equation (19) to obtain the calculation formula of solid skeleton displacement at step $k+1$, written as follows:

$${\mathit{u}}_{\mathit{k}+1}={\mathit{u}}_{\mathit{k}}+\Delta t{\dot{\mathit{u}}}_{\mathit{k}}+{\displaystyle \frac{\Delta t^2}{2}}{\mathit{M}}_{\mathit{l}}^{-1}({\mathit{f}}_{\mathit{u}\mathit{k}}-\mathit{K}{\mathit{u}}_{\mathit{k}}+\mathit{Q}{\mathit{p}}_{\mathit{k}})$$

(22)

Assuming a linear transformation of fluid pressure $\mathit{p}$ in a time $\Delta t$, the rate of change with respect to time can be approximated by backward difference method as follows:

$${\dot{\mathit{p}}}_{\mathit{k}}={\displaystyle \frac{{\mathit{p}}_{\mathit{k}+1}-{\mathit{p}}_{\mathit{k}}}{\Delta t}}$$

(23)

By substituting Equation (23) into Equation (20), the following calculation formula of pore fluid pressure at step $k+1$ is obtained:

$${\mathit{p}}_{\mathit{k}+1}={\mathit{p}}_{\mathit{k}}+\Delta t{\mathit{S}}_{\mathit{l}}^{-1}({\mathit{f}}_{\mathit{p}\mathit{k}}-\mathit{J}{\mathit{p}}_{\mathit{k}}-{\mathit{Q}}^{\mathit{T}}{\dot{\mathit{u}}}_{\mathit{k}})$$

(24)

The number of time iteration steps of Formula (17) is $k$, which is superimposed with the formula whose number of iteration steps is $k+1$ to obtain the following:

$${\mathit{M}}_{\mathit{l}}({\ddot{\mathit{u}}}_{\mathit{k}}{+\ddot{\mathit{u}}}_{\mathit{k}+1})+\mathit{K}{(\mathit{u}}_{\mathit{k}}+{\mathit{u}}_{\mathit{k}+1})-\mathit{Q}({\mathit{p}}_{\mathit{k}}+{\mathit{p}}_{\mathit{k}+1})=({\mathit{f}}_{\mathit{u}\mathit{k}}+{\mathit{f}}_{\mathit{u}(\mathit{k}+1)})$$

(25)

The Newmark constant average acceleration method adopts the following assumptions:

$${\displaystyle \frac{{\ddot{\mathit{u}}}_{\mathit{k}}{+\ddot{\mathit{u}}}_{\mathit{k}+1}}{2}}={\displaystyle \frac{{\mathit{u}}_{\mathit{k}+1}-{\mathit{u}}_{\mathit{k}}}{\Delta t}}$$

(26)

Substituting Equation (26) into (25) gives the following expression for solid skeleton velocity at step $k+1$:

$${\dot{\mathit{u}}}_{\mathit{k}+1}={\dot{\mathit{u}}}_{\mathit{k}}+{\displaystyle \frac{\Delta t}{2}}{\mathit{M}}_{\mathit{l}}^{-1}\left[\left({\mathit{f}}_{\mathit{u}\mathit{k}}+{\mathit{f}}_{\mathit{u}\left(\mathit{k}+1\right)}\right)-\mathit{K}{(\mathit{u}}_{\mathit{k}}+{\mathit{u}}_{\mathit{k}+1}\right)+\mathit{Q}({\mathit{p}}_{\mathit{k}}+{\mathit{p}}_{\mathit{k}+1})]$$

(27)

Formulas (22), (24) and (27) together form the time-domain explicit step-by-step explicit time-domain solution.

The time domain explicit numerical calculation method has conditional stability. The literature [27] deduced the stability condition of the algorithm as follows:

0 < Ω < 2, 0 < ϕ < 2, (1 + α)Ω² < 2ϕ

(28)

$$\Delta \mathrm{t}<\mathrm{m}\mathrm{i}\mathrm{n}\left\{\sqrt{{\displaystyle \frac{2\mathrm{m}}{\mathrm{k}}}},{\displaystyle \frac{2\mathrm{s}}{\mathrm{j}}},{\displaystyle \frac{2\mathrm{j}\mathrm{m}}{\mathrm{k}\mathrm{s}+{\mathrm{q}}^2}}\right\}$$

(29)

where ${\Omega }^2={\displaystyle \frac{k}{m}}{\Delta t}^2,\mathsf{\varphi }={\displaystyle \frac{j}{s}}\Delta t,\mathsf{\alpha }{\Omega }^2={\displaystyle \frac{q^2}{ms}}{\Delta t}^2(\mathsf{\alpha }={\displaystyle \frac{q^2}{ks}})$.

When material and calculation parameters are determined, Equations (28) and (29) can be used to evaluate the stability of the calculation results.

3. Implementation of the Kernel Function on a GPU

The program presents a process for calculating the near-field wave propagation in a saturated two-phase medium, where the main part of the calculation is performed on the GPU, with only the initial data loading from the CPU and the final iteration solution data exported from the GPU to the CPU, which can greatly save the time for data transmission (Figure 1).

The initial step of the finite element program lies in discretizing the structure, establishing a two-dimensional or three-dimensional model, discretizing the model, choosing the appropriate element size and type, and subsequently generating the original set of data points. For a two-dimensional planar four-node finite element model, the information regarding cells and nodes in its data file comprises five parts, namely element ID, node 1 ID, node 2 ID, node 3 ID, and node 4 ID (Figure 2 left). In this manner, a concrete graph is converted into an abstract data set, which can subsequently be read and processed by a computer to generate a new data matrix.

The challenge in assembling the global matrix resides in the circumstance that a node might be shared by distinct elements, yet the calculation of an element matrix demands a base-unit element, signifying that the identical node could be called multiple times (the frequency of its being called depends on the number of elements it is shared by, as depicted in Figure 2 left). The subsequent explanations of the kernel functions are all elaborated through the FEM in Figure 2. Once the calculation of the element matrix is accomplished, the element matrices need to be assembled into the global matrix, where there is overlapping data, as indicated by the red blocks in Figure 2 right. In serial computation, upon the completion of the calculation of one element matrix, its value is written into the corresponding position in the global matrix, and this procedure is repeated, where the overlapping parts are added. Nevertheless, for parallel computation, an issue will emerge. Parallel computation does not write data in a sequential fashion but will concurrently calculate all the element matrices and write them into the corresponding positions in the global matrix simultaneously. However, this will lead to the same position being written with diverse element matrix data, which gives rise to the problem of either data error or thread deadlock.

This paper uses the atomic operation method to solve the problem of thread conflict. Although the atomic operation is used to solve the problem of thread conflict among threads, the problem is solved in a more subtle way so that the program does not take a lot of time. Only 10⁻² s is required for a testing model with a mesh size of 1 million, and 10⁻¹ s for a testing model with a mesh size of 10 million.

PNAM initially processes the relationship between nodes and elements, generates preprocessed data that is more conducive to processing, and subsequently generates the result of the global matrix using the preprocessed data. When generating the results of global matrices, it is unnecessary to generate the global matrices in the interim but to generate the data directly in the CSR storage format in one step. The advantage of this process lies in that the global matrix in the middle contains numerous invalid zero values, and the useful data occupies less than 1% and decreases as the matrix expands. The CSR storage format is a compressed format that can significantly save storage and GPU memory space. This enables the computation of matrices with tens or even hundreds of millions of degrees of freedom using very limited GPU memory. It is worth noting that both the compressed sparse column (CSC) storage format and the CSR storage format are generated because the CSC storage format is the transposed form of the CSR storage format when employing the CUDA built-in function. This remarkable property saves the time required for matrix transposition.

The preprocessed data are generated from the data file of FEM and obtained step by step through 3 kernel functions, as shown in Figure 3. The first step involves generating eleNloc and ele_count from kernel function 1. In the second step, eleNloc and ele_count are employed as arguments to generate node_count from kernel function 3 (part; refer to the first 10 lines of Algorithm 3). In the third step, acc_ele_count is produced from kernel function 2 with node_count as an argument. In the fourth step, eleNloc, ele_count, and acc_ele_count are utilized as arguments to generate node_list from kernel function 3 (full; refer to the all codes of Algorithm 3). The preprocessed data will be used to acquire the CSR and CSC storage format data for global matrices.

3.1. Resolve Thread Conflicts and Determine Relationships Between Nodes and Elements

The kernel function 1 is utilized to calculate in which elements the node is contained and the location of those elements. The default eleNloc size is employed here. Generally, for a two-dimensional model, it is most prevalent for a node to be encompassed by four elements, so it is hypothesized that a node is contained by up to four elements (of course, real cases may be greater than 4; simply set a larger number). The array eleNloc is used to store in which elements the node is contained and its location within the element. The array ele_count counts the number of elements that each node encompasses. For instance, node 2 is encompassed by element 1 and element 2, and it is 1 and 0, respectively, in their positions (counting from 0), and the corresponding number in ele_count is 2, as shown in Figure 3. Since the size of the eleNloc is presumed in advance, this requires additional memory space. A more superior approach is to calculate the ele_count and acc_ele_count data first, thereby enabling one to precisely know the size of eleNloc and the corresponding data location. This scheme will be adopted in the future.

Kernel function 1 deals with the entire data file and uses atomicAdd() to avoid collisions between threads (see Algorithm 1 for details). Each thread is used to process the data of each row, that is, the node data contained in each element, a total of four. For each node in this row, the element ID and the location of the node in the elements are informed. However, due to thread conflicts, data cannot be directly written to eleNloc, and atomicAdd() is needed to solve thread conflicts. Specifically, for the node, when it is its turn, add one to the corresponding position of ele_count and assign the value to temp_count so that the location of the data to be written is known. For example, for the node 1 in the first row of data in Figure 3, we know that it is in the location of 0 of the element 1. This is done by entering a 1 at element 1 and a 0 next to it. The number entered in ele_count is 1.

Algorithm 1: Use atomicAdd() to resolve thread conflicts and obtain the data of eleNloc and ele_count

1: initial eleNloc to −1 and ele_count to 0  // The −1 is inserted to indicate that these numbers are not available. 2: MAX_ELEMENT←4  // MAX_ELEMENT represents the maximum number of elements shared by a node, which is set to 4. 3: x ← threadIdx.x + blockIdx.x * blockDim.x  // x controls the serial number of the element ID. 4: for i = 0 to 3 do 5:   node ← element(i) // element is the data files of FEM, and i shows 4 nodes. 6:   temp_count ← atomicAdd(&ele_count[node],1)  // This step resolves thread conflicts. 7:   eleNloc[2 * (node-1) * MAX_ELEMENT + 2 * temp_count] ← x + 1  // element numbers start at 1, so +1 8:   eleNloc[2 * (node-1) * MAX_ELEMENT + 2 * temp_count + 1] ← i   // i is the location of node in the element, counting from 0 9: end for

3.2. Obtaining the Data of the Prefix Sum

Kernel function 2 computes the prefix sum for node_count and represents the most time-consuming component of the entire program, aside from data transmission; however, it possesses a high potential for optimization. This article systematically introduces the complete algorithm, with plans for further optimization of kernel function 2 in future work. The process to achieve the result involves two steps (see Figure 4). First, node_count is partitioned into several blocks; subsequently, within each block, the Hillis/Steele scan algorithm is employed to compute the prefix sum of the data contained in that block, resulting in acc_ele_count0 for the whole data. The rationale behind dividing it into multiple thread blocks lies in the limitations on thread block size and the inability to share data across different thread blocks. Additionally, it should be emphasized that acc_ele_count0 initializes at zero while subsequent entries automatically extend back by one position. Utilizing this acc_ele_count0 allows each thread to add every last number from its preceding thread block onto itself, thereby yielding an updated value for acc_ele_count. The method of obtaining acc_ele_count0 is shown in Algorithm 2.

Algorithm 2: From node_count to acc_ele_count0

1: buffer_idx ← threadIdx.x + blockDim.x * blockIdx.x  // buffer_idx controls the serial number of the node. 2: BLOCK_SIZE←32  // BLOCK_SIZE is the size of kernel block and should not larger than 32, otherwise the data will be wrong. 3: __shared__ int cache[BLOCK_SIZE] 4: if buffer_idx< BLOCK_SIZE then // The first data is 0, and the other data is moved one position later. 5:   cache[0] ← 0 6:   cache[threadIdx.x+1] ← node_count[buffer_idx] 7: else 8:   cache[threadIdx.x] ← node_count[buffer_idx-1] 9: end if 10: __syncthreads(); 11: for (int stride = 1; stride < blockDim.x; stride * = 2) { 12:    if (threadIdx.x >= stride) { 13:      cache[threadIdx.x] += cache[threadIdx.x − stride];} 14:     __syncthreads();} 15: acc_ele_count0[buffer_idx] ← cache[threadIdx.x]

3.3. Obtaining the Relationship Between Nodes

Kernel function 3 counts information about each node and its associated nodes (including itself). Algorithm 3 shows how to obtain node_list and node_count from ele_count and acc_ele_count. Figure 5 shows the relationship between the four preprocessed data. node_count[x] represents the number of nodes each node is associated with (including itself). It is always less than 4*ele_count[x] because nodes between two different elements may overlap. acc_ele_count specifies the position of the first data in each row in the node_list.

Algorithm 3: Obtain information about each node and its associated nodes (including itself)

1: initial node_list to 0 2: x←threadIdx.x + blockIdx.x * blockDim.x  // x controls the serial number of the node 3: element_num←4 * ele_count[x]  // The number of nodes associated with a node can never exceed 4 * ele_count[x] 5: temp_node_list[element_num] ← −1  // Initialize the array temp_node_list to any number less than 0 (this is to make it easier to compare value of nodes with values of temp_node_list and sort them). 6: for element in eleNloc(x) do  // Get all the elements of this node from eleNloc. 7:   Get the nodes in element, sort them, and place them in the array temp_node_list. 8: end for 9: The value of node_count(x) is the amount in temp_node_list that is not −1. 10: element_num←node_count(x) 11: for i = 0 to element_num do 12:    node_list[i + acc_ele_count[x]] ← temp_node_list[i]  // Store the data from temp_node_list into node_list 13: end for

3.4. Generate CSR and CSC Storage Format Data for Each Global Matrix

The subsequent step involves generating three sets of data in CSR and CSC storage formats for each global matrix, specifically: CSR_value/CSC_value, CSR_y/CSC_x, and CSR_offset/CSC_offset. In this context, CSR_value corresponds to the spmv_csr function parameter row_ptr, while CSR_y aligns with col_ind, and again, CSR_value pertains to values (see Algorithm 7). This distinction is necessary because both of the CSR data are utilized when invoking the spmv_csr function and the CSC data are required for computing the transpose of the matrix; thus, these terms are designated here for clarity. Given that there are five matrices associated with the u-p formulation equation—each differing in size. The element stiffness matrix and element mass matrix have dimensions of 8 × 8, indicating that a single node generates 2 × 2 data. Conversely, both the element seepage matrix of pore fluid and the element compressibility matrix of pore fluid possess dimensions of 4 × 4, meaning one node produces a singular piece of data. For the element coupling matrix related to dynamic responses between solid skeletons and pore fluids, its size measures at 8 × 4; hence a node yields 2 × 1 data. Since the stiffness matrix represents the most intricate case among these matrices, it will be used as an example to elucidate algorithm implementation.

3.4.1. Generation of CSR_offset Data

Kernel function 4 is used to obtain the data of CSR_offset. CSR_offset data store the row index of the non-zero element. For the element stiffness matrix, one node generates 2 × 2 data (two rows and two columns). Figure 6 shows the generation of CSR_offset for the global stiffness matrix in detail. The first index always starts at 0. The first node is shared by an element, which means that there are four nodes associated with it (including itself), so the first row has eight data points, so the second index starts at eight. There are also eight data points in the second row, so there are 16 data points before the third row. And then all the indexes are acquired in the same way. One thing to note is that the last index is the sum of all the previous elements plus 1 (see Algorithm 4 for details).

Algorithm 4: Generation of CSR_offset data for global stiffness matrix

1: initial CSR_offset to 0 2: x ← threadIdx.x + blockIdx.x * blockDim.x  // x controls the serial number of the node 3: element_num←node_count(x) 4: CSR_offset[2 * x] ← acc_ele_count[x] * 2 * 2 5: CSR_offset[2 * x + 1] ← acc_ele_count[x] * 2 * 2 + 2 * temp_count 6: if x = COORDS_NUM then  // COORDS_NUM is the total number of nodes, and it’s 25 in this case 7:   CSR_offset[2 * x] = acc_ele_count[x] * 2 * 2 + 1  // The last value of CSR_offset adds one 8: end if

3.4.2. Generation of CSR_y Data

Kernel function 5 is used to obtain the data of CSR_y. CSR_y data are used to store column index information for non-zero elements. A node of the stiffness matrix generates 2 rows of data, and the CSR_y value of the two rows is the same. A node generates two columns of data, and the CSR_y value for these two data can be generated from node_list, as shown in Figure 7. The detailed calculation process is shown in Algorithm 5.

Algorithm 5: Generation of CSR_y data for global stiffness matrix

1: initial CSR_offset to 0 2: x ← threadIdx.x + blockIdx.x * blockDim.x  / x controls the serial number of the node 3: element_num←node_count(x) 4: acc_ele_num ← acc_ele_count[x] 5: for i = 0 to temp_count do 6:   CSR_y[2 * i + CSR_offset[2 * x]] ← 2 * node_list[acc_element_num + i] − 2    // The first element in the first row for a node. 7:   CSR_y[2 * i + 1 + CSR_offset[2 * x]] ← 2 * node_list[acc_element_num + i] − 1  // The second element in the first row for a node. 8:   CSR_y[2 * i + CSR_offset[2 * x + 1]] ← 2 * node_list[acc_element_num + i] − 2    // The first element in the second row for a node. 9:   CSR_y[2 * i + 1 + CSR_offset[2 * x + 1]] ← 2 * node_list[acc_element_num + i] − 1  // The second element in the second row for a node. 10: end for

3.4.3. Generation of CSR_value Data

Kernel function 6 is used to obtain the data of CSR_value. CSR_value data are the non-zero data in the matrix (in general). As shown in Figure 8, the generation of CSR_value data for the global stiffness matrix through node 2 is illustrated. eleNloc is used to obtain which elements node 2 is contained in and where it is located. Then node information can be obtained from data files of FEM, and the element stiffness matrix can be generated. The obtained element stiffness matrices can be indexed to the required number of rows by location in eleNloc. For instance, the first element matrix of node 2, the data in the second row is indexed, and then the column in which it is located is extracted and compared with node_list so that the position of these nodes in K_CSR_value can be known. From CSR_offset we know that the index of the first data from node 2 in K_CSR_value (16 in this case). The function is shown in Algorithm 6.

Algorithm 6: Generation of CSR_value data for global stiffness matrix

1: initial CSR_value to 0 2: x ← threadIdx.x + blockIdx.x * blockDim.x  / x controls the serial number of the node 3: element_num←node_count(x) 4: for K = 0 to ele_count[x] do  // K controls element and location 5:   element_N←eleNloc(K,x)  // Gets the number of elment from eleNloc and assigns it to element_N. 6:   element_loc←eleNloc(K,x) // Gets the elment location assignment from eleNloc to element_loc 7:   Get the four nodes of each element, i,j,m,n, from data files of FEM 8:   calculate ke(i,j,m,n) 10: for K = 0 to 4 do  // Control four nodes. 11: if element_loc = K do  // Control the position of nodes in the elements. 12:   for M = 0 to element_num do  // here is to compare the position of node_list with i,j,m,n. 13:   Compare i,j,m,n with node_list(M,x) respectively, and then fill the data into the corresponding location. //The x in node_list(M,x) is to locate the data in node_list, for instance, node_list(node 2) in Figure 8. //for example, //if i = node_list(M,x) do      CSR_value[2 * M + CSR_offset[2 * x]] += ke [2 * K][0];      CSR_value[2 * M + 1 + CSR_offset[2 * x]] += ke [2 * K][1];      CSR_value[2 * M + CSR_offset[2 * x + 1]] += ke [2 * K + 1][0];      CSR_value[2 * M + 1 + CSR_offset[2 * x + 1]] += ke [2 * K + 1][1];      end if 14:   end for 15: end if 16: end for

3.5. Description of Iterative Solution

The algorithm for matrix diagonalization is extremely straightforward and will not be elaborated on in detail. The iterative solution involves a lot of dot product, vector addition, and number product, and the time-consuming one is the dot product. The CUDA spmv_csr kernel function was employed for the dot product between the matrix and the vector (for transposed matrix times vector, CSC data are used instead of CSR data).

Algorithm 7: CUDA built-in spmv_csr function

1: __global__ void spmv_csr(const int * row_ptr, const int * col_ind, const double * values, const int * num_rows, const double * x, double * y) { 2:  // Uses a grid-stride loop to perform dot product 3:    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < num_rows; i += blockDim.x * gridDim.x) { 4:      ACCURACY dotProduct = 0; 5:      const int row_start = row_ptr[i]; 6:      const int row_end = row_ptr[i + 1]; 7:      for (int j = row_start; j < row_end; j++) { 8:        dotProduct += values[j] * x[col_ind[j]]; 9:      } 10:      y[i] = dotProduct; 11:    } 12: }

This paper presents a new algorithm based on the above function. This algorithm computes Equation (20), Equation (22), or Equation (25) in one step. The advantage is that the calculation time is less, and the disadvantage is that it involves a lot of data at once, requiring large GPU memory. This is explained in detail by the calculation Formula (20), as shown in Algorithm 8. The first and second row calculations obtain Kuk and Qpk, which correspond to the data of a single row rather than the overall data, because a thread will only calculate the data of a single row. The third and fourth rows are the U_k+1 result of computing a row at a time, where M_l⁻¹ has already been processed before.

Algorithm 8: Calculate uk+1

1: Kuk←K * Uk;  // Using spmv_csr to calculate K * Uk, Kuk is the value of a row 2: Qpk←Q * pk;  // Using spmv_csr to calculate Q * Pk, Qpk is the value of a row 3: for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < TOTAL_COORDS_NUM * 2; i += blockDim.x * gridDim.x) { // TOTAL_COORDS_NUM is the number of all nodes. 4: uk+1[i]←uk[i] + Δt * $\dot{u}$k[i] + 0.5 * Δt2 * Ml−1[i] * (fuk[i] − Kuk + Qpk);}  // Get the result

The vectors $u_{k+1}$, $p_{k+1}$, and ${\dot{u}}_{k+1}$ obtained from each calculation also require some processing. The final output from the GPU must be transferred to the CPU, a process that requires a certain amount of time and can also be optimized. Combining these three vectors into one vector and then operating on it can save a lot of time.

4. Results and Analysis

4.1. Example Simulation

The program is employed to address the two-dimensional saturated soil dynamic issue under sinusoidal loading. The model and load are depicted in Figure 9. The material parameters are shown in

Table 1. Material parameters.

Parameters	Value	Parameters	Value
Young’s modulus Es/MPa	70	Bulk modulus of pore fluid Kf/Gpa	2.24
Bulk modulus of solid Ks/GPa	1 × 1014	Density of pore fluid ρf/(kg·m3)	1000
Density of solid and pore fluid ρs/(kg·m3)	1700	porosity n	0.4
Poisson’s ratio ν	0.3	Hydraulic conductivity k/(m·s−1)	0.001

. The surface of the model constitutes a free permeable boundary subjected to full sinusoidal load, while the other three sides are impermeable to water. The boundary on the left and right sides limits lateral displacement, and the boundary on the bottom limits horizontal and vertical displacement. Set three points, A(15,3), B(15,15), and C(15,27), as observation points. Observe the displacement in the y direction and the pore pressure response (with the coordinate origin in the lower left corner). The material properties of the soil are presented in Table 1. The initial pore pressure is 0.

Each element is 0.03 m × 0.03 m square, with a total of 1 million elements, 1,002,001 nodes, and 2,004,002 degrees of freedom. The time increment of each step is 10⁻⁴ s, the precision is double, and the thread block size is 16.

Figure 10 shows the vertical displacement and pore pressure of the three observation points. It can be found that the absolute value of displacement decreases with the increase of depth, while pore pressure increases with the increase of depth. The CPU’s calculation results are oscillating since its grid is not fine enough.

Table 2. Calculation time of each kernel function in example simulation.

Operation	Time/s	Description
Kernel function 1	0.019	Determine relationships between nodes and elements
Kernel function 2	0.570	Get the data of the prefix sum
Kernel function 3 (full)	0.018	Gets the relationship between nodes
Kernel function 4	0.001	Generation of CSR_offset data for global stiffness matrix
Kernel function 5	0.029	Generation of CSR_y data for global stiffness matrix
Kernel function 6	0.070	Generation of CSR_value data for global stiffness matrix
displacement	0.009	Calculate Formula (20) once
pressure	0.009	Calculate Formula (22) once
velocity	0.012	Calculate Formula (25) once
Sperated_DtoH	0.026	Transfer three vectors of data from the GPU to the CPU separately (once)
Centralized_DtoH	0.020	Combine three vectors into one vector and transfer it to the CPU (once)
Sperated_storage	5.556	Sperated_DtoH and then write these three vectors to the text file separately (once)
Centralized_storage	4.623	Centralized_DtoH and then write it to text file (once)

presents the computation time in the above example simulation (the average time needed for each operation to compute once). Generally, for data with 2 million degrees of freedom, the calculation time of the kernel function is no more than 0.02 s. The calculation time is extremely short. Nevertheless, the time of data storage occupies the majority of the time. It takes less than 0.02 s to calculate the three vectors $u_{k+1}$, $p_{k+1}$ and ${\dot{u}}_{k+1}$ separately, and the time of data transmission needs to be 0.026 s. However, if the three vectors are combined into one vector and transferred to the CPU, the time taken is only 0.020 s, saving 22% of the time. Storing the three vectors individually takes an average of 5.556 s, but it takes only 4.623 s if the three vectors are packed into one vector (including calling the kernel function to package the three vectors), saving 17% of the time. This step is significant because data transfer takes up the majority of the computing time. Especially for seismic simulation, there are numerous steps required, and the amount of data that needs to be transmitted is also substantial.

4.2. Influence of Accuracy on Calculation Efficiency

The aforementioned model is employed to compute the program using both single precision and double precision, yielding the time it takes for these operations to compute once (refer to Figure 11). For the initial five kernel functions, precision is not of great significance; however, for kernel function 6 and the computation of three vectors, precision has a significant influence on them, and the result of double precision is 1.5 to 2 times that of single precision. Storage is Centralized_storage in Table 2, which takes up most of the time of the entire program. For a model with 1 million elements, its time does not vary significantly with accuracy.

4.3. Influence of Model Size on Computing Efficiency

Four models are utilized: Model 1 with a mesh size of 0.64 million, Model 2 at 1 million, Model 3 at 4 million, and Model 4 at 6.24 million. The calculation accuracy was set to double precision, with a thread block size of 16. As illustrated in Figure 12 (left), among the four models, model size exerts the most significant influence on Kernel function 2, which exhibits exponential growth and necessitates further optimization. Additionally, model size substantially affects data storage; however, this is limited by CPU performance. As illustrated in Figure 12 (right), Kernel function 6 also has a large influence on model size, but its computation time is small, and it is only computed once. In the iterative solution, the computation time of the three vectors escalates marginally with the expansion of the model size. This escalation becomes conspicuous when there are numerous iterative steps in the calculation. DtoH (Centralized_DtoH in Table 2) also increases exponentially with the rise of the model size, and this growth will rise sharply with the increase of iteration steps. However, this is contingent upon computer hardware specifications.

Table 3. The computation time, in seconds, taken to obtain data in three CSR storage formats for six global matrices.

	KF 1	KF 2	KF 3	$\mathit{K}$	$\mathit{M}$	$\mathit{J}$	$\mathit{S}$	$\mathit{Q}$	${\mathit{Q}}^{\mathit{T}}$
Model 1	0.0140	0.2126	0.0059	0.0562	0.0573	0.0103	0.0090	0.0313	0.0299
Model 2	0.0189	0.5702	0.0184	0.0999	0.0984	0.0141	0.0148	0.0462	0.0445
Model 3	0.0850	9.1537	0.0286	0.3893	0.4332	0.1080	0.1219	0.3854	0.4229
Model 4	0.2433	22.4836	0.0430	0.8828	0.9836	0.2562	0.3140	0.8340	0.8929

presents the time needed for generating the three CSR format data of the global matrices under varying model sizes. The degrees of freedom of the four models are 1.28 million, 2 million, 8 million, and 12.48 million, respectively, and the total time consumed for computing the six global matrices is 0.4265 s, 0.9254 s, 12.1280 s, and 26.9334 s, respectively. Kernel function 2 accounts for the majority of the time expended on acquiring global matrix data, accounting for 50%, 62%, 82%, and 83%, respectively. Global matrix $\mathit{J}$ and $\mathit{S}$ consume the least time among the six matrices as they have less data, and their CSR_y and CSR_offset correspond to node_list and acc_ele_count, respectively, eliminating the need for double calculation.

4.4. The Speedup Ratio of a GPU Program to a CPU Program

It should be noted that the speedup ratio is largely dependent on the optimization of the original CPU program; the less optimized the original CPU program is, the higher the speedup ratio.

Table 4. The computing time of the GPU version and CPU version.

Element Size	CPU	GPU
100	3.26 s	5.82 s
400	21.06 s	6.10 s
2500	677.67 s	6.47 s
6400	3380.35 s	9.17 s
100,000	Not enough memory	107.22 s
640,000	Not enough memory	559.81 s
1,000,000	Not enough memory	898.02 s

shows the computation time of CPU programs and GPU programs at different element sizes. The model in Figure 9 is adopted, with steps of 1 × 10⁻³ s and a total time of 30 s, totaling 30,000 steps. The original CPU program is written in Matlab, using eight physical cores. In the assembly of the global matrix, the CPU program generates an empty global matrix and then fills the data of the element matrix into the global matrix in turn (as shown in Figure 2). This step consumes a considerable amount of time and unnecessary memory space. The computing time of the CPU version rises dramatically as the element size grows; the computing time of the GPU version increases at a slower pace compared to that of the CPU version. Although each core of the CPU computes very rapidly compared to the GPU, the advantage of the GPU lies in its numerous cores, making it suitable for handling very large-scale parallel data computations. Additionally, since the CPU version needs to calculate the global matrices in the middle, it quickly reaches the upper limit of memory; the GPU version directly acquires the data in CSR format of the global matrices, omitting most of the zero value, thus enabling it to compute larger models.

5. Conclusions

A comprehensive set of GPU-based algorithms for solving the dynamic response of fluid-saturated porous media is proposed. The algorithms are based on the fully explicit time-domain integration method with second-order precision. They directly acquire the CSR format data of the entire matrix in the matrix assembly part; in the iterative solution part, a novel calculation method is proposed based on CUDA’s built-in kernel function spmv_csr. Compared with the original CPU version, the calculation speed is significantly increased, and the problem that the traditional computing method requires a large amount of CPU memory space is solved. In this paper, an example is employed to demonstrate that the GPU version has higher computational accuracy than the CPU version due to the finer grid.

Given the increased number of matrices involved, PNAM demonstrates a significant advantage by requiring only a single calculation across the three algorithms, KF1, KF2, and KF3, to generate preprocessed data. For the optimization of KF2, the literature [32] is suggested for reference.