18 #include <type_traits>
20 #ifndef __SYCL_DEVICE_ONLY__
24 #define __SYCL_STATIC_ASSERT_NOT_FLOAT(T) \
25 static_assert(!std::is_same<T, float>::value, \
26 "SYCL atomic function not available for float type")
29 inline namespace _V1 {
42 (std::is_same_v<T, int> || std::is_same_v<T, unsigned int> ||
43 std::is_same_v<T, long> || std::is_same_v<T, unsigned long> ||
44 std::is_same_v<T, long long> || std::is_same_v<T, unsigned long long> ||
45 std::is_same_v<T, float>);
63 access::address_space::ext_intel_global_device_space> {
74 #ifndef __SYCL_DEVICE_ONLY__
77 inline namespace _V1 {
103 template <
typename T>
109 template <
typename T>
115 template <
typename T>
121 template <
typename T>
127 template <
typename T>
133 template <
typename T>
139 template <
typename T>
143 T Val = Ptr->load(MemoryOrder);
145 if (Ptr->compare_exchange_strong(Val, V, MemoryOrder, MemoryOrder))
147 Val = Ptr->load(MemoryOrder);
152 template <
typename T>
156 T Val = Ptr->load(MemoryOrder);
158 if (Ptr->compare_exchange_strong(Val, V, MemoryOrder, MemoryOrder))
160 Val = Ptr->load(MemoryOrder);
168 inline namespace _V1 {
174 friend class atomic<T, access::address_space::global_space>;
175 static_assert(detail::IsValidAtomicType<T>::value,
176 "Invalid SYCL atomic type. Valid types are: int, "
177 "unsigned int, long, unsigned long, long long, unsigned "
179 static_assert(detail::IsValidAtomicAddressSpace<addressSpace>::value,
180 "Invalid SYCL atomic address_space. Valid address spaces are: "
181 "global_space, local_space, ext_intel_global_device_space");
182 static constexpr auto SpirvScope =
183 detail::GetSpirvMemoryScope<addressSpace>::scope;
185 template <typename pointerT, access::decorated IsDecorated>
187 GetDecoratedPtr(multi_ptr<pointerT, addressSpace, IsDecorated> ptr) {
188 if constexpr (IsDecorated == access::decorated::legacy)
191 return ptr.get_decorated();
195 template <
typename po
interT, access::decorated IsDecorated>
196 #ifdef __SYCL_DEVICE_ONLY__
197 atomic(multi_ptr<pointerT, addressSpace, IsDecorated> ptr)
198 : Ptr(GetDecoratedPtr(ptr))
204 static_assert(
sizeof(T) ==
sizeof(pointerT),
205 "T and pointerT must be same size");
208 #ifdef __ENABLE_USM_ADDR_SPACE__
211 typename =
typename std::enable_if_t<
212 _Space == addressSpace &&
214 atomic(
const atomic<T, access::address_space::ext_intel_global_device_space>
220 typename =
typename std::enable_if_t<
221 _Space == addressSpace &&
224 atomic<T, access::address_space::ext_intel_global_device_space> &&RHS) {
234 #ifdef __SYCL_DEVICE_ONLY__
235 template <
typename T2 = T>
236 std::enable_if_t<!std::is_same<cl_float, T2>::value, T>
241 template <
typename T2 = T>
242 std::enable_if_t<std::is_same<cl_float, T2>::value, T>
244 auto *TmpPtr =
reinterpret_cast<typename
multi_ptr<
248 cl_float ResVal = sycl::bit_cast<cl_float>(TmpVal);
264 compare_exchange_strong(T &Expected, T Desired,
268 #ifdef __SYCL_DEVICE_ONLY__
269 T Value = __spirv_AtomicCompareExchange(
273 if (Value == Expected)
279 return Ptr->compare_exchange_strong(Expected, Desired,
328 #ifdef __SYCL_DEVICE_ONLY__
329 typename detail::DecoratedType<T, addressSpace>::type *Ptr;
335 template <
typename T, access::address_space addressSpace>
338 Object.store(Operand, MemoryOrder);
341 template <
typename T, access::address_space addressSpace>
344 return Object.load(MemoryOrder);
347 template <
typename T, access::address_space addressSpace>
350 return Object.exchange(Operand, MemoryOrder);
353 template <
typename T, access::address_space addressSpace>
355 atomic<T, addressSpace> Object, T &Expected, T Desired,
358 return Object.compare_exchange_strong(Expected, Desired, SuccessOrder,
362 template <
typename T, access::address_space addressSpace>
365 return Object.fetch_add(Operand, MemoryOrder);
368 template <
typename T, access::address_space addressSpace>
371 return Object.fetch_sub(Operand, MemoryOrder);
374 template <
typename T, access::address_space addressSpace>
377 return Object.fetch_and(Operand, MemoryOrder);
380 template <
typename T, access::address_space addressSpace>
383 return Object.fetch_or(Operand, MemoryOrder);
386 template <
typename T, access::address_space addressSpace>
389 return Object.fetch_xor(Operand, MemoryOrder);
392 template <
typename T, access::address_space addressSpace>
395 return Object.fetch_min(Operand, MemoryOrder);
398 template <
typename T, access::address_space addressSpace>
401 return Object.fetch_max(Operand, MemoryOrder);
407 #undef __SYCL_STATIC_ASSERT_NOT_FLOAT
@ ext_intel_global_device_space
sycl::memory_order memory_order
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
std::memory_order getStdMemoryOrder(__spv::MemorySemanticsMask::Flag)
T atomic_fetch_xor(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
constexpr auto memory_order_relaxed
T atomic_load(atomic< T, addressSpace > Object, memory_order MemoryOrder=memory_order::relaxed)
T atomic_exchange(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
T atomic_fetch_and(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
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)
signed char __SYCL2020_DEPRECATED
T atomic_fetch_min(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
T atomic_fetch_add(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
std::conditional_t< is_decorated, decorated_type *, std::add_pointer_t< value_type > > pointer
T atomic_fetch_sub(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
T atomic_fetch_max(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
void atomic_store(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
T atomic_fetch_or(atomic< T, addressSpace > Object, T Operand, memory_order MemoryOrder=memory_order::relaxed)
static constexpr bool value
static constexpr bool value
T __spirv_AtomicOr(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
void __spirv_AtomicStore(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
T __spirv_AtomicExchange(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
#define __SYCL_STATIC_ASSERT_NOT_FLOAT(T)
T __spirv_AtomicLoad(const std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS)
T __spirv_AtomicMax(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
T __spirv_AtomicIAdd(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
T __spirv_AtomicAnd(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
T __spirv_AtomicXor(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
T __spirv_AtomicISub(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)
T __spirv_AtomicMin(std::atomic< T > *Ptr, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag MS, T V)