atomic.hpp

/*
 * VecMem project, part of the ACTS project (R&D line)
 *
 * (c) 2021-2026 CERN for the benefit of the ACTS project
 *
 * Mozilla Public License Version 2.0
 */
#pragma once
 
// VecMem include(s).
#include "vecmem/utils/types.hpp"
 
// 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 {
 
template <typename T>
class [
    [deprecated("vecmem::atomic is deprecated, use vecmem::device_atomic_ref "
                "instead")]] atomic
    : public ::sycl::atomic_ref<T, ::sycl::memory_order::relaxed,
                                ::sycl::memory_scope::device,
                                ::sycl::access::address_space::global_space> {
 
public:
 
    using value_type = T;
    using difference_type = value_type;
    using pointer = std::add_pointer_t<value_type>;
    using volatile_pointer =
        std::add_pointer_t<std::add_volatile_t<value_type>>;
    using reference = std::add_lvalue_reference_t<value_type>;
 
 
    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 {
 
template <typename T>
class [
    [deprecated("vecmem::atomic is deprecated, use vecmem::device_atomic_ref "
                "instead")]] atomic {
 
public:
 
    using value_type = T;
    using difference_type = value_type;
    using pointer = std::add_pointer_t<value_type>;
    using volatile_pointer =
        std::add_pointer_t<std::add_volatile_t<value_type>>;
    using reference = std::add_lvalue_reference_t<value_type>;
 
 
 
    static_assert(std::is_integral<value_type>::value,
                  "vecmem::atomic only accepts built-in integral types");
 
 
    VECMEM_HOST_AND_DEVICE
    atomic(pointer ptr);
 
 
    VECMEM_HOST_AND_DEVICE
    void store(value_type data);
    VECMEM_HOST_AND_DEVICE
    value_type load() const;
 
    VECMEM_HOST_AND_DEVICE
    value_type exchange(value_type data);
 
    VECMEM_HOST_AND_DEVICE
    bool compare_exchange_strong(reference expected, value_type desired);
 
 
 
    VECMEM_HOST_AND_DEVICE
    value_type fetch_add(value_type data);
    VECMEM_HOST_AND_DEVICE
    value_type fetch_sub(value_type data);
 
    VECMEM_HOST_AND_DEVICE
    value_type fetch_and(value_type data);
    VECMEM_HOST_AND_DEVICE
    value_type fetch_or(value_type data);
    VECMEM_HOST_AND_DEVICE
    value_type fetch_xor(value_type data);
 
 
private:
    pointer m_ptr;
 
};  // class atomic
 
}  // namespace vecmem
 
// Include the implementation.
#include "vecmem/memory/impl/atomic.ipp"
 
#endif  // sycl::atomic_ref