Skip to content

Commit c6fbac5

Browse files
mfrancepilloisEwanC
mfrancepillois
and
EwanC
authored
[SYCL][Graph] Support for Prefetch and memory advise (#11474)
Adds support for `prefetch` and `advise` memory hints. Adds e2e tests that verify that backend functions are called as they should be. --------- Co-authored-by: Ewan Crawford <ewan@codeplay.com>
1 parent a6d5b98 commit c6fbac5

27 files changed

+536
-9
lines changed

‎sycl/doc/design/CommandGraph.md

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -351,6 +351,8 @@ The types of commands which are unsupported, and lead to this exception are:
351351
`dest` are USM pointers. This corresponds to a USM copy command.
352352
* `handler::memset(ptr, value, numBytes)` - This corresponds to a USM memory
353353
fill command.
354+
* `handler::prefetch()`.
355+
* `handler::mem_advise()`.
354356

355357
Note that `handler::copy(src, dest)` where both `src` and `dest` are an accessor
356358
is supported, as a memory buffer copy command exists in the OpenCL extension.
@@ -377,6 +379,8 @@ adapter where there is matching support for each function in the list.
377379
| urCommandBufferAppendMemBufferWriteRectExp | | No |
378380
| urCommandBufferAppendMemBufferReadRectExp | | No |
379381
| urCommandBufferAppendMemBufferFillExp | clCommandFillBufferKHR | Yes |
382+
| urCommandBufferAppendUSMPrefetchExp | | No |
383+
| urCommandBufferAppendUSMAdviseExp | | No |
380384
| urCommandBufferEnqueueExp | clEnqueueCommandBufferKHR | Yes |
381385
| | clCommandBarrierWithWaitListKHR | No |
382386
| | clCommandCopyImageKHR | No |

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -178,6 +178,8 @@ _PI_API(piextCommandBufferMemBufferRead)
178178
_PI_API(piextCommandBufferMemBufferReadRect)
179179
_PI_API(piextCommandBufferMemBufferFill)
180180
_PI_API(piextCommandBufferFillUSM)
181+
_PI_API(piextCommandBufferPrefetchUSM)
182+
_PI_API(piextCommandBufferAdviseUSM)
181183
_PI_API(piextEnqueueCommandBuffer)
182184

183185
_PI_API(piextUSMPitchedAlloc)

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

Lines changed: 34 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,9 +148,10 @@
148148
// 14.39 Added PI_EXT_INTEL_DEVICE_INFO_ESIMD_SUPPORT device info query.
149149
// 14.40 Add HIP _pi_mem_advice alises to match the PI_MEM_ADVICE_CUDA* ones.
150150
// 14.41 Added piextCommandBufferMemBufferFill & piextCommandBufferFillUSM
151+
// 14.42 Added piextCommandBufferPrefetchUSM and piextCommandBufferAdviseUSM
151152

152153
#define _PI_H_VERSION_MAJOR 14
153-
#define _PI_H_VERSION_MINOR 41
154+
#define _PI_H_VERSION_MINOR 42
154155

155156
#define _PI_STRING_HELPER(a) #a
156157
#define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -2524,6 +2525,38 @@ __SYCL_EXPORT pi_result piextCommandBufferFillUSM(
25242525
const pi_ext_sync_point *sync_point_wait_list,
25252526
pi_ext_sync_point *sync_point);
25262527

2528+
/// API to append a USM Prefetch command to the command-buffer.
2529+
/// \param command_buffer The command-buffer to append onto.
2530+
/// \param ptr points to the memory to migrate.
2531+
/// \param size is the number of bytes to migrate.
2532+
/// \param flags is a bitfield used to specify memory migration options.
2533+
/// \param num_sync_points_in_wait_list The number of sync points in the
2534+
/// provided wait list.
2535+
/// \param sync_point_wait_list A list of sync points that this command must
2536+
/// wait on.
2537+
/// \param sync_point The sync_point associated with this memory operation.
2538+
__SYCL_EXPORT pi_result piextCommandBufferPrefetchUSM(
2539+
pi_ext_command_buffer command_buffer, const void *ptr, size_t size,
2540+
pi_usm_migration_flags flags, pi_uint32 num_sync_points_in_wait_list,
2541+
const pi_ext_sync_point *sync_point_wait_list,
2542+
pi_ext_sync_point *sync_point);
2543+
2544+
/// API to append a USM Advise command to the command-buffer.
2545+
/// \param command_buffer The command-buffer to append onto.
2546+
/// \param ptr is the data to be advised.
2547+
/// \param length is the size in bytes of the memory to advise.
2548+
/// \param advice is device specific advice.
2549+
/// \param num_sync_points_in_wait_list The number of sync points in the
2550+
/// provided wait list.
2551+
/// \param sync_point_wait_list A list of sync points that this command must
2552+
/// wait on.
2553+
/// \param sync_point The sync_point associated with this memory operation.
2554+
__SYCL_EXPORT pi_result piextCommandBufferAdviseUSM(
2555+
pi_ext_command_buffer command_buffer, const void *ptr, size_t length,
2556+
pi_mem_advice advice, pi_uint32 num_sync_points_in_wait_list,
2557+
const pi_ext_sync_point *sync_point_wait_list,
2558+
pi_ext_sync_point *sync_point);
2559+
25272560
/// API to submit the command-buffer to queue for execution, returns an error if
25282561
/// the command-buffer is not finalized or another instance of the same
25292562
/// command-buffer is currently executing.

‎sycl/plugins/cuda/pi_cuda.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1158,6 +1158,24 @@ pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
11581158
SyncPointWaitList, SyncPoint);
11591159
}
11601160

1161+
pi_result piextCommandBufferPrefetchUSM(
1162+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
1163+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
1164+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1165+
return pi2ur::piextCommandBufferPrefetchUSM(CommandBuffer, Ptr, Size, Flags,
1166+
NumSyncPointsInWaitList,
1167+
SyncPointWaitList, SyncPoint);
1168+
}
1169+
1170+
pi_result piextCommandBufferAdviseUSM(
1171+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
1172+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
1173+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1174+
return pi2ur::piextCommandBufferAdviseUSM(CommandBuffer, Ptr, Length, Advice,
1175+
NumSyncPointsInWaitList,
1176+
SyncPointWaitList, SyncPoint);
1177+
}
1178+
11611179
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11621180
pi_queue Queue,
11631181
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/hip/pi_hip.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1166,6 +1166,24 @@ pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
11661166
SyncPointWaitList, SyncPoint);
11671167
}
11681168

1169+
pi_result piextCommandBufferPrefetchUSM(
1170+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
1171+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
1172+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1173+
return pi2ur::piextCommandBufferPrefetchUSM(CommandBuffer, Ptr, Size, Flags,
1174+
NumSyncPointsInWaitList,
1175+
SyncPointWaitList, SyncPoint);
1176+
}
1177+
1178+
pi_result piextCommandBufferAdviseUSM(
1179+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
1180+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
1181+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1182+
return pi2ur::piextCommandBufferAdviseUSM(CommandBuffer, Ptr, Length, Advice,
1183+
NumSyncPointsInWaitList,
1184+
SyncPointWaitList, SyncPoint);
1185+
}
1186+
11691187
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11701188
pi_queue Queue,
11711189
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/level_zero/pi_level_zero.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1324,6 +1324,24 @@ pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
13241324
SyncPointWaitList, SyncPoint);
13251325
}
13261326

1327+
pi_result piextCommandBufferPrefetchUSM(
1328+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
1329+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
1330+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1331+
return pi2ur::piextCommandBufferPrefetchUSM(CommandBuffer, Ptr, Size, Flags,
1332+
NumSyncPointsInWaitList,
1333+
SyncPointWaitList, SyncPoint);
1334+
}
1335+
1336+
pi_result piextCommandBufferAdviseUSM(
1337+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
1338+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
1339+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1340+
return pi2ur::piextCommandBufferAdviseUSM(CommandBuffer, Ptr, Length, Advice,
1341+
NumSyncPointsInWaitList,
1342+
SyncPointWaitList, SyncPoint);
1343+
}
1344+
13271345
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
13281346
pi_queue Queue,
13291347
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/native_cpu/pi_native_cpu.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1212,6 +1212,24 @@ pi_result piextPeerAccessGetInfo(pi_device command_device,
12121212
ParamValueSizeRet);
12131213
}
12141214

1215+
pi_result piextCommandBufferPrefetchUSM(
1216+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
1217+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
1218+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1219+
return pi2ur::piextCommandBufferPrefetchUSM(CommandBuffer, Ptr, Size, Flags,
1220+
NumSyncPointsInWaitList,
1221+
SyncPointWaitList, SyncPoint);
1222+
}
1223+
1224+
pi_result piextCommandBufferAdviseUSM(
1225+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
1226+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
1227+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1228+
return pi2ur::piextCommandBufferAdviseUSM(CommandBuffer, Ptr, Length, Advice,
1229+
NumSyncPointsInWaitList,
1230+
SyncPointWaitList, SyncPoint);
1231+
}
1232+
12151233
// Initialize function table with stubs.
12161234
#define _PI_API(api) \
12171235
(PluginInit->PiFunctionTable).api = (decltype(&::api))(&api);

‎sycl/plugins/opencl/pi_opencl.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1102,6 +1102,24 @@ pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
11021102
SyncPointWaitList, SyncPoint);
11031103
}
11041104

1105+
pi_result piextCommandBufferPrefetchUSM(
1106+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
1107+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
1108+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1109+
return pi2ur::piextCommandBufferPrefetchUSM(CommandBuffer, Ptr, Size, Flags,
1110+
NumSyncPointsInWaitList,
1111+
SyncPointWaitList, SyncPoint);
1112+
}
1113+
1114+
pi_result piextCommandBufferAdviseUSM(
1115+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
1116+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
1117+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1118+
return pi2ur::piextCommandBufferAdviseUSM(CommandBuffer, Ptr, Length, Advice,
1119+
NumSyncPointsInWaitList,
1120+
SyncPointWaitList, SyncPoint);
1121+
}
1122+
11051123
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11061124
pi_queue Queue,
11071125
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/unified_runtime/pi2ur.hpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4611,6 +4611,41 @@ inline pi_result piextCommandBufferFillUSM(
46114611
return PI_SUCCESS;
46124612
}
46134613

4614+
inline pi_result piextCommandBufferPrefetchUSM(
4615+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
4616+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
4617+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
4618+
4619+
// flags is currently unused so fail if set
4620+
PI_ASSERT(Flags == 0, PI_ERROR_INVALID_VALUE);
4621+
4622+
ur_exp_command_buffer_handle_t UrCommandBuffer =
4623+
reinterpret_cast<ur_exp_command_buffer_handle_t>(CommandBuffer);
4624+
4625+
// TODO: to map from pi_usm_migration_flags to
4626+
// ur_usm_migration_flags_t
4627+
// once we have those defined
4628+
ur_usm_migration_flags_t UrFlags{};
4629+
HANDLE_ERRORS(urCommandBufferAppendUSMPrefetchExp(
4630+
UrCommandBuffer, Ptr, Size, UrFlags, NumSyncPointsInWaitList,
4631+
SyncPointWaitList, SyncPoint));
4632+
return PI_SUCCESS;
4633+
}
4634+
4635+
inline pi_result piextCommandBufferAdviseUSM(
4636+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
4637+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
4638+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
4639+
ur_exp_command_buffer_handle_t UrCommandBuffer =
4640+
reinterpret_cast<ur_exp_command_buffer_handle_t>(CommandBuffer);
4641+
4642+
ur_usm_advice_flags_t UrAdvice{};
4643+
HANDLE_ERRORS(urCommandBufferAppendUSMAdviseExp(
4644+
UrCommandBuffer, Ptr, Length, UrAdvice, NumSyncPointsInWaitList,
4645+
SyncPointWaitList, SyncPoint));
4646+
return PI_SUCCESS;
4647+
}
4648+
46144649
inline pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
46154650
pi_queue Queue,
46164651
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/unified_runtime/pi_unified_runtime.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1123,6 +1123,24 @@ pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
11231123
SyncPointWaitList, SyncPoint);
11241124
}
11251125

1126+
pi_result piextCommandBufferPrefetchUSM(
1127+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Size,
1128+
pi_usm_migration_flags Flags, pi_uint32 NumSyncPointsInWaitList,
1129+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1130+
return pi2ur::piextCommandBufferPrefetchUSM(CommandBuffer, Ptr, Size, Flags,
1131+
NumSyncPointsInWaitList,
1132+
SyncPointWaitList, SyncPoint);
1133+
}
1134+
1135+
pi_result piextCommandBufferAdviseUSM(
1136+
pi_ext_command_buffer CommandBuffer, const void *Ptr, size_t Length,
1137+
pi_mem_advice Advice, pi_uint32 NumSyncPointsInWaitList,
1138+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1139+
return pi2ur::piextCommandBufferAdviseUSM(CommandBuffer, Ptr, Length, Advice,
1140+
NumSyncPointsInWaitList,
1141+
SyncPointWaitList, SyncPoint);
1142+
}
1143+
11261144
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11271145
pi_queue Queue,
11281146
pi_uint32 NumEventsInWaitList,

‎sycl/source/detail/memory_manager.cpp

Lines changed: 35 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -194,15 +194,15 @@ void memBufferMapHelper(const PluginPtr &Plugin, pi_queue Queue, pi_mem Buffer,
194194
// We only want to instrument piEnqueueMemBufferMap
195195

196196
#ifdef XPTI_ENABLE_INSTRUMENTATION
197-
CorrID = emitMemAllocBeginTrace(MemObjID, Size, 0 /* guard zone */);
198-
xpti::utils::finally _{[&] {
199-
emitMemAllocEndTrace(MemObjID, (uintptr_t)(*RetMap), Size,
200-
0 /* guard zone */, CorrID);
201-
}};
197+
CorrID = emitMemAllocBeginTrace(MemObjID, Size, 0 /* guard zone */);
198+
xpti::utils::finally _{[&] {
199+
emitMemAllocEndTrace(MemObjID, (uintptr_t)(*RetMap), Size,
200+
0 /* guard zone */, CorrID);
201+
}};
202202
#endif
203-
Plugin->call<PiApiKind::piEnqueueMemBufferMap>(
204-
Queue, Buffer, Blocking, Flags, Offset, Size, NumEvents, WaitList,
205-
Event, RetMap);
203+
Plugin->call<PiApiKind::piEnqueueMemBufferMap>(Queue, Buffer, Blocking, Flags,
204+
Offset, Size, NumEvents,
205+
WaitList, Event, RetMap);
206206
}
207207

208208
void memUnmapHelper(const PluginPtr &Plugin, pi_queue Queue, pi_mem Mem,
@@ -1711,6 +1711,33 @@ void MemoryManager::ext_oneapi_fill_cmd_buffer(
17111711
PI_ERROR_INVALID_OPERATION);
17121712
}
17131713

1714+
void MemoryManager::ext_oneapi_prefetch_usm_cmd_buffer(
1715+
sycl::detail::ContextImplPtr Context,
1716+
sycl::detail::pi::PiExtCommandBuffer CommandBuffer, void *Mem,
1717+
size_t Length, std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
1718+
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
1719+
assert(!Context->is_host() && "Host queue not supported in prefetch_usm.");
1720+
1721+
const PluginPtr &Plugin = Context->getPlugin();
1722+
Plugin->call<PiApiKind::piextCommandBufferPrefetchUSM>(
1723+
CommandBuffer, Mem, Length, _pi_usm_migration_flags(0), Deps.size(),
1724+
Deps.data(), OutSyncPoint);
1725+
}
1726+
1727+
void MemoryManager::ext_oneapi_advise_usm_cmd_buffer(
1728+
sycl::detail::ContextImplPtr Context,
1729+
sycl::detail::pi::PiExtCommandBuffer CommandBuffer, const void *Mem,
1730+
size_t Length, pi_mem_advice Advice,
1731+
std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
1732+
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
1733+
assert(!Context->is_host() && "Host queue not supported in advise_usm.");
1734+
1735+
const PluginPtr &Plugin = Context->getPlugin();
1736+
Plugin->call<PiApiKind::piextCommandBufferAdviseUSM>(
1737+
CommandBuffer, Mem, Length, Advice, Deps.size(), Deps.data(),
1738+
OutSyncPoint);
1739+
}
1740+
17141741
void MemoryManager::copy_image_bindless(
17151742
void *Src, QueueImplPtr Queue, void *Dst,
17161743
const sycl::detail::pi::PiMemImageDesc &Desc,

‎sycl/source/detail/memory_manager.hpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -334,6 +334,19 @@ class __SYCL_EXPORT MemoryManager {
334334
std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
335335
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint);
336336

337+
static void ext_oneapi_prefetch_usm_cmd_buffer(
338+
sycl::detail::ContextImplPtr Context,
339+
sycl::detail::pi::PiExtCommandBuffer CommandBuffer, void *Mem,
340+
size_t Length, std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
341+
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint);
342+
343+
static void ext_oneapi_advise_usm_cmd_buffer(
344+
sycl::detail::ContextImplPtr Context,
345+
sycl::detail::pi::PiExtCommandBuffer CommandBuffer, const void *Mem,
346+
size_t Length, pi_mem_advice Advice,
347+
std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
348+
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint);
349+
337350
static void
338351
copy_image_bindless(void *Src, QueueImplPtr Queue, void *Dst,
339352
const sycl::detail::pi::PiMemImageDesc &Desc,

‎sycl/source/detail/queue_impl.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -249,6 +249,17 @@ event queue_impl::mem_advise(const std::shared_ptr<detail::queue_impl> &Self,
249249
const void *Ptr, size_t Length,
250250
pi_mem_advice Advice,
251251
const std::vector<event> &DepEvents) {
252+
// If we have a command graph set we need to capture the advise through normal
253+
// queue submission.
254+
if (MGraph.lock()) {
255+
return submit(
256+
[&](handler &CGH) {
257+
CGH.depends_on(DepEvents);
258+
CGH.mem_advise(Ptr, Length, Advice);
259+
},
260+
Self, {});
261+
}
262+
252263
if (MHasDiscardEventsSupport) {
253264
MemoryManager::advise_usm(Ptr, Self, Length, Advice,
254265
getOrWaitEvents(DepEvents, MContext), nullptr);

‎sycl/source/detail/scheduler/commands.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2786,6 +2786,24 @@ pi_int32 ExecCGCommand::enqueueImpCommandBuffer() {
27862786
MEvent->setSyncPoint(OutSyncPoint);
27872787
return PI_SUCCESS;
27882788
}
2789+
case CG::CGTYPE::PrefetchUSM: {
2790+
CGPrefetchUSM *Prefetch = (CGPrefetchUSM *)MCommandGroup.get();
2791+
MemoryManager::ext_oneapi_prefetch_usm_cmd_buffer(
2792+
MQueue->getContextImplPtr(), MCommandBuffer, Prefetch->getDst(),
2793+
Prefetch->getLength(), std::move(MSyncPointDeps), &OutSyncPoint);
2794+
MEvent->setSyncPoint(OutSyncPoint);
2795+
return PI_SUCCESS;
2796+
}
2797+
case CG::CGTYPE::AdviseUSM: {
2798+
CGAdviseUSM *Advise = (CGAdviseUSM *)MCommandGroup.get();
2799+
MemoryManager::ext_oneapi_advise_usm_cmd_buffer(
2800+
MQueue->getContextImplPtr(), MCommandBuffer, Advise->getDst(),
2801+
Advise->getLength(), Advise->getAdvice(), std::move(MSyncPointDeps),
2802+
&OutSyncPoint);
2803+
MEvent->setSyncPoint(OutSyncPoint);
2804+
return PI_SUCCESS;
2805+
}
2806+
27892807
default:
27902808
throw runtime_error("CG type not implemented for command buffers.",
27912809
PI_ERROR_INVALID_OPERATION);

0 commit comments

Comments
 (0)