vecmem 1.30.0
Loading...
Searching...
No Matches
atomic.ipp
1/*
2 * VecMem project, part of the ACTS project (R&D line)
3 *
4 * (c) 2021-2024 CERN for the benefit of the ACTS project
5 *
6 * Mozilla Public License Version 2.0
7 */
8#pragma once
9
10// HIP include(s).
11#if defined(__HIP_DEVICE_COMPILE__)
12#include <hip/hip_runtime.h>
13#endif
14
15// SYCL include(s).
16#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
17#include <sycl/sycl.hpp>
18#endif
19
21#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
22#define __VECMEM_SYCL_ATOMIC_CALL0(FNAME, PTR) \
23 ::sycl::atomic_##FNAME<value_type>( \
24 ::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)))
25#define __VECMEM_SYCL_ATOMIC_CALL1(FNAME, PTR, ARG1) \
26 ::sycl::atomic_##FNAME<value_type>( \
27 ::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)), ARG1)
28#define __VECMEM_SYCL_ATOMIC_CALL2(FNAME, PTR, ARG1, ARG2) \
29 ::sycl::atomic_##FNAME<value_type>( \
30 ::sycl::atomic<value_type>(::sycl::global_ptr<value_type>(PTR)), ARG1, \
31 ARG2)
32#endif
33
34namespace vecmem {
35
36template <typename T>
37VECMEM_HOST_AND_DEVICE atomic<T>::atomic(pointer ptr) : m_ptr(ptr) {}
38
39template <typename T>
40VECMEM_HOST_AND_DEVICE void atomic<T>::store(value_type data) {
41
42#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
43 (!defined(SYCL_LANGUAGE_VERSION))
44 volatile_pointer addr = m_ptr;
45 __threadfence();
46 *addr = data;
47#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
48 __VECMEM_SYCL_ATOMIC_CALL1(store, m_ptr, data);
49#else
50 *m_ptr = data;
51#endif
52}
53
54template <typename T>
55VECMEM_HOST_AND_DEVICE auto atomic<T>::load() const -> value_type {
56
57#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
58 (!defined(SYCL_LANGUAGE_VERSION))
59 volatile_pointer addr = m_ptr;
60 __threadfence();
61 const value_type value = *addr;
62 __threadfence();
63 return value;
64#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
65 return __VECMEM_SYCL_ATOMIC_CALL0(load, m_ptr);
66#else
67 return *m_ptr;
68#endif
69}
70
71template <typename T>
72VECMEM_HOST_AND_DEVICE auto atomic<T>::exchange(value_type data) -> value_type {
73
74#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
75 (!defined(SYCL_LANGUAGE_VERSION))
76 return atomicExch(m_ptr, data);
77#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
78 return __VECMEM_SYCL_ATOMIC_CALL1(exchange, m_ptr, data);
79#else
80 value_type current_value = *m_ptr;
81 *m_ptr = data;
82 return current_value;
83#endif
84}
85
86template <typename T>
87VECMEM_HOST_AND_DEVICE bool atomic<T>::compare_exchange_strong(
88 reference expected, value_type desired) {
89
90#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
91 (!defined(SYCL_LANGUAGE_VERSION))
92 return atomicCAS(m_ptr, expected, desired);
93#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
94 return __VECMEM_SYCL_ATOMIC_CALL2(compare_exchange_strong, m_ptr, expected,
95 desired);
96#else
97 if (*m_ptr == expected) {
98 *m_ptr = desired;
99 return true;
100 } else {
101 expected = *m_ptr;
102 return false;
103 }
104#endif
105}
106
107template <typename T>
108VECMEM_HOST_AND_DEVICE auto atomic<T>::fetch_add(value_type data)
109 -> value_type {
110
111#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
112 (!defined(SYCL_LANGUAGE_VERSION))
113 return atomicAdd(m_ptr, data);
114#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
115 return __VECMEM_SYCL_ATOMIC_CALL1(fetch_add, m_ptr, data);
116#else
117 const value_type result = *m_ptr;
118 *m_ptr += data;
119 return result;
120#endif
121}
122
123template <typename T>
124VECMEM_HOST_AND_DEVICE auto atomic<T>::fetch_sub(value_type data)
125 -> value_type {
126
127#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
128 (!defined(SYCL_LANGUAGE_VERSION))
129 return atomicSub(m_ptr, data);
130#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
131 return __VECMEM_SYCL_ATOMIC_CALL1(fetch_sub, m_ptr, data);
132#else
133 const value_type result = *m_ptr;
134 *m_ptr -= data;
135 return result;
136#endif
137}
138
139template <typename T>
140VECMEM_HOST_AND_DEVICE auto atomic<T>::fetch_and(value_type data)
141 -> value_type {
142
143#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
144 (!defined(SYCL_LANGUAGE_VERSION))
145 return atomicAnd(m_ptr, data);
146#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
147 return __VECMEM_SYCL_ATOMIC_CALL1(fetch_and, m_ptr, data);
148#else
149 const value_type result = *m_ptr;
150 *m_ptr &= data;
151 return result;
152#endif
153}
154
155template <typename T>
156VECMEM_HOST_AND_DEVICE auto atomic<T>::fetch_or(value_type data) -> value_type {
157
158#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
159 (!defined(SYCL_LANGUAGE_VERSION))
160 return atomicOr(m_ptr, data);
161#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
162 return __VECMEM_SYCL_ATOMIC_CALL1(fetch_or, m_ptr, data);
163#else
164 const value_type result = *m_ptr;
165 *m_ptr |= data;
166 return result;
167#endif
168}
169
170template <typename T>
171VECMEM_HOST_AND_DEVICE auto atomic<T>::fetch_xor(value_type data)
172 -> value_type {
173
174#if (defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)) && \
175 (!defined(SYCL_LANGUAGE_VERSION))
176 return atomicXor(m_ptr, data);
177#elif defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
178 return __VECMEM_SYCL_ATOMIC_CALL1(fetch_xor, m_ptr, data);
179#else
180 const value_type result = *m_ptr;
181 *m_ptr ^= data;
182 return result;
183#endif
184}
185
186} // namespace vecmem
vecmem
Main namespace for the vecmem classes/functions.
Definition atomic_ref.hpp:16
Generated by doxygen 1.9.8