raja-adaptivecpp.patch

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 0aa305e9e..c68a90c96 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -231,6 +231,16 @@ endif ()
 if (RAJA_ENABLE_SYCL)
   set (RAJA_ENABLE_DESUL_ATOMICS "On")
   set (ENABLE_SYCL "On") # Enable SYCL atomics in Desul
+
+  # BLT has no SYCL setup, so "sycl" is not a CMake target: with the Intel
+  # oneAPI toolchain it falls through to the linker and resolves as -lsycl.
+  # AdaptiveCpp has no such library -- acpp supplies the SYCL runtime itself --
+  # so stand in an empty interface target to keep the dependency name valid.
+  # RAJAMacros.cmake adds "sycl" to every test/example it builds as well.
+  if (RAJA_USING_ADAPTIVECPP AND NOT TARGET sycl)
+    add_library (sycl INTERFACE IMPORTED GLOBAL)
+  endif ()
+
   set (raja_depends
     ${raja_depends}
     sycl)
@@ -243,6 +253,12 @@ if (RAJA_ENABLE_DESUL_ATOMICS)
   if (NOT RAJA_ENABLE_TARGET_OPENMP)
     set (DESUL_ENABLE_OPENMP Off CACHE BOOL "Disable Desul OpenMP offload atomics to select builtin host atomics")
   endif()
+  # Desul's SYCL atomics are written against oneAPI-only extensions
+  # (sycl::ext::oneapi this_sub_group/group_ballot/sub_group_mask), which
+  # AdaptiveCpp does not provide. Fall back to the builtin host atomics.
+  if (RAJA_USING_ADAPTIVECPP)
+    set (DESUL_ENABLE_SYCL Off CACHE BOOL "Disable Desul SYCL atomics; they require oneAPI extensions AdaptiveCpp lacks")
+  endif()
   add_subdirectory(tpl/desul)
   set(raja_depends
     ${raja_depends}
diff --git a/cmake/SetupCompilers.cmake b/cmake/SetupCompilers.cmake
index bc9367cfc..b3e376a8d 100644
--- a/cmake/SetupCompilers.cmake
+++ b/cmake/SetupCompilers.cmake
@@ -84,6 +84,33 @@ if (RAJA_ENABLE_HIP)
 endif()
 # end RAJA_ENABLE_HIP section
 
+if (RAJA_ENABLE_SYCL)
+  # AdaptiveCpp (acpp) does not predefine SYCL_LANGUAGE_VERSION on the command
+  # line; it is only defined once <sycl/sycl.hpp> has been included. RAJA and
+  # camp both key their SYCL support off that macro at configuration points
+  # that are reached before any SYCL header is pulled in, so it has to be
+  # provided up front. RAJA_CXX_FLAGS lets a host-config/command line supply
+  # that (and any other compiler flag) for the whole build.
+  get_filename_component(RAJA_SYCL_COMPILER_NAME "${CMAKE_CXX_COMPILER}" NAME)
+  if (RAJA_SYCL_COMPILER_NAME MATCHES "acpp|syclcc|hipsycl")
+    set(RAJA_USING_ADAPTIVECPP On)
+    if (ACPP_TARGETS)
+      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --acpp-targets=${ACPP_TARGETS}")
+    endif()
+    if (NOT RAJA_CXX_FLAGS MATCHES "SYCL_LANGUAGE_VERSION")
+      set(RAJA_CXX_FLAGS "${RAJA_CXX_FLAGS} -DSYCL_LANGUAGE_VERSION=202012")
+    endif()
+  endif()
+endif()
+# end RAJA_ENABLE_SYCL section
+
+# Extra user-supplied C++ flags, applied to RAJA and everything it builds
+# (tests, tpls) after the compiler defaults above.
+if (RAJA_CXX_FLAGS)
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${RAJA_CXX_FLAGS}")
+  message(STATUS "Appending RAJA_CXX_FLAGS to CMAKE_CXX_FLAGS: ${RAJA_CXX_FLAGS}")
+endif()
+
 set(RAJA_RANGE_ALIGN 4 CACHE STRING "")
 set(RAJA_RANGE_MIN_LENGTH 32 CACHE STRING "")
 set(RAJA_DATA_ALIGN 64 CACHE STRING "")
diff --git a/include/RAJA/pattern/launch/launch_core.hpp b/include/RAJA/pattern/launch/launch_core.hpp
index 574440eb5..11cf0d2f1 100644
--- a/include/RAJA/pattern/launch/launch_core.hpp
+++ b/include/RAJA/pattern/launch/launch_core.hpp
@@ -277,6 +277,23 @@ public:
     // SGS ODR Issue
 #if defined(RAJA_GPU_DEVICE_COMPILE_PASS_ACTIVE) && defined(RAJA_SYCL_ACTIVE)
     itm->barrier(::sycl::access::fence_space::local_space);
+#elif defined(RAJA_SYCL_ACTIVE) && defined(__ACPP__) &&                        \
+    (defined(__ACPP_ENABLE_LLVM_SSCP_TARGET__) ||                             \
+     defined(__ACPP_USE_ACCELERATED_CPU__))
+    // AdaptiveCpp compiles kernel and host code in a single pass and never
+    // defines __SYCL_DEVICE_ONLY__, so the guard above is always false and the
+    // barrier silently disappears -- leaving shared memory races (see the
+    // launch shared-memory tests). itm is only non-null inside a SYCL kernel,
+    // so branch on that instead.
+    //
+    // Restricted to the targets that can actually execute a work-group
+    // barrier. The library-only CPU target runs work items sequentially and
+    // has no loop splitting, so a barrier there deadlocks rather than
+    // synchronising; kernels needing teamSync are unsupported on it.
+    if (itm != nullptr)
+    {
+      itm->barrier(::sycl::access::fence_space::local_space);
+    }
 #endif
 
 #if defined(RAJA_GPU_DEVICE_COMPILE_PASS_ACTIVE) && !defined(RAJA_SYCL_ACTIVE)
diff --git a/include/RAJA/policy/sycl/params/kernel_name.hpp b/include/RAJA/policy/sycl/params/kernel_name.hpp
index e4faf54ff..6d2bf55da 100644
--- a/include/RAJA/policy/sycl/params/kernel_name.hpp
+++ b/include/RAJA/policy/sycl/params/kernel_name.hpp
@@ -12,6 +12,9 @@
 
 #include "RAJA/pattern/params/kernel_name.hpp"
 
+// Provides RAJA_SYCL_EXTERNAL.
+#include "RAJA/util/sycl_compat.hpp"
+
 namespace RAJA
 {
 namespace expt
@@ -31,8 +34,9 @@ param_init(EXEC_POL const&, RAJA::detail::Name&)
 
 // Combine
 template<typename EXEC_POL, typename T>
-camp::concepts::enable_if<RAJA::type_traits::is_sycl_policy<EXEC_POL>>
-    SYCL_EXTERNAL param_combine(EXEC_POL const&, RAJA::detail::Name&, T)
+RAJA_SYCL_EXTERNAL camp::concepts::enable_if<
+    RAJA::type_traits::is_sycl_policy<EXEC_POL>>
+param_combine(EXEC_POL const&, RAJA::detail::Name&, T)
 {}
 
 // Resolve
diff --git a/include/RAJA/policy/sycl/reduce.hpp b/include/RAJA/policy/sycl/reduce.hpp
index c30e65e0a..7673c2f68 100644
--- a/include/RAJA/policy/sycl/reduce.hpp
+++ b/include/RAJA/policy/sycl/reduce.hpp
@@ -25,6 +25,7 @@
 #if defined(RAJA_ENABLE_SYCL)
 
 #include <algorithm>
+#include <mutex>
 #include <stdexcept>
 #include <string>
 
@@ -247,6 +248,12 @@ struct Shared_Host_Data
   T hostVal;
   T identity;
   const void* rootToken;
+  // Guards hostVal on the host reduction path. That path is not only taken by
+  // serial host code: on a CPU SYCL backend that runs kernels as host code
+  // (e.g. AdaptiveCpp's library-only OpenMP target) __SYCL_DEVICE_ONLY__ is
+  // never defined, so every work item combines through here concurrently.
+  // Declared last so the aggregate initializers above stay valid.
+  std::mutex hostMutex;
 };
 
 template<typename T, typename IndexType>
@@ -257,6 +264,9 @@ struct Shared_Host_Loc_Data
   T identity;
   IndexType identityLoc;
   const void* rootToken;
+  // See Shared_Host_Data::hostMutex. The value and its location have to be
+  // updated together, so this pair needs a lock rather than an atomic.
+  std::mutex hostMutex;
 };
 
 }  // end namespace sycl
@@ -376,6 +386,7 @@ struct TargetReduce
         val.device[i]);
     Reducer {}(atm, rhsVal);
 #else
+    const std::lock_guard<std::mutex> guard(hostData->hostMutex);
     Reducer {}(hostData->hostVal, rhsVal);
 #endif
     return *this;
@@ -576,6 +587,7 @@ public:
     ::sycl::atomic_fence(::sycl::memory_order_release,
                          ::sycl::memory_scope::device);
 #else
+    const std::lock_guard<std::mutex> guard(hostData->hostMutex);
     Reducer {}(hostData->hostVal, hostData->hostLoc, rhsVal, rhsLoc);
 #endif
     return *this;
diff --git a/include/RAJA/util/sycl_compat.hpp b/include/RAJA/util/sycl_compat.hpp
index f54fd5b03..9bcbd6814 100644
--- a/include/RAJA/util/sycl_compat.hpp
+++ b/include/RAJA/util/sycl_compat.hpp
@@ -30,4 +30,16 @@
 #endif
 #endif
 
+//
+// SYCL_EXTERNAL is not universally available: AdaptiveCpp deliberately leaves
+// it undefined on its accelerated-CPU compilation path
+// (__ACPP_USE_ACCELERATED_CPU__), which is what an -O2/-O3 build for a host
+// target uses. Marking a function is a no-op there, so fall back to nothing.
+//
+#if defined(SYCL_EXTERNAL)
+#define RAJA_SYCL_EXTERNAL SYCL_EXTERNAL
+#else
+#define RAJA_SYCL_EXTERNAL
+#endif
+
 #endif  // RAJA_util_sycl_compat_HPP