From af8d5553b37786e4a04352b09a560531f1c0b876 Mon Sep 17 00:00:00 2001
From: James Osborn <osborn@alcf.anl.gov>
Date: Sun, 17 Sep 2023 19:30:10 -0500
Subject: [PATCH] add missing file

---
 include/targets/cuda/special_ops_target.h | 229 ----------------------
 include/targets/hip/special_ops_target.h  |  29 +++
 2 files changed, 29 insertions(+), 229 deletions(-)
 create mode 100644 include/targets/hip/special_ops_target.h

diff --git a/include/targets/cuda/special_ops_target.h b/include/targets/cuda/special_ops_target.h
index 352d99e6b7..dbc4703d83 100644
--- a/include/targets/cuda/special_ops_target.h
+++ b/include/targets/cuda/special_ops_target.h
@@ -11,137 +11,6 @@ namespace quda {
     }
   };
 
-#if 0
-  // blockSync
-  template <typename ...T>
-  inline void blockSync(SpecialOps<T...> *ops) {
-    static_assert(hasBlockSync<T...>);
-    //if (ops->ndi == nullptr) {
-    //  errorQuda("SpecialOps not set");
-    //}
-#ifdef __SYCL_DEVICE_ONLY__
-    sycl::group_barrier(ops->ndi->get_group());
-#endif
-  }
-  template <typename ...T> inline void blockSync(SpecialOps<T...> ops) { blockSync(&ops); }
-
-  template <typename ...T> static constexpr bool isOpConcurrent = false;
-  template <typename ...T> static constexpr bool isOpConcurrent<op_Concurrent<T...>> = true;
-
-  template <typename T, typename ...U> static constexpr int getOpIndex = 0;
-  template <typename T, typename ...U> static constexpr int getOpIndex<T,op_Concurrent<U...>> = getOpIndex<T,U...>;
-  template <typename T, typename U, typename ...V> static constexpr int getOpIndex<T, U, V...> =
-    std::is_same_v<T,U> ? 0 : (1 + getOpIndex<T,V...>);
-
-  // getSpecialOp
-  template <typename U, int n = 0, typename ...T>
-  inline SpecialOpsType<U,n> getSpecialOp(const SpecialOps<T...> &ops) {
-    if constexpr (!isOpConcurrent<U> && sizeof...(T) == 1 && isOpConcurrent<T...>) {
-      static constexpr int i = getOpIndex<U, T...>;
-      return getSpecialOp<T...,i>(ops);
-    } else {
-      static_assert(hasSpecialOpType<U,T...>);
-      //if (ops->ndi == nullptr || ops->smem == nullptr) {
-      //	errorQuda("SpecialOps not set");
-      //}
-      SpecialOpsType<U,n> s;
-      s.ndi = ops.ndi;
-      //s.smem = ops->smem + sharedMemOffset<U,n>()(ops->ndi->get_local_range());  // FIXME: need to pass arg
-      s.smem = ops.smem + sharedMemOffset<U,n>()(getBlockDim());  // FIXME: need to pass arg
-      return s;
-    }
-  }
-  template <typename U, int n = 0, typename ...T>
-    inline SpecialOpsType<U,n> getSpecialOp(const SpecialOps<T...> *ops) { return getSpecialOp<U,n>(*ops); }
-  template <typename U, int n = 0> struct getSpecialOpF {
-    template <typename T> inline SpecialOpsType<U,n> operator()(const T &ops) { return getSpecialOp<U,n>(ops); }
-  };
-
-  // getDependentOps
-  template <typename U, int n = 0, typename ...T>
-  inline SpecialOpDependencies<SpecialOpsType<U,n>> getDependentOps(const SpecialOps<T...> &ops) {
-    static_assert(hasSpecialOpType<U,T...>);
-    //if (ops->ndi == nullptr || ops->smem == nullptr) {
-    //errorQuda("SpecialOps not set");
-    //}
-    SpecialOpDependencies<SpecialOpsType<U,n>> s;
-    s.ndi = ops.ndi;
-    //s.smem = ops->smem + sharedMemOffset<U,n>()(ops->ndi->get_local_range());  // FIXME: need to pass arg
-    s.smem = ops.smem + sharedMemOffset<U,n>()(getBlockDim());  // FIXME: need to pass arg
-    return s;
-  }
-
-  // getSharedMemPtr
-#if 0
-  template <typename ...T>
-  //SpecialOpsElemType<T...> *getSharedMemPtr(SpecialOps<T...> *ops) {
-  sycl::local_ptr<SpecialOpsElemType<T...>> getSharedMemPtr(SpecialOps<T...> *ops) {
-    static_assert(!std::is_same_v<SpecialOpsElemType<T...>,void>);
-    //return reinterpret_cast<SpecialOpsElemType<T...>*>(ops->smem);
-    //return reinterpret_cast<SpecialOpsElemType<T...>*>(ops->smem.get());
-    //sycl::local_ptr<SpecialOpsElemType<T...>> smem = ops->smem.get();
-    //return smem.get();
-    //auto p = ops->smem.get();
-    sycl::local_ptr<void> v(ops->smem);
-    sycl::local_ptr<SpecialOpsElemType<T...>> p(v);
-    return p;
-    //sycl::local_ptr<SpecialOpsElemType<T...>> smem;
-    //using LT = decltype(smem.get());
-    //LT pt = reinterpret_cast<LT>(p);
-    //sycl::local_ptr<SpecialOpsElemType<T...>> smem2(pt);
-    //return smem2;
-    //return reinterpret_cast<SpecialOpsElemType<T...>*>(0);
-  }
-  template <typename ...T>
-  inline SpecialOpsElemType<T...> *getSharedMemPtr(SpecialOps<T...> ops) { return getSharedMemPtr(&ops); }
-#endif
-
-  template <typename T, typename S, typename O = op_SharedMemory<T,S>>
-  inline sycl::local_ptr<T> getSharedMemPtr(const only_SharedMemory<T,S> &ops) {
-    //if (ops->ndi == nullptr || ops->smem == nullptr) {
-    //errorQuda("SpecialOps not set");
-    //}
-    sycl::local_ptr<void> v(ops.smem);
-    sycl::local_ptr<T> p(v);
-    return p;
-  }
-  //template <typename T, typename S>
-  //inline sycl::local_ptr<T> getSharedMemPtr(only_SharedMemory<T,S> ops) { return getSharedMemPtr(&ops); }
-  template <typename O, typename T, typename U, typename ...V>
-  inline auto getSharedMemPtr(const SpecialOps<T,U,V...> &ops) {
-    SpecialOps<O> op = getSpecialOp<O>(ops);
-    return getSharedMemPtr(op);
-  }
-
-  template <typename T, typename O>
-  inline auto getSharedMemory(O *ops)
-  {
-    auto s = getSpecialOp<T>(ops);
-    return getSharedMemPtr(s);
-  }
-
-  // base operation dependencies
-  struct depNone {};
-  template <> struct sharedMemSizeS<depNone> {
-    template <typename ...Arg>
-    static constexpr unsigned int size(dim3 block, Arg &...arg) { return 0; }
-  };
-
-  struct depFullBlock {};
-  template <> struct sharedMemSizeS<depFullBlock> {
-    template <typename ...Arg>
-    static constexpr unsigned int size(dim3 block, Arg &...arg) { return 0; }
-  };
-
-  template <typename T, typename S>
-  struct depSharedMem {};
-  template <typename T, typename S> struct sharedMemSizeS<depSharedMem<T,S>> {
-    template <typename ...Arg>
-    static constexpr unsigned int size(dim3 block, Arg &...arg) { return S().template size<T>(block, arg...); }
-  };
-
-#endif
-
   // op implementations
   struct op_blockSync : op_BaseT<void> {
     //using dependencies = depFullBlock;
@@ -157,102 +26,4 @@ namespace quda {
     static constexpr unsigned int shared_mem_size(dim3, Arg &...arg) { return 0; }
   };
 
-#if 0
-
-  template <typename T, int N>
-  struct op_thread_array : op_BaseT<T,N> {
-    //using dependencies = depNone;
-    using dependencies = op_SharedMemory<array<T,N>,opSizeBlock>;
-  };
-
-  template <typename T>
-  struct op_BlockReduce : op_BaseT<T> {
-    using concurrentOps = op_Concurrent<op_blockSync,op_SharedMemory<T,opSizeBlockDivWarp>>;
-    using opBlockSync = getSpecialOpF<concurrentOps,0>;
-    using opSharedMem = getSpecialOpF<concurrentOps,1>;
-    //using specialOps = SpecialOps<concurrentOps>;
-    using dependencies = concurrentOps;
-  };
-
-  template <typename T, typename D>
-  struct op_SharedMemoryCache : op_BaseT<T> {
-    template <typename ...Arg> static constexpr dim3 dims(dim3 block, Arg &...arg) { return D::dims(block, arg...); }
-    using dependencies = op_Sequential<op_blockSync,op_SharedMemory<T,opSizeDims<D>>>;
-  };
-
-  template <typename T, typename S>
-  struct op_SharedMemory : op_BaseT<T> {
-    using dependencies = depSharedMem<T,S>;
-    template <typename ...Arg>
-    static constexpr unsigned int shared_mem_size(dim3 block, Arg &...arg) { return S::template size<T>(block, arg...); }
-  };
-
-  // needsFullWarp?
-
-  // needsFullBlock
-#if 0
-  template <typename T> static constexpr bool needsFullBlock = needsFullBlock<getSpecialOps<T>>;
-  template <typename ...T> static constexpr bool needsFullBlockImpl = (needsFullBlockImpl<T> || ...);
-  template <> static constexpr bool needsFullBlockImpl<depNone> = false;
-  template <> static constexpr bool needsFullBlockImpl<depFullBlock> = true;
-  template <typename T, typename S> static constexpr bool needsFullBlockImpl<depSharedMem<T,S>> = false;
-  template <typename ...T> static constexpr bool needsFullBlockImpl<op_Concurrent<T...>> = needsFullBlockImpl<T...>;
-  template <typename ...T> static constexpr bool needsFullBlockImpl<op_Sequential<T...>> = needsFullBlockImpl<T...>;
-  template <typename T> static constexpr bool needsFullBlockF() {
-    if constexpr (std::is_base_of<op_Base,T>::value) {
-      return needsFullBlockImpl<typename T::dependencies>;
-    } else {
-      //if constexpr (hasSpecialOps<T>) {
-      //return needsFullBlock<getSpecialOps<T>>;
-      //} else {
-      //return false;
-      return needsFullBlock<typename T::dependentOps>;
-      //}
-    }
-  }
-  template <typename T> static constexpr bool needsFullBlockImpl<T> = needsFullBlockF<T>();
-  template <> static constexpr bool needsFullBlock<NoSpecialOps> = false;
-  template <typename ...T> static constexpr bool needsFullBlock<SpecialOps<T...>> = needsFullBlockImpl<T...>;
-#else
-  template <typename T> static constexpr bool needsFullBlock = needsFullBlock<getSpecialOps<T>>;
-  template <typename ...T> static constexpr bool needsFullBlock<SpecialOps<T...>> = (needsFullBlock<T> || ...);
-  template <> static constexpr bool needsFullBlock<NoSpecialOps> = false;
-#endif
-
-
-  // needsSharedMem
-#if 0
-  template <typename T> static constexpr bool needsSharedMem = needsSharedMem<getSpecialOps<T>>;
-  template <typename ...T> static constexpr bool needsSharedMemImpl = (needsSharedMemImpl<T> || ...);
-  template <> static constexpr bool needsSharedMemImpl<depNone> = false;
-  template <> static constexpr bool needsSharedMemImpl<depFullBlock> = false;
-  template <typename T, typename S> static constexpr bool needsSharedMemImpl<depSharedMem<T,S>> = true;
-  template <typename ...T> static constexpr bool needsSharedMemImpl<op_Concurrent<T...>> = needsSharedMemImpl<T...>;
-  template <typename ...T> static constexpr bool needsSharedMemImpl<op_Sequential<T...>> = needsSharedMemImpl<T...>;
-  template <typename T> static constexpr bool needsSharedMemF() {
-    if constexpr (std::is_base_of<op_Base,T>::value) {
-    //if constexpr (is_instance<T,op_Base>) {
-      return needsSharedMemImpl<typename T::dependencies>;
-    } else {
-      //if constexpr (hasSpecialOps<T>) {
-      //return needsSharedMem<getSpecialOps<T>>;
-      //} else {
-      //return false;
-      return needsSharedMem<typename T::dependentOps>;
-      //}
-    }
-  }
-  template <typename T> static constexpr bool needsSharedMemImpl<T> = needsSharedMemF<T>();
-  template <> static constexpr bool needsSharedMem<NoSpecialOps> = false;
-  template <typename ...T> static constexpr bool needsSharedMem<SpecialOps<T...>> = needsSharedMemImpl<T...>;
-#else
-  //template <typename ...T> static constexpr bool needsSharedMemImpl = (needsSharedMemImpl<T> || ...);
-  template <typename T> static constexpr bool needsSharedMemImpl = (T::shared_mem_size(dim3{8,8,8}) > 0);
-  template <typename... T> static constexpr bool needsSharedMemImpl<SpecialOps<T...>> = (needsSharedMemImpl<T> || ...);
-  template <typename T> static constexpr bool needsSharedMem = needsSharedMem<getSpecialOps<T>>;
-  template <typename... T> static constexpr bool needsSharedMem<SpecialOps<T...>> = (needsSharedMemImpl<T> || ...);
-  template <> static constexpr bool needsSharedMem<NoSpecialOps> = false;
-#endif
-
-#endif
 }
diff --git a/include/targets/hip/special_ops_target.h b/include/targets/hip/special_ops_target.h
new file mode 100644
index 0000000000..dbc4703d83
--- /dev/null
+++ b/include/targets/hip/special_ops_target.h
@@ -0,0 +1,29 @@
+#pragma once
+#include <special_ops.h>
+
+namespace quda {
+
+  // SpecialOps
+  template <typename ...T>
+  struct SpecialOps : SpecialOps_Base<T...> {
+    template <typename ...U> constexpr void setSpecialOps(const SpecialOps<U...> &ops) {
+      static_assert(std::is_same_v<SpecialOps<T...>,SpecialOps<U...>>);
+    }
+  };
+
+  // op implementations
+  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; }
+  };
+
+  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; }
+  };
+
+}