Skip to content

Commit 8ea0229

Browse files
BensuomfrancepilloisEwanC
authored
[SYCL][Graph] Add support for fill and memset nodes in graphs (#11472)
- Adds support for fill and memset nodes in graphs. - Supported on Level Zero only for now. - Adds E2E and unit tests for these new node types. - Minor modifications due to renaming of some UR functions. --------- Co-authored-by: Maxime France-Pillois <[email protected]> Co-authored-by: Ewan Crawford <[email protected]>
1 parent 218d9fe commit 8ea0229

26 files changed

+524
-4
lines changed

‎sycl/doc/design/CommandGraph.md

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -37,12 +37,14 @@ with the following entry-points:
3737
| `urCommandBufferFinalizeExp` | No more commands can be appended, makes command-buffer ready to enqueue on a command-queue. |
3838
| `urCommandBufferAppendKernelLaunchExp` | Append a kernel execution command to command-buffer. |
3939
| `urCommandBufferAppendUSMMemcpyExp` | Append a USM memcpy command to the command-buffer. |
40+
| `urCommandBufferAppendUSMFillExp` | Append a USM fill command to the command-buffer. |
4041
| `urCommandBufferAppendMemBufferCopyExp` | Append a mem buffer copy command to the command-buffer. |
4142
| `urCommandBufferAppendMemBufferWriteExp` | Append a memory write command to a command-buffer object. |
4243
| `urCommandBufferAppendMemBufferReadExp` | Append a memory read command to a command-buffer object. |
4344
| `urCommandBufferAppendMemBufferCopyRectExp` | Append a rectangular memory copy command to a command-buffer object. |
4445
| `urCommandBufferAppendMemBufferWriteRectExp` | Append a rectangular memory write command to a command-buffer object. |
4546
| `urCommandBufferAppendMemBufferReadRectExp` | Append a rectangular memory read command to a command-buffer object. |
47+
| `urCommandBufferAppendMemBufferFillExp` | Append a memory fill command to a command-buffer object. |
4648
| `urCommandBufferEnqueueExp` | Submit command-buffer to a command-queue for execution. |
4749

4850
See the [UR EXP-COMMAND-BUFFER](https://oneapi-src.github.io/unified-runtime/core/EXP-COMMAND-BUFFER.html)
@@ -347,6 +349,8 @@ The types of commands which are unsupported, and lead to this exception are:
347349
This corresponds to a memory buffer write command.
348350
* `handler::copy(src, dest)` or `handler::memcpy(dest, src)` - Where both `src` and
349351
`dest` are USM pointers. This corresponds to a USM copy command.
352+
* `handler::memset(ptr, value, numBytes)` - This corresponds to a USM memory
353+
fill command.
350354

351355
Note that `handler::copy(src, dest)` where both `src` and `dest` are an accessor
352356
is supported, as a memory buffer copy command exists in the OpenCL extension.

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -176,6 +176,8 @@ _PI_API(piextCommandBufferMemBufferWrite)
176176
_PI_API(piextCommandBufferMemBufferWriteRect)
177177
_PI_API(piextCommandBufferMemBufferRead)
178178
_PI_API(piextCommandBufferMemBufferReadRect)
179+
_PI_API(piextCommandBufferMemBufferFill)
180+
_PI_API(piextCommandBufferFillUSM)
179181
_PI_API(piextEnqueueCommandBuffer)
180182

181183
_PI_API(piextUSMPitchedAlloc)

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

Lines changed: 41 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -147,9 +147,10 @@
147147
// 14.38 Change PI_MEM_ADVICE_* values to flags for use in bitwise operations.
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.
150+
// 14.41 Added piextCommandBufferMemBufferFill & piextCommandBufferFillUSM
150151

151152
#define _PI_H_VERSION_MAJOR 14
152-
#define _PI_H_VERSION_MINOR 40
153+
#define _PI_H_VERSION_MINOR 41
153154

154155
#define _PI_STRING_HELPER(a) #a
155156
#define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -2441,7 +2442,7 @@ __SYCL_EXPORT pi_result piextCommandBufferMemBufferReadRect(
24412442
pi_buff_rect_offset buffer_offset, pi_buff_rect_offset host_offset,
24422443
pi_buff_rect_region region, size_t buffer_row_pitch,
24432444
size_t buffer_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch,
2444-
void *ptr, pi_uint32 num_events_in_wait_list,
2445+
void *ptr, pi_uint32 num_sync_points_in_wait_list,
24452446
const pi_ext_sync_point *sync_point_wait_list,
24462447
pi_ext_sync_point *sync_point);
24472448

@@ -2458,7 +2459,7 @@ __SYCL_EXPORT pi_result piextCommandBufferMemBufferReadRect(
24582459
/// \param sync_point The sync_point associated with this memory operation.
24592460
__SYCL_EXPORT pi_result piextCommandBufferMemBufferWrite(
24602461
pi_ext_command_buffer command_buffer, pi_mem buffer, size_t offset,
2461-
size_t size, const void *ptr, pi_uint32 num_events_in_wait_list,
2462+
size_t size, const void *ptr, pi_uint32 num_sync_points_in_wait_list,
24622463
const pi_ext_sync_point *sync_point_wait_list,
24632464
pi_ext_sync_point *sync_point);
24642465

@@ -2483,7 +2484,43 @@ __SYCL_EXPORT pi_result piextCommandBufferMemBufferWriteRect(
24832484
pi_buff_rect_offset buffer_offset, pi_buff_rect_offset host_offset,
24842485
pi_buff_rect_region region, size_t buffer_row_pitch,
24852486
size_t buffer_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch,
2486-
const void *ptr, pi_uint32 num_events_in_wait_list,
2487+
const void *ptr, pi_uint32 num_sync_points_in_wait_list,
2488+
const pi_ext_sync_point *sync_point_wait_list,
2489+
pi_ext_sync_point *sync_point);
2490+
2491+
/// API to append a mem buffer fill command to the command-buffer.
2492+
/// \param command_buffer The command-buffer to append onto.
2493+
/// \param buffer is the location to fill the data.
2494+
/// \param pattern pointer to the pattern to fill the buffer with.
2495+
/// \param pattern_size size of the pattern in bytes.
2496+
/// \param offset Offset into the buffer to fill from.
2497+
/// \param size fill size in bytes.
2498+
/// \param num_sync_points_in_wait_list The number of sync points in the
2499+
/// provided wait list.
2500+
/// \param sync_point_wait_list A list of sync points that this command must
2501+
/// wait on.
2502+
/// \param sync_point The sync_point associated with this memory operation.
2503+
__SYCL_EXPORT pi_result piextCommandBufferMemBufferFill(
2504+
pi_ext_command_buffer command_buffer, pi_mem buffer, const void *pattern,
2505+
size_t pattern_size, size_t offset, size_t size,
2506+
pi_uint32 num_sync_points_in_wait_list,
2507+
const pi_ext_sync_point *sync_point_wait_list,
2508+
pi_ext_sync_point *sync_point);
2509+
2510+
/// API to append a USM fill command to the command-buffer.
2511+
/// \param command_buffer The command-buffer to append onto.
2512+
/// \param ptr pointer to the USM allocation to fill.
2513+
/// \param pattern pointer to the pattern to fill ptr with.
2514+
/// \param pattern_size size of the pattern in bytes.
2515+
/// \param size fill size in bytes.
2516+
/// \param num_sync_points_in_wait_list The number of sync points in the
2517+
/// provided wait list.
2518+
/// \param sync_point_wait_list A list of sync points that this command must
2519+
/// wait on.
2520+
/// \param sync_point The sync_point associated with this memory operation.
2521+
__SYCL_EXPORT pi_result piextCommandBufferFillUSM(
2522+
pi_ext_command_buffer command_buffer, void *ptr, const void *pattern,
2523+
size_t pattern_size, size_t size, pi_uint32 num_sync_points_in_wait_list,
24872524
const pi_ext_sync_point *sync_point_wait_list,
24882525
pi_ext_sync_point *sync_point);
24892526

‎sycl/plugins/cuda/pi_cuda.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1137,6 +1137,27 @@ pi_result piextCommandBufferMemBufferWriteRect(
11371137
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
11381138
}
11391139

1140+
pi_result piextCommandBufferMemBufferFill(
1141+
pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
1142+
size_t PatternSize, size_t Offset, size_t Size,
1143+
pi_uint32 NumSyncPointsInWaitList,
1144+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1145+
return pi2ur::piextCommandBufferMemBufferFill(
1146+
CommandBuffer, Buffer, Pattern, PatternSize, Offset, Size,
1147+
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
1148+
}
1149+
1150+
pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
1151+
void *Ptr, const void *Pattern,
1152+
size_t PatternSize, size_t Size,
1153+
pi_uint32 NumSyncPointsInWaitList,
1154+
const pi_ext_sync_point *SyncPointWaitList,
1155+
pi_ext_sync_point *SyncPoint) {
1156+
return pi2ur::piextCommandBufferFillUSM(
1157+
CommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
1158+
SyncPointWaitList, SyncPoint);
1159+
}
1160+
11401161
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11411162
pi_queue Queue,
11421163
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/hip/pi_hip.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1145,6 +1145,27 @@ pi_result piextCommandBufferMemBufferWriteRect(
11451145
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
11461146
}
11471147

1148+
pi_result piextCommandBufferMemBufferFill(
1149+
pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
1150+
size_t PatternSize, size_t Offset, size_t Size,
1151+
pi_uint32 NumSyncPointsInWaitList,
1152+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1153+
return pi2ur::piextCommandBufferMemBufferFill(
1154+
CommandBuffer, Buffer, Pattern, PatternSize, Offset, Size,
1155+
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
1156+
}
1157+
1158+
pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
1159+
void *Ptr, const void *Pattern,
1160+
size_t PatternSize, size_t Size,
1161+
pi_uint32 NumSyncPointsInWaitList,
1162+
const pi_ext_sync_point *SyncPointWaitList,
1163+
pi_ext_sync_point *SyncPoint) {
1164+
return pi2ur::piextCommandBufferFillUSM(
1165+
CommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
1166+
SyncPointWaitList, SyncPoint);
1167+
}
1168+
11481169
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11491170
pi_queue Queue,
11501171
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/level_zero/pi_level_zero.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1303,6 +1303,27 @@ pi_result piextCommandBufferMemBufferWriteRect(
13031303
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
13041304
}
13051305

1306+
pi_result piextCommandBufferMemBufferFill(
1307+
pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
1308+
size_t PatternSize, size_t Offset, size_t Size,
1309+
pi_uint32 NumSyncPointsInWaitList,
1310+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1311+
return pi2ur::piextCommandBufferMemBufferFill(
1312+
CommandBuffer, Buffer, Pattern, PatternSize, Offset, Size,
1313+
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
1314+
}
1315+
1316+
pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
1317+
void *Ptr, const void *Pattern,
1318+
size_t PatternSize, size_t Size,
1319+
pi_uint32 NumSyncPointsInWaitList,
1320+
const pi_ext_sync_point *SyncPointWaitList,
1321+
pi_ext_sync_point *SyncPoint) {
1322+
return pi2ur::piextCommandBufferFillUSM(
1323+
CommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
1324+
SyncPointWaitList, SyncPoint);
1325+
}
1326+
13061327
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
13071328
pi_queue Queue,
13081329
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/opencl/pi_opencl.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1081,6 +1081,27 @@ pi_result piextCommandBufferMemBufferWriteRect(
10811081
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
10821082
}
10831083

1084+
pi_result piextCommandBufferMemBufferFill(
1085+
pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
1086+
size_t PatternSize, size_t Offset, size_t Size,
1087+
pi_uint32 NumSyncPointsInWaitList,
1088+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1089+
return pi2ur::piextCommandBufferMemBufferFill(
1090+
CommandBuffer, Buffer, Pattern, PatternSize, Offset, Size,
1091+
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
1092+
}
1093+
1094+
pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
1095+
void *Ptr, const void *Pattern,
1096+
size_t PatternSize, size_t Size,
1097+
pi_uint32 NumSyncPointsInWaitList,
1098+
const pi_ext_sync_point *SyncPointWaitList,
1099+
pi_ext_sync_point *SyncPoint) {
1100+
return pi2ur::piextCommandBufferFillUSM(
1101+
CommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
1102+
SyncPointWaitList, SyncPoint);
1103+
}
1104+
10841105
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
10851106
pi_queue Queue,
10861107
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/unified_runtime/pi2ur.hpp

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4580,6 +4580,37 @@ inline pi_result piextCommandBufferMemBufferWrite(
45804580
return PI_SUCCESS;
45814581
}
45824582

4583+
inline pi_result piextCommandBufferMemBufferFill(
4584+
pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
4585+
size_t PatternSize, size_t Offset, size_t Size,
4586+
pi_uint32 NumSyncPointsInWaitList,
4587+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
4588+
PI_ASSERT(Buffer, PI_ERROR_INVALID_MEM_OBJECT);
4589+
4590+
ur_exp_command_buffer_handle_t UrCommandBuffer =
4591+
reinterpret_cast<ur_exp_command_buffer_handle_t>(CommandBuffer);
4592+
ur_mem_handle_t UrBuffer = reinterpret_cast<ur_mem_handle_t>(Buffer);
4593+
4594+
HANDLE_ERRORS(urCommandBufferAppendMemBufferFillExp(
4595+
UrCommandBuffer, UrBuffer, Pattern, PatternSize, Offset, Size,
4596+
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint));
4597+
return PI_SUCCESS;
4598+
}
4599+
4600+
inline pi_result piextCommandBufferFillUSM(
4601+
pi_ext_command_buffer CommandBuffer, void *Ptr, const void *Pattern,
4602+
size_t PatternSize, size_t Size, pi_uint32 NumSyncPointsInWaitList,
4603+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
4604+
4605+
ur_exp_command_buffer_handle_t UrCommandBuffer =
4606+
reinterpret_cast<ur_exp_command_buffer_handle_t>(CommandBuffer);
4607+
4608+
HANDLE_ERRORS(urCommandBufferAppendUSMFillExp(
4609+
UrCommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
4610+
SyncPointWaitList, SyncPoint));
4611+
return PI_SUCCESS;
4612+
}
4613+
45834614
inline pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
45844615
pi_queue Queue,
45854616
pi_uint32 NumEventsInWaitList,

‎sycl/plugins/unified_runtime/pi_unified_runtime.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1102,6 +1102,27 @@ pi_result piextCommandBufferMemBufferWriteRect(
11021102
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
11031103
}
11041104

1105+
pi_result piextCommandBufferMemBufferFill(
1106+
pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
1107+
size_t PatternSize, size_t Offset, size_t Size,
1108+
pi_uint32 NumSyncPointsInWaitList,
1109+
const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
1110+
return pi2ur::piextCommandBufferMemBufferFill(
1111+
CommandBuffer, Buffer, Pattern, PatternSize, Offset, Size,
1112+
NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
1113+
}
1114+
1115+
pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
1116+
void *Ptr, const void *Pattern,
1117+
size_t PatternSize, size_t Size,
1118+
pi_uint32 NumSyncPointsInWaitList,
1119+
const pi_ext_sync_point *SyncPointWaitList,
1120+
pi_ext_sync_point *SyncPoint) {
1121+
return pi2ur::piextCommandBufferFillUSM(
1122+
CommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
1123+
SyncPointWaitList, SyncPoint);
1124+
}
1125+
11051126
pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
11061127
pi_queue Queue,
11071128
pi_uint32 NumEventsInWaitList,

‎sycl/source/detail/memory_manager.cpp

Lines changed: 44 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1666,6 +1666,50 @@ void MemoryManager::ext_oneapi_copy_usm_cmd_buffer(
16661666
}
16671667
}
16681668

1669+
void MemoryManager::ext_oneapi_fill_usm_cmd_buffer(
1670+
sycl::detail::ContextImplPtr Context,
1671+
sycl::detail::pi::PiExtCommandBuffer CommandBuffer, void *DstMem,
1672+
size_t Len, int Pattern, std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
1673+
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
1674+
1675+
if (!DstMem)
1676+
throw runtime_error("NULL pointer argument in memory fill operation.",
1677+
PI_ERROR_INVALID_VALUE);
1678+
1679+
const PluginPtr &Plugin = Context->getPlugin();
1680+
// Pattern is interpreted as an unsigned char so pattern size is always 1.
1681+
size_t PatternSize = 1;
1682+
Plugin->call<PiApiKind::piextCommandBufferFillUSM>(
1683+
CommandBuffer, DstMem, &Pattern, PatternSize, Len, Deps.size(),
1684+
Deps.data(), OutSyncPoint);
1685+
}
1686+
1687+
void MemoryManager::ext_oneapi_fill_cmd_buffer(
1688+
sycl::detail::ContextImplPtr Context,
1689+
sycl::detail::pi::PiExtCommandBuffer CommandBuffer, SYCLMemObjI *SYCLMemObj,
1690+
void *Mem, size_t PatternSize, const char *Pattern, unsigned int Dim,
1691+
sycl::range<3> Size, sycl::range<3> AccessRange, sycl::id<3> AccessOffset,
1692+
unsigned int ElementSize,
1693+
std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
1694+
sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
1695+
assert(SYCLMemObj && "The SYCLMemObj is nullptr");
1696+
1697+
const PluginPtr &Plugin = Context->getPlugin();
1698+
if (SYCLMemObj->getType() != detail::SYCLMemObjI::MemObjType::Buffer) {
1699+
throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
1700+
"Images are not supported in Graphs");
1701+
}
1702+
if (Dim <= 1) {
1703+
Plugin->call<PiApiKind::piextCommandBufferMemBufferFill>(
1704+
CommandBuffer, pi::cast<sycl::detail::pi::PiMem>(Mem), Pattern,
1705+
PatternSize, AccessOffset[0] * ElementSize,
1706+
AccessRange[0] * ElementSize, Deps.size(), Deps.data(), OutSyncPoint);
1707+
return;
1708+
}
1709+
throw runtime_error("Not supported configuration of fill requested",
1710+
PI_ERROR_INVALID_OPERATION);
1711+
}
1712+
16691713
void MemoryManager::copy_image_bindless(
16701714
void *Src, QueueImplPtr Queue, void *Dst,
16711715
const sycl::detail::pi::PiMemImageDesc &Desc,

0 commit comments

Comments
 (0)