DPC++ Runtime
Runtime libraries for oneAPI DPC++
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>
172  "sycl::atomic is deprecated since SYCL 2020") atomic {
173  friend class atomic<T, access::address_space::global_space>;
174  static_assert(detail::IsValidAtomicType<T>::value,
175  "Invalid SYCL atomic type. Valid types are: int, "
176  "unsigned int, long, unsigned long, long long, unsigned "
177  "long long, float");
178  static_assert(detail::IsValidAtomicAddressSpace<addressSpace>::value,
179  "Invalid SYCL atomic address_space. Valid address spaces are: "
180  "global_space, local_space, global_device_space");
181  static constexpr auto SpirvScope =
182  detail::GetSpirvMemoryScope<addressSpace>::scope;
183 
184 public:
185  template <typename pointerT>
186 #ifdef __SYCL_DEVICE_ONLY__
187  atomic(multi_ptr<pointerT, addressSpace> ptr)
188  : Ptr(ptr.get())
189 #else
190  atomic(multi_ptr<pointerT, addressSpace> ptr)
191  : Ptr(reinterpret_cast<std::atomic<T> *>(ptr.get()))
192 #endif
193  {
194  static_assert(sizeof(T) == sizeof(pointerT),
195  "T and pointerT must be same size");
196  }
197 
198 #ifdef __ENABLE_USM_ADDR_SPACE__
199  // Create atomic in global_space with one from global_device_space
200  template <access::address_space _Space = addressSpace,
201  typename = typename detail::enable_if_t<
202  _Space == addressSpace &&
203  addressSpace == access::address_space::global_space>>
204  atomic(const atomic<T, access::address_space::global_device_space> &RHS) {
205  Ptr = RHS.Ptr;
206  }
207 
208  template <access::address_space _Space = addressSpace,
209  typename = typename detail::enable_if_t<
210  _Space == addressSpace &&
211  addressSpace == access::address_space::global_space>>
212  atomic(atomic<T, access::address_space::global_device_space> &&RHS) {
213  Ptr = RHS.Ptr;
214  }
215 #endif // __ENABLE_USM_ADDR_SPACE__
216 
217  void store(T Operand, memory_order Order = memory_order::relaxed) {
219  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
220  }
221 
222 #ifdef __SYCL_DEVICE_ONLY__
223  template <typename T2 = T>
225  load(memory_order Order = memory_order::relaxed) const {
226  return __spirv_AtomicLoad(Ptr, SpirvScope,
228  }
229  template <typename T2 = T>
231  load(memory_order Order = memory_order::relaxed) const {
232  auto *TmpPtr =
233  reinterpret_cast<typename multi_ptr<cl_int, addressSpace>::pointer_t>(
234  Ptr);
235  cl_int TmpVal = __spirv_AtomicLoad(
236  TmpPtr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order));
237  cl_float ResVal = bit_cast<cl_float>(TmpVal);
238  return ResVal;
239  }
240 #else
241  T load(memory_order Order = memory_order::relaxed) const {
242  return __spirv_AtomicLoad(Ptr, SpirvScope,
244  }
245 #endif
246 
247  T exchange(T Operand, memory_order Order = memory_order::relaxed) {
248  return __spirv_AtomicExchange(
249  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
250  }
251 
252  bool
253  compare_exchange_strong(T &Expected, T Desired,
254  memory_order SuccessOrder = memory_order::relaxed,
255  memory_order FailOrder = memory_order::relaxed) {
257 #ifdef __SYCL_DEVICE_ONLY__
258  T Value = __spirv_AtomicCompareExchange(
259  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(SuccessOrder),
260  detail::getSPIRVMemorySemanticsMask(FailOrder), Desired, Expected);
261 
262  if (Value == Expected)
263  return true;
264 
265  Expected = Value;
266  return false;
267 #else
268  return Ptr->compare_exchange_strong(Expected, Desired,
269  detail::getStdMemoryOrder(SuccessOrder),
270  detail::getStdMemoryOrder(FailOrder));
271 #endif
272  }
273 
274  T fetch_add(T Operand, memory_order Order = memory_order::relaxed) {
276  return __spirv_AtomicIAdd(
277  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
278  }
279 
280  T fetch_sub(T Operand, memory_order Order = memory_order::relaxed) {
282  return __spirv_AtomicISub(
283  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
284  }
285 
286  T fetch_and(T Operand, memory_order Order = memory_order::relaxed) {
288  return __spirv_AtomicAnd(
289  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
290  }
291 
292  T fetch_or(T Operand, memory_order Order = memory_order::relaxed) {
294  return __spirv_AtomicOr(
295  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
296  }
297 
298  T fetch_xor(T Operand, memory_order Order = memory_order::relaxed) {
300  return __spirv_AtomicXor(
301  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
302  }
303 
304  T fetch_min(T Operand, memory_order Order = memory_order::relaxed) {
306  return __spirv_AtomicMin(
307  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
308  }
309 
310  T fetch_max(T Operand, memory_order Order = memory_order::relaxed) {
312  return __spirv_AtomicMax(
313  Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
314  }
315 
316 private:
317 #ifdef __SYCL_DEVICE_ONLY__
319 #else
320  std::atomic<T> *Ptr;
321 #endif
322 };
323 
324 template <typename T, access::address_space addressSpace>
325 void atomic_store(atomic<T, addressSpace> Object, T Operand,
326  memory_order MemoryOrder = memory_order::relaxed) {
327  Object.store(Operand, MemoryOrder);
328 }
329 
330 template <typename T, access::address_space addressSpace>
331 T atomic_load(atomic<T, addressSpace> Object,
332  memory_order MemoryOrder = memory_order::relaxed) {
333  return Object.load(MemoryOrder);
334 }
335 
336 template <typename T, access::address_space addressSpace>
337 T atomic_exchange(atomic<T, addressSpace> Object, T Operand,
338  memory_order MemoryOrder = memory_order::relaxed) {
339  return Object.exchange(Operand, MemoryOrder);
340 }
341 
342 template <typename T, access::address_space addressSpace>
344  atomic<T, addressSpace> Object, T &Expected, T Desired,
345  memory_order SuccessOrder = memory_order::relaxed,
346  memory_order FailOrder = memory_order::relaxed) {
347  return Object.compare_exchange_strong(Expected, Desired, SuccessOrder,
348  FailOrder);
349 }
350 
351 template <typename T, access::address_space addressSpace>
352 T atomic_fetch_add(atomic<T, addressSpace> Object, T Operand,
353  memory_order MemoryOrder = memory_order::relaxed) {
354  return Object.fetch_add(Operand, MemoryOrder);
355 }
356 
357 template <typename T, access::address_space addressSpace>
358 T atomic_fetch_sub(atomic<T, addressSpace> Object, T Operand,
359  memory_order MemoryOrder = memory_order::relaxed) {
360  return Object.fetch_sub(Operand, MemoryOrder);
361 }
362 
363 template <typename T, access::address_space addressSpace>
364 T atomic_fetch_and(atomic<T, addressSpace> Object, T Operand,
365  memory_order MemoryOrder = memory_order::relaxed) {
366  return Object.fetch_and(Operand, MemoryOrder);
367 }
368 
369 template <typename T, access::address_space addressSpace>
370 T atomic_fetch_or(atomic<T, addressSpace> Object, T Operand,
371  memory_order MemoryOrder = memory_order::relaxed) {
372  return Object.fetch_or(Operand, MemoryOrder);
373 }
374 
375 template <typename T, access::address_space addressSpace>
376 T atomic_fetch_xor(atomic<T, addressSpace> Object, T Operand,
377  memory_order MemoryOrder = memory_order::relaxed) {
378  return Object.fetch_xor(Operand, MemoryOrder);
379 }
380 
381 template <typename T, access::address_space addressSpace>
382 T atomic_fetch_min(atomic<T, addressSpace> Object, T Operand,
383  memory_order MemoryOrder = memory_order::relaxed) {
384  return Object.fetch_min(Operand, MemoryOrder);
385 }
386 
387 template <typename T, access::address_space addressSpace>
388 T atomic_fetch_max(atomic<T, addressSpace> Object, T Operand,
389  memory_order MemoryOrder = memory_order::relaxed) {
390  return Object.fetch_max(Operand, MemoryOrder);
391 }
392 
393 } // namespace sycl
394 } // __SYCL_INLINE_NAMESPACE(cl)
395 
396 #undef __SYCL_STATIC_ASSERT_NOT_FLOAT
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:370
__spv::Scope::Workgroup
@ Workgroup
Definition: spirv_types.hpp:30
cl::sycl::memory_order
memory_order
Definition: memory_enums.hpp:16
T
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:56
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:364
cl::sycl::multi_ptr::pointer_t
typename detail::DecoratedType< ElementType, Space >::type * pointer_t
Definition: multi_ptr.hpp:39
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:343
helpers.hpp
sycl
Definition: invoke_simd.hpp:68
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
access.hpp
cl::sycl::atomic_store
void atomic_store(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:325
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:352
cl::sycl::atomic_exchange
T atomic_exchange(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:337
cl::sycl::atomic_fetch_max
T atomic_fetch_max(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:388
__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::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::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:331
__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:358
__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:376
cl::sycl::atomic_fetch_min
T atomic_fetch_min(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:382
__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
__spv::Scope::Device
@ Device
Definition: spirv_types.hpp:29
cl::sycl::detail::DecoratedType
Definition: access.hpp:159
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12