DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
atomic.hpp
Go to the documentation of this file.
1 //==---------------- atomic.hpp - SYCL atomics -----------------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <CL/__spirv/spirv_ops.hpp>
14 #include <CL/sycl/memory_enums.hpp>
15 
16 #ifndef __SYCL_DEVICE_ONLY__
17 #include <atomic>
18 #else
19 #include <cstring>
20 #endif
21 #include <type_traits>
22 
23 #define __SYCL_STATIC_ASSERT_NOT_FLOAT(T) \
24  static_assert(!std::is_same<T, float>::value, \
25  "SYCL atomic function not available for float type")
26 
28 namespace sycl {
29 
30 // Forward declaration
31 template <typename pointerT, access::address_space addressSpace>
32 class multi_ptr;
33 
34 namespace detail {
35 
37 
38 template <typename T> struct IsValidAtomicType {
39  static constexpr bool value =
40  (std::is_same<T, int>::value || std::is_same<T, unsigned int>::value ||
41  std::is_same<T, long>::value || std::is_same<T, unsigned long>::value ||
42  std::is_same<T, long long>::value ||
43  std::is_same<T, unsigned long long>::value ||
44  std::is_same<T, float>::value);
45 };
46 
47 template <cl::sycl::access::address_space AS> struct IsValidAtomicAddressSpace {
48  static constexpr bool value =
49  (AS == access::address_space::global_space ||
50  AS == access::address_space::local_space ||
51  AS == access::address_space::global_device_space);
52 };
53 
54 // Type trait to translate a cl::sycl::access::address_space to
55 // a SPIR-V memory scope
56 template <access::address_space AS> struct GetSpirvMemoryScope {};
57 template <> struct GetSpirvMemoryScope<access::address_space::global_space> {
58  static constexpr auto scope = __spv::Scope::Device;
59 };
60 template <>
61 struct GetSpirvMemoryScope<access::address_space::global_device_space> {
62  static constexpr auto scope = __spv::Scope::Device;
63 };
64 template <> struct GetSpirvMemoryScope<access::address_space::local_space> {
65  static constexpr auto scope = __spv::Scope::Workgroup;
66 };
67 
68 } // namespace detail
69 } // namespace sycl
70 } // __SYCL_INLINE_NAMESPACE(cl)
71 
72 #ifndef __SYCL_DEVICE_ONLY__
73 // host implementation of SYCL atomics
75 namespace sycl {
76 namespace detail {
77 // Translate cl::sycl::memory_order or __spv::MemorySemanticsMask::Flag
78 // into std::memory_order
79 // Only relaxed memory semantics are supported currently
80 static inline std::memory_order
82  return std::memory_order_relaxed;
83 }
84 } // namespace detail
85 } // namespace sycl
86 } // __SYCL_INLINE_NAMESPACE(cl)
87 
88 // std::atomic version of atomic SPIR-V builtins
89 
90 template <typename T>
91 void __spirv_AtomicStore(std::atomic<T> *Ptr, __spv::Scope::Flag,
93  Ptr->store(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
94 }
95 
96 template <typename T>
97 T __spirv_AtomicLoad(const std::atomic<T> *Ptr, __spv::Scope::Flag,
99  return Ptr->load(::cl::sycl::detail::getStdMemoryOrder(MS));
100 }
101 
102 template <typename T>
105  return Ptr->exchange(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
106 }
107 
108 template <typename T>
109 extern T __spirv_AtomicIAdd(std::atomic<T> *Ptr, __spv::Scope::Flag,
111  return Ptr->fetch_add(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
112 }
113 
114 template <typename T>
115 extern T __spirv_AtomicISub(std::atomic<T> *Ptr, __spv::Scope::Flag,
117  return Ptr->fetch_sub(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
118 }
119 
120 template <typename T>
121 extern T __spirv_AtomicAnd(std::atomic<T> *Ptr, __spv::Scope::Flag,
123  return Ptr->fetch_and(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
124 }
125 
126 template <typename T>
127 extern T __spirv_AtomicOr(std::atomic<T> *Ptr, __spv::Scope::Flag,
129  return Ptr->fetch_or(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
130 }
131 
132 template <typename T>
133 extern T __spirv_AtomicXor(std::atomic<T> *Ptr, __spv::Scope::Flag,
135  return Ptr->fetch_xor(V, ::cl::sycl::detail::getStdMemoryOrder(MS));
136 }
137 
138 template <typename T>
139 extern T __spirv_AtomicMin(std::atomic<T> *Ptr, __spv::Scope::Flag,
142  T Val = Ptr->load(MemoryOrder);
143  while (V < Val) {
144  if (Ptr->compare_exchange_strong(Val, V, MemoryOrder, MemoryOrder))
145  break;
146  Val = Ptr->load(MemoryOrder);
147  }
148  return Val;
149 }
150 
151 template <typename T>
152 extern T __spirv_AtomicMax(std::atomic<T> *Ptr, __spv::Scope::Flag,
155  T Val = Ptr->load(MemoryOrder);
156  while (V > Val) {
157  if (Ptr->compare_exchange_strong(Val, V, MemoryOrder, MemoryOrder))
158  break;
159  Val = Ptr->load(MemoryOrder);
160  }
161  return Val;
162 }
163 
164 #endif // !defined(__SYCL_DEVICE_ONLY__)
165 
167 namespace sycl {
168 
169 template <typename T, access::address_space addressSpace =
170  access::address_space::global_space>
171 class atomic {
172  friend class atomic<T, access::address_space::global_space>;
173  static_assert(detail::IsValidAtomicType<T>::value,
174  "Invalid SYCL atomic type. Valid types are: int, "
175  "unsigned int, long, unsigned long, long long, unsigned "
176  "long long, float");
177  static_assert(detail::IsValidAtomicAddressSpace<addressSpace>::value,
178  "Invalid SYCL atomic address_space. Valid address spaces are: "
179  "global_space, local_space, global_device_space");
180  static constexpr auto SpirvScope =
181  detail::GetSpirvMemoryScope<addressSpace>::scope;
182 
183 public:
184  template <typename pointerT>
185 #ifdef __SYCL_DEVICE_ONLY__
186  atomic(multi_ptr<pointerT, addressSpace> ptr)
187  : Ptr(ptr.get())
188 #else
189  atomic(multi_ptr<pointerT, addressSpace> ptr)
190  : Ptr(reinterpret_cast<std::atomic<T> *>(ptr.get()))
191 #endif
192  {
193  static_assert(sizeof(T) == sizeof(pointerT),
194  "T and pointerT must be same size");
195  }
196 
197 #ifdef __ENABLE_USM_ADDR_SPACE__
198  // Create atomic in global_space with one from global_device_space
199  template <access::address_space _Space = addressSpace,
200  typename = typename detail::enable_if_t<
201  _Space == addressSpace &&
202  addressSpace == access::address_space::global_space>>
204  Ptr = RHS.Ptr;
205  }
206 
207  template <access::address_space _Space = addressSpace,
208  typename = typename detail::enable_if_t<
209  _Space == addressSpace &&
210  addressSpace == access::address_space::global_space>>
211  atomic(atomic<T, access::address_space::global_device_space> &&RHS) {
212  Ptr = RHS.Ptr;
213  }
214 #endif // __ENABLE_USM_ADDR_SPACE__
215 
216  void store(T Operand, memory_order Order = memory_order::relaxed) {
218  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
219  }
220 
221 #ifdef __SYCL_DEVICE_ONLY__
222  template <typename T2 = T>
224  load(memory_order Order = memory_order::relaxed) const {
225  return __spirv_AtomicLoad(Ptr, SpirvScope,
227  }
228  template <typename T2 = T>
229  detail::enable_if_t<std::is_same<cl_float, T2>::value, T>
230  load(memory_order Order = memory_order::relaxed) const {
231  auto *TmpPtr =
232  reinterpret_cast<typename multi_ptr<cl_int, addressSpace>::pointer_t>(
233  Ptr);
234  cl_int TmpVal = __spirv_AtomicLoad(
235  TmpPtr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order));
236  cl_float ResVal = bit_cast<cl_float>(TmpVal);
237  return ResVal;
238  }
239 #else
240  T load(memory_order Order = memory_order::relaxed) const {
241  return __spirv_AtomicLoad(Ptr, SpirvScope,
243  }
244 #endif
245 
246  T exchange(T Operand, memory_order Order = memory_order::relaxed) {
247  return __spirv_AtomicExchange(
248  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
249  }
250 
251  bool
252  compare_exchange_strong(T &Expected, T Desired,
253  memory_order SuccessOrder = memory_order::relaxed,
254  memory_order FailOrder = memory_order::relaxed) {
256 #ifdef __SYCL_DEVICE_ONLY__
257  T Value = __spirv_AtomicCompareExchange(
258  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(SuccessOrder),
259  detail::getSPIRVMemorySemanticsMask(FailOrder), Desired, Expected);
260 
261  if (Value == Expected)
262  return true;
263 
264  Expected = Value;
265  return false;
266 #else
267  return Ptr->compare_exchange_strong(Expected, Desired,
268  detail::getStdMemoryOrder(SuccessOrder),
269  detail::getStdMemoryOrder(FailOrder));
270 #endif
271  }
272 
273  T fetch_add(T Operand, memory_order Order = memory_order::relaxed) {
275  return __spirv_AtomicIAdd(
276  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
277  }
278 
279  T fetch_sub(T Operand, memory_order Order = memory_order::relaxed) {
281  return __spirv_AtomicISub(
282  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
283  }
284 
285  T fetch_and(T Operand, memory_order Order = memory_order::relaxed) {
287  return __spirv_AtomicAnd(
288  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
289  }
290 
291  T fetch_or(T Operand, memory_order Order = memory_order::relaxed) {
293  return __spirv_AtomicOr(
294  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
295  }
296 
297  T fetch_xor(T Operand, memory_order Order = memory_order::relaxed) {
299  return __spirv_AtomicXor(
300  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
301  }
302 
303  T fetch_min(T Operand, memory_order Order = memory_order::relaxed) {
305  return __spirv_AtomicMin(
306  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
307  }
308 
309  T fetch_max(T Operand, memory_order Order = memory_order::relaxed) {
311  return __spirv_AtomicMax(
312  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
313  }
314 
315 private:
316 #ifdef __SYCL_DEVICE_ONLY__
318 #else
319  std::atomic<T> *Ptr;
320 #endif
321 };
322 
323 template <typename T, access::address_space addressSpace>
324 void atomic_store(atomic<T, addressSpace> Object, T Operand,
325  memory_order MemoryOrder = memory_order::relaxed) {
326  Object.store(Operand, MemoryOrder);
327 }
328 
329 template <typename T, access::address_space addressSpace>
331  memory_order MemoryOrder = memory_order::relaxed) {
332  return Object.load(MemoryOrder);
333 }
334 
335 template <typename T, access::address_space addressSpace>
337  memory_order MemoryOrder = memory_order::relaxed) {
338  return Object.exchange(Operand, MemoryOrder);
339 }
340 
341 template <typename T, access::address_space addressSpace>
343  atomic<T, addressSpace> Object, T &Expected, T Desired,
344  memory_order SuccessOrder = memory_order::relaxed,
345  memory_order FailOrder = memory_order::relaxed) {
346  return Object.compare_exchange_strong(Expected, Desired, SuccessOrder,
347  FailOrder);
348 }
349 
350 template <typename T, access::address_space addressSpace>
352  memory_order MemoryOrder = memory_order::relaxed) {
353  return Object.fetch_add(Operand, MemoryOrder);
354 }
355 
356 template <typename T, access::address_space addressSpace>
358  memory_order MemoryOrder = memory_order::relaxed) {
359  return Object.fetch_sub(Operand, MemoryOrder);
360 }
361 
362 template <typename T, access::address_space addressSpace>
364  memory_order MemoryOrder = memory_order::relaxed) {
365  return Object.fetch_and(Operand, MemoryOrder);
366 }
367 
368 template <typename T, access::address_space addressSpace>
370  memory_order MemoryOrder = memory_order::relaxed) {
371  return Object.fetch_or(Operand, MemoryOrder);
372 }
373 
374 template <typename T, access::address_space addressSpace>
376  memory_order MemoryOrder = memory_order::relaxed) {
377  return Object.fetch_xor(Operand, MemoryOrder);
378 }
379 
380 template <typename T, access::address_space addressSpace>
382  memory_order MemoryOrder = memory_order::relaxed) {
383  return Object.fetch_min(Operand, MemoryOrder);
384 }
385 
386 template <typename T, access::address_space addressSpace>
388  memory_order MemoryOrder = memory_order::relaxed) {
389  return Object.fetch_max(Operand, MemoryOrder);
390 }
391 
392 } // namespace sycl
393 } // __SYCL_INLINE_NAMESPACE(cl)
394 
395 #undef __SYCL_STATIC_ASSERT_NOT_FLOAT
cl::sycl::atomic::fetch_max
T fetch_max(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:309
spirv_ops.hpp
__spirv_AtomicStore
void __spirv_AtomicStore(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:91
cl::sycl::atomic_fetch_or
T atomic_fetch_or(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:369
__spv::Scope::Workgroup
@ Workgroup
Definition: spirv_types.hpp:30
cl::sycl::memory_order
memory_order
Definition: memory_enums.hpp:16
T
cl::sycl::detail::IsValidAtomicAddressSpace
Definition: atomic.hpp:47
cl::sycl::atomic_fetch_and
T atomic_fetch_and(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:363
cl::sycl::atomic::exchange
T exchange(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:246
cl::sycl::atomic_compare_exchange_strong
bool atomic_compare_exchange_strong(atomic< T, addressSpace > Object, T &Expected, T Desired, memory_order SuccessOrder=memory_order::relaxed, memory_order FailOrder=memory_order::relaxed)
Definition: atomic.hpp:342
cl::sycl::atomic::fetch_or
T fetch_or(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:291
helpers.hpp
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
cl::sycl::atomic::fetch_xor
T fetch_xor(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:297
access.hpp
cl::sycl::atomic::fetch_and
T fetch_and(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:285
cl::sycl::atomic::store
void store(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:216
cl::sycl::atomic::fetch_add
T fetch_add(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:273
cl::sycl::atomic_store
void atomic_store(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:324
cl::sycl::detail::IsValidAtomicType
Definition: atomic.hpp:38
cl::sycl::detail::memory_order
cl::sycl::memory_order memory_order
Definition: atomic.hpp:36
__spv::Scope::Flag
Flag
Definition: spirv_types.hpp:27
cl::sycl::detail::getSPIRVMemorySemanticsMask
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
Definition: helpers.hpp:200
__SYCL_STATIC_ASSERT_NOT_FLOAT
#define __SYCL_STATIC_ASSERT_NOT_FLOAT(T)
Definition: atomic.hpp:23
cl::sycl::detail::getStdMemoryOrder
static std::memory_order getStdMemoryOrder(__spv::MemorySemanticsMask::Flag)
Definition: atomic.hpp:81
cl::sycl::atomic_fetch_add
T atomic_fetch_add(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:351
cl::sycl::atomic_exchange
T atomic_exchange(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:336
cl::sycl::atomic_fetch_max
T atomic_fetch_max(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:387
__spirv_AtomicExchange
T __spirv_AtomicExchange(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:103
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
__spirv_AtomicISub
T __spirv_AtomicISub(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:115
__spirv_AtomicMax
T __spirv_AtomicMax(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:152
__spirv_AtomicIAdd
T __spirv_AtomicIAdd(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:109
cl::sycl::cl_float
float cl_float
Definition: aliases.hpp:87
cl::sycl::atomic::fetch_min
T fetch_min(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:303
cl::sycl::access::address_space
address_space
Definition: access.hpp:45
__spirv_AtomicMin
T __spirv_AtomicMin(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:139
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:82
cl::sycl::atomic::fetch_sub
T fetch_sub(T Operand, memory_order Order=memory_order::relaxed)
Definition: atomic.hpp:279
cl::sycl::detail::GetSpirvMemoryScope
Definition: atomic.hpp:56
cl::sycl::atomic_load
T atomic_load(atomic< T, addressSpace > Object, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:330
__spirv_AtomicAnd
T __spirv_AtomicAnd(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:121
__spv::MemorySemanticsMask::Flag
Flag
Definition: spirv_types.hpp:84
__spirv_AtomicLoad
T __spirv_AtomicLoad(const std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS)
Definition: atomic.hpp:97
memory_enums.hpp
cl::sycl::atomic_fetch_sub
T atomic_fetch_sub(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:357
cl::sycl::atomic::compare_exchange_strong
bool compare_exchange_strong(T &Expected, T Desired, memory_order SuccessOrder=memory_order::relaxed, memory_order FailOrder=memory_order::relaxed)
Definition: atomic.hpp:252
__spirv_AtomicXor
T __spirv_AtomicXor(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:133
cl::sycl::atomic_fetch_xor
T atomic_fetch_xor(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:375
cl::sycl::atomic_fetch_min
T atomic_fetch_min(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:381
__spirv_AtomicOr
T __spirv_AtomicOr(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
Definition: atomic.hpp:127
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::atomic::load
T load(memory_order Order=memory_order::relaxed) const
Definition: atomic.hpp:240
__spv::Scope::Device
@ Device
Definition: spirv_types.hpp:29
cl::sycl::atomic
Definition: atomic.hpp:171
cl::sycl::detail::DecoratedType
Definition: access.hpp:158
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12