A multi‐GPU shallow‐water simulation with transport of contaminants

Abstract

This work presents cost-effective multi-graphics processing unit (GPU) parallel implementations of a finitevolume numerical scheme for solving pollutant transport problems in bidimensional domains. The fluid is modeled by 2D shallow-water equations, whereas the transport of pollutant is modeled by a transport equation. The 2D domain is discretized using a first-order Roe finite-volume scheme. Specifically, this paper presents multi-GPU implementations of both a solution that exploits recomputation on the GPU and an optimized solution that is based on a ghost cell decoupling approach. Our multi-GPU implementations have been optimized using nonblocking communications, overlapping communications and computations and the application of ghost cell expansion to minimize communications. The fastest one reached a speedup of 78 using four GPUs on an InfiniBand network with respect to a parallel execution on a multicore CPU with six cores and two-way hyperthreading per core. Such performance, measured using a realistic problem, enabled the calculation of solutions not only in real time but also in orders of magnitude faster than the simulated time.synchronization between thread blocks). For this purpose, the 2D domain is split into 2D subdomains that include several ghost cells. The ghost cells represent flux contributions that are recomputed in two neighbor thread blocks. In our shallow-water problem, these ghost cells are a row and a column of each 2D subdomain. This way, this memory region (ghost region from now on) is read by two thread blocks, although it is only updated by one. The reason is that the ghost region, together with the rows and columns whose updating is the responsibility of the thread block, provides the information the thread block needs to perform its computations, therefore making the block self-sufficient. Overall, the ghost cell decoupling technique removes the replicated computations for most of the cells, the exception being the ghost cells of each thread block.The algorithm shown in Figure 4 shows the implementation details of the ghost cell decoupling technique. The thread responsible for volume v computes the flow from the neighbor volumes on the right (flR in line 8) and bottom (flD in line 9). Next, a partial flow M OEv is calculated as flROEv C flDOEv (line 10). The same procedure is followed to obtain the partial tOEv (lines 12-14). These partial values, flROEv, flDOEv, dtROEv and dtDOEv, are stored in the shared memory, so that in the second phase (lines 19-22), the thread responsible for volume v only has to add to its partial flow the opposite contribution of its left and up neighbors (line 20; correspondingly, tOEv in line 21). A synchronization barrier is needed between the first and second phases (line 17) because in the second phase each thread reads the partial flows (lines 20-21) stored in the shared memory by another thread in the first phase (lines 8-14). In CUDA, a synchronization barrier (__syncthreads()) stops all warps within a given thread block until all the warps have rea...

Facilitating the development of stencil applications using the Heterogeneous Programming Library

“…Each cell can be dead or alive and following several behavior rules each cell changes its state during the game. Finally, SHWA is a shallow water simulator with pollutant transport presented in Viñas et al This iterative benchmark needs three stages/kernels to compute the evolution of a mesh of finite volumes: in the first one, the flux variation among the elements of the mesh is computed. Then, the global time step is computed for each iteration.…”

“…For this reason, the kernel does not involve any template. Regarding SHWA, although it can benefit from local memory, the implementation using the generic stencil templates we developed would be less efficient than a manual implementation, and thus, the kernel used was written manually. There are two main reasons why the stencil templates are not a good match for the stencil kernel of this problem.…”

“…In our example, this involves manually updating the ghost regions of input in each iteration, as lines 36 to 40 in Figure 4 show. Another issue is the correct computation of the limits of each subarray needed in any multi-device problem, which is performed in function obtain_subarrays_size() (lines [1][2][3][4][5][6][7][8][9][10][11][12][13][14][15][16][17][18][19][20]).…”

“…The management of the memory coherence of these ghost regions among the different devices gives place to an exchange of information from the owner device to the one where the ghost region resides. In general, stencil computations are performed repetitively in a loop and the result of an array in the current iteration depends on the result of this same array in the previous one. Thus, the update of the ghost regions has to be done each iteration, before starting the next stencil computation.…”

Developing adaptive multi-device applications with the Heterogeneous Programming Library

“…These structures need to be refreshed with the most recent version of the data they replicate after each update and before the next round of computations begins. When accelerators are used, this gives place to a data exchange between them [25]. Another example are pipelines, in which data proceed through a series of tasks that transform them and handle the results to the next task in the sequence.…”