DPC++ Runtime
Runtime libraries for oneAPI DPC++
atomic_intrin.hpp
Go to the documentation of this file.
1 //==-------- atomic_intrin.hpp - Atomic intrinsic definition file ----------==//
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 #pragma once
9 
11 
12 #include <sycl/exception.hpp>
13 
14 namespace __ESIMD_DNS {
15 
16 // This function implements atomic update of pre-existing variable in the
17 // absense of C++ 20's atomic_ref.
18 
19 template <typename Ty> Ty atomic_load(Ty *ptr) {
20 #ifdef _WIN32
21  // TODO: Windows will be supported soon
22  __ESIMD_UNSUPPORTED_ON_HOST;
23 #else
24  return __atomic_load(ptr, __ATOMIC_SEQ_CST);
25 #endif
26 }
27 
28 template <typename Ty> Ty atomic_store(Ty *ptr, Ty val) {
29 #ifdef _WIN32
30  // TODO: Windows will be supported soon
31  __ESIMD_UNSUPPORTED_ON_HOST;
32 #else
33  __atomic_store(ptr, val, __ATOMIC_SEQ_CST);
34 #endif
35 }
36 
37 template <typename Ty> Ty atomic_add_fetch(Ty *ptr, Ty val) {
38 #ifdef _WIN32
39  // TODO: Windows will be supported soon
40  __ESIMD_UNSUPPORTED_ON_HOST;
41 #else
42  return __atomic_add_fetch(ptr, val, __ATOMIC_SEQ_CST);
43 #endif
44 }
45 
46 template <typename Ty> Ty atomic_sub_fetch(Ty *ptr, Ty val) {
47 #ifdef _WIN32
48  // TODO: Windows will be supported soon
49  __ESIMD_UNSUPPORTED_ON_HOST;
50 #else
51  return __atomic_sub_fetch(ptr, val, __ATOMIC_SEQ_CST);
52 #endif
53 }
54 
55 template <typename Ty> Ty atomic_and_fetch(Ty *ptr, Ty val) {
56 #ifdef _WIN32
57  // TODO: Windows will be supported soon
58  __ESIMD_UNSUPPORTED_ON_HOST;
59 #else
60  return __atomic_and_fetch(ptr, val, __ATOMIC_SEQ_CST);
61 #endif
62 }
63 
64 template <typename Ty> Ty atomic_or_fetch(Ty *ptr, Ty val) {
65 #ifdef _WIN32
66  // TODO: Windows will be supported soon
67  __ESIMD_UNSUPPORTED_ON_HOST;
68 #else
69  return __atomic_or_fetch(ptr, val, __ATOMIC_SEQ_CST);
70 #endif
71 }
72 
73 template <typename Ty> Ty atomic_xor_fetch(Ty *ptr, Ty val) {
74 #ifdef _WIN32
75  // TODO: Windows will be supported soon
76  __ESIMD_UNSUPPORTED_ON_HOST;
77 #else
78  return __atomic_xor_fetch(ptr, val, __ATOMIC_SEQ_CST);
79 #endif
80 }
81 
82 template <typename Ty> Ty atomic_min(Ty *ptr, Ty val) {
83 #ifdef _WIN32
84  // TODO: Windows will be supported soon
85  __ESIMD_UNSUPPORTED_ON_HOST;
86 #else
87  // TODO FIXME: fix implementation for FP types.
88  if constexpr (std::is_integral_v<Ty>) {
89  Ty _old, _new;
90  do {
91  _old = *ptr;
92  _new = std::min<Ty>(_old, val);
93  } while (!__atomic_compare_exchange_n(ptr, &_old, _new, false,
94  __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST));
95  return _new;
96  } else {
97  __ESIMD_UNSUPPORTED_ON_HOST;
98  }
99 #endif
100 }
101 
102 template <typename Ty> Ty atomic_max(Ty *ptr, Ty val) {
103 #ifdef _WIN32
104  // TODO: Windows will be supported soon
105  __ESIMD_UNSUPPORTED_ON_HOST;
106 #else
107  // TODO FIXME: fix implementation for FP types.
108  if constexpr (std::is_integral_v<Ty>) {
109  Ty _old, _new;
110  do {
111  _old = *ptr;
112  _new = std::max<Ty>(_old, val);
113  } while (!__atomic_compare_exchange_n(ptr, &_old, _new, false,
114  __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST));
115  return _new;
116  } else {
117  __ESIMD_UNSUPPORTED_ON_HOST;
118  }
119 #endif
120 }
121 
122 template <typename Ty> Ty atomic_cmpxchg(Ty *ptr, Ty expected, Ty desired) {
123 #ifdef _WIN32
124  // TODO: Windows will be supported soon
125  __ESIMD_UNSUPPORTED_ON_HOST;
126 #else
127  // TODO FIXME: fix implementation for FP types.
128  if constexpr (std::is_integral_v<Ty>) {
129  Ty _old = expected;
130  __atomic_compare_exchange_n(ptr, &_old, desired, false, __ATOMIC_SEQ_CST,
131  __ATOMIC_SEQ_CST);
132  return *ptr;
133  } else {
134  __ESIMD_UNSUPPORTED_ON_HOST;
135  }
136 #endif
137 }
138 
139 } // namespace __ESIMD_DNS
140 
sycl::_V1::atomic_load
T atomic_load(atomic< T, addressSpace > Object, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:335
exception.hpp
sycl::_V1::atomic_store
void atomic_store(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
Definition: atomic.hpp:329