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

@@ -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
         }
     }
 

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   =