diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index b31392b0..1d5e482a 100644 --- a/core/CMakeLists.txt +++ b/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() diff --git a/core/cmake/atomic_ref_test.sycl b/core/cmake/atomic_ref_test.sycl new file mode 100644 index 00000000..20675263 --- /dev/null +++ b/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 + +int main() { + // Try to use sycl::atomic_ref. + int dummy = 0; + sycl::atomic_ref + atomic_dummy(dummy); + atomic_dummy.store(3); + atomic_dummy.fetch_add(1); + return 0; +} diff --git a/core/include/vecmem/memory/atomic.hpp b/core/include/vecmem/memory/atomic.hpp index 8548f32b..002a10ee 100644 --- a/core/include/vecmem/memory/atomic.hpp +++ b/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 +// 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 +class atomic + : public ::sycl::atomic_ref { + +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( + *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