From 612dd81b57e35a659a657ccb8570a617b826466a Mon Sep 17 00:00:00 2001
From: Roland Haas <rhaas@illinois.edu>
Date: Fri, 10 Jan 2025 10:27:54 -0600
Subject: [PATCH] CarpetX: manually implement custom OpenMP reductions for
 nvc++

---
 CarpetX/src/reduction.cxx | 12 ++++++++++++
 CarpetX/src/reduction.hxx |  2 ++
 2 files changed, 14 insertions(+)

diff --git a/CarpetX/src/reduction.cxx b/CarpetX/src/reduction.cxx
index f75684b35..3b7d39122 100644
--- 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
index e1b03121c..b6180afb1 100644
--- 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();