Add more const to Arg &arg; Fix compile when MMA is not available.

lattice · Oct 31, 2024 · 013aa50 · 013aa50
1 parent 88bedc9
commit 013aa50
Show file tree

Hide file tree

Showing 4 changed files with 9 additions and 9 deletions.
diff --git a/include/kernels/dslash_mdw_fused.cuh b/include/kernels/dslash_mdw_fused.cuh
@@ -199,7 +199,7 @@ namespace quda {
     -> Everything should be understood in a 4d checkboarding sense.
     */
     template <class storage_type, bool dagger, bool halo, bool back, class Vector, class Arg>
-    __device__ inline void apply_wilson_5d(Vector &out, int coordinate[4], Arg &arg, int s)
+    __device__ inline void apply_wilson_5d(Vector &out, int coordinate[4], const Arg &arg, int s)
     {
       typedef typename mapper<storage_type>::type compute_type;
       typedef Matrix<complex<compute_type>, 3> Link;

diff --git a/include/targets/cuda/mdw_dslash5_tensor_core.cuh b/include/targets/cuda/mdw_dslash5_tensor_core.cuh
@@ -66,7 +66,7 @@ namespace quda
   // one(spin).
   // x by y
   template <int M_sm, bool dagger, class Arg>
-  __device__ inline void construct_matrix_a_m5inv(Arg &arg, half *sm_a, const float *mp = nullptr,
+  __device__ inline void construct_matrix_a_m5inv(const Arg &arg, half *sm_a, const float *mp = nullptr,
                                                   const float *mm = nullptr)
   {
     constexpr int Ls = Arg::Ls;
@@ -137,7 +137,7 @@ namespace quda
   // one(spin).
   // x by y
   template <int M_sm, bool dagger, class Arg>
-  __device__ inline void construct_matrix_a_d5(Arg &arg, half *sm_a)
+  __device__ inline void construct_matrix_a_d5(const Arg &arg, half *sm_a)
   {
     constexpr int Ls = Arg::Ls;
     // if we rescale, then the actual matrix is alpha*m5inv+beta.

diff --git a/lib/prolongator_mma.in.cu b/lib/prolongator_mma.in.cu
@@ -1,18 +1,18 @@
 #include <color_spinor_field.h>
 #include <multigrid.h>
 #include <tunable_nd.h>
-#include <kernels/prolongator_mma.cuh>
 #include <device.hpp>
 #include <expand_list.hpp>
 
-#if QUDA_MMA_AVAILABLE
+#ifdef QUDA_MMA_AVAILABLE
+#include <kernels/prolongator_mma.cuh>
 #include <mma_tensor_op/smma_m16n8k8_sm70.cuh>
 #endif
 
 namespace quda
 {
 
-#if QUDA_MMA_AVAILABLE
+#ifdef QUDA_MMA_AVAILABLE
 
   template <typename Float, typename vFloat, int fineSpin, int fineColor, int coarseSpin, int coarseColor, int nVec>
   class ProlongateLaunchMma : public TunableKernel
@@ -217,7 +217,7 @@ namespace quda
                                                    const ColorSpinorField &v, const int *fine_to_coarse,
                                                    const int *const *spin_map, int parity)
   {
-#if QUDA_MMA_AVAILABLE
+#ifdef QUDA_MMA_AVAILABLE
     if constexpr (is_enabled_multigrid()) {
       QudaPrecision precision = checkPrecision(out, in);
 

diff --git a/lib/restrictor_mma.in.cu b/lib/restrictor_mma.in.cu
@@ -2,11 +2,11 @@
 #include <multigrid.h>
 #include <power_of_two_array.h>
 #include <tunable_block_reduction.h>
-#include <kernels/restrictor_mma.cuh>
 #include <device.hpp>
 #include <expand_list.hpp>
 
 #ifdef QUDA_MMA_AVAILABLE
+#include <kernels/restrictor_mma.cuh>
 #include <mma_tensor_op/smma_m16n8k8_sm70.cuh>
 #endif
 
@@ -253,7 +253,7 @@ namespace quda
                                                  const ColorSpinorField &v, const int *fine_to_coarse,
                                                  const int *coarse_to_fine, const int *const *spin_map, int parity)
   {
-#if QUDA_MMA_AVAILABLE
+#ifdef QUDA_MMA_AVAILABLE
     if constexpr (is_enabled_multigrid()) {
 
       checkLocation(out, in, v);