sycl_custom_device_atomic_ref.hpp

/*
 * VecMem project, part of the ACTS project (R&D line)
 *
 * (c) 2022-2024 CERN for the benefit of the ACTS project
 *
 * Mozilla Public License Version 2.0
 */
#pragma once
 
// Local include(s).
#include "vecmem/memory/device_address_space.hpp"
#include "vecmem/memory/memory_order.hpp"
 
// System include(s).
#include <type_traits>
 
namespace vecmem {
namespace sycl {
 
template <typename T,
          device_address_space address = device_address_space::global>
class custom_device_atomic_ref {
 
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::sycl::custom_device_atomic_ref only accepts built-in "
        "integral types");
 
 
    explicit custom_device_atomic_ref(reference ref);
    custom_device_atomic_ref(const custom_device_atomic_ref& parent);
 
    custom_device_atomic_ref& operator=(const custom_device_atomic_ref&) =
        delete;
 
 
    value_type operator=(value_type data) const;
 
    void store(value_type data,
               memory_order order = memory_order::seq_cst) const;
    value_type load(memory_order order = memory_order::seq_cst) const;
 
    value_type exchange(value_type data,
                        memory_order order = memory_order::seq_cst) const;
 
    bool compare_exchange_strong(reference expected, value_type desired,
                                 memory_order success,
                                 memory_order failure) const;
    bool compare_exchange_strong(
        reference expected, value_type desired,
        memory_order order = memory_order::seq_cst) const;
 
 
 
    value_type fetch_add(value_type data,
                         memory_order order = memory_order::seq_cst) const;
    value_type fetch_sub(value_type data,
                         memory_order order = memory_order::seq_cst) const;
 
    value_type fetch_and(value_type data,
                         memory_order order = memory_order::seq_cst) const;
    value_type fetch_or(value_type data,
                        memory_order order = memory_order::seq_cst) const;
    value_type fetch_xor(value_type data,
                         memory_order order = memory_order::seq_cst) const;
 
 
private:
    pointer m_ptr;
 
};  // class custom_device_atomic_ref
 
}  // namespace sycl
}  // namespace vecmem
 
// Include the implementation.
#include "vecmem/memory/impl/sycl_custom_device_atomic_ref.ipp"