From 5117d21af7902a8032a004f09ca495b5be4fd752 Mon Sep 17 00:00:00 2001
From: Weiqun Zhang <WeiqunZhang@lbl.gov>
Date: Wed, 12 Jul 2023 15:50:28 -0700
Subject: [PATCH] Use atomic load in SYCL version of the single-pass scan
 (#3419)

Otherwise the single-pass scan of long ints will hang because it does
not load the most recent data even though the member function is
volatile.
---
 Src/Base/AMReX_Scan.H         | 15 +++++++++++++++
 Tools/GNUMake/comps/dpcpp.mak |  3 +++
 2 files changed, 18 insertions(+)

diff --git a/Src/Base/AMReX_Scan.H b/Src/Base/AMReX_Scan.H
index 1d39adaee65..bee1d910ec9 100644
--- a/Src/Base/AMReX_Scan.H
+++ b/Src/Base/AMReX_Scan.H
@@ -139,12 +139,27 @@ struct BlockStatus<T, false>
 
     AMREX_GPU_DEVICE AMREX_FORCE_INLINE
     STVA<T> read () volatile {
+#if defined(AMREX_USE_SYCL)
+        constexpr auto mo = sycl::memory_order::relaxed;
+        constexpr auto ms = sycl::memory_scope::device;
+        constexpr auto as = sycl::access::address_space::global_space;
+#endif
         if (status == 'x') {
             return {'x', 0};
         } else if (status == 'a') {
+#if defined(AMREX_USE_SYCL)
+            sycl::atomic_ref<T,mo,ms,as> ar{const_cast<T&>(aggregate)};
+            return {'a', ar.load()};
+#else
             return {'a', aggregate};
+#endif
         } else {
+#if defined(AMREX_USE_SYCL)
+            sycl::atomic_ref<T,mo,ms,as> ar{const_cast<T&>(inclusive)};
+            return {'p', ar.load()};
+#else
             return {'p', inclusive};
+#endif
         }
     }
 
diff --git a/Tools/GNUMake/comps/dpcpp.mak b/Tools/GNUMake/comps/dpcpp.mak
index b82546a9f70..3ae78e39433 100644
--- a/Tools/GNUMake/comps/dpcpp.mak
+++ b/Tools/GNUMake/comps/dpcpp.mak
@@ -6,6 +6,9 @@ CC  = icx
 FC  = ifx
 F90 = ifx
 
+amrex_oneapi_version = $(shell $(CXX) --version | head -1)
+$(info oneAPI version: $(amrex_oneapi_version))
+
 CXXFLAGS =
 CFLAGS   =
 FFLAGS   =