feature/omptarget

.gitignore

@@ -1,14 +1,18 @@
 *.o
+*.d
 *.f90
 *.mod
 *.a
 *~
 tests/*_test
+tests/*_ctest
+make.inc
 milc_interface/*
 *#*
 *.pyc
 tunecache.tsv
 profile.tsv
+profile_*.tsv
 config.log
 CMakeCache.txt
 CMakeFiles

CMakeLists.txt

@@ -72,11 +72,11 @@ else()
   set(DEFTARGET "CUDA")
 endif()
 
-set(VALID_TARGET_TYPES CUDA HIP SYCL)
+set(VALID_TARGET_TYPES CUDA HIP SYCL OMPTARGET)
 set(QUDA_TARGET_TYPE
   "${DEFTARGET}"
   CACHE STRING "Choose the type of target, options are: ${VALID_TARGET_TYPES}")
-set_property(CACHE QUDA_TARGET_TYPE PROPERTY STRINGS CUDA HIP SYCL)
+set_property(CACHE QUDA_TARGET_TYPE PROPERTY STRINGS CUDA HIP SYCL OMPTARGET)
 
 string(TOUPPER ${QUDA_TARGET_TYPE} CHECK_TARGET_TYPE)
 list(FIND VALID_TARGET_TYPES ${CHECK_TARGET_TYPE} TARGET_TYPE_VALID)
@@ -240,10 +240,10 @@ option(QUDA_ALTERNATIVE_I_TO_F "enable using alternative integer-to-float conver
 
 option(QUDA_OPENMP "enable OpenMP" OFF)
 set(QUDA_CXX_STANDARD
-    17
-    CACHE STRING "set the CXX Standard (14 or 17)")
+    20
+    CACHE STRING "set the CXX Standard (14, 17, 20)")
 
-set_property(CACHE QUDA_CXX_STANDARD PROPERTY STRINGS 14 17)
+set_property(CACHE QUDA_CXX_STANDARD PROPERTY STRINGS 14 17 20)
 
 option(QUDA_BACKWARDS "Enable stacktrace generation using backwards-cpp")
 

include/array.h

@@ -17,13 +17,6 @@ namespace quda
     constexpr T &operator[](int i) { return data[i]; }
     constexpr const T &operator[](int i) const { return data[i]; }
     constexpr int size() const { return n; }
-
-    array() = default;
-    array(const array<T, n> &) = default;
-    array(array<T, n> &&) = default;
-
-    array<T, n> &operator=(const array<T, n> &) = default;
-    array<T, n> &operator=(array<T, n> &&) = default;
   };
 
   template <typename T, int n> std::ostream &operator<<(std::ostream &output, const array<T, n> &a)

include/comm_quda.h

@@ -1,5 +1,6 @@
 #pragma once
 #include <cstdint>
+#include <cstdlib>
 #include <vector>
 #include <quda_constants.h>
 #include <quda_api.h>

include/communicator_quda.h

@@ -544,7 +544,7 @@ namespace quda
 
     if (gpuid < 0) {
       int device_count = device::get_device_count();
-      if (device_count == 0) { errorQuda("No devices found"); }
+      if (device_count == 0) { warningQuda("No devices found"); }
 
       // We initialize gpuid if it's still negative.
       gpuid = 0;
@@ -558,7 +558,7 @@ namespace quda
           gpuid = gpuid % device_count;
           printf("MPS enabled, rank=%3d -> gpu=%d\n", comm_rank(), gpuid);
         } else {
-          errorQuda("Too few GPUs available on %s", comm_hostname());
+          warningQuda("Too few GPUs available on %s", comm_hostname());
         }
       }
     } // -ve gpuid

include/complex_quda.h

@@ -27,6 +27,8 @@
 #include <type_traits>
 #include <quda_arch.h> // for double2 / float2
 
+#include <quda_api.h>
+
 namespace quda {
   namespace gauge {
     template<typename Float, typename storeFloat> struct fieldorder_wrapper;

include/dslash_helper.cuh

@@ -12,7 +12,7 @@
 #include <kernel_helper.h>
 #include <tune_quda.h>
 
-#if defined(_NVHPC_CUDA)
+#if defined(_NVHPC_CUDA) || defined(QUDA_TARGET_OMPTARGET)
 #include <constant_kernel_arg.h>
 constexpr quda::use_kernel_arg_p use_kernel_arg = quda::use_kernel_arg_p::FALSE;
 #else

include/kernel_helper.h

@@ -5,6 +5,17 @@
 namespace quda
 {
 
+  enum ThreadsSync {
+    ThreadsSyncNo = 0,
+    ThreadsSyncX = 1,
+    ThreadsSyncY = 2,
+    ThreadsSyncXY = 3,
+    ThreadsSyncZ = 4,
+    ThreadsSyncXZ = 5,
+    ThreadsSyncYZ = 6,
+    ThreadsSyncAll = 7
+  };
+
   struct kernel_t {
     const void *func;
     const std::string name;

include/kernels/coarse_op_kernel.cuh

@@ -1029,6 +1029,7 @@ namespace quda {
 
   template <typename Arg>
   __device__ __host__ inline int virtualThreadIdx(const Arg &arg) {
+    QUDA_RT_CONSTS;
     int warp_id = threadIdx.x / device::warp_size();
     int warp_lane = threadIdx.x % device::warp_size();
     int tx = warp_id * (device::warp_size() / arg.aggregates_per_block) + warp_lane / arg.aggregates_per_block;
@@ -1037,12 +1038,14 @@ namespace quda {
 
   template <typename Arg>
   __device__ __host__ inline int virtualBlockDim(const Arg &arg) {
+    QUDA_RT_CONSTS;
     int block_dim_x = blockDim.x / arg.aggregates_per_block;
     return block_dim_x;
   }
 
   template <typename Arg>
   __device__ __host__ inline int coarseIndex(const Arg &arg) {
+    QUDA_RT_CONSTS;
     int warp_lane = threadIdx.x % device::warp_size();
     int x_coarse = (arg.coarse_color_wave ? blockIdx.y : blockIdx.x) * arg.aggregates_per_block + warp_lane % arg.aggregates_per_block;
     return x_coarse;
@@ -1396,6 +1399,7 @@ namespace quda {
     template <typename VUV, typename Pack, typename Arg>
     inline __device__ void operator()(VUV &vuv, bool isDiagonal, int coarse_x_cb, int coarse_parity, int i0, int j0, int parity, const Pack &pack, const Arg &arg)
     {
+      QUDA_RT_CONSTS;
       using real = typename Arg::Float;
       using TileType = typename Arg::vuvTileType;
       const int dim_index = arg.dim_index % arg.Y_atomic.geometry;
@@ -1638,6 +1642,7 @@ namespace quda {
     template <typename Arg> __device__ inline void operator()(int &parity_coarse, int &x_coarse_cb, int &parity, int &x_cb,
                                                               int &parity_c_row, int &c_row, int &c_col, const Arg &arg)
     {
+      QUDA_RT_CONSTS;
       if (arg.coarse_color_wave) {
         int parity_c_row_block_idx_z = blockDim.y*blockIdx.x + threadIdx.y;
         int c_row_block_idx_z = arg.parity_flip ? (parity_c_row_block_idx_z % arg.coarse_color_grid_z ) : (parity_c_row_block_idx_z / 2); // coarse color row index

include/kernels/color_spinor_pack.cuh

@@ -64,6 +64,7 @@ namespace quda {
 
   template <typename store_t, typename ghost_store_t, int nSpin_, int nColor_, int nDim_, QudaFieldOrder order>
   struct PackGhostArg : kernel_param<> {
+    static constexpr ThreadsSync requires_threads_sync = ThreadsSyncX;
     static constexpr bool block_float = sizeof(store_t) == QUDA_SINGLE_PRECISION && isFixed<ghost_store_t>::value;
 
     // ensure we only compile supported block-float kernels

include/kernels/dslash_coarse.cuh

@@ -31,6 +31,8 @@ namespace quda {
   template <bool dslash_, bool clover_, bool dagger_, DslashType type_, int color_stride_, int dim_stride_, typename Float,
             typename yFloat, typename ghostFloat, int nSpin_, int nColor_, bool native>
   struct DslashCoarseArg : kernel_param<> {
+    static constexpr ThreadsSync requires_threads_sync = ThreadsSyncAll;
+
     static constexpr bool dslash = dslash_;
     static constexpr bool clover = clover_;
     static constexpr bool dagger = dagger_;
@@ -364,11 +366,18 @@ namespace quda {
 
       if (doBulk<Arg::type>() && Arg::clover && dir==0 && dim==0) applyClover<Mc>(out, arg, x_cb, src_idx, parity, s, color_block, color_offset);
 
+#ifdef QUDA_TARGET_OMPTARGET
+      // reduce down to the first group of column-split threads, do it for every threads in openmp.
+      out = warp_combine<Arg::color_stride>(out);
+#endif
+
       if (dir==0 && dim==0) {
         const int my_spinor_parity = (arg.nParity == 2) ? parity : 0;
 
+#ifndef QUDA_TARGET_OMPTARGET
         // reduce down to the first group of column-split threads
         out = warp_combine<Arg::color_stride>(out);
+#endif
 
 #pragma unroll
         for (int color_local=0; color_local<Mc; color_local++) {

include/kernels/dslash_domain_wall_m5.cuh

@@ -212,6 +212,7 @@ namespace quda
   template <bool sync, bool dagger, bool shared, class Vector, class Arg, Dslash5Type type = Arg::type>
   __device__ __host__ inline Vector d5(const Arg &arg, const Vector &in, int parity, int x_cb, int s, int src_idx)
   {
+    QUDA_RT_CONSTS;
 
     using real = typename Arg::real;
     constexpr bool is_variable = true;
@@ -378,6 +379,7 @@ namespace quda
   __device__ __host__ inline Vector constantInv(const Arg &arg, const Vector &in, int parity, int x_cb, int s_,
                                                 int src_idx)
   {
+    QUDA_RT_CONSTS;
     using real = typename Arg::real;
     const auto k = arg.kappa;
     const auto inv = arg.inv;
@@ -436,6 +438,7 @@ namespace quda
   __device__ __host__ inline Vector variableInv(const Arg &arg, const Vector &in, int parity, int x_cb, int s_,
                                                 int src_idx)
   {
+    QUDA_RT_CONSTS;
     constexpr int nSpin = 4;
     using real = typename Arg::real;
     typedef ColorSpinor<real, Arg::nColor, nSpin / 2> HalfVector;

include/kernels/dslash_mobius_eofa.cuh

@@ -106,6 +106,7 @@ namespace quda
 
       __device__ __host__ inline void operator()(int x_cb, int src_s, int parity)
       {
+        QUDA_RT_CONSTS;
         using real = typename Arg::real;
         typedef ColorSpinor<real, Arg::nColor, 4> Vector;
 
@@ -186,6 +187,7 @@ namespace quda
 
       __device__ __host__ inline void operator()(int x_cb, int src_s, int parity)
       {
+        QUDA_RT_CONSTS;
         using real = typename Arg::real;
         typedef ColorSpinor<real, Arg::nColor, 4> Vector;
 

include/kernels/gauge_fix_ovr.cuh

@@ -101,6 +101,7 @@ namespace quda {
    */
   template <typename store_t, QudaReconstructType recon, int gauge_dir_, bool halo_, int type_>
   struct GaugeFixArg : kernel_param<> {
+    static constexpr ThreadsSync requires_threads_sync = ThreadsSyncAll;
     using real = typename mapper<store_t>::type;
     static constexpr int gauge_dir = gauge_dir_;
     static constexpr bool halo = halo_;

include/kernels/momentum.cuh

@@ -90,6 +90,17 @@ namespace quda {
     constexpr MomUpdate(const Arg &arg) : arg(arg) {}
     static constexpr const char *filename() { return KERNEL_FILE; }
 
+#ifdef QUDA_TARGET_OMPTARGET
+    static reduce_t reduce_omp(const reduce_t &a, const reduce_t &b)
+    {
+      auto c = a;
+      if (b[0] > a[0]) c[0] = b[0];
+      if (b[1] > a[1]) c[1] = b[1];
+      return c;
+    }
+    static reduce_t init_omp() { return reduce_t(); }  // see UpdateMomArg::init().
+#endif
+
     // calculate the momentum contribution to the action.  This uses the
     // MILC convention where we subtract 4.0 from each matrix norm in
     // order to increase stability

include/kernels/multi_blas_core.cuh

@@ -35,6 +35,7 @@ namespace quda
     struct MultiBlasArg : kernel_param<>,
       SpinorXZ<NXZ_, store_t, N, Functor_::use_z>,
       SpinorYW<max_YW_size<NXZ_, store_t, y_store_t, Functor_>(), store_t, N, y_store_t, Ny, Functor_::use_w> {
+      static constexpr ThreadsSync requires_threads_sync = ThreadsSyncAll;
       using real = real_;
       using Functor = Functor_;
       static constexpr int warp_split = warp_split_;

include/kernels/random_init.cuh

@@ -9,6 +9,7 @@
 namespace quda {
 
   struct rngArg : kernel_param<> {
+    static constexpr ThreadsSync requires_threads_sync = ThreadsSyncNo;
     int commCoord[QUDA_MAX_DIM];
     int X[QUDA_MAX_DIM];
     int X_global[QUDA_MAX_DIM];

include/kernels/reduce_init.cuh

@@ -6,6 +6,7 @@ namespace quda {
   namespace reducer {
 
     template <typename T_> struct init_arg : kernel_param<> {
+      static constexpr ThreadsSync requires_threads_sync = ThreadsSyncNo;
       using T = T_;
       T *count;
       init_arg(T *count, int n_reduce) :

include/quda_api.h

@@ -5,6 +5,13 @@
 #include <enum_quda.h>
 #include <quda_ptr.h>
 
+/* We have to overwrite some cuda-ism here even for public interface,
+   other wise we can't compile tests.
+ */
+#ifdef QUDA_TARGET_OMPTARGET
+#include "targets/omptarget/quda_api.h"
+#endif
+
 /**
    @file quda_api.h
 

include/quda_arch.h

@@ -14,6 +14,9 @@
 
 #elif defined(QUDA_TARGET_SYCL)
 #include <targets/sycl/quda_sycl.h>
+
+#elif defined(QUDA_TARGET_OMPTARGET)
+#include <omp.h>
 #endif
 
 #ifdef QUDA_OPENMP

include/quda_define.h.in

@@ -238,6 +238,12 @@ static_assert(QUDA_ORDER_FP_MG == 2 || QUDA_ORDER_FP_MG == 4 || QUDA_ORDER_FP_MG
  */
 #cmakedefine QUDA_TARGET_SYCL @QUDA_TARGET_SYCL@
 
-#if !defined(QUDA_TARGET_CUDA) && !defined(QUDA_TARGET_HIP) && !defined(QUDA_TARGET_SYCL)
+/**
+ * @def QUDA_TARGET_OMPTARGET
+ * @brief This macro is set by CMake if the OMPTARGET Build Target is selected
+ */
+#cmakedefine QUDA_TARGET_OMPTARGET @QUDA_TARGET_OMPTARGET@
+
+#if !defined(QUDA_TARGET_CUDA) && !defined(QUDA_TARGET_HIP) && !defined(QUDA_TARGET_SYCL) && !defined(QUDA_TARGET_OMPTARGET)
 #error "No QUDA_TARGET selected"
 #endif

include/quda_ptr.h

@@ -1,6 +1,7 @@
 #pragma once
 
 #include <ostream>
+#include <utility>
 #include "malloc_quda.h"
 
 namespace quda