Merge PPMP and PPMC, add MHD support to PPMP

cholla-hydro · Apr 29, 2024 · 0654285 · 0654285
1 parent 3b727fb
commit 0654285
Show file tree

Hide file tree

Showing 15 changed files with 264 additions and 979 deletions.
diff --git a/src/integrators/VL_1D_cuda.cu b/src/integrators/VL_1D_cuda.cu
@@ -13,8 +13,7 @@
   #include "../integrators/VL_1D_cuda.h"
   #include "../io/io.h"
   #include "../reconstruction/plm_cuda.h"
-  #include "../reconstruction/ppmc_cuda.h"
-  #include "../reconstruction/ppmp_cuda.h"
+  #include "../reconstruction/ppm_cuda.h"
   #include "../reconstruction/reconstruction.h"
   #include "../riemann_solvers/exact_cuda.h"
   #include "../riemann_solvers/hllc_cuda.h"
@@ -85,12 +84,8 @@ void VL_Algorithm_1D_CUDA(Real *d_conserved, int nx, int x_off, int n_ghost, Rea
   #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
-  #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_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
+  #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();
 

diff --git a/src/integrators/VL_2D_cuda.cu b/src/integrators/VL_2D_cuda.cu
@@ -11,8 +11,7 @@
   #include "../hydro/hydro_cuda.h"
   #include "../integrators/VL_2D_cuda.h"
   #include "../reconstruction/plm_cuda.h"
-  #include "../reconstruction/ppmc_cuda.h"
-  #include "../reconstruction/ppmp_cuda.h"
+  #include "../reconstruction/ppm_cuda.h"
   #include "../reconstruction/reconstruction.h"
   #include "../riemann_solvers/exact_cuda.h"
   #include "../riemann_solvers/hllc_cuda.h"
@@ -97,16 +96,10 @@ void VL_Algorithm_2D_CUDA(Real *d_conserved, int nx, int ny, int x_off, int y_of
   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
-  #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_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Lx, Q_Rx, nx, ny, nz, dx, dt,
+  #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(PPMC_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved_half, Q_Ly, Q_Ry, nx, ny, nz, dy, dt,
+  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();

diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu
@@ -18,8 +18,7 @@
   #include "../mhd/ct_electric_fields.h"
   #include "../mhd/magnetic_update.h"
   #include "../reconstruction/plm_cuda.h"
-  #include "../reconstruction/ppmc_cuda.h"
-  #include "../reconstruction/ppmp_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"
@@ -245,23 +244,14 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int
   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
-  #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_cuda<0>, n_cells);
-  hipLaunchKernelGGL(PPMC_cuda<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, dx, dt, gama);
-  hipLaunchKernelGGL(PPMC_cuda<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, dy, dt, gama);
-  hipLaunchKernelGGL(PPMC_cuda<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, dz, dt, gama);
+  #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();
 

diff --git a/src/integrators/simple_1D_cuda.cu b/src/integrators/simple_1D_cuda.cu
@@ -11,8 +11,7 @@
 #include "../integrators/simple_1D_cuda.h"
 #include "../io/io.h"
 #include "../reconstruction/plm_cuda.h"
-#include "../reconstruction/ppmc_cuda.h"
-#include "../reconstruction/ppmp_cuda.h"
+#include "../reconstruction/ppm_cuda.h"
 #include "../reconstruction/reconstruction.h"
 #include "../riemann_solvers/exact_cuda.h"
 #include "../riemann_solvers/hllc_cuda.h"
@@ -56,13 +55,8 @@ void Simple_Algorithm_1D_CUDA(Real *d_conserved, int nx, int x_off, int n_ghost,
   hipLaunchKernelGGL(PLM_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
   GPU_Error_Check();
 #endif  // PLMP or PLMC
-#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_cuda<0>, dimGrid, dimBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
+#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
 

diff --git a/src/integrators/simple_2D_cuda.cu b/src/integrators/simple_2D_cuda.cu
@@ -9,8 +9,7 @@
 #include "../hydro/hydro_cuda.h"
 #include "../integrators/simple_2D_cuda.h"
 #include "../reconstruction/plm_cuda.h"
-#include "../reconstruction/ppmc_cuda.h"
-#include "../reconstruction/ppmp_cuda.h"
+#include "../reconstruction/ppm_cuda.h"
 #include "../reconstruction/reconstruction.h"
 #include "../riemann_solvers/exact_cuda.h"
 #include "../riemann_solvers/hllc_cuda.h"
@@ -57,15 +56,9 @@ void Simple_Algorithm_2D_CUDA(Real *d_conserved, int nx, int ny, int x_off, int
   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
-#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_cuda<0>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
-  hipLaunchKernelGGL(PPMC_cuda<1>, dim2dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
+#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();
 

diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu
@@ -13,8 +13,7 @@
   #include "../integrators/simple_3D_cuda.h"
   #include "../io/io.h"
   #include "../reconstruction/plm_cuda.h"
-  #include "../reconstruction/ppmc_cuda.h"
-  #include "../reconstruction/ppmp_cuda.h"
+  #include "../reconstruction/ppm_cuda.h"
   #include "../reconstruction/reconstruction.h"
   #include "../riemann_solvers/exact_cuda.h"
   #include "../riemann_solvers/hll_cuda.h"
@@ -89,18 +88,10 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx,
   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
-  #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_cuda<0>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lx, Q_Rx, nx, ny, nz, dx, dt, gama);
-  hipLaunchKernelGGL(PPMC_cuda<1>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Ly, Q_Ry, nx, ny, nz, dy, dt, gama);
-  hipLaunchKernelGGL(PPMC_cuda<2>, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, Q_Lz, Q_Rz, nx, ny, nz, dz, dt, gama);
+  #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
 

diff --git a/src/reconstruction/plm_cuda.h b/src/reconstruction/plm_cuda.h
@@ -2,8 +2,8 @@
  *  \brief Declarations of the cuda plm kernels
  */
 
-#ifndef PLMC_CUDA_H
-#define PLMC_CUDA_H
+#ifndef PLM_CUDA_H
+#define PLM_CUDA_H
 
 #include "../global/global.h"
 #include "../grid/grid_enum.h"
@@ -314,4 +314,4 @@ auto __device__ __inline__ PLM_Reconstruction(Real *dev_conserved, int const xid
 }
 }  // namespace reconstruction
 
-#endif  // PLMC_CUDA_H
+#endif  // PLM_CUDA_H
diff --git a/src/reconstruction/ppmc_cuda.cu → src/reconstruction/ppm_cuda.cu b/src/reconstruction/ppmc_cuda.cu → src/reconstruction/ppm_cuda.cu
@@ -1,12 +1,12 @@
-/*! \file ppmc_cuda.cu
+/*! \file ppm_cuda.cu
  *  \brief Functions definitions for the ppm kernels, using characteristic
  tracing. Written following Stone et al. 2008. */
 
 #include <math.h>
 
 #include "../global/global.h"
 #include "../global/global_cuda.h"
-#include "../reconstruction/ppmc_cuda.h"
+#include "../reconstruction/ppm_cuda.h"
 #include "../reconstruction/reconstruction_internals.h"
 #include "../utils/gpu.hpp"
 #include "../utils/hydro_utilities.h"
@@ -17,8 +17,8 @@
 
 // =====================================================================================================================
 template <int dir>
-__global__ __launch_bounds__(TPB) void PPMC_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx,
-                                                 int ny, int nz, Real dx, Real dt, Real gamma)
+__global__ __launch_bounds__(TPB) void PPM_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx,
+                                                int ny, int nz, Real dx, Real dt, Real gamma)
 {
   // get a thread ID
   int const thread_id = threadIdx.x + blockIdx.x * blockDim.x;
@@ -74,6 +74,7 @@ __global__ __launch_bounds__(TPB) void PPMC_cuda(Real *dev_conserved, Real *dev_
   hydro_utilities::Primitive const cell_ip2 = hydro_utilities::Load_Cell_Primitive<dir>(
       dev_conserved, xid + 2 * int(dir == 0), yid + 2 * int(dir == 1), zid + 2 * int(dir == 2), nx, ny, n_cells, gamma);
 
+#ifdef PPMC
   // Compute the eigenvectors
   reconstruction::EigenVecs const eigenvectors = reconstruction::Compute_Eigenvectors(cell_i, gamma);
 
@@ -98,32 +99,9 @@ __global__ __launch_bounds__(TPB) void PPMC_cuda(Real *dev_conserved, Real *dev_
       reconstruction::Primitive_To_Characteristic(cell_i, cell_ip2, eigenvectors, gamma);
 
   // Compute the interface states for each field
-  reconstruction::Characteristic interface_R_imh_characteristic, interface_L_iph_characteristic;
-
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a0, cell_im1_characteristic.a0, cell_i_characteristic.a0,
-                                      cell_ip1_characteristic.a0, cell_ip2_characteristic.a0,
-                                      interface_L_iph_characteristic.a0, interface_R_imh_characteristic.a0);
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a1, cell_im1_characteristic.a1, cell_i_characteristic.a1,
-                                      cell_ip1_characteristic.a1, cell_ip2_characteristic.a1,
-                                      interface_L_iph_characteristic.a1, interface_R_imh_characteristic.a1);
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a2, cell_im1_characteristic.a2, cell_i_characteristic.a2,
-                                      cell_ip1_characteristic.a2, cell_ip2_characteristic.a2,
-                                      interface_L_iph_characteristic.a2, interface_R_imh_characteristic.a2);
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a3, cell_im1_characteristic.a3, cell_i_characteristic.a3,
-                                      cell_ip1_characteristic.a3, cell_ip2_characteristic.a3,
-                                      interface_L_iph_characteristic.a3, interface_R_imh_characteristic.a3);
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a4, cell_im1_characteristic.a4, cell_i_characteristic.a4,
-                                      cell_ip1_characteristic.a4, cell_ip2_characteristic.a4,
-                                      interface_L_iph_characteristic.a4, interface_R_imh_characteristic.a4);
-
-#ifdef MHD
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a5, cell_im1_characteristic.a5, cell_i_characteristic.a5,
-                                      cell_ip1_characteristic.a5, cell_ip2_characteristic.a5,
-                                      interface_L_iph_characteristic.a5, interface_R_imh_characteristic.a5);
-  reconstruction::PPM_Single_Variable(cell_im2_characteristic.a6, cell_im1_characteristic.a6, cell_i_characteristic.a6,
-                                      cell_ip1_characteristic.a6, cell_ip2_characteristic.a6,
-                                      interface_L_iph_characteristic.a6, interface_R_imh_characteristic.a6);
-#endif  // MHD
+  auto const [interface_L_iph_characteristic, interface_R_imh_characteristic] =
+      reconstruction::PPM_Interfaces(cell_im2_characteristic, cell_im1_characteristic, cell_i_characteristic,
+                                     cell_ip1_characteristic, cell_ip2_characteristic);
 
   // Convert back to primitive variables
   hydro_utilities::Primitive interface_L_iph =
@@ -132,20 +110,24 @@ __global__ __launch_bounds__(TPB) void PPMC_cuda(Real *dev_conserved, Real *dev_
       reconstruction::Characteristic_To_Primitive(cell_i, interface_R_imh_characteristic, eigenvectors, gamma);
 
   // Compute the interfaces for the variables that don't have characteristics
-#ifdef DE
+  #ifdef DE
   reconstruction::PPM_Single_Variable(cell_im2.gas_energy_specific, cell_im1.gas_energy_specific,
                                       cell_i.gas_energy_specific, cell_ip1.gas_energy_specific,
                                       cell_ip2.gas_energy_specific, interface_L_iph.gas_energy_specific,
                                       interface_R_imh.gas_energy_specific);
-#endif  // DE
-#ifdef SCALAR
+  #endif  // DE
+  #ifdef SCALAR
   for (int i = 0; i < NSCALARS; i++) {
     reconstruction::PPM_Single_Variable(cell_im2.scalar_specific[i], cell_im1.scalar_specific[i],
                                         cell_i.scalar_specific[i], cell_ip1.scalar_specific[i],
                                         cell_ip2.scalar_specific[i], interface_L_iph.scalar_specific[i],
                                         interface_R_imh.scalar_specific[i]);
   }
-#endif  // SCALAR
+  #endif  // SCALAR
+#else     // PPMC
+  auto [interface_L_iph, interface_R_imh] =
+      reconstruction::PPM_Interfaces(cell_im2, cell_im1, cell_i, cell_ip1, cell_ip2);
+#endif    // PPMC
 
   // Do the characteristic tracing
 #ifndef VL
@@ -169,13 +151,10 @@ __global__ __launch_bounds__(TPB) void PPMC_cuda(Real *dev_conserved, Real *dev_
   reconstruction::Write_Data(interface_R_imh, dev_bounds_R, dev_conserved, id, n_cells, o1, o2, o3, gamma);
 }
 // Instantiate the relevant template specifications
-template __global__ __launch_bounds__(TPB) void PPMC_cuda<0>(Real *dev_conserved, Real *dev_bounds_L,
-                                                             Real *dev_bounds_R, int nx, int ny, int nz, Real dx,
-                                                             Real dt, Real gamma);
-template __global__ __launch_bounds__(TPB) void PPMC_cuda<1>(Real *dev_conserved, Real *dev_bounds_L,
-                                                             Real *dev_bounds_R, int nx, int ny, int nz, Real dx,
-                                                             Real dt, Real gamma);
-template __global__ __launch_bounds__(TPB) void PPMC_cuda<2>(Real *dev_conserved, Real *dev_bounds_L,
-                                                             Real *dev_bounds_R, int nx, int ny, int nz, Real dx,
-                                                             Real dt, Real gamma);
+template __global__ __launch_bounds__(TPB) void PPM_cuda<0>(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R,
+                                                            int nx, int ny, int nz, Real dx, Real dt, Real gamma);
+template __global__ __launch_bounds__(TPB) void PPM_cuda<1>(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R,
+                                                            int nx, int ny, int nz, Real dx, Real dt, Real gamma);
+template __global__ __launch_bounds__(TPB) void PPM_cuda<2>(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R,
+                                                            int nx, int ny, int nz, Real dx, Real dt, Real gamma);
 // =====================================================================================================================