add missing file

remove some unused arguments

include/kernels/restrictor.cuh

@@ -119,7 +119,7 @@ namespace quda {
     static constexpr const char *filename() { return KERNEL_FILE; }
 
     template <bool allthreads = false>
-    __device__ __host__ inline void operator()(dim3 block, dim3 thread, bool active = true)
+    __device__ __host__ inline void operator()(dim3 block, dim3 thread, bool /*active*/ = true)
     {
       int x_fine_offset = thread.x;
       const int x_coarse = block.x;

include/targets/generic/helpers.h

@@ -45,7 +45,7 @@ namespace quda
 
   template <int x, int y, int z>
   struct DimsStatic {
-    static constexpr dim3 dims(dim3 block) {
+    static constexpr dim3 dims(dim3) {
       return dim3(x,y,z);
     }
   };

include/targets/generic/special_ops.h

@@ -6,25 +6,25 @@ namespace quda {
 
   // dimensions functors for SharedMemoryCache
   struct opDimsBlock {
-    template <typename ...Arg> static constexpr dim3 dims(dim3 b, const Arg &...arg) { return b; }
+    template <typename ...Arg> static constexpr dim3 dims(dim3 b, const Arg &...) { return b; }
   };
   template <int bx, int by, int bz> struct opDimsStatic {
-    template <typename ...Arg> static constexpr dim3 dims(dim3 b, const Arg &...arg) { return dim3(bx,by,bz); }
+    template <typename ...Arg> static constexpr dim3 dims(dim3, const Arg &...) { return dim3(bx,by,bz); }
   };
 
   // size functors for determining shared memory size
   struct opSizeBlock {
-    template <typename T, typename ...Arg> static constexpr unsigned int size(dim3 b, const Arg &...arg) {
+    template <typename T, typename ...Arg> static constexpr unsigned int size(dim3 b, const Arg &...) {
       return b.x * b.y * b.z * sizeof(T);
     }
   };
   struct opSizeBlockDivWarp {
-    template <typename T, typename ...Arg> static constexpr unsigned int size(dim3 b, const Arg &...arg) {
+    template <typename T, typename ...Arg> static constexpr unsigned int size(dim3 b, const Arg &...) {
       return ((b.x * b.y * b.z + device::warp_size() - 1)/device::warp_size()) * sizeof(T);
     }
   };
   template <unsigned int S> struct opSizeStatic {
-    template <typename T, typename ...Arg> static constexpr unsigned int size(dim3 b, const Arg &...arg) {
+    template <typename T, typename ...Arg> static constexpr unsigned int size(dim3 b, const Arg &...) {
       return S * sizeof(T);
     }
   };
@@ -54,7 +54,7 @@ namespace quda {
   static constexpr bool hasSpecialOp<T,SpecialOps<U...>> = ( std::is_same_v<T,U> || ... );
 
   //template <typename T, typename Ops> void checkSpecialOps() { static_assert(hasSpecialOp<T,Ops>); }
-  template <typename T, typename Ops> void checkSpecialOps(const Ops &ops) {
+  template <typename T, typename Ops> void checkSpecialOps(const Ops &) {
     static_assert(hasSpecialOp<T,typename Ops::SpecialOpsT>);
   }
 
@@ -248,7 +248,7 @@ namespace quda {
   };
   template <> struct sharedMemSizeS<NoSpecialOps> {
     template <typename ...Arg>
-    static constexpr unsigned int size(dim3, Arg &...arg) {
+    static constexpr unsigned int size(dim3, Arg &...) {
       return 0;
     }
   };
@@ -266,15 +266,15 @@ namespace quda {
   // sharedMemOffset
   template <typename T, int n> struct sharedMemOffset {
     template <typename ...Arg>
-    inline int operator()(dim3 block, Arg &...arg) { return 0; }
+    inline int operator()(dim3 block, Arg &...) { return 0; }
   };
   template <typename T, int n> struct sharedMemOffset<SpecialOps<T>,n> {
     template <typename ...Arg>
     inline int operator()(dim3 block, Arg &...arg) { return sharedMemOffset<T,n>()(block, arg...); }
   };
   template <typename ...T> struct sharedMemOffset<op_Concurrent<T...>,0> {
     template <typename ...Arg>
-    inline int operator()(dim3 block, Arg &...arg) { return 0; }
+    inline int operator()(dim3 block, Arg &...) { return 0; }
   };
   template <typename ...T, int n> struct sharedMemOffset<op_Concurrent<T...>,n> {
     template <typename ...Arg>

include/targets/hip/shared_memory_helper.h

@@ -0,0 +1,93 @@
+#pragma once
+
+#include <target_device.h>
+#include <special_ops.h>
+
+/**
+   @file shared_memory_helper.h
+
+   Target specific helper for allocating and accessing shared memory.
+ */
+
+namespace quda
+{
+
+  /**
+     @brief Class which is used to allocate and access shared memory.
+     The shared memory is treated as an array of type T, with the
+     number of elements given by a call to the static member
+     S::size(target::block_dim()).  The byte offset from the beginning
+     of the total shared memory block is given by the static member
+     O::shared_mem_size(target::block_dim()), or 0 if O is void.
+   */
+  template <typename T, typename S, typename O = void> class SharedMemory
+  {
+  public:
+    using value_type = T;
+
+  private:
+    T *data;
+
+    /**
+       @brief This is a dummy instantiation for the host compiler
+    */
+    template <bool, typename dummy = void> struct cache_dynamic {
+      T *operator()(unsigned int)
+      {
+        static T *cache_;
+        return cache_;
+      }
+    };
+
+    /**
+       @brief This is the handle to the dynamic shared memory
+       @return Shared memory pointer
+     */
+    template <typename dummy> struct cache_dynamic<true, dummy> {
+      __device__ inline T *operator()(unsigned int offset)
+      {
+        extern __shared__ int cache_[];
+        return reinterpret_cast<T *>(reinterpret_cast<char *>(cache_) + offset);
+      }
+    };
+
+    __device__ __host__ inline T *cache(unsigned int offset) const { return target::dispatch<cache_dynamic>(offset); }
+
+  public:
+    /**
+       @brief Byte offset for this shared memory object.
+    */
+    static constexpr unsigned int get_offset(dim3 block)
+    {
+      unsigned int o = 0;
+      if constexpr (!std::is_same_v<O, void>) { o = O::shared_mem_size(block); }
+      return o;
+    }
+
+    /**
+       @brief Shared memory size in bytes.
+    */
+    static constexpr unsigned int shared_mem_size(dim3 block) { return get_offset(block) + S::size(block) * sizeof(T); }
+
+    /**
+       @brief Constructor for SharedMemory object.
+    */
+    HostDevice constexpr SharedMemory() : data(cache(get_offset(target::block_dim()))) { }
+
+    template <typename ...U>
+    HostDevice constexpr SharedMemory(const SpecialOps<U...> &ops) : data(cache(get_offset(target::block_dim()))) { }
+
+    /**
+       @brief Return this SharedMemory object.
+    */
+    HostDevice constexpr auto sharedMem() const { return *this; }
+
+    /**
+       @brief Subscripting operator returning a reference to element.
+       @param[in] i The index to use.
+       @return Reference to value stored at that index.
+     */
+    HostDevice T &operator[](const int i) const { return data[i]; }
+  };
+
+} // namespace quda

include/targets/hip/special_ops_target.h

@@ -6,7 +6,7 @@ namespace quda {
   // SpecialOps
   template <typename ...T>
   struct SpecialOps : SpecialOps_Base<T...> {
-    template <typename ...U> constexpr void setSpecialOps(const SpecialOps<U...> &ops) {
+    template <typename ...U> constexpr void setSpecialOps(const SpecialOps<U...> &) {
       static_assert(std::is_same_v<SpecialOps<T...>,SpecialOps<U...>>);
     }
   };
@@ -15,15 +15,15 @@ namespace quda {
   struct op_blockSync : op_BaseT<void> {
     //using dependencies = depFullBlock;
     template <typename ...Arg>
-    static constexpr unsigned int shared_mem_size(dim3 block, Arg &...arg) { return 0; }
+    static constexpr unsigned int shared_mem_size(dim3 block, Arg &...) { return 0; }
   };
 
   template <typename T>
   struct op_warp_combine : op_BaseT<T> {
     //using dependencies = depNone;
     //using dependencies = depFullBlock;
     template <typename ...Arg>
-    static constexpr unsigned int shared_mem_size(dim3, Arg &...arg) { return 0; }
+    static constexpr unsigned int shared_mem_size(dim3, Arg &...) { return 0; }
   };
 
 }