atomic.hpp

/*
 * VecMem project, part of the ACTS project (R&D line)
 *
 * (c) 2021-2022 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 atomic
    : public ::sycl::atomic_ref<T, ::sycl::memory_order::relaxed,
                                ::sycl::memory_scope::device,
                                ::sycl::access::address_space::global_space> {
 
public:
 
    typedef T value_type;
    typedef value_type difference_type;
    typedef value_type* pointer;
    typedef value_type& reference;
 
 
    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 atomic {
 
public:
 
    typedef T value_type;
    typedef value_type difference_type;
    typedef value_type* pointer;
    typedef value_type& reference;
 
 
 
    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