make special ops initialization mandatory at construction time

lattice · Oct 4, 2023 · a5d3463 · a5d3463
1 parent e9cd06d
commit a5d3463
Show file tree

Hide file tree

Showing 48 changed files with 321 additions and 133 deletions.
diff --git a/include/dslash_helper.cuh b/include/dslash_helper.cuh
@@ -662,15 +662,18 @@ namespace quda
     static constexpr bool dagger = Arg::dagger;
     static constexpr KernelType kernel_type = Arg::kernel_type;
     static constexpr const char *filename() { return Arg::D::filename(); }
-    constexpr dslash_functor(const Arg &arg) : arg(arg.arg) { }
+    using typename getSpecialOps<typename Arg::D>::KernelOpsT;
+    template <typename ...Ops>
+    constexpr dslash_functor(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg.arg) { }
 
     template <bool allthreads = false>
     __forceinline__ __device__ void operator()(int, int s, int parity, bool active = true)
     {
-      typename Arg::D dslash(arg);
-      if constexpr (hasSpecialOps<typename Arg::D>) {
-	dslash.setSpecialOps(*this);
-      }
+      //typename Arg::D dslash(arg);
+      //if constexpr (hasSpecialOps<typename Arg::D>) {
+      //dslash.setSpecialOps(*this);
+      //}
+      typename Arg::D dslash(*this);
       // for full fields set parity from z thread index else use arg setting
       if (nParity == 1) parity = arg.parity;
 

diff --git a/include/kernels/block_orthogonalize.cuh b/include/kernels/block_orthogonalize.cuh
@@ -114,7 +114,9 @@ namespace quda {
     using dot_t = typename BlockOrtho_Params<Arg>::dot_t;
     using real = typename Arg::real;
 
-    constexpr BlockOrtho_(const Arg &arg) : arg(arg) {}
+    using typename BlockOrtho_Params<Arg>::Ops::KernelOpsT;
+    template <typename ...Ops>
+    constexpr BlockOrtho_(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     __device__ __host__ inline void load(ColorSpinor<real, nColor, spinBlock> &v, int parity, int x_cb, int chirality, int i)

diff --git a/include/kernels/block_transpose.cuh b/include/kernels/block_transpose.cuh
@@ -57,7 +57,9 @@ namespace quda
 
   template <typename Arg> struct BlockTransposeKernel : BlockTransposeKernelOps<Arg>::Ops {
     const Arg &arg;
-    constexpr BlockTransposeKernel(const Arg &arg) : arg(arg) { }
+    using typename BlockTransposeKernelOps<Arg>::Ops::KernelOpsT;
+    template <typename ...OpsArgs>
+    constexpr BlockTransposeKernel(const Arg &arg, const OpsArgs &...ops) : KernelOpsT(ops...), arg(arg) { }
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     /**

diff --git a/include/kernels/clover_deriv.cuh b/include/kernels/clover_deriv.cuh
@@ -205,7 +205,8 @@ namespace quda
   template <typename Arg> struct CloverDerivative : computeForceOps
   {
     const Arg &arg;
-    constexpr CloverDerivative(const Arg &arg) : arg(arg) {}
+    template <typename ...Ops>
+    constexpr CloverDerivative(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     __host__ __device__ void operator()(int x_cb, int parity, int mu)

diff --git a/include/kernels/coarse_op_kernel.cuh b/include/kernels/coarse_op_kernel.cuh
@@ -1702,7 +1702,9 @@ namespace quda {
     static constexpr int nFace = 1;
     const Arg &arg;
     static constexpr const char *filename() { return KERNEL_FILE; }
-    constexpr compute_vuv(const Arg &arg) : arg(arg) { }
+    using typename storeCoarseSharedAtomic_impl<true>::Ops<Arg>::KernelOpsT;
+    template <typename ...Ops>
+    constexpr compute_vuv(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) { }
 
     /**
        3-d parallelism
@@ -1735,7 +1737,9 @@ namespace quda {
     static constexpr int nFace = 3;
     const Arg &arg;
     static constexpr const char *filename() { return KERNEL_FILE; }
-    constexpr compute_vlv(const Arg &arg) : arg(arg) { }
+    using typename storeCoarseSharedAtomic_impl<true>::Ops<Arg_>::KernelOpsT;
+    template <typename ...Ops>
+    constexpr compute_vlv(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) { }
 
     /**
        3-d parallelism

diff --git a/include/kernels/color_spinor_pack.cuh b/include/kernels/color_spinor_pack.cuh
@@ -292,11 +292,15 @@ namespace quda {
     }
   }
 
-  template <typename Arg_> struct GhostPacker :
-    std::conditional_t<Arg_::block_float, site_max<true>::Ops<Arg_>, NoSpecialOps> {
+  template <typename Arg_> using GhostPackerOps =
+    std::conditional_t<Arg_::block_float, site_max<true>::Ops<Arg_>, NoSpecialOps>;
+
+  template <typename Arg_> struct GhostPacker : GhostPackerOps<Arg_> {
     using Arg = Arg_;
     const Arg &arg;
-    constexpr GhostPacker(const Arg &arg) : arg(arg) {}
+    using typename GhostPackerOps<Arg>::KernelOpsT;
+    template <typename ...Ops>
+    constexpr GhostPacker(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     template <bool allthreads = false>

diff --git a/include/kernels/covDev.cuh b/include/kernels/covDev.cuh
@@ -124,7 +124,8 @@ namespace quda
     dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr covDev(const Arg &arg) : arg(arg) {}
+    //constexpr covDev(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr covDev(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     template <KernelType mykernel_type = kernel_type>

diff --git a/include/kernels/dslash_clover_helper.cuh b/include/kernels/dslash_clover_helper.cuh
@@ -172,19 +172,23 @@ namespace quda {
       arg.out(x_cb, spinor_parity) = out;
     }
   };
-
+
+  template <typename Arg> using NdegTwistCloverApplyOps =
+    SpecialOps<SharedMemoryCache<ColorSpinor<typename Arg::real, Arg::nColor, Arg::nSpin / 2>>>;
+
   // if (!inverse) apply (Clover + i*a*gamma_5*tau_3 + b*epsilon*tau_1) to the input spinor
   // else apply (Clover + i*a*gamma_5*tau_3 + b*epsilon*tau_1)/(Clover^2 + a^2 - b^2) to the input spinor
   // noting that appropriate signs are carried by a and b depending on inverse
-  template <typename Arg> struct NdegTwistCloverApply :
-    SpecialOps<SharedMemoryCache<ColorSpinor<typename Arg::real, Arg::nColor, Arg::nSpin / 2>>> {
+  template <typename Arg> struct NdegTwistCloverApply : NdegTwistCloverApplyOps<Arg> {
     static constexpr int N = Arg::nColor * Arg::nSpin / 2;
     using real = typename Arg::real;
     using fermion = ColorSpinor<typename Arg::real, Arg::nColor, Arg::nSpin>;
     using half_fermion = ColorSpinor<typename Arg::real, Arg::nColor, Arg::nSpin / 2>;
     using Mat = HMatrix<typename Arg::real, N>;
     const Arg &arg;
-    constexpr NdegTwistCloverApply(const Arg &arg) : arg(arg) {}
+    using typename NdegTwistCloverApplyOps<Arg>::KernelOpsT;
+    template <typename ...Ops>
+    constexpr NdegTwistCloverApply(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
     static constexpr const char* filename() { return KERNEL_FILE; }
 
     template <bool allthreads = false>

diff --git a/include/kernels/dslash_coarse.cuh b/include/kernels/dslash_coarse.cuh
@@ -338,7 +338,9 @@ namespace quda {
   template <typename Arg_> struct CoarseDslash : CoarseDslashParams<Arg_>::Ops {
     using Arg = Arg_;
     const Arg &arg;
-    constexpr CoarseDslash(const Arg &arg) : arg(arg) {}
+    using typename CoarseDslashParams<Arg>::Ops::KernelOpsT;
+    template <typename ...Ops>
+    constexpr CoarseDslash(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     template <bool allthreads = false>

diff --git a/include/kernels/dslash_domain_wall_4d.cuh b/include/kernels/dslash_domain_wall_4d.cuh
@@ -28,7 +28,8 @@ namespace quda
   struct domainWall4D : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr domainWall4D(const Arg &arg) : arg(arg) {}
+    //constexpr domainWall4D(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr domainWall4D(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     template <KernelType mykernel_type = kernel_type>

diff --git a/include/kernels/dslash_domain_wall_4d_fused_m5.cuh b/include/kernels/dslash_domain_wall_4d_fused_m5.cuh
@@ -61,7 +61,9 @@ namespace quda
     static constexpr Dslash5Type dslash5_type = Arg::type;
 
     const Arg &arg;
-    constexpr domainWall4DFusedM5(const Arg &arg) : arg(arg) { }
+    using typename d5Params<Arg_>::Ops::KernelOpsT;
+    //constexpr domainWall4DFusedM5(const Arg &arg) : arg(arg) { }
+    template <typename Ftor> constexpr domainWall4DFusedM5(const Ftor &ftor) : KernelOpsT(ftor), arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     template <KernelType mykernel_type = kernel_type, bool allthreads = false>

diff --git a/include/kernels/dslash_domain_wall_5d.cuh b/include/kernels/dslash_domain_wall_5d.cuh
@@ -26,7 +26,8 @@ namespace quda
   struct domainWall5D : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr domainWall5D(const Arg &arg) : arg(arg) {}
+    //constexpr domainWall5D(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr domainWall5D(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
     static constexpr QudaPCType pc_type() { return QUDA_5D_PC; }
 

diff --git a/include/kernels/dslash_domain_wall_m5.cuh b/include/kernels/dslash_domain_wall_m5.cuh
@@ -333,7 +333,9 @@ namespace quda
   template <typename Arg_> struct dslash5 : d5Params<Arg_>::Ops {
     using Arg = Arg_;
     const Arg &arg;
-    constexpr dslash5(const Arg &arg) : arg(arg) { }
+    using typename d5Params<Arg_>::Ops::KernelOpsT;
+    template <typename ...OpsArgs>
+    constexpr dslash5(const Arg &arg, const OpsArgs &...ops) : KernelOpsT(ops...), arg(arg) { }
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     /**
@@ -589,7 +591,9 @@ namespace quda
   template <typename Arg_> struct dslash5inv : dslash5invParams<Arg_>::Ops {
     using Arg = Arg_;
     const Arg &arg;
-    constexpr dslash5inv(const Arg &arg) : arg(arg) {}
+    using typename dslash5invParams<Arg>::Ops::KernelOpsT;
+    template <typename ...OpsArgs>
+    constexpr dslash5inv(const Arg &arg, const OpsArgs &...ops) : KernelOpsT(ops...), arg(arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     /**

diff --git a/include/kernels/dslash_mobius_eofa.cuh b/include/kernels/dslash_mobius_eofa.cuh
@@ -90,17 +90,20 @@ namespace quda
       }
     };
 
+    template <typename Arg> using eofa_dslash5Ops =
+      SpecialOps<SharedMemoryCache<ColorSpinor<typename Arg::real, Arg::nColor, 4>>>;
     /**
       @brief Apply the D5 operator at given site
       @param[in] arg    Argument struct containing any meta data and accessors
       @param[in] parity Parity we are on
       @param[in] x_cb   Checkerboarded 4-d space-time index
       @param[in] s      Ls dimension coordinate
      */
-    template <typename Arg> struct eofa_dslash5 :
-      SpecialOps<SharedMemoryCache<ColorSpinor<typename Arg::real, Arg::nColor, 4>>> {
+    template <typename Arg> struct eofa_dslash5 : eofa_dslash5Ops<Arg> {
       const Arg &arg;
-      constexpr eofa_dslash5(const Arg &arg) : arg(arg) {}
+      using typename eofa_dslash5Ops<Arg>::KernelOpsT;
+      template <typename ...Ops>
+      constexpr eofa_dslash5(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
       static constexpr const char *filename() { return KERNEL_FILE; }
 
       template <bool allthreads = false>
@@ -170,6 +173,8 @@ namespace quda
       }
     };
 
+    template <typename Arg> using eofa_dslash5invOps =
+      SpecialOps<SharedMemoryCache<ColorSpinor<typename Arg::real, Arg::nColor, 4>>>;
     /**
       @brief Apply the M5 inverse operator at a given site on the
       lattice.  This is the original algorithm as described in Kim and
@@ -182,10 +187,11 @@ namespace quda
       @param[in] x_cb   Checkerboarded 4-d space-time index
       @param[in] s      Ls dimension coordinate
      */
-    template <typename Arg> struct eofa_dslash5inv :
-      SpecialOps<SharedMemoryCache<ColorSpinor<typename Arg::real, Arg::nColor, 4>>> {
+    template <typename Arg> struct eofa_dslash5inv : eofa_dslash5invOps<Arg> {
       const Arg &arg;
-      constexpr eofa_dslash5inv(const Arg &arg) : arg(arg) {}
+      using typename eofa_dslash5invOps<Arg>::KernelOpsT;
+      template <typename ...Ops>
+      constexpr eofa_dslash5inv(const Arg &arg, const Ops &...ops) : KernelOpsT(ops...), arg(arg) {}
       static constexpr const char *filename() { return KERNEL_FILE; }
 
       template <bool allthreads = false>

diff --git a/include/kernels/dslash_ndeg_twisted_clover.cuh b/include/kernels/dslash_ndeg_twisted_clover.cuh
@@ -40,16 +40,18 @@ namespace quda
     using real = typename mapper<typename Arg::Float>::type;
     using Vec = ColorSpinor<real, Arg::nColor, 4>;
     using Cache = SharedMemoryCache<Vec>;
-    using Ops = SpecialOps<Cache>;
+    //using Ops = SpecialOps<Cache>;
     //template <KernelType kernel_type>
-    //using Ops = conditional_t<kernel_type == INTERIOR_KERNEL,SpecialOps<Cache>,NoSpecialOps>;
+    using Ops = std::conditional_t<kernel_type == INTERIOR_KERNEL,SpecialOps<Cache>,NoSpecialOps>;
   };
 
   template <int nParity, bool dagger, bool xpay, KernelType kernel_type, typename Arg>
   struct nDegTwistedClover : dslash_default, nDegTwistedCloverParams<kernel_type,Arg>::Ops {
 
     const Arg &arg;
-    constexpr nDegTwistedClover(const Arg &arg) : arg(arg) {}
+    using typename nDegTwistedCloverParams<kernel_type,Arg>::Ops::KernelOpsT;
+    //constexpr nDegTwistedClover(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr nDegTwistedClover(const Ftor &ftor) : KernelOpsT(ftor), arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_ndeg_twisted_clover_preconditioned.cuh b/include/kernels/dslash_ndeg_twisted_clover_preconditioned.cuh
@@ -52,7 +52,9 @@ namespace quda
   struct nDegTwistedCloverPreconditioned : dslash_default, nDegTwistedCloverPreconditionedParams<Arg>::Ops {
 
     const Arg &arg;
-    constexpr nDegTwistedCloverPreconditioned(const Arg &arg) : arg(arg) {}
+    using typename nDegTwistedCloverPreconditionedParams<Arg>::Ops::KernelOpsT;
+    //constexpr nDegTwistedCloverPreconditioned(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr nDegTwistedCloverPreconditioned(const Ftor &ftor) : KernelOpsT(ftor), arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_ndeg_twisted_mass.cuh b/include/kernels/dslash_ndeg_twisted_mass.cuh
@@ -26,7 +26,8 @@ namespace quda
   struct nDegTwistedMass : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr nDegTwistedMass(const Arg &arg) : arg(arg) {}
+    //constexpr nDegTwistedMass(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr nDegTwistedMass(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_ndeg_twisted_mass_preconditioned.cuh b/include/kernels/dslash_ndeg_twisted_mass_preconditioned.cuh
@@ -48,7 +48,9 @@ namespace quda
   struct nDegTwistedMassPreconditioned : dslash_default, nDegTwistedMassPreconditionedParams<dagger, Arg>::Ops {
 
     const Arg &arg;
-    constexpr nDegTwistedMassPreconditioned(const Arg &arg) : arg(arg) {}
+    using typename nDegTwistedMassPreconditionedParams<dagger,Arg>::Ops::KernelOpsT;
+    //constexpr nDegTwistedMassPreconditioned(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr nDegTwistedMassPreconditioned(const Ftor &ftor) : KernelOpsT(ftor), arg(ftor.arg) {}
     constexpr int twist_pack() const { return (!Arg::asymmetric && dagger) ? 2 : 0; }
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 

diff --git a/include/kernels/dslash_staggered.cuh b/include/kernels/dslash_staggered.cuh
@@ -173,7 +173,8 @@ namespace quda
   struct staggered : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr staggered(const Arg &arg) : arg(arg) {}
+    //constexpr staggered(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr staggered(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     template <KernelType mykernel_type = kernel_type>

diff --git a/include/kernels/dslash_twisted_clover_preconditioned.cuh b/include/kernels/dslash_twisted_clover_preconditioned.cuh
@@ -40,7 +40,8 @@ namespace quda
   struct twistedCloverPreconditioned : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr twistedCloverPreconditioned(const Arg &arg) : arg(arg) {}
+    //constexpr twistedCloverPreconditioned(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr twistedCloverPreconditioned(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_twisted_mass.cuh b/include/kernels/dslash_twisted_mass.cuh
@@ -24,7 +24,8 @@ namespace quda
   struct twistedMass : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr twistedMass(const Arg &arg) : arg(arg) {}
+    //constexpr twistedMass(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr twistedMass(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_twisted_mass_preconditioned.cuh b/include/kernels/dslash_twisted_mass_preconditioned.cuh
@@ -135,7 +135,8 @@ namespace quda
   struct twistedMassPreconditioned : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr twistedMassPreconditioned(const Arg &arg) : arg(arg) {}
+    //constexpr twistedMassPreconditioned(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr twistedMassPreconditioned(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
     constexpr int twist_pack() const { return (!Arg::asymmetric && dagger) ? 1 : 0; }
 

diff --git a/include/kernels/dslash_wilson.cuh b/include/kernels/dslash_wilson.cuh
@@ -137,7 +137,8 @@ namespace quda
     dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr wilson(const Arg &arg) : arg(arg) {}
+    //constexpr wilson(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr wilson(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     // out(x) = M*in = (-D + m) * in(x-mu)

diff --git a/include/kernels/dslash_wilson_clover.cuh b/include/kernels/dslash_wilson_clover.cuh
@@ -36,7 +36,8 @@ namespace quda
   struct wilsonClover : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr wilsonClover(const Arg &arg) : arg(arg) {}
+    //constexpr wilsonClover(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr wilsonClover(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_wilson_clover_hasenbusch_twist.cuh b/include/kernels/dslash_wilson_clover_hasenbusch_twist.cuh
@@ -36,7 +36,8 @@ namespace quda
   struct cloverHasenbusch : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr cloverHasenbusch(const Arg &arg) : arg(arg) {}
+    //constexpr cloverHasenbusch(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr cloverHasenbusch(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_wilson_clover_hasenbusch_twist_preconditioned.cuh b/include/kernels/dslash_wilson_clover_hasenbusch_twist_preconditioned.cuh
@@ -38,7 +38,8 @@ namespace quda
   struct cloverHasenbuschPreconditioned : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr cloverHasenbuschPreconditioned(const Arg &arg) : arg(arg) {}
+    //constexpr cloverHasenbuschPreconditioned(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr cloverHasenbuschPreconditioned(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**

diff --git a/include/kernels/dslash_wilson_clover_preconditioned.cuh b/include/kernels/dslash_wilson_clover_preconditioned.cuh
@@ -34,7 +34,8 @@ namespace quda
   struct wilsonCloverPreconditioned : dslash_default, NoSpecialOps {
 
     const Arg &arg;
-    constexpr wilsonCloverPreconditioned(const Arg &arg) : arg(arg) {}
+    //constexpr wilsonCloverPreconditioned(const Arg &arg) : arg(arg) {}
+    template <typename Ftor> constexpr wilsonCloverPreconditioned(const Ftor &ftor) : arg(ftor.arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; } // this file name - used for run-time compilation
 
     /**