Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merge All Three PPM Kernels Into One #392

Draft
wants to merge 29 commits into
base: dev
Choose a base branch
from
Draft

Merge All Three PPM Kernels Into One #392

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
29 commits
Select commit Hold shift + click to select a range
873e580
Add a __device__ function version of PCM
bcaddy Feb 20, 2024
423ca52
Fuse PCM reconstruction and HLLC solver kernels
bcaddy Feb 22, 2024
a8e3146
Fuse PCM reconstruction and Roe solver kernels
bcaddy Feb 26, 2024
627ddad
Remove standalone PCM kernels
bcaddy Feb 26, 2024
ddd4560
Refactor Monotonize_Characteristic_Return_Primitive
bcaddy Mar 4, 2024
eee20a5
Move sound_speed into EigenVecs
bcaddy Mar 4, 2024
068d537
Convert PLMC into a device function
bcaddy Mar 4, 2024
8735833
Add documentation, standard type of direction
bcaddy Mar 5, 2024
cdfd4c0
Fix a few errors from bad merges
bcaddy Apr 17, 2024
8827644
Move PLMC characteristic tracing into its own function
bcaddy Mar 5, 2024
7ff04e1
Add missing namespace qualification
bcaddy Mar 7, 2024
a9d5e4a
Fix bugs in Riemann_Thread_Guard
bcaddy Apr 18, 2024
8fa52eb
Add doxygen comments to refactored PLMC code
bcaddy Apr 18, 2024
c08007c
Move minimum value enforcement in PLMC into the proper function
bcaddy Apr 18, 2024
7e2df9a
Fix typo
bcaddy Apr 18, 2024
67fc328
Add switch to PLMC to turn it into PLMP
bcaddy Apr 18, 2024
284203f
Merge PLMC & PLMP, MHD support in PLMP
bcaddy Apr 18, 2024
0f3f1a0
Update PLMC tests to work with the new PLM kernel
bcaddy Apr 18, 2024
8749c1d
Move the eigenvector computation so it's not in PLMP
bcaddy Apr 18, 2024
170f8a8
Rename PLMC_Characteristic_Tracing to PLM_...
bcaddy Apr 24, 2024
24b8082
Fix a bad rebase
bcaddy Apr 29, 2024
f79db63
Remove standalone PCM kernels
bcaddy Feb 26, 2024
502f30a
Refactor Monotonize_Characteristic_Return_Primitive
bcaddy Mar 4, 2024
35ce512
Move sound_speed into EigenVecs
bcaddy Mar 4, 2024
761098a
Merge PPMC_CTU and PPMC_VL
bcaddy Apr 24, 2024
3e851d1
Remove now unused PPMC related device functions
bcaddy Apr 26, 2024
d555bab
Merge PPMP and PPMC, add MHD support to PPMP
bcaddy Apr 26, 2024
a028b12
Fix bad rebase
bcaddy Apr 29, 2024
f07c6b7
Restore fuction that got lost in rebase
bcaddy Apr 29, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 4 additions & 10 deletions src/grid/grid3D.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,18 +53,12 @@ Grid3D::Grid3D(void)
#ifdef PCM
H.n_ghost = 2;
#endif // PCM
#ifdef PLMP
#if defined(PLMP) or defined(PLMC)
H.n_ghost = 3;
#endif // PLMP
#ifdef PLMC
H.n_ghost = 3;
#endif // PLMC
#ifdef PPMP
H.n_ghost = 4;
#endif // PPMP
#ifdef PPMC
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
H.n_ghost = 4;
#endif // PPMC
#endif // PPMP or PLMC

#ifdef GRAVITY
H.n_ghost_potential_offset = H.n_ghost - N_GHOST_POTENTIAL;
Expand Down
25 changes: 7 additions & 18 deletions src/integrators/VL_1D_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,10 +12,8 @@
#include "../hydro/hydro_cuda.h"
#include "../integrators/VL_1D_cuda.h"
#include "../io/io.h"
#include "../reconstruction/plmc_cuda.h"
#include "../reconstruction/plmp_cuda.h"
#include "../reconstruction/ppmc_cuda.h"
#include "../reconstruction/ppmp_cuda.h"
#include "../reconstruction/plm_cuda.h"
#include "../reconstruction/ppm_cuda.h"
#include "../reconstruction/reconstruction.h"
#include "../riemann_solvers/exact_cuda.h"
#include "../riemann_solvers/hllc_cuda.h"
Expand Down Expand Up @@ -83,20 +81,11 @@ void VL_Algorithm_1D_CUDA(Real *d_conserved, int nx, int x_off, int n_ghost, Rea
GPU_Error_Check();

// Step 4: Construct left and right interface values using updated conserved variables
#ifdef PLMC
hipLaunchKernelGGL(PLMC_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama,
n_fields);
#endif
#ifdef PLMP
hipLaunchKernelGGL(PLMP_cuda, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt,
gama, 0, n_fields);
#endif
#ifdef PPMP
hipLaunchKernelGGL(PPMP_cuda, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt,
gama, 0, n_fields);
#endif
#ifdef PPMC
hipLaunchKernelGGL(PPMC_VL<0>, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, gama);
#if defined(PLMP) or defined(PLMC)
hipLaunchKernelGGL(PLM_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
hipLaunchKernelGGL(PPM_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
#endif
GPU_Error_Check();

Expand Down
38 changes: 13 additions & 25 deletions src/integrators/VL_2D_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,10 +10,8 @@
#include "../global/global_cuda.h"
#include "../hydro/hydro_cuda.h"
#include "../integrators/VL_2D_cuda.h"
#include "../reconstruction/plmc_cuda.h"
#include "../reconstruction/plmp_cuda.h"
#include "../reconstruction/ppmc_cuda.h"
#include "../reconstruction/ppmp_cuda.h"
#include "../reconstruction/plm_cuda.h"
#include "../reconstruction/ppm_cuda.h"
#include "../reconstruction/reconstruction.h"
#include "../riemann_solvers/exact_cuda.h"
#include "../riemann_solvers/hllc_cuda.h"
Expand Down Expand Up @@ -92,27 +90,17 @@ void VL_Algorithm_2D_CUDA(Real *d_conserved, int nx, int ny, int x_off, int y_of

// Step 4: Construct left and right interface values using updated conserved
// variables
#ifdef PLMP
hipLaunchKernelGGL(PLMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx,
dt, gama, 0, n_fields);
hipLaunchKernelGGL(PLMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy,
dt, gama, 1, n_fields);
#endif
#ifdef PLMC
hipLaunchKernelGGL(PLMC_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt,
gama, n_fields);
hipLaunchKernelGGL(PLMC_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt,
gama, n_fields);
#endif
#ifdef PPMP
hipLaunchKernelGGL(PPMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx,
dt, gama, 0, n_fields);
hipLaunchKernelGGL(PPMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy,
dt, gama, 1, n_fields);
#endif // PPMP
#ifdef PPMC
hipLaunchKernelGGL(PPMC_VL<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, gama);
hipLaunchKernelGGL(PPMC_VL<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, gama);
#if defined(PLMP) or defined(PLMC)
hipLaunchKernelGGL(PLM_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt,
gama);
hipLaunchKernelGGL(PLM_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt,
gama);
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
hipLaunchKernelGGL(PPM_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt,
gama);
hipLaunchKernelGGL(PPM_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt,
gama);
#endif // PPMC
GPU_Error_Check();

Expand Down
62 changes: 21 additions & 41 deletions src/integrators/VL_3D_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,10 +17,8 @@
#include "../io/io.h"
#include "../mhd/ct_electric_fields.h"
#include "../mhd/magnetic_update.h"
#include "../reconstruction/plmc_cuda.h"
#include "../reconstruction/plmp_cuda.h"
#include "../reconstruction/ppmc_cuda.h"
#include "../reconstruction/ppmp_cuda.h"
#include "../reconstruction/plm_cuda.h"
#include "../reconstruction/ppm_cuda.h"
#include "../riemann_solvers/exact_cuda.h"
#include "../riemann_solvers/hll_cuda.h"
#include "../riemann_solvers/hllc_cuda.h"
Expand Down Expand Up @@ -158,10 +156,10 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int
hipLaunchKernelGGL(HIP_KERNEL_NAME(Calculate_Roe_Fluxes_CUDA<reconstruction::Kind::pcm, 0>),
roe_pcm_launch_params.get_numBlocks(), roe_pcm_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved, Q_Lx, Q_Rx, F_x, nx, ny, nz, n_cells, gama, n_fields);
hipLaunchKernelGGL(HIP_KERNEL_NAME(Calculate_Roe_Fluxes_CUDA<reconstruction::Kind::pcm, 1>),
hipLaunchKernelGGL(HIP_KERNEL_NAME(Calculate_Roe_Fluxes_CUDA<reconstruction::Kind::pcm, 0>),
roe_pcm_launch_params.get_numBlocks(), roe_pcm_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved, Q_Ly, Q_Ry, F_y, nx, ny, nz, n_cells, gama, n_fields);
hipLaunchKernelGGL(HIP_KERNEL_NAME(Calculate_Roe_Fluxes_CUDA<reconstruction::Kind::pcm, 2>),
hipLaunchKernelGGL(HIP_KERNEL_NAME(Calculate_Roe_Fluxes_CUDA<reconstruction::Kind::pcm, 0>),
roe_pcm_launch_params.get_numBlocks(), roe_pcm_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved, Q_Lz, Q_Rz, F_z, nx, ny, nz, n_cells, gama, n_fields);
#endif // ROE
Expand Down Expand Up @@ -237,41 +235,23 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int
#endif // MHD

// Step 4: Construct left and right interface values using updated conserved variables
#ifdef PLMP
cuda_utilities::AutomaticLaunchParams static const plmp_launch_params(PLMP_cuda, n_cells);
hipLaunchKernelGGL(PLMP_cuda, plmp_launch_params.get_numBlocks(), plmp_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt, gama, 0, n_fields);
hipLaunchKernelGGL(PLMP_cuda, plmp_launch_params.get_numBlocks(), plmp_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy, dt, gama, 1, n_fields);
hipLaunchKernelGGL(PLMP_cuda, plmp_launch_params.get_numBlocks(), plmp_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Lz, Q_Rz, nx, ny, nz, n_ghost, dz, dt, gama, 2, n_fields);
#endif // PLMP
#ifdef PLMC
cuda_utilities::AutomaticLaunchParams static const plmc_vl_launch_params(PLMC_cuda<0>, n_cells);
hipLaunchKernelGGL(PLMC_cuda<0>, plmc_vl_launch_params.get_numBlocks(), plmc_vl_launch_params.get_threadsPerBlock(),
0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama, n_fields);
hipLaunchKernelGGL(PLMC_cuda<1>, plmc_vl_launch_params.get_numBlocks(), plmc_vl_launch_params.get_threadsPerBlock(),
0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama, n_fields);
hipLaunchKernelGGL(PLMC_cuda<2>, plmc_vl_launch_params.get_numBlocks(), plmc_vl_launch_params.get_threadsPerBlock(),
0, 0, dev_conserved_half, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama, n_fields);
#endif // PLMC
#ifdef PPMP
cuda_utilities::AutomaticLaunchParams static const ppmp_launch_params(PPMP_cuda, n_cells);
hipLaunchKernelGGL(PPMP_cuda, ppmp_launch_params.get_numBlocks(), ppmp_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt, gama, 0, n_fields);
hipLaunchKernelGGL(PPMP_cuda, ppmp_launch_params.get_numBlocks(), ppmp_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy, dt, gama, 1, n_fields);
hipLaunchKernelGGL(PPMP_cuda, ppmp_launch_params.get_numBlocks(), ppmp_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Lz, Q_Rz, nx, ny, nz, n_ghost, dz, dt, gama, 2, n_fields);
#endif // PPMP
#ifdef PPMC
cuda_utilities::AutomaticLaunchParams static const ppmc_vl_launch_params(PPMC_VL<0>, n_cells);
hipLaunchKernelGGL(PPMC_VL<0>, ppmc_vl_launch_params.get_numBlocks(), ppmc_vl_launch_params.get_threadsPerBlock(), 0,
0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, gama);
hipLaunchKernelGGL(PPMC_VL<1>, ppmc_vl_launch_params.get_numBlocks(), ppmc_vl_launch_params.get_threadsPerBlock(), 0,
0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, gama);
hipLaunchKernelGGL(PPMC_VL<2>, ppmc_vl_launch_params.get_numBlocks(), ppmc_vl_launch_params.get_threadsPerBlock(), 0,
0, dev_conserved_half, Q_Lz, Q_Rz, nx, ny, nz, gama);
#if defined(PLMP) or defined(PLMC)
cuda_utilities::AutomaticLaunchParams static const plm_vl_launch_params(PLM_cuda<0>, n_cells);
hipLaunchKernelGGL(PLM_cuda<0>, plm_vl_launch_params.get_numBlocks(), plm_vl_launch_params.get_threadsPerBlock(), 0,
0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PLM_cuda<1>, plm_vl_launch_params.get_numBlocks(), plm_vl_launch_params.get_threadsPerBlock(), 0,
0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
hipLaunchKernelGGL(PLM_cuda<2>, plm_vl_launch_params.get_numBlocks(), plm_vl_launch_params.get_threadsPerBlock(), 0,
0, dev_conserved_half, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama);
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
cuda_utilities::AutomaticLaunchParams static const ppm_launch_params(PPM_cuda<0>, n_cells);
hipLaunchKernelGGL(PPM_cuda<0>, ppm_launch_params.get_numBlocks(), ppm_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PPM_cuda<1>, ppm_launch_params.get_numBlocks(), ppm_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
hipLaunchKernelGGL(PPM_cuda<2>, ppm_launch_params.get_numBlocks(), ppm_launch_params.get_threadsPerBlock(), 0, 0,
dev_conserved_half, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama);
#endif // PPMC
GPU_Error_Check();

Expand Down
27 changes: 7 additions & 20 deletions src/integrators/simple_1D_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,10 +10,8 @@
#include "../hydro/hydro_cuda.h"
#include "../integrators/simple_1D_cuda.h"
#include "../io/io.h"
#include "../reconstruction/plmc_cuda.h"
#include "../reconstruction/plmp_cuda.h"
#include "../reconstruction/ppmc_cuda.h"
#include "../reconstruction/ppmp_cuda.h"
#include "../reconstruction/plm_cuda.h"
#include "../reconstruction/ppm_cuda.h"
#include "../reconstruction/reconstruction.h"
#include "../riemann_solvers/exact_cuda.h"
#include "../riemann_solvers/hllc_cuda.h"
Expand Down Expand Up @@ -53,23 +51,12 @@ void Simple_Algorithm_1D_CUDA(Real *d_conserved, int nx, int x_off, int n_ghost,
}

// Step 1: Do the reconstruction
#ifdef PLMP
hipLaunchKernelGGL(PLMP_cuda, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt, gama,
0, n_fields);
#if defined(PLMP) or defined(PLMC)
hipLaunchKernelGGL(PLM_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
GPU_Error_Check();
#endif
#ifdef PLMC
hipLaunchKernelGGL(PLMC_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama,
n_fields);
GPU_Error_Check();
#endif
#ifdef PPMP
hipLaunchKernelGGL(PPMP_cuda, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt, gama,
0, n_fields);
GPU_Error_Check();
#endif
#ifdef PPMC
hipLaunchKernelGGL(PPMC_CTU<0>, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
hipLaunchKernelGGL(PPM_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
GPU_Error_Check();
#endif

Expand Down
34 changes: 9 additions & 25 deletions src/integrators/simple_2D_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,10 +8,8 @@
#include "../global/global_cuda.h"
#include "../hydro/hydro_cuda.h"
#include "../integrators/simple_2D_cuda.h"
#include "../reconstruction/plmc_cuda.h"
#include "../reconstruction/plmp_cuda.h"
#include "../reconstruction/ppmc_cuda.h"
#include "../reconstruction/ppmp_cuda.h"
#include "../reconstruction/plm_cuda.h"
#include "../reconstruction/ppm_cuda.h"
#include "../reconstruction/reconstruction.h"
#include "../riemann_solvers/exact_cuda.h"
#include "../riemann_solvers/hllc_cuda.h"
Expand Down Expand Up @@ -54,27 +52,13 @@ void Simple_Algorithm_2D_CUDA(Real *d_conserved, int nx, int ny, int x_off, int
}

// Step 1: Do the reconstruction
#ifdef PLMP
hipLaunchKernelGGL(PLMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt,
gama, 0, n_fields);
hipLaunchKernelGGL(PLMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy, dt,
gama, 1, n_fields);
#endif
#ifdef PLMC
hipLaunchKernelGGL(PLMC_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama,
n_fields);
hipLaunchKernelGGL(PLMC_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama,
n_fields);
#endif
#ifdef PPMP
hipLaunchKernelGGL(PPMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt,
gama, 0, n_fields);
hipLaunchKernelGGL(PPMP_cuda, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy, dt,
gama, 1, n_fields);
#endif
#ifdef PPMC
hipLaunchKernelGGL(PPMC_CTU<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PPMC_CTU<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
#if defined(PLMP) or defined(PLMC)
hipLaunchKernelGGL(PLM_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PLM_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
hipLaunchKernelGGL(PPM_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PPM_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
#endif
GPU_Error_Check();

Expand Down
43 changes: 11 additions & 32 deletions src/integrators/simple_3D_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,10 +12,8 @@
#include "../hydro/hydro_cuda.h"
#include "../integrators/simple_3D_cuda.h"
#include "../io/io.h"
#include "../reconstruction/plmc_cuda.h"
#include "../reconstruction/plmp_cuda.h"
#include "../reconstruction/ppmc_cuda.h"
#include "../reconstruction/ppmp_cuda.h"
#include "../reconstruction/plm_cuda.h"
#include "../reconstruction/ppm_cuda.h"
#include "../reconstruction/reconstruction.h"
#include "../riemann_solvers/exact_cuda.h"
#include "../riemann_solvers/hll_cuda.h"
Expand Down Expand Up @@ -85,34 +83,15 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx,
#endif

// Step 1: Construct left and right interface values using updated conserved variables
#ifdef PLMP
hipLaunchKernelGGL(PLMP_cuda, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt,
gama, 0, n_fields);
hipLaunchKernelGGL(PLMP_cuda, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy, dt,
gama, 1, n_fields);
hipLaunchKernelGGL(PLMP_cuda, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, n_ghost, dz, dt,
gama, 2, n_fields);
#endif // PLMP
#ifdef PLMC
hipLaunchKernelGGL(PLMC_cuda<0>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama,
n_fields);
hipLaunchKernelGGL(PLMC_cuda<1>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama,
n_fields);
hipLaunchKernelGGL(PLMC_cuda<2>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama,
n_fields);
#endif
#ifdef PPMP
hipLaunchKernelGGL(PPMP_cuda, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, n_ghost, dx, dt,
gama, 0, n_fields);
hipLaunchKernelGGL(PPMP_cuda, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, n_ghost, dy, dt,
gama, 1, n_fields);
hipLaunchKernelGGL(PPMP_cuda, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, n_ghost, dz, dt,
gama, 2, n_fields);
#endif // PPMP
#ifdef PPMC
hipLaunchKernelGGL(PPMC_CTU<0>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PPMC_CTU<1>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
hipLaunchKernelGGL(PPMC_CTU<2>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama);
#if defined(PLMP) or defined(PLMC)
hipLaunchKernelGGL(PLM_cuda<0>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PLM_cuda<1>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
hipLaunchKernelGGL(PLM_cuda<2>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama);
#endif // PLMP or PLMC
#if defined(PPMP) or defined(PPMC)
hipLaunchKernelGGL(PPM_cuda<0>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
hipLaunchKernelGGL(PPM_cuda<1>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
hipLaunchKernelGGL(PPM_cuda<2>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama);
GPU_Error_Check();
#endif // PPMC

Expand Down
2 changes: 1 addition & 1 deletion src/reconstruction/pcm_cuda.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ namespace reconstruction
* \param[in] gamma The adiabatic index
* \return reconstruction::InterfaceState The interface state at xid, yid, zid
*/
template <size_t direction>
template <uint direction>
reconstruction::InterfaceState __device__ __host__ inline PCM_Reconstruction(Real const *dev_conserved,
size_t const xid, size_t const yid,
size_t const zid, size_t const nx,
Expand Down
Loading
Loading