Skip to content

Commit 9b04f41

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Adds info query for atomic_memory_order_capabilities on device and context (#4105)
Implements support for `device::info::atomic_memory_order_capabilities` and `context::info::atomic_memory_order_capabilities`. The corresponding PI query is only implemented for the PI CUDA backend and host. Additionally, these changes reverse the dependency between ` sycl/include/CL/sycl/ext/oneapi/atomic_enums.hpp` and `sycl/include/CL/sycl/memory_enums.hpp` and replace all uses of `sycl::ext::oneapi::memory_order` and `sycl::ext::oneapi::memory_scope` with `sycl::memory_order` and `sycl::memory_scope` respectively.
1 parent fb08adf commit 9b04f41

23 files changed

+197
-122
lines changed

‎sycl/include/CL/sycl/detail/pi.h

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -292,7 +292,8 @@ typedef enum {
292292
PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE = 0x10025,
293293
PI_DEVICE_INFO_MAX_MEM_BANDWIDTH = 0x10026,
294294
PI_DEVICE_INFO_IMAGE_SRGB = 0x10027,
295-
PI_DEVICE_INFO_ATOMIC_64 = 0x10110
295+
PI_DEVICE_INFO_ATOMIC_64 = 0x10110,
296+
PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES = 0x10111
296297
} _pi_device_info;
297298

298299
typedef enum {
@@ -312,6 +313,8 @@ typedef enum {
312313
PI_CONTEXT_INFO_NUM_DEVICES = CL_CONTEXT_NUM_DEVICES,
313314
PI_CONTEXT_INFO_PROPERTIES = CL_CONTEXT_PROPERTIES,
314315
PI_CONTEXT_INFO_REFERENCE_COUNT = CL_CONTEXT_REFERENCE_COUNT,
316+
// Atomics capabilities extensions
317+
PI_CONTEXT_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES = 0x10010
315318
} _pi_context_info;
316319

317320
typedef enum {
@@ -509,6 +512,13 @@ constexpr pi_sampler_properties PI_SAMPLER_PROPERTIES_ADDRESSING_MODE =
509512
constexpr pi_sampler_properties PI_SAMPLER_PROPERTIES_FILTER_MODE =
510513
CL_SAMPLER_FILTER_MODE;
511514

515+
using pi_memory_order_capabilities = pi_bitfield;
516+
constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_RELAXED = 0x01;
517+
constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_ACQUIRE = 0x02;
518+
constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_RELEASE = 0x04;
519+
constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_ACQ_REL = 0x08;
520+
constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_SEQ_CST = 0x10;
521+
512522
typedef enum {
513523
PI_PROFILING_INFO_COMMAND_QUEUED = CL_PROFILING_COMMAND_QUEUED,
514524
PI_PROFILING_INFO_COMMAND_SUBMIT = CL_PROFILING_COMMAND_SUBMIT,

‎sycl/include/CL/sycl/detail/spirv.hpp

Lines changed: 46 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -224,11 +224,10 @@ EnableIfGenericBroadcast<T> GroupBroadcast(T x, id<Dimensions> local_id) {
224224
// Single happens-before means semantics should always apply to all spaces
225225
// Although consume is unsupported, forwarding to acquire is valid
226226
template <typename T>
227-
static inline constexpr typename std::enable_if<
228-
std::is_same<T, sycl::ext::oneapi::memory_order>::value ||
229-
std::is_same<T, sycl::memory_order>::value,
230-
__spv::MemorySemanticsMask::Flag>::type
231-
getMemorySemanticsMask(T Order) {
227+
static inline constexpr
228+
typename std::enable_if<std::is_same<T, sycl::memory_order>::value,
229+
__spv::MemorySemanticsMask::Flag>::type
230+
getMemorySemanticsMask(T Order) {
232231
__spv::MemorySemanticsMask::Flag SpvOrder = __spv::MemorySemanticsMask::None;
233232
switch (Order) {
234233
case T::relaxed:
@@ -254,28 +253,25 @@ getMemorySemanticsMask(T Order) {
254253
__spv::MemorySemanticsMask::CrossWorkgroupMemory);
255254
}
256255

257-
static inline constexpr __spv::Scope::Flag
258-
getScope(ext::oneapi::memory_scope Scope) {
256+
static inline constexpr __spv::Scope::Flag getScope(memory_scope Scope) {
259257
switch (Scope) {
260-
case ext::oneapi::memory_scope::work_item:
258+
case memory_scope::work_item:
261259
return __spv::Scope::Invocation;
262-
case ext::oneapi::memory_scope::sub_group:
260+
case memory_scope::sub_group:
263261
return __spv::Scope::Subgroup;
264-
case ext::oneapi::memory_scope::work_group:
262+
case memory_scope::work_group:
265263
return __spv::Scope::Workgroup;
266-
case ext::oneapi::memory_scope::device:
264+
case memory_scope::device:
267265
return __spv::Scope::Device;
268-
case ext::oneapi::memory_scope::system:
266+
case memory_scope::system:
269267
return __spv::Scope::CrossDevice;
270268
}
271269
}
272270

273271
template <typename T, access::address_space AddressSpace>
274272
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
275-
AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
276-
ext::oneapi::memory_scope Scope,
277-
ext::oneapi::memory_order Success,
278-
ext::oneapi::memory_order Failure, T Desired,
273+
AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
274+
memory_order Success, memory_order Failure, T Desired,
279275
T Expected) {
280276
auto SPIRVSuccess = getMemorySemanticsMask(Success);
281277
auto SPIRVFailure = getMemorySemanticsMask(Failure);
@@ -287,10 +283,8 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
287283

288284
template <typename T, access::address_space AddressSpace>
289285
inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
290-
AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
291-
ext::oneapi::memory_scope Scope,
292-
ext::oneapi::memory_order Success,
293-
ext::oneapi::memory_order Failure, T Desired,
286+
AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
287+
memory_order Success, memory_order Failure, T Desired,
294288
T Expected) {
295289
using I = detail::make_unsinged_integer_t<T>;
296290
auto SPIRVSuccess = getMemorySemanticsMask(Success);
@@ -308,8 +302,8 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
308302

309303
template <typename T, access::address_space AddressSpace>
310304
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
311-
AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
312-
ext::oneapi::memory_order Order) {
305+
AtomicLoad(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
306+
memory_order Order) {
313307
auto *Ptr = MPtr.get();
314308
auto SPIRVOrder = getMemorySemanticsMask(Order);
315309
auto SPIRVScope = getScope(Scope);
@@ -318,8 +312,8 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
318312

319313
template <typename T, access::address_space AddressSpace>
320314
inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
321-
AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
322-
ext::oneapi::memory_order Order) {
315+
AtomicLoad(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
316+
memory_order Order) {
323317
using I = detail::make_unsinged_integer_t<T>;
324318
auto *PtrInt =
325319
reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -332,8 +326,8 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
332326

333327
template <typename T, access::address_space AddressSpace>
334328
inline typename detail::enable_if_t<std::is_integral<T>::value>
335-
AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
336-
ext::oneapi::memory_order Order, T Value) {
329+
AtomicStore(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
330+
memory_order Order, T Value) {
337331
auto *Ptr = MPtr.get();
338332
auto SPIRVOrder = getMemorySemanticsMask(Order);
339333
auto SPIRVScope = getScope(Scope);
@@ -342,8 +336,8 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
342336

343337
template <typename T, access::address_space AddressSpace>
344338
inline typename detail::enable_if_t<std::is_floating_point<T>::value>
345-
AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
346-
ext::oneapi::memory_order Order, T Value) {
339+
AtomicStore(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
340+
memory_order Order, T Value) {
347341
using I = detail::make_unsinged_integer_t<T>;
348342
auto *PtrInt =
349343
reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -356,8 +350,8 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
356350

357351
template <typename T, access::address_space AddressSpace>
358352
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
359-
AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
360-
ext::oneapi::memory_order Order, T Value) {
353+
AtomicExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
354+
memory_order Order, T Value) {
361355
auto *Ptr = MPtr.get();
362356
auto SPIRVOrder = getMemorySemanticsMask(Order);
363357
auto SPIRVScope = getScope(Scope);
@@ -366,8 +360,8 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
366360

367361
template <typename T, access::address_space AddressSpace>
368362
inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
369-
AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
370-
ext::oneapi::memory_order Order, T Value) {
363+
AtomicExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
364+
memory_order Order, T Value) {
371365
using I = detail::make_unsinged_integer_t<T>;
372366
auto *PtrInt =
373367
reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -382,8 +376,8 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
382376

383377
template <typename T, access::address_space AddressSpace>
384378
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
385-
AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
386-
ext::oneapi::memory_order Order, T Value) {
379+
AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
380+
memory_order Order, T Value) {
387381
auto *Ptr = MPtr.get();
388382
auto SPIRVOrder = getMemorySemanticsMask(Order);
389383
auto SPIRVScope = getScope(Scope);
@@ -392,8 +386,8 @@ AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
392386

393387
template <typename T, access::address_space AddressSpace>
394388
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
395-
AtomicISub(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
396-
ext::oneapi::memory_order Order, T Value) {
389+
AtomicISub(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
390+
memory_order Order, T Value) {
397391
auto *Ptr = MPtr.get();
398392
auto SPIRVOrder = getMemorySemanticsMask(Order);
399393
auto SPIRVScope = getScope(Scope);
@@ -402,8 +396,8 @@ AtomicISub(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
402396

403397
template <typename T, access::address_space AddressSpace>
404398
inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
405-
AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
406-
ext::oneapi::memory_order Order, T Value) {
399+
AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
400+
memory_order Order, T Value) {
407401
auto *Ptr = MPtr.get();
408402
auto SPIRVOrder = getMemorySemanticsMask(Order);
409403
auto SPIRVScope = getScope(Scope);
@@ -412,8 +406,8 @@ AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
412406

413407
template <typename T, access::address_space AddressSpace>
414408
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
415-
AtomicAnd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
416-
ext::oneapi::memory_order Order, T Value) {
409+
AtomicAnd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
410+
memory_order Order, T Value) {
417411
auto *Ptr = MPtr.get();
418412
auto SPIRVOrder = getMemorySemanticsMask(Order);
419413
auto SPIRVScope = getScope(Scope);
@@ -422,8 +416,8 @@ AtomicAnd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
422416

423417
template <typename T, access::address_space AddressSpace>
424418
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
425-
AtomicOr(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
426-
ext::oneapi::memory_order Order, T Value) {
419+
AtomicOr(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
420+
memory_order Order, T Value) {
427421
auto *Ptr = MPtr.get();
428422
auto SPIRVOrder = getMemorySemanticsMask(Order);
429423
auto SPIRVScope = getScope(Scope);
@@ -432,8 +426,8 @@ AtomicOr(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
432426

433427
template <typename T, access::address_space AddressSpace>
434428
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
435-
AtomicXor(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
436-
ext::oneapi::memory_order Order, T Value) {
429+
AtomicXor(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
430+
memory_order Order, T Value) {
437431
auto *Ptr = MPtr.get();
438432
auto SPIRVOrder = getMemorySemanticsMask(Order);
439433
auto SPIRVScope = getScope(Scope);
@@ -442,8 +436,8 @@ AtomicXor(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
442436

443437
template <typename T, access::address_space AddressSpace>
444438
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
445-
AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
446-
ext::oneapi::memory_order Order, T Value) {
439+
AtomicMin(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
440+
memory_order Order, T Value) {
447441
auto *Ptr = MPtr.get();
448442
auto SPIRVOrder = getMemorySemanticsMask(Order);
449443
auto SPIRVScope = getScope(Scope);
@@ -452,8 +446,8 @@ AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
452446

453447
template <typename T, access::address_space AddressSpace>
454448
inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
455-
AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
456-
ext::oneapi::memory_order Order, T Value) {
449+
AtomicMin(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
450+
memory_order Order, T Value) {
457451
auto *Ptr = MPtr.get();
458452
auto SPIRVOrder = getMemorySemanticsMask(Order);
459453
auto SPIRVScope = getScope(Scope);
@@ -462,8 +456,8 @@ AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
462456

463457
template <typename T, access::address_space AddressSpace>
464458
inline typename detail::enable_if_t<std::is_integral<T>::value, T>
465-
AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
466-
ext::oneapi::memory_order Order, T Value) {
459+
AtomicMax(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
460+
memory_order Order, T Value) {
467461
auto *Ptr = MPtr.get();
468462
auto SPIRVOrder = getMemorySemanticsMask(Order);
469463
auto SPIRVScope = getScope(Scope);
@@ -472,8 +466,8 @@ AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
472466

473467
template <typename T, access::address_space AddressSpace>
474468
inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
475-
AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
476-
ext::oneapi::memory_order Order, T Value) {
469+
AtomicMax(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
470+
memory_order Order, T Value) {
477471
auto *Ptr = MPtr.get();
478472
auto SPIRVOrder = getMemorySemanticsMask(Order);
479473
auto SPIRVScope = getScope(Scope);

‎sycl/include/CL/sycl/info/context_traits.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
11
__SYCL_PARAM_TRAITS_SPEC(context, reference_count, cl_uint)
22
__SYCL_PARAM_TRAITS_SPEC(context, platform, cl::sycl::platform)
33
__SYCL_PARAM_TRAITS_SPEC(context, devices, std::vector<cl::sycl::device>)
4+
__SYCL_PARAM_TRAITS_SPEC(context, atomic_memory_order_capabilities, std::vector<cl::sycl::memory_order>)

‎sycl/include/CL/sycl/info/device_traits.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,8 @@ __SYCL_PARAM_TRAITS_SPEC(device, address_bits, pi_uint32)
2323
__SYCL_PARAM_TRAITS_SPEC(device, max_mem_alloc_size, pi_uint64)
2424
__SYCL_PARAM_TRAITS_SPEC(device, image_support, bool)
2525
__SYCL_PARAM_TRAITS_SPEC(device, atomic64, bool)
26+
__SYCL_PARAM_TRAITS_SPEC(device, atomic_memory_order_capabilities,
27+
std::vector<cl::sycl::memory_order>)
2628
__SYCL_PARAM_TRAITS_SPEC(device, max_read_image_args, pi_uint32)
2729
__SYCL_PARAM_TRAITS_SPEC(device, max_write_image_args, pi_uint32)
2830
__SYCL_PARAM_TRAITS_SPEC(device, image2d_max_width, size_t)

‎sycl/include/CL/sycl/info/info_desc.hpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,8 @@ enum class context : cl_context_info {
3737
reference_count = CL_CONTEXT_REFERENCE_COUNT,
3838
platform = CL_CONTEXT_PLATFORM,
3939
devices = CL_CONTEXT_DEVICES,
40+
atomic_memory_order_capabilities =
41+
PI_CONTEXT_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES,
4042
};
4143

4244
// A.3 Device information descriptors
@@ -149,7 +151,9 @@ enum class device : cl_device_info {
149151
ext_intel_mem_channel = PI_MEM_PROPERTIES_CHANNEL,
150152
ext_oneapi_srgb = PI_DEVICE_INFO_IMAGE_SRGB,
151153
ext_intel_device_info_uuid = PI_DEVICE_INFO_UUID,
152-
atomic64 = PI_DEVICE_INFO_ATOMIC_64
154+
atomic64 = PI_DEVICE_INFO_ATOMIC_64,
155+
atomic_memory_order_capabilities =
156+
PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES,
153157
};
154158

155159
enum class device_type : pi_uint64 {

‎sycl/include/CL/sycl/memory_enums.hpp

Lines changed: 26 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -8,11 +8,8 @@
88

99
#pragma once
1010

11-
#include <sycl/ext/oneapi/atomic_enums.hpp>
12-
1311
__SYCL_INLINE_NAMESPACE(cl) {
1412
namespace sycl {
15-
using ext::oneapi::memory_scope;
1613

1714
enum class memory_order : int {
1815
relaxed = 0,
@@ -24,6 +21,14 @@ enum class memory_order : int {
2421
seq_cst = 5
2522
};
2623

24+
enum class memory_scope : int {
25+
work_item = 0,
26+
sub_group = 1,
27+
work_group = 2,
28+
device = 3,
29+
system = 4
30+
};
31+
2732
#if __cplusplus >= 201703L
2833
inline constexpr auto memory_scope_work_item = memory_scope::work_item;
2934
inline constexpr auto memory_scope_sub_group = memory_scope::sub_group;
@@ -38,9 +43,25 @@ inline constexpr auto memory_order_acq_rel = memory_order::acq_rel;
3843
inline constexpr auto memory_order_seq_cst = memory_order::seq_cst;
3944
#endif
4045

41-
#ifndef __SYCL_DEVICE_ONLY__
4246
namespace detail {
4347

48+
inline std::vector<memory_order>
49+
readMemoryOrderBitfield(pi_memory_order_capabilities bits) {
50+
std::vector<memory_order> result;
51+
if (bits & PI_MEMORY_ORDER_RELAXED)
52+
result.push_back(memory_order::relaxed);
53+
if (bits & PI_MEMORY_ORDER_ACQUIRE)
54+
result.push_back(memory_order::acquire);
55+
if (bits & PI_MEMORY_ORDER_RELEASE)
56+
result.push_back(memory_order::release);
57+
if (bits & PI_MEMORY_ORDER_ACQ_REL)
58+
result.push_back(memory_order::acq_rel);
59+
if (bits & PI_MEMORY_ORDER_SEQ_CST)
60+
result.push_back(memory_order::seq_cst);
61+
return result;
62+
}
63+
64+
#ifndef __SYCL_DEVICE_ONLY__
4465
static constexpr std::memory_order getStdMemoryOrder(sycl::memory_order order) {
4566
switch (order) {
4667
case memory_order::relaxed:
@@ -57,8 +78,8 @@ static constexpr std::memory_order getStdMemoryOrder(sycl::memory_order order) {
5778
return std::memory_order_seq_cst;
5879
}
5980
}
81+
#endif // __SYCL_DEVICE_ONLY__
6082

6183
} // namespace detail
62-
#endif // __SYCL_DEVICE_ONLY__
6384
} // namespace sycl
6485
} // __SYCL_INLINE_NAMESPACE(cl)

0 commit comments

Comments
 (0)