CarpetX: manually implement custom OpenMP reductions for nvc++

EinsteinToolkit · Jan 10, 2025 · 612dd81 · 612dd81
1 parent ad15372
commit 612dd81
Show file tree

Hide file tree

Showing 2 changed files with 14 additions and 0 deletions.
diff --git a/CarpetX/src/reduction.cxx b/CarpetX/src/reduction.cxx
@@ -201,7 +201,14 @@ reduction<CCTK_REAL, dim> reduce(int gi, int vi, int tl) {
       // TODO: check that multi-threading actually helps (and we are
       // not dominated by memory latency)
       // TODO: document required version of OpenMP to use custom reductions
+#ifdef __NVCOMPILER
+#pragma omp parallel
+      {
+      auto& outer(red);
+      reduction<CCTK_REAL, dim> red;
+#else
 #pragma omp parallel reduction(reduction : red)
+#endif
       for (amrex::MFIter mfi(mfab, mfitinfo); mfi.isValid(); ++mfi) {
         const amrex::Box &bx = mfi.tilebox(); // current tile (without ghosts)
         const vect<int, dim> tmin{bx.smallEnd(0), bx.smallEnd(1),
@@ -233,6 +240,11 @@ reduction<CCTK_REAL, dim> reduce(int gi, int vi, int tl) {
         red += reduce_array(vars, vi, tmin, tmax, indextype, imin, imax,
                             finemask.get(), x0, dx);
       }
+#ifdef __NVCOMPILER
+#pragma omp critical
+      outer += red;
+      }
+#endif
     }
   }
 

diff --git a/CarpetX/src/reduction.hxx b/CarpetX/src/reduction.hxx
@@ -121,7 +121,9 @@ template <typename T, int D> struct combine {
 };
 
 typedef reduction<CCTK_REAL, dim> reduction_CCTK_REAL;
+#ifndef __NVCOMPILER
 #pragma omp declare reduction(reduction:reduction_CCTK_REAL : omp_out += omp_in)
+#endif
 
 MPI_Datatype reduction_mpi_datatype_CCTK_REAL();
 MPI_Op reduction_mpi_op();