diff --git a/src/reconstruction/plmc_cuda.cu b/src/reconstruction/plmc_cuda.cu
index e45bbf771..46fcbfd89 100644
--- a/src/reconstruction/plmc_cuda.cu
+++ b/src/reconstruction/plmc_cuda.cu
@@ -21,8 +21,8 @@
  gamma, int dir)
  *  \brief When passed a stencil of conserved variables, returns the left and
  right boundary values for the interface calculated using plm. */
-__global__ void PLMC_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx, int ny, int nz, Real dx,
-                          Real dt, Real gamma, int dir, int n_fields)
+__global__ __launch_bounds__(TPB) void PLMC_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx,
+                                                 int ny, int nz, Real dx, Real dt, Real gamma, int dir, int n_fields)
 {
   // get a thread ID
   int const thread_id = threadIdx.x + blockIdx.x * blockDim.x;
diff --git a/src/reconstruction/plmc_cuda.h b/src/reconstruction/plmc_cuda.h
index 4a1ca322b..c2d25df84 100644
--- a/src/reconstruction/plmc_cuda.h
+++ b/src/reconstruction/plmc_cuda.h
@@ -15,7 +15,7 @@
  gamma, int dir)
  *  \brief When passed a stencil of conserved variables, returns the left and
  right boundary values for the interface calculated using plm. */
-__global__ void PLMC_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx, int ny, int nz, Real dx,
-                          Real dt, Real gamma, int dir, int n_fields);
+__global__ __launch_bounds__(TPB) void PLMC_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx,
+                                                 int ny, int nz, Real dx, Real dt, Real gamma, int dir, int n_fields);
 
 #endif  // PLMC_CUDA_H
diff --git a/src/reconstruction/ppmc_cuda.cu b/src/reconstruction/ppmc_cuda.cu
index 1dec4bbc2..af40af921 100644
--- a/src/reconstruction/ppmc_cuda.cu
+++ b/src/reconstruction/ppmc_cuda.cu
@@ -539,8 +539,8 @@ __global__ void PPMC_CTU(Real *dev_conserved, Real *dev_bounds_L, Real *dev_boun
 // =====================================================================================================================
 
 // =====================================================================================================================
-__global__ void PPMC_VL(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx, int ny, int nz, Real gamma,
-                        int dir)
+__global__ __launch_bounds__(TPB) void PPMC_VL(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx,
+                                               int ny, int nz, Real gamma, int dir)
 {
   // get a thread ID
   int const thread_id = threadIdx.x + blockIdx.x * blockDim.x;
diff --git a/src/reconstruction/ppmc_cuda.h b/src/reconstruction/ppmc_cuda.h
index 033f2505b..916853874 100644
--- a/src/reconstruction/ppmc_cuda.h
+++ b/src/reconstruction/ppmc_cuda.h
@@ -47,7 +47,7 @@ __global__ void PPMC_CTU(Real *dev_conserved, Real *dev_bounds_L, Real *dev_boun
  * \param[in] gamma The adiabatic index
  * \param[in] dir The direction to reconstruct. 0=X, 1=Y, 2=Z
  */
-__global__ void PPMC_VL(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx, int ny, int nz, Real gamma,
-                        int dir);
+__global__ __launch_bounds__(TPB) void PPMC_VL(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bounds_R, int nx,
+                                               int ny, int nz, Real gamma, int dir);
 
 #endif  // PPMC_CUDA_H