Skip to content

Commit 0ccb0b7

Browse files
konradkusiak97konradkusiak97
authored
[SYCL][ABI-Break] Improve Queue fill (#13788)
Changed the `queue.fill()` implementation to make use of the native functions for a specific backend. Also, unified the implementation with the one for memset, since it is just an 8-bit subset operation of fill. In the CUDA case, both memset and fill are currently calling `urEnqueueUSMFill` which depending on the size of the filling pattern calls either `cuMemsetD8Async`, `cuMemsetD16Async`, `cuMemsetD32Async` or `commonMemSetLargePattern`. Before this patch memset was using the same thing, just beforehand setting patternSize always to 1 byte which resulted in calling `cuMemsetD8Async`. In other backends, the behaviour is analogous. The fill method was just invoking a `parallel_for` to fill the memory with the pattern which was making this operation quite slow.
1 parent b026de4 commit 0ccb0b7

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

44 files changed

+266
-198
lines changed

‎sycl/doc/design/CommandGraph.md

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -537,6 +537,8 @@ The types of commands which are unsupported, and lead to this exception are:
537537
This corresponds to a memory buffer write command.
538538
* `handler::copy(src, dest)` or `handler::memcpy(dest, src)` - Where both `src` and
539539
`dest` are USM pointers. This corresponds to a USM copy command.
540+
* `handler::fill(ptr, pattern, count)` - This corresponds to a USM memory
541+
fill command.
540542
* `handler::memset(ptr, value, numBytes)` - This corresponds to a USM memory
541543
fill command.
542544
* `handler::prefetch()`.

‎sycl/include/sycl/detail/cg.hpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -147,7 +147,7 @@ class CG {
147147
getAuxiliaryResources() const {
148148
return {};
149149
}
150-
virtual void clearAuxiliaryResources(){};
150+
virtual void clearAuxiliaryResources() {};
151151

152152
virtual ~CG() = default;
153153

@@ -247,11 +247,11 @@ class CGCopy : public CG {
247247
/// "Fill memory" command group class.
248248
class CGFill : public CG {
249249
public:
250-
std::vector<char> MPattern;
250+
std::vector<unsigned char> MPattern;
251251
AccessorImplHost *MPtr;
252252

253-
CGFill(std::vector<char> Pattern, void *Ptr, CG::StorageInitHelper CGData,
254-
detail::code_location loc = {})
253+
CGFill(std::vector<unsigned char> Pattern, void *Ptr,
254+
CG::StorageInitHelper CGData, detail::code_location loc = {})
255255
: CG(Fill, std::move(CGData), std::move(loc)),
256256
MPattern(std::move(Pattern)), MPtr((AccessorImplHost *)Ptr) {}
257257
AccessorImplHost *getReqToFill() { return MPtr; }
@@ -289,18 +289,18 @@ class CGCopyUSM : public CG {
289289

290290
/// "Fill USM" command group class.
291291
class CGFillUSM : public CG {
292-
std::vector<char> MPattern;
292+
std::vector<unsigned char> MPattern;
293293
void *MDst;
294294
size_t MLength;
295295

296296
public:
297-
CGFillUSM(std::vector<char> Pattern, void *DstPtr, size_t Length,
297+
CGFillUSM(std::vector<unsigned char> Pattern, void *DstPtr, size_t Length,
298298
CG::StorageInitHelper CGData, detail::code_location loc = {})
299299
: CG(FillUSM, std::move(CGData), std::move(loc)),
300300
MPattern(std::move(Pattern)), MDst(DstPtr), MLength(Length) {}
301301
void *getDst() { return MDst; }
302302
size_t getLength() { return MLength; }
303-
int getFill() { return MPattern[0]; }
303+
const std::vector<unsigned char> &getPattern() { return MPattern; }
304304
};
305305

306306
/// "Prefetch USM" command group class.
@@ -378,14 +378,14 @@ class CGCopy2DUSM : public CG {
378378

379379
/// "Fill 2D USM" command group class.
380380
class CGFill2DUSM : public CG {
381-
std::vector<char> MPattern;
381+
std::vector<unsigned char> MPattern;
382382
void *MDst;
383383
size_t MPitch;
384384
size_t MWidth;
385385
size_t MHeight;
386386

387387
public:
388-
CGFill2DUSM(std::vector<char> Pattern, void *DstPtr, size_t Pitch,
388+
CGFill2DUSM(std::vector<unsigned char> Pattern, void *DstPtr, size_t Pitch,
389389
size_t Width, size_t Height, CG::StorageInitHelper CGData,
390390
detail::code_location loc = {})
391391
: CG(Fill2DUSM, std::move(CGData), std::move(loc)),
@@ -395,7 +395,7 @@ class CGFill2DUSM : public CG {
395395
size_t getPitch() const { return MPitch; }
396396
size_t getWidth() const { return MWidth; }
397397
size_t getHeight() const { return MHeight; }
398-
const std::vector<char> &getPattern() const { return MPattern; }
398+
const std::vector<unsigned char> &getPattern() const { return MPattern; }
399399
};
400400

401401
/// "Memset 2D USM" command group class.

‎sycl/include/sycl/detail/pi.def

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -129,7 +129,7 @@ _PI_API(piextUSMHostAlloc)
129129
_PI_API(piextUSMDeviceAlloc)
130130
_PI_API(piextUSMSharedAlloc)
131131
_PI_API(piextUSMFree)
132-
_PI_API(piextUSMEnqueueMemset)
132+
_PI_API(piextUSMEnqueueFill)
133133
_PI_API(piextUSMEnqueueMemcpy)
134134
_PI_API(piextUSMEnqueuePrefetch)
135135
_PI_API(piextUSMEnqueueMemAdvise)

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

Lines changed: 14 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -196,9 +196,10 @@
196196
// _pi_virtual_mem_granularity_info enum, _pi_virtual_mem_info enum and
197197
// pi_virtual_access_flags bit flags.
198198
// 15.55 Added piextEnqueueNativeCommand as well as associated types and enums
199+
// 16.56 Replaced piextUSMEnqueueMemset with piextUSMEnqueueFill
199200

200-
#define _PI_H_VERSION_MAJOR 15
201-
#define _PI_H_VERSION_MINOR 55
201+
#define _PI_H_VERSION_MAJOR 16
202+
#define _PI_H_VERSION_MINOR 56
202203

203204
#define _PI_STRING_HELPER(a) #a
204205
#define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -2174,22 +2175,22 @@ __SYCL_EXPORT pi_result piextUSMPitchedAlloc(
21742175
/// \param ptr is the memory to be freed
21752176
__SYCL_EXPORT pi_result piextUSMFree(pi_context context, void *ptr);
21762177

2177-
/// USM Memset API
2178+
/// USM Fill API
21782179
///
21792180
/// \param queue is the queue to submit to
2180-
/// \param ptr is the ptr to memset
2181-
/// \param value is value to set. It is interpreted as an 8-bit value and the
2182-
/// upper
2183-
/// 24 bits are ignored
2184-
/// \param count is the size in bytes to memset
2181+
/// \param ptr is the ptr to fill
2182+
/// \param pattern is the ptr with the bytes of the pattern to set
2183+
/// \param patternSize is the size in bytes of the pattern to set
2184+
/// \param count is the size in bytes to fill
21852185
/// \param num_events_in_waitlist is the number of events to wait on
21862186
/// \param events_waitlist is an array of events to wait on
21872187
/// \param event is the event that represents this operation
2188-
__SYCL_EXPORT pi_result piextUSMEnqueueMemset(pi_queue queue, void *ptr,
2189-
pi_int32 value, size_t count,
2190-
pi_uint32 num_events_in_waitlist,
2191-
const pi_event *events_waitlist,
2192-
pi_event *event);
2188+
__SYCL_EXPORT pi_result piextUSMEnqueueFill(pi_queue queue, void *ptr,
2189+
const void *pattern,
2190+
size_t patternSize, size_t count,
2191+
pi_uint32 num_events_in_waitlist,
2192+
const pi_event *events_waitlist,
2193+
pi_event *event);
21932194

21942195
/// USM Memcpy API
21952196
///

‎sycl/include/sycl/handler.hpp

Lines changed: 16 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2826,10 +2826,14 @@ class __SYCL_EXPORT handler {
28262826
setUserFacingNodeType(ext::oneapi::experimental::node_type::memfill);
28272827
static_assert(is_device_copyable<T>::value,
28282828
"Pattern must be device copyable");
2829-
parallel_for<__usmfill<T>>(range<1>(Count), [=](id<1> Index) {
2830-
T *CastedPtr = static_cast<T *>(Ptr);
2831-
CastedPtr[Index] = Pattern;
2832-
});
2829+
if (getDeviceBackend() == backend::ext_oneapi_level_zero) {
2830+
parallel_for<__usmfill<T>>(range<1>(Count), [=](id<1> Index) {
2831+
T *CastedPtr = static_cast<T *>(Ptr);
2832+
CastedPtr[Index] = Pattern;
2833+
});
2834+
} else {
2835+
this->fill_impl(Ptr, &Pattern, sizeof(T), Count);
2836+
}
28332837
}
28342838

28352839
/// Prevents any commands submitted afterward to this queue from executing
@@ -3297,7 +3301,7 @@ class __SYCL_EXPORT handler {
32973301
/// Length to copy or fill (for USM operations).
32983302
size_t MLength = 0;
32993303
/// Pattern that is used to fill memory object in case command type is fill.
3300-
std::vector<char> MPattern;
3304+
std::vector<unsigned char> MPattern;
33013305
/// Storage for a lambda or function object.
33023306
std::unique_ptr<detail::HostKernelBase> MHostKernel;
33033307
/// Storage for lambda/function when using HostTask
@@ -3442,6 +3446,10 @@ class __SYCL_EXPORT handler {
34423446
// Helper function for getting a loose bound on work-items.
34433447
id<2> computeFallbackKernelBounds(size_t Width, size_t Height);
34443448

3449+
// Function to get information about the backend for which the code is
3450+
// compiled for
3451+
backend getDeviceBackend() const;
3452+
34453453
// Common function for launching a 2D USM memcpy kernel to avoid redefinitions
34463454
// of the kernel from copy and memcpy.
34473455
template <typename T>
@@ -3553,6 +3561,9 @@ class __SYCL_EXPORT handler {
35533561
});
35543562
}
35553563

3564+
// Implementation of USM fill using command for native fill.
3565+
void fill_impl(void *Dest, const void *Value, size_t ValueSize, size_t Count);
3566+
35563567
// Implementation of ext_oneapi_memcpy2d using command for native 2D memcpy.
35573568
void ext_oneapi_memcpy2d_impl(void *Dest, size_t DestPitch, const void *Src,
35583569
size_t SrcPitch, size_t Width, size_t Height);

‎sycl/plugins/cuda/pi_cuda.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -930,12 +930,12 @@ pi_result piextKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
930930
return pi2ur::piextKernelSetArgPointer(Kernel, ArgIndex, ArgSize, ArgValue);
931931
}
932932

933-
pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr, pi_int32 Value,
934-
size_t Count, pi_uint32 NumEventsInWaitlist,
935-
const pi_event *EventsWaitlist,
936-
pi_event *Event) {
937-
return pi2ur::piextUSMEnqueueMemset(
938-
Queue, Ptr, Value, Count, NumEventsInWaitlist, EventsWaitlist, Event);
933+
pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr, const void *Pattern,
934+
size_t PatternSize, size_t Count,
935+
pi_uint32 NumEventsInWaitlist,
936+
const pi_event *EventsWaitlist, pi_event *Event) {
937+
return pi2ur::piextUSMEnqueueFill(Queue, Ptr, Pattern, PatternSize, Count,
938+
NumEventsInWaitlist, EventsWaitlist, Event);
939939
}
940940

941941
pi_result piextUSMEnqueueMemcpy(pi_queue Queue, pi_bool Blocking, void *DstPtr,

‎sycl/plugins/hip/pi_hip.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -933,12 +933,12 @@ pi_result piextKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
933933
return pi2ur::piextKernelSetArgPointer(Kernel, ArgIndex, ArgSize, ArgValue);
934934
}
935935

936-
pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr, pi_int32 Value,
937-
size_t Count, pi_uint32 NumEventsInWaitlist,
938-
const pi_event *EventsWaitlist,
939-
pi_event *Event) {
940-
return pi2ur::piextUSMEnqueueMemset(
941-
Queue, Ptr, Value, Count, NumEventsInWaitlist, EventsWaitlist, Event);
936+
pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr, const void *Pattern,
937+
size_t PatternSize, size_t Count,
938+
pi_uint32 NumEventsInWaitlist,
939+
const pi_event *EventsWaitlist, pi_event *Event) {
940+
return pi2ur::piextUSMEnqueueFill(Queue, Ptr, Pattern, PatternSize, Count,
941+
NumEventsInWaitlist, EventsWaitlist, Event);
942942
}
943943

944944
pi_result piextUSMEnqueueMemcpy(pi_queue Queue, pi_bool Blocking, void *DstPtr,

‎sycl/plugins/level_zero/pi_level_zero.cpp

Lines changed: 11 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -957,23 +957,22 @@ pi_result piextKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
957957
return pi2ur::piextKernelSetArgPointer(Kernel, ArgIndex, ArgSize, ArgValue);
958958
}
959959

960-
/// USM Memset API
960+
/// USM Fill API
961961
///
962962
/// @param Queue is the queue to submit to
963-
/// @param Ptr is the ptr to memset
964-
/// @param Value is value to set. It is interpreted as an 8-bit value and the
965-
/// upper
966-
/// 24 bits are ignored
967-
/// @param Count is the size in bytes to memset
963+
/// @param Ptr is the ptr to fill
964+
/// \param Pattern is the ptr with the bytes of the pattern to set
965+
/// \param PatternSize is the size in bytes of the pattern to set
966+
/// @param Count is the size in bytes to fill
968967
/// @param NumEventsInWaitlist is the number of events to wait on
969968
/// @param EventsWaitlist is an array of events to wait on
970969
/// @param Event is the event that represents this operation
971-
pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr, pi_int32 Value,
972-
size_t Count, pi_uint32 NumEventsInWaitlist,
973-
const pi_event *EventsWaitlist,
974-
pi_event *Event) {
975-
return pi2ur::piextUSMEnqueueMemset(
976-
Queue, Ptr, Value, Count, NumEventsInWaitlist, EventsWaitlist, Event);
970+
pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr, const void *Pattern,
971+
size_t PatternSize, size_t Count,
972+
pi_uint32 NumEventsInWaitlist,
973+
const pi_event *EventsWaitlist, pi_event *Event) {
974+
return pi2ur::piextUSMEnqueueFill(Queue, Ptr, Pattern, PatternSize, Count,
975+
NumEventsInWaitlist, EventsWaitlist, Event);
977976
}
978977

979978
pi_result piextUSMEnqueueMemcpy(pi_queue Queue, pi_bool Blocking, void *DstPtr,

‎sycl/plugins/native_cpu/pi_native_cpu.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -933,12 +933,12 @@ pi_result piextKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
933933
return pi2ur::piextKernelSetArgPointer(Kernel, ArgIndex, ArgSize, ArgValue);
934934
}
935935

936-
pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr, pi_int32 Value,
937-
size_t Count, pi_uint32 NumEventsInWaitlist,
938-
const pi_event *EventsWaitlist,
939-
pi_event *Event) {
940-
return pi2ur::piextUSMEnqueueMemset(
941-
Queue, Ptr, Value, Count, NumEventsInWaitlist, EventsWaitlist, Event);
936+
pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr, const void *Pattern,
937+
size_t PatternSize, size_t Count,
938+
pi_uint32 NumEventsInWaitlist,
939+
const pi_event *EventsWaitlist, pi_event *Event) {
940+
return pi2ur::piextUSMEnqueueFill(Queue, Ptr, Pattern, PatternSize, Count,
941+
NumEventsInWaitlist, EventsWaitlist, Event);
942942
}
943943

944944
pi_result piextUSMEnqueueMemcpy(pi_queue Queue, pi_bool Blocking, void *DstPtr,

‎sycl/plugins/opencl/pi_opencl.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -889,12 +889,12 @@ pi_result piextKernelSetArgPointer(pi_kernel Kernel, pi_uint32 ArgIndex,
889889
return pi2ur::piextKernelSetArgPointer(Kernel, ArgIndex, ArgSize, ArgValue);
890890
}
891891

892-
pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr, pi_int32 Value,
893-
size_t Count, pi_uint32 NumEventsInWaitlist,
894-
const pi_event *EventsWaitlist,
895-
pi_event *Event) {
896-
return pi2ur::piextUSMEnqueueMemset(
897-
Queue, Ptr, Value, Count, NumEventsInWaitlist, EventsWaitlist, Event);
892+
pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr, const void *Pattern,
893+
size_t PatternSize, size_t Count,
894+
pi_uint32 NumEventsInWaitlist,
895+
const pi_event *EventsWaitlist, pi_event *Event) {
896+
return pi2ur::piextUSMEnqueueFill(Queue, Ptr, Pattern, PatternSize, Count,
897+
NumEventsInWaitlist, EventsWaitlist, Event);
898898
}
899899

900900
pi_result piextUSMEnqueueMemcpy(pi_queue Queue, pi_bool Blocking, void *DstPtr,

‎sycl/plugins/unified_runtime/pi2ur.hpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -3913,11 +3913,12 @@ inline pi_result piEnqueueMemBufferFill(pi_queue Queue, pi_mem Buffer,
39133913
return PI_SUCCESS;
39143914
}
39153915

3916-
inline pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr,
3917-
pi_int32 Value, size_t Count,
3918-
pi_uint32 NumEventsInWaitList,
3919-
const pi_event *EventsWaitList,
3920-
pi_event *OutEvent) {
3916+
inline pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr,
3917+
const void *Pattern, size_t PatternSize,
3918+
size_t Count,
3919+
pi_uint32 NumEventsInWaitList,
3920+
const pi_event *EventsWaitList,
3921+
pi_event *OutEvent) {
39213922
PI_ASSERT(Queue, PI_ERROR_INVALID_QUEUE);
39223923
if (!Ptr) {
39233924
return PI_ERROR_INVALID_VALUE;
@@ -3929,8 +3930,7 @@ inline pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr,
39293930

39303931
ur_event_handle_t *UREvent = reinterpret_cast<ur_event_handle_t *>(OutEvent);
39313932

3932-
size_t PatternSize = 1;
3933-
HANDLE_ERRORS(urEnqueueUSMFill(UrQueue, Ptr, PatternSize, &Value, Count,
3933+
HANDLE_ERRORS(urEnqueueUSMFill(UrQueue, Ptr, PatternSize, Pattern, Count,
39343934
NumEventsInWaitList, UrEventsWaitList,
39353935
UREvent));
39363936

‎sycl/plugins/unified_runtime/pi_unified_runtime.cpp

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -442,24 +442,24 @@ __SYCL_EXPORT pi_result piQueueGetInfo(pi_queue Queue, pi_queue_info ParamName,
442442
ParamValueSizeRet);
443443
}
444444

445-
/// USM Memset API
445+
/// USM Fill API
446446
///
447-
/// @param Queue is the queue to submit to
448-
/// @param Ptr is the ptr to memset
449-
/// @param Value is value to set. It is interpreted as an 8-bit value and the
450-
/// upper
451-
/// 24 bits are ignored
452-
/// @param Count is the size in bytes to memset
453-
/// @param NumEventsInWaitlist is the number of events to wait on
454-
/// @param EventsWaitlist is an array of events to wait on
455-
/// @param Event is the event that represents this operation
456-
__SYCL_EXPORT pi_result piextUSMEnqueueMemset(pi_queue Queue, void *Ptr,
457-
pi_int32 Value, size_t Count,
458-
pi_uint32 NumEventsInWaitlist,
459-
const pi_event *EventsWaitlist,
460-
pi_event *Event) {
461-
return pi2ur::piextUSMEnqueueMemset(
462-
Queue, Ptr, Value, Count, NumEventsInWaitlist, EventsWaitlist, Event);
447+
/// \param queue is the queue to submit to
448+
/// \param ptr is the ptr to fill
449+
/// \param pattern is the ptr with the bytes of the pattern to set
450+
/// \param patternSize is the size in bytes of the pattern to set
451+
/// \param count is the size in bytes to fill
452+
/// \param num_events_in_waitlist is the number of events to wait on
453+
/// \param events_waitlist is an array of events to wait on
454+
/// \param event is the event that represents this operation
455+
__SYCL_EXPORT pi_result piextUSMEnqueueFill(pi_queue Queue, void *Ptr,
456+
const void *Pattern,
457+
size_t PatternSize, size_t Count,
458+
pi_uint32 NumEventsInWaitlist,
459+
const pi_event *EventsWaitlist,
460+
pi_event *Event) {
461+
return pi2ur::piextUSMEnqueueFill(Queue, Ptr, Pattern, PatternSize, Count,
462+
NumEventsInWaitlist, EventsWaitlist, Event);
463463
}
464464

465465
__SYCL_EXPORT pi_result piEnqueueMemBufferCopyRect(
@@ -1598,7 +1598,7 @@ __SYCL_EXPORT pi_result piPluginInit(pi_plugin *PluginInit) {
15981598
_PI_API(piEnqueueMemBufferMap)
15991599
_PI_API(piEnqueueMemUnmap)
16001600
_PI_API(piEnqueueMemBufferFill)
1601-
_PI_API(piextUSMEnqueueMemset)
1601+
_PI_API(piextUSMEnqueueFill)
16021602
_PI_API(piEnqueueMemBufferCopyRect)
16031603
_PI_API(piEnqueueMemBufferCopy)
16041604
_PI_API(piextUSMEnqueueMemcpy)

‎sycl/source/detail/graph_impl.hpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -703,8 +703,10 @@ class node_impl {
703703
sycl::detail::CGFillUSM *FillUSM =
704704
static_cast<sycl::detail::CGFillUSM *>(MCommandGroup.get());
705705
Stream << "Dst: " << FillUSM->getDst()
706-
<< " Length: " << FillUSM->getLength()
707-
<< " Pattern: " << FillUSM->getFill() << "\\n";
706+
<< " Length: " << FillUSM->getLength() << " Pattern: ";
707+
for (auto byte : FillUSM->getPattern())
708+
Stream << byte;
709+
Stream << "\\n";
708710
}
709711
break;
710712
case sycl::detail::CG::CGTYPE::PrefetchUSM:

0 commit comments

Comments
 (0)