Merge pull request #169 from krasznaa/AtomicRef-main-20220307

Add vecmem::device_atomic_ref, main branch (2022.03.07.)

core/CMakeLists.txt

@@ -52,6 +52,8 @@ vecmem_add_library( vecmem_core core
    "src/memory/details/memory_resource_base.cpp"
    "include/vecmem/memory/atomic.hpp"
    "include/vecmem/memory/impl/atomic.ipp"
+   "include/vecmem/memory/device_atomic_ref.hpp"
+   "include/vecmem/memory/impl/device_atomic_ref.ipp"
    "include/vecmem/memory/polymorphic_allocator.hpp"
    "include/vecmem/memory/memory_resource.hpp"
    "src/memory/alignment.hpp"

core/include/vecmem/containers/device_vector.hpp

@@ -1,6 +1,6 @@
 /** 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
  */
@@ -9,7 +9,7 @@
 // Local include(s).
 #include "vecmem/containers/data/vector_view.hpp"
 #include "vecmem/containers/details/reverse_iterator.hpp"
-#include "vecmem/memory/atomic.hpp"
+#include "vecmem/memory/device_atomic_ref.hpp"
 #include "vecmem/utils/type_traits.hpp"
 #include "vecmem/utils/types.hpp"
 
@@ -145,7 +145,7 @@ class device_vector {
         // Create copies of all of the elements one-by-one. It's very
         // inefficient, but we can't make any assumptions about the type of the
         // input iterator received by this function.
-        atomic<size_type> asize(m_size);
+        device_atomic_ref<size_type> asize(*m_size);
         for (InputIt itr = other_begin; itr != other_end; ++itr) {
             construct(asize.fetch_add(1), *itr);
         }

core/include/vecmem/containers/impl/device_vector.ipp

@@ -1,6 +1,6 @@
 /** 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
  */
@@ -175,7 +175,7 @@ VECMEM_HOST_AND_DEVICE void device_vector<TYPE>::assign(size_type count,
     // Remove all previous elements.
     clear();
     // Set the assigned size of the vector.
-    atomic<size_type> asize(m_size);
+    device_atomic_ref<size_type> asize(*m_size);
     asize.store(count);
 
     // Create the required number of identical elements.
@@ -194,7 +194,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::emplace_back(Args&&... args)
 
     // Increment the size of the vector at first. So that we would "claim" the
     // index from other threads.
-    atomic<size_type> asize(m_size);
+    device_atomic_ref<size_type> asize(*m_size);
     const size_type index = asize.fetch_add(1);
     assert(index < m_capacity);
 
@@ -214,7 +214,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::push_back(
 
     // Increment the size of the vector at first. So that we would "claim" the
     // index from other threads.
-    atomic<size_type> asize(m_size);
+    device_atomic_ref<size_type> asize(*m_size);
     const size_type index = asize.fetch_add(1);
     assert(index < m_capacity);
 
@@ -232,7 +232,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::pop_back() -> size_type {
     assert(m_size != nullptr);
 
     // Decrement the size of the vector, and remember this new size.
-    atomic<size_type> asize(m_size);
+    device_atomic_ref<size_type> asize(*m_size);
     const size_type new_size = asize.fetch_sub(1) - 1;
 
     // Remove the last element.
@@ -249,7 +249,7 @@ VECMEM_HOST_AND_DEVICE void device_vector<TYPE>::clear() {
     assert(m_size != nullptr);
 
     // Destruct all of the elements that the vector has "at the moment".
-    atomic<size_type> asize(m_size);
+    device_atomic_ref<size_type> asize(*m_size);
     const size_type current_size = asize.load();
     for (size_type i = 0; i < current_size; ++i) {
         destruct(i);
@@ -273,7 +273,7 @@ VECMEM_HOST_AND_DEVICE void device_vector<TYPE>::resize(size_type new_size,
     assert(m_size != nullptr);
 
     // Get the current size of the vector.
-    atomic<size_type> asize(m_size);
+    device_atomic_ref<size_type> asize(*m_size);
     const size_type current_size = asize.load();
 
     // Check if anything needs to be done.
@@ -395,7 +395,7 @@ VECMEM_HOST_AND_DEVICE auto device_vector<TYPE>::size() const -> size_type {
         // SYCL we must pass a non-const pointer to the sycl::atomic object
         // that performs the load operation. And for that we need a non-const
         // pointer...
-        atomic<size_type> asize(const_cast<size_type*>(m_size));
+        device_atomic_ref<size_type> asize(*(const_cast<size_type*>(m_size)));
         return asize.load();
     }
 }

core/include/vecmem/memory/device_atomic_ref.hpp

@@ -0,0 +1,187 @@
+/*
+ * 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
+ */
+#pragma once
+
+// System include(s).
+#include <atomic>
+
+#if (defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)) && \
+    defined(VECMEM_HAVE_SYCL_ATOMIC_REF)
+
+// SYCL include(s).
+#include <CL/sycl.hpp>
+
+namespace vecmem {
+
+/// Define @c vecmem::memory_order as @c sycl::memory_order
+using memory_order = ::sycl::memory_order;
+
+/// @c vecmem::atomic_ref equals @c sycl::atomic_ref with "modern SYCL"
+template <typename T>
+using device_atomic_ref =
+    ::sycl::atomic_ref<T, ::sycl::memory_order::relaxed,
+                       ::sycl::memory_scope::device,
+                       ::sycl::access::address_space::global_space>;
+
+}  // namespace vecmem
+
+#elif ((!defined(__CUDA_ARCH__)) && (!defined(__HIP_DEVICE_COMPILE__)) && \
+       (!defined(CL_SYCL_LANGUAGE_VERSION)) &&                            \
+       (!defined(SYCL_LANGUAGE_VERSION)) && __cpp_lib_atomic_ref)
+
+namespace vecmem {
+
+/// Define @c vecmem::memory_order as @c std::memory_order
+using memory_order = std::memory_order;
+
+/// @c vecmem::atomic_ref equals @c std::atomic_ref in host code with C++20
+template <typename T>
+using device_atomic_ref = std::atomic_ref<T>;
+
+}  // namespace vecmem
+
+#else
+
+// VecMem include(s).
+#include "vecmem/utils/types.hpp"
+
+// System include(s).
+#include <type_traits>
+
+namespace vecmem {
+
+/// Custom (dummy) definition for the memory order
+enum class memory_order {
+    relaxed = 0,
+    consume = 1,
+    acquire = 2,
+    release = 3,
+    acq_rel = 4,
+    seq_cst = 5
+};
+
+/// Class providing atomic operations for the VecMem code
+///
+/// It is only meant to be used with primitive types. Ones that CUDA, HIP and
+/// SYCL built-in functions exist for. So no structs, or even pointers.
+///
+/// Note that it does not perform atomic operations in host code! That is only
+/// implemented with @c std::atomic_ref in C++20. With earlier C++ standards all
+/// operations in host code are performed as "regular" operations.
+///
+template <typename T>
+class device_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;
+
+    /// @}
+
+    /// @name Check(s) on the value type
+    /// @{
+
+    static_assert(std::is_integral<value_type>::value,
+                  "vecmem::atomic_ref only accepts built-in integral types");
+
+    /// @}
+
+    /// Constructor, with a pointer to the managed variable
+    VECMEM_HOST_AND_DEVICE
+    device_atomic_ref(reference ref);
+    /// Copy constructor
+    VECMEM_HOST_AND_DEVICE
+    device_atomic_ref(const device_atomic_ref& parent);
+
+    /// Disable the assignment operator
+    device_atomic_ref& operator=(const device_atomic_ref&) = delete;
+
+    /// @name Value setter/getter functions
+    /// @{
+
+    /// Assigns a value desired to the referenced object
+    ///
+    /// @see vecmem::device_atomic_ref::store
+    ///
+    VECMEM_HOST_AND_DEVICE
+    value_type operator=(value_type data) const;
+
+    /// Set the variable to the desired value
+    VECMEM_HOST_AND_DEVICE
+    void store(value_type data,
+               memory_order order = memory_order::seq_cst) const;
+    /// Get the value of the variable
+    VECMEM_HOST_AND_DEVICE
+    value_type load(memory_order order = memory_order::seq_cst) const;
+
+    /// Exchange the current value of the variable with a different one
+    VECMEM_HOST_AND_DEVICE
+    value_type exchange(value_type data,
+                        memory_order order = memory_order::seq_cst) const;
+
+    /// Compare against the current value, and exchange only if different
+    VECMEM_HOST_AND_DEVICE
+    bool compare_exchange_strong(reference expected, value_type desired,
+                                 memory_order success,
+                                 memory_order failure) const;
+    /// Compare against the current value, and exchange only if different
+    VECMEM_HOST_AND_DEVICE
+    bool compare_exchange_strong(
+        reference expected, value_type desired,
+        memory_order order = memory_order::seq_cst) const;
+
+    /// @}
+
+    /// @name Value modifier functions
+    /// @{
+
+    /// Add a chosen value to the stored variable
+    VECMEM_HOST_AND_DEVICE
+    value_type fetch_add(value_type data,
+                         memory_order order = memory_order::seq_cst) const;
+    /// Substitute a chosen value from the stored variable
+    VECMEM_HOST_AND_DEVICE
+    value_type fetch_sub(value_type data,
+                         memory_order order = memory_order::seq_cst) const;
+
+    /// Replace the current value with the specified value AND-ed to it
+    VECMEM_HOST_AND_DEVICE
+    value_type fetch_and(value_type data,
+                         memory_order order = memory_order::seq_cst) const;
+    /// Replace the current value with the specified value OR-d to it
+    VECMEM_HOST_AND_DEVICE
+    value_type fetch_or(value_type data,
+                        memory_order order = memory_order::seq_cst) const;
+    /// Replace the current value with the specified value XOR-d to it
+    VECMEM_HOST_AND_DEVICE
+    value_type fetch_xor(value_type data,
+                         memory_order order = memory_order::seq_cst) const;
+
+    /// @}
+
+private:
+    /// Pointer to the value to perform atomic operations on
+    pointer m_ptr;
+
+};  // class device_atomic_ref
+
+}  // namespace vecmem
+
+// Include the implementation.
+#include "vecmem/memory/impl/device_atomic_ref.ipp"
+
+#endif  // Platform selection