Skip to content

Commit 4ae6867

Browse files
author
Steffen Larsen
committed
[SYCL] Adds info query for atomic_memory_order_capabilities on device and context
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. Signed-off-by: Steffen Larsen <[email protected]>
1 parent ea4b8a9 commit 4ae6867

23 files changed

+200
-121
lines changed

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

Lines changed: 13 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,15 @@ 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+
typedef enum {
517+
PI_MEMORY_ORDER_RELAXED = 0x01,
518+
PI_MEMORY_ORDER_ACQUIRE = 0x02,
519+
PI_MEMORY_ORDER_RELEASE = 0x04,
520+
PI_MEMORY_ORDER_ACQ_REL = 0x08,
521+
PI_MEMORY_ORDER_SEQ_CST = 0x10,
522+
} _pi_memory_order_capability;
523+
512524
typedef enum {
513525
PI_PROFILING_INFO_COMMAND_QUEUED = CL_PROFILING_COMMAND_QUEUED,
514526
PI_PROFILING_INFO_COMMAND_SUBMIT = CL_PROFILING_COMMAND_SUBMIT,

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

Lines changed: 47 additions & 51 deletions
Original file line numberDiff line numberDiff line change
@@ -224,11 +224,11 @@ 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+
std::is_same<T, sycl::memory_order>::value,
230+
__spv::MemorySemanticsMask::Flag>::type
231+
getMemorySemanticsMask(T Order) {
232232
__spv::MemorySemanticsMask::Flag SpvOrder = __spv::MemorySemanticsMask::None;
233233
switch (Order) {
234234
case T::relaxed:
@@ -255,27 +255,25 @@ getMemorySemanticsMask(T Order) {
255255
}
256256

257257
static inline constexpr __spv::Scope::Flag
258-
getScope(ext::oneapi::memory_scope Scope) {
258+
getScope(memory_scope Scope) {
259259
switch (Scope) {
260-
case ext::oneapi::memory_scope::work_item:
260+
case memory_scope::work_item:
261261
return __spv::Scope::Invocation;
262-
case ext::oneapi::memory_scope::sub_group:
262+
case memory_scope::sub_group:
263263
return __spv::Scope::Subgroup;
264-
case ext::oneapi::memory_scope::work_group:
264+
case memory_scope::work_group:
265265
return __spv::Scope::Workgroup;
266-
case ext::oneapi::memory_scope::device:
266+
case memory_scope::device:
267267
return __spv::Scope::Device;
268-
case ext::oneapi::memory_scope::system:
268+
case memory_scope::system:
269269
return __spv::Scope::CrossDevice;
270270
}
271271
}
272272

273273
template <typename T, access::address_space AddressSpace>
274274
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,
275+
AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
276+
memory_order Success, memory_order Failure, T Desired,
279277
T Expected) {
280278
auto SPIRVSuccess = getMemorySemanticsMask(Success);
281279
auto SPIRVFailure = getMemorySemanticsMask(Failure);
@@ -287,10 +285,8 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
287285

288286
template <typename T, access::address_space AddressSpace>
289287
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,
288+
AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
289+
memory_order Success, memory_order Failure, T Desired,
294290
T Expected) {
295291
using I = detail::make_unsinged_integer_t<T>;
296292
auto SPIRVSuccess = getMemorySemanticsMask(Success);
@@ -308,8 +304,8 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
308304

309305
template <typename T, access::address_space AddressSpace>
310306
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) {
307+
AtomicLoad(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
308+
memory_order Order) {
313309
auto *Ptr = MPtr.get();
314310
auto SPIRVOrder = getMemorySemanticsMask(Order);
315311
auto SPIRVScope = getScope(Scope);
@@ -318,8 +314,8 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
318314

319315
template <typename T, access::address_space AddressSpace>
320316
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) {
317+
AtomicLoad(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
318+
memory_order Order) {
323319
using I = detail::make_unsinged_integer_t<T>;
324320
auto *PtrInt =
325321
reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -332,8 +328,8 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
332328

333329
template <typename T, access::address_space AddressSpace>
334330
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) {
331+
AtomicStore(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
332+
memory_order Order, T Value) {
337333
auto *Ptr = MPtr.get();
338334
auto SPIRVOrder = getMemorySemanticsMask(Order);
339335
auto SPIRVScope = getScope(Scope);
@@ -342,8 +338,8 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
342338

343339
template <typename T, access::address_space AddressSpace>
344340
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) {
341+
AtomicStore(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
342+
memory_order Order, T Value) {
347343
using I = detail::make_unsinged_integer_t<T>;
348344
auto *PtrInt =
349345
reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -356,8 +352,8 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
356352

357353
template <typename T, access::address_space AddressSpace>
358354
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) {
355+
AtomicExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
356+
memory_order Order, T Value) {
361357
auto *Ptr = MPtr.get();
362358
auto SPIRVOrder = getMemorySemanticsMask(Order);
363359
auto SPIRVScope = getScope(Scope);
@@ -366,8 +362,8 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
366362

367363
template <typename T, access::address_space AddressSpace>
368364
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) {
365+
AtomicExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
366+
memory_order Order, T Value) {
371367
using I = detail::make_unsinged_integer_t<T>;
372368
auto *PtrInt =
373369
reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -382,8 +378,8 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
382378

383379
template <typename T, access::address_space AddressSpace>
384380
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) {
381+
AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
382+
memory_order Order, T Value) {
387383
auto *Ptr = MPtr.get();
388384
auto SPIRVOrder = getMemorySemanticsMask(Order);
389385
auto SPIRVScope = getScope(Scope);
@@ -392,8 +388,8 @@ AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
392388

393389
template <typename T, access::address_space AddressSpace>
394390
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) {
391+
AtomicISub(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
392+
memory_order Order, T Value) {
397393
auto *Ptr = MPtr.get();
398394
auto SPIRVOrder = getMemorySemanticsMask(Order);
399395
auto SPIRVScope = getScope(Scope);
@@ -402,8 +398,8 @@ AtomicISub(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
402398

403399
template <typename T, access::address_space AddressSpace>
404400
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) {
401+
AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
402+
memory_order Order, T Value) {
407403
auto *Ptr = MPtr.get();
408404
auto SPIRVOrder = getMemorySemanticsMask(Order);
409405
auto SPIRVScope = getScope(Scope);
@@ -412,8 +408,8 @@ AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
412408

413409
template <typename T, access::address_space AddressSpace>
414410
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) {
411+
AtomicAnd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
412+
memory_order Order, T Value) {
417413
auto *Ptr = MPtr.get();
418414
auto SPIRVOrder = getMemorySemanticsMask(Order);
419415
auto SPIRVScope = getScope(Scope);
@@ -422,8 +418,8 @@ AtomicAnd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
422418

423419
template <typename T, access::address_space AddressSpace>
424420
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) {
421+
AtomicOr(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
422+
memory_order Order, T Value) {
427423
auto *Ptr = MPtr.get();
428424
auto SPIRVOrder = getMemorySemanticsMask(Order);
429425
auto SPIRVScope = getScope(Scope);
@@ -432,8 +428,8 @@ AtomicOr(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
432428

433429
template <typename T, access::address_space AddressSpace>
434430
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) {
431+
AtomicXor(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
432+
memory_order Order, T Value) {
437433
auto *Ptr = MPtr.get();
438434
auto SPIRVOrder = getMemorySemanticsMask(Order);
439435
auto SPIRVScope = getScope(Scope);
@@ -442,8 +438,8 @@ AtomicXor(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
442438

443439
template <typename T, access::address_space AddressSpace>
444440
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) {
441+
AtomicMin(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
442+
memory_order Order, T Value) {
447443
auto *Ptr = MPtr.get();
448444
auto SPIRVOrder = getMemorySemanticsMask(Order);
449445
auto SPIRVScope = getScope(Scope);
@@ -452,8 +448,8 @@ AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
452448

453449
template <typename T, access::address_space AddressSpace>
454450
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) {
451+
AtomicMin(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
452+
memory_order Order, T Value) {
457453
auto *Ptr = MPtr.get();
458454
auto SPIRVOrder = getMemorySemanticsMask(Order);
459455
auto SPIRVScope = getScope(Scope);
@@ -462,8 +458,8 @@ AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
462458

463459
template <typename T, access::address_space AddressSpace>
464460
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) {
461+
AtomicMax(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
462+
memory_order Order, T Value) {
467463
auto *Ptr = MPtr.get();
468464
auto SPIRVOrder = getMemorySemanticsMask(Order);
469465
auto SPIRVScope = getScope(Scope);
@@ -472,8 +468,8 @@ AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
472468

473469
template <typename T, access::address_space AddressSpace>
474470
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) {
471+
AtomicMax(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
472+
memory_order Order, T Value) {
477473
auto *Ptr = MPtr.get();
478474
auto SPIRVOrder = getMemorySemanticsMask(Order);
479475
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_capability 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)