Using sycl::atomic_ref for vecmem::atomic in case it is available.

When sycl::atomic_ref is available, vecmem::atomic is set up simply as
a thin layer above that class.

Since oneAPI only added support for sycl::atomic_ref relatively recently,
the CMake configuration needs to check whether that type is available
with the particular SYCL compiler that we use, or not.

core/CMakeLists.txt

@@ -168,7 +168,7 @@ if(NOT VECMEM_HAVE_DEFAULT_RESOURCE)
       FILES "src/memory/default_resource_polyfill.cpp" )
 endif()
 
-# Figure out how to produce SYCL debug printouts.
+# Figure out how to use various SYCL features.
 if( VECMEM_BUILD_SYCL_LIBRARY )
 
    # Test which printf function(s) is/are available.
@@ -190,4 +190,12 @@ if( VECMEM_BUILD_SYCL_LIBRARY )
       target_compile_definitions( vecmem_core PUBLIC
          VECMEM_ONEAPI_PRINTF_FUNCTION=printf )
    endif()
+
+   # Test whether sycl::atomic_ref is available.
+   vecmem_check_sycl_code_compiles( VECMEM_HAVE_SYCL_ATOMIC_REF
+      "${CMAKE_CURRENT_SOURCE_DIR}/cmake/atomic_ref_test.sycl" )
+   if( VECMEM_HAVE_SYCL_ATOMIC_REF )
+      target_compile_definitions( vecmem_core PUBLIC
+         VECMEM_HAVE_SYCL_ATOMIC_REF )
+   endif()
 endif()

core/cmake/atomic_ref_test.sycl

@@ -0,0 +1,21 @@
+/** VecMem project, part of the ACTS project (R&D line)
+ *
+ * (c) 2022 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// SYCL include(s).
+#include <CL/sycl.hpp>
+
+int main() {
+    // Try to use sycl::atomic_ref.
+    int dummy = 0;
+    sycl::atomic_ref<int, sycl::memory_order::relaxed,
+                     sycl::memory_scope::device,
+                     sycl::access::address_space::global_space>
+        atomic_dummy(dummy);
+    atomic_dummy.store(3);
+    atomic_dummy.fetch_add(1);
+    return 0;
+}

core/include/vecmem/memory/atomic.hpp

@@ -1,7 +1,7 @@
 /*
  * VecMem project, part of the ACTS project (R&D line)
  *
- * (c) 2021 CERN for the benefit of the ACTS project
+ * (c) 2021-2022 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
@@ -13,6 +13,53 @@
 // System include(s).
 #include <type_traits>
 
+// Provide a different implementation for modern SYCL and everything else
+#if (defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)) && \
+    defined(VECMEM_HAVE_SYCL_ATOMIC_REF)
+
+namespace vecmem {
+
+/// In modern SYCL code @c vecmem::atomic is an alias for @c sycl::atomic_ref
+///
+/// It has to be an actual class and not just a typedef, because
+/// @c sycl::atomic_ref, as the name implies, is created on top of references,
+/// and not on top of pointers.
+///
+template <typename T>
+class atomic
+    : public ::sycl::atomic_ref<T, ::sycl::memory_order::relaxed,
+                                ::sycl::memory_scope::device,
+                                ::sycl::access::address_space::global_space> {
+
+public:
+    /// @name Type definitions
+    /// @{
+
+    /// Type managed by the object
+    typedef T value_type;
+    /// Difference between two objects
+    typedef value_type difference_type;
+    /// Pointer to the value in global memory
+    typedef value_type* pointer;
+    /// Reference to a value given by the user
+    typedef value_type& reference;
+
+    /// @}
+
+    /// Constructor, with a pointer to the managed variable
+    VECMEM_HOST_AND_DEVICE
+    atomic(pointer ptr)
+        : ::sycl::atomic_ref<T, ::sycl::memory_order::relaxed,
+                             ::sycl::memory_scope::device,
+                             ::sycl::access::address_space::global_space>(
+              *ptr) {}
+
+};  // class atomic
+
+}  // namespace vecmem
+
+#else
+
 namespace vecmem {
 
 /// Class providing atomic operations for the VecMem code
@@ -107,3 +154,5 @@ class atomic {
 
 // Include the implementation.
 #include "vecmem/memory/impl/atomic.ipp"
+
+#endif  // sycl::atomic_ref