[SYCL][Graph] Add support for fill and memset nodes in graphs

- Add support for fill/memset nodes in command graphs
- Add tests for buffer fills and memset
- Changes to UR Append* naming

Author: Bensuo

sycl/doc/design/CommandGraph.md

@@ -37,12 +37,14 @@ with the following entry-points:
 | `urCommandBufferFinalizeExp`                 | No more commands can be appended, makes command-buffer ready to enqueue on a command-queue. |
 | `urCommandBufferAppendKernelLaunchExp`       | Append a kernel execution command to command-buffer. |
 | `urCommandBufferAppendUSMMemcpyExp`          | Append a USM memcpy command to the command-buffer. |
+| `urCommandBufferAppendUSMFillExp`          | Append a USM fill command to the command-buffer. |
 | `urCommandBufferAppendMemBufferCopyExp`      | Append a mem buffer copy command to the command-buffer. |
 | `urCommandBufferAppendMemBufferWriteExp`     | Append a memory write command to a command-buffer object. |
 | `urCommandBufferAppendMemBufferReadExp`      | Append a memory read command to a command-buffer object. |
 | `urCommandBufferAppendMemBufferCopyRectExp`  | Append a rectangular memory copy command to a command-buffer object. |
 | `urCommandBufferAppendMemBufferWriteRectExp` | Append a rectangular memory write command to a command-buffer object. |
 | `urCommandBufferAppendMemBufferReadRectExp`  | Append a rectangular memory read command to a command-buffer object. |
+| `urCommandBufferAppendMemBufferFillExp`  | Append a memory fill command to a command-buffer object. |
 | `urCommandBufferEnqueueExp`                  | Submit command-buffer to a command-queue for execution. |
 
 See the [UR EXP-COMMAND-BUFFER](https://oneapi-src.github.io/unified-runtime/core/EXP-COMMAND-BUFFER.html)

sycl/include/sycl/detail/pi.def

@@ -176,6 +176,8 @@ _PI_API(piextCommandBufferMemBufferWrite)
 _PI_API(piextCommandBufferMemBufferWriteRect)
 _PI_API(piextCommandBufferMemBufferRead)
 _PI_API(piextCommandBufferMemBufferReadRect)
+_PI_API(piextCommandBufferMemBufferFill)
+_PI_API(piextCommandBufferFillUSM)
 _PI_API(piextEnqueueCommandBuffer)
 
 _PI_API(piextUSMPitchedAlloc)

sycl/include/sycl/detail/pi.h

@@ -2417,7 +2417,7 @@ __SYCL_EXPORT pi_result piextCommandBufferMemBufferReadRect(
     pi_buff_rect_offset buffer_offset, pi_buff_rect_offset host_offset,
     pi_buff_rect_region region, size_t buffer_row_pitch,
     size_t buffer_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch,
-    void *ptr, pi_uint32 num_events_in_wait_list,
+    void *ptr, pi_uint32 num_sync_points_in_wait_list,
     const pi_ext_sync_point *sync_point_wait_list,
     pi_ext_sync_point *sync_point);
 
@@ -2434,7 +2434,7 @@ __SYCL_EXPORT pi_result piextCommandBufferMemBufferReadRect(
 /// \param sync_point The sync_point associated with this memory operation.
 __SYCL_EXPORT pi_result piextCommandBufferMemBufferWrite(
     pi_ext_command_buffer command_buffer, pi_mem buffer, size_t offset,
-    size_t size, const void *ptr, pi_uint32 num_events_in_wait_list,
+    size_t size, const void *ptr, pi_uint32 num_sync_points_in_wait_list,
     const pi_ext_sync_point *sync_point_wait_list,
     pi_ext_sync_point *sync_point);
 
@@ -2459,7 +2459,43 @@ __SYCL_EXPORT pi_result piextCommandBufferMemBufferWriteRect(
     pi_buff_rect_offset buffer_offset, pi_buff_rect_offset host_offset,
     pi_buff_rect_region region, size_t buffer_row_pitch,
     size_t buffer_slice_pitch, size_t host_row_pitch, size_t host_slice_pitch,
-    const void *ptr, pi_uint32 num_events_in_wait_list,
+    const void *ptr, pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to append a mem buffer fill command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param buffer is the location to fill the data
+/// \param pattern pointer to the pattern to fill the buffer with.
+/// \param pattern_size size of the pattern in bytes.
+/// \param offset Offset into the buffer to fill from.
+/// \param size fill size in bytes.
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this command must
+/// wait on.
+/// \param sync_point The sync_point associated with this memory operation.
+__SYCL_EXPORT pi_result piextCommandBufferMemBufferFill(
+    pi_ext_command_buffer command_buffer, pi_mem buffer, const void *pattern,
+    size_t pattern_size, size_t offset, size_t size,
+    pi_uint32 num_sync_points_in_wait_list,
+    const pi_ext_sync_point *sync_point_wait_list,
+    pi_ext_sync_point *sync_point);
+
+/// API to append a USM fill command to the command-buffer.
+/// \param command_buffer The command-buffer to append onto.
+/// \param ptr pointer to the USM allocation to fill.
+/// \param pattern pointer to the pattern to fill ptr with.
+/// \param pattern_size size of the pattern in bytes.
+/// \param size fill size in bytes.
+/// \param num_sync_points_in_wait_list The number of sync points in the
+/// provided wait list.
+/// \param sync_point_wait_list A list of sync points that this command must
+/// wait on.
+/// \param sync_point The sync_point associated with this memory operation.
+__SYCL_EXPORT pi_result piextCommandBufferFillUSM(
+    pi_ext_command_buffer command_buffer, void *ptr, const void *pattern,
+    size_t pattern_size, size_t size, pi_uint32 num_sync_points_in_wait_list,
     const pi_ext_sync_point *sync_point_wait_list,
     pi_ext_sync_point *sync_point);
 

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -1304,6 +1304,27 @@ pi_result piextCommandBufferMemBufferWriteRect(
       NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
 }
 
+pi_result piextCommandBufferMemBufferFill(
+    pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
+    size_t PatternSize, size_t Offset, size_t Size,
+    pi_uint32 NumSyncPointsInWaitList,
+    const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
+  return pi2ur::piextCommandBufferMemBufferFill(
+      CommandBuffer, Buffer, Pattern, PatternSize, Offset, Size,
+      NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint);
+}
+
+pi_result piextCommandBufferFillUSM(pi_ext_command_buffer CommandBuffer,
+                                    void *Ptr, const void *Pattern,
+                                    size_t PatternSize, size_t Size,
+                                    pi_uint32 NumSyncPointsInWaitList,
+                                    const pi_ext_sync_point *SyncPointWaitList,
+                                    pi_ext_sync_point *SyncPoint) {
+  return pi2ur::piextCommandBufferFillUSM(
+      CommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
+      SyncPointWaitList, SyncPoint);
+}
+
 pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
                                     pi_queue Queue,
                                     pi_uint32 NumEventsInWaitList,

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -56,14 +56,8 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 47af3ee296ae0517213114332ffd3ac54a456b16
-  # Merge: bd76c510 f2ca7a91
-  # Author: Omar Ahmed <[email protected]>
-  # Date:   Thu Nov 30 16:11:56 2023 +0000
-  # - Merge pull request #1072 from omarahmed1111/merge-some-main-changes-into-adapters-third-patch
-  # - Merge main into adapters branch
-  set(UNIFIED_RUNTIME_TAG 47af3ee296ae0517213114332ffd3ac54a456b16)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/bensuo/unified-runtime.git")
+  set(UNIFIED_RUNTIME_TAG cmdbuf-fill-memset-l0)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -4521,6 +4521,37 @@ inline pi_result piextCommandBufferMemBufferWrite(
   return PI_SUCCESS;
 }
 
+inline pi_result piextCommandBufferMemBufferFill(
+    pi_ext_command_buffer CommandBuffer, pi_mem Buffer, const void *Pattern,
+    size_t PatternSize, size_t Offset, size_t Size,
+    pi_uint32 NumSyncPointsInWaitList,
+    const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
+  PI_ASSERT(Buffer, PI_ERROR_INVALID_MEM_OBJECT);
+
+  ur_exp_command_buffer_handle_t UrCommandBuffer =
+      reinterpret_cast<ur_exp_command_buffer_handle_t>(CommandBuffer);
+  ur_mem_handle_t UrBuffer = reinterpret_cast<ur_mem_handle_t>(Buffer);
+
+  HANDLE_ERRORS(urCommandBufferAppendMemBufferFillExp(
+      UrCommandBuffer, UrBuffer, Pattern, PatternSize, Offset, Size,
+      NumSyncPointsInWaitList, SyncPointWaitList, SyncPoint));
+  return PI_SUCCESS;
+}
+
+inline pi_result piextCommandBufferFillUSM(
+    pi_ext_command_buffer CommandBuffer, void *Ptr, const void *Pattern,
+    size_t PatternSize, size_t Size, pi_uint32 NumSyncPointsInWaitList,
+    const pi_ext_sync_point *SyncPointWaitList, pi_ext_sync_point *SyncPoint) {
+
+  ur_exp_command_buffer_handle_t UrCommandBuffer =
+      reinterpret_cast<ur_exp_command_buffer_handle_t>(CommandBuffer);
+
+  HANDLE_ERRORS(urCommandBufferAppendUSMFillExp(
+      UrCommandBuffer, Ptr, Pattern, PatternSize, Size, NumSyncPointsInWaitList,
+      SyncPointWaitList, SyncPoint));
+  return PI_SUCCESS;
+}
+
 inline pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
                                            pi_queue Queue,
                                            pi_uint32 NumEventsInWaitList,

sycl/source/detail/memory_manager.cpp

@@ -1666,6 +1666,50 @@ void MemoryManager::ext_oneapi_copy_usm_cmd_buffer(
   }
 }
 
+void MemoryManager::ext_oneapi_fill_usm_cmd_buffer(
+    sycl::detail::ContextImplPtr Context,
+    sycl::detail::pi::PiExtCommandBuffer CommandBuffer, void *DstMem,
+    size_t Len, int Pattern, std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
+    sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
+
+  if (!DstMem)
+    throw runtime_error("NULL pointer argument in memory fill operation.",
+                        PI_ERROR_INVALID_VALUE);
+
+  const PluginPtr &Plugin = Context->getPlugin();
+  // Pattern is interpreted as an unsigned char so pattern size is always 1.
+  size_t PatternSize = 1;
+  Plugin->call<PiApiKind::piextCommandBufferFillUSM>(
+      CommandBuffer, DstMem, &Pattern, PatternSize, Len, Deps.size(),
+      Deps.data(), OutSyncPoint);
+}
+
+void MemoryManager::ext_oneapi_fill_cmd_buffer(
+    sycl::detail::ContextImplPtr Context,
+    sycl::detail::pi::PiExtCommandBuffer CommandBuffer, SYCLMemObjI *SYCLMemObj,
+    void *Mem, size_t PatternSize, const char *Pattern, unsigned int Dim,
+    sycl::range<3> Size, sycl::range<3> AccessRange, sycl::id<3> AccessOffset,
+    unsigned int ElementSize,
+    std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
+    sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
+  assert(SYCLMemObj && "The SYCLMemObj is nullptr");
+
+  const PluginPtr &Plugin = Context->getPlugin();
+  if (SYCLMemObj->getType() != detail::SYCLMemObjI::MemObjType::Buffer) {
+    throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
+                          "Images are not supported in Graphs");
+  }
+  if (Dim <= 1) {
+    Plugin->call<PiApiKind::piextCommandBufferMemBufferFill>(
+        CommandBuffer, pi::cast<sycl::detail::pi::PiMem>(Mem), Pattern,
+        PatternSize, AccessOffset[0] * ElementSize,
+        AccessRange[0] * ElementSize, Deps.size(), Deps.data(), OutSyncPoint);
+    return;
+  }
+  throw runtime_error("Not supported configuration of fill requested",
+                      PI_ERROR_INVALID_OPERATION);
+}
+
 void MemoryManager::copy_image_bindless(
     void *Src, QueueImplPtr Queue, void *Dst,
     const sycl::detail::pi::PiMemImageDesc &Desc,

sycl/source/detail/memory_manager.hpp

@@ -316,6 +316,24 @@ class __SYCL_EXPORT MemoryManager {
       void *DstMem, std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
       sycl::detail::pi::PiExtSyncPoint *OutSyncPoint);
 
+  static void ext_oneapi_fill_usm_cmd_buffer(
+      sycl::detail::ContextImplPtr Context,
+      sycl::detail::pi::PiExtCommandBuffer CommandBuffer, void *DstMem,
+      size_t Len, int Pattern,
+      std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
+      sycl::detail::pi::PiExtSyncPoint *OutSyncPoint);
+
+  static void
+  ext_oneapi_fill_cmd_buffer(sycl::detail::ContextImplPtr Context,
+                             sycl::detail::pi::PiExtCommandBuffer CommandBuffer,
+                             SYCLMemObjI *SYCLMemObj, void *Mem,
+                             size_t PatternSize, const char *Pattern,
+                             unsigned int Dim, sycl::range<3> Size,
+                             sycl::range<3> AccessRange,
+                             sycl::id<3> AccessOffset, unsigned int ElementSize,
+                             std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
+                             sycl::detail::pi::PiExtSyncPoint *OutSyncPoint);
+
   static void
   copy_image_bindless(void *Src, QueueImplPtr Queue, void *Dst,
                       const sycl::detail::pi::PiMemImageDesc &Desc,

sycl/source/detail/scheduler/commands.cpp

@@ -2768,6 +2768,28 @@ pi_int32 ExecCGCommand::enqueueImpCommandBuffer() {
     MEvent->setSyncPoint(OutSyncPoint);
     return PI_SUCCESS;
   }
+  case CG::CGTYPE::Fill: {
+    CGFill *Fill = (CGFill *)MCommandGroup.get();
+    Requirement *Req = (Requirement *)(Fill->getReqToFill());
+    AllocaCommandBase *AllocaCmd = getAllocaForReq(Req);
+
+    MemoryManager::ext_oneapi_fill_cmd_buffer(
+        MQueue->getContextImplPtr(), MCommandBuffer, AllocaCmd->getSYCLMemObj(),
+        AllocaCmd->getMemAllocation(), Fill->MPattern.size(),
+        Fill->MPattern.data(), Req->MDims, Req->MMemoryRange, Req->MAccessRange,
+        Req->MOffset, Req->MElemSize, std::move(MSyncPointDeps), &OutSyncPoint);
+
+    return PI_SUCCESS;
+  }
+  case CG::CGTYPE::FillUSM: {
+    CGFillUSM *Fill = (CGFillUSM *)MCommandGroup.get();
+    MemoryManager::ext_oneapi_fill_usm_cmd_buffer(
+        MQueue->getContextImplPtr(), MCommandBuffer, Fill->getDst(),
+        Fill->getLength(), Fill->getFill(), std::move(MSyncPointDeps),
+        &OutSyncPoint);
+
+    return PI_SUCCESS;
+  }
   default:
     throw runtime_error("CG type not implemented for command buffers.",
                         PI_ERROR_INVALID_OPERATION);

sycl/test-e2e/Graph/Explicit/buffer_fill.cpp

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_fill.cpp"

sycl/test-e2e/Graph/Explicit/usm_memset.cpp

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/usm_memset.cpp"

sycl/test-e2e/Graph/Inputs/buffer_fill.cpp

@@ -0,0 +1,34 @@
+// Tests adding a Buffer fill operation as a graph node.
+
+#include "../graph_common.hpp"
+
+int main() {
+
+  queue Queue;
+  const size_t N = 10;
+  const float Pattern = 3.14f;
+  std::vector<float> Data(N);
+  buffer<float> Buffer(Data);
+  Buffer.set_write_back(false);
+  {
+    exp_ext::command_graph Graph{
+        Queue.get_context(),
+        Queue.get_device(),
+        {exp_ext::property::graph::assume_buffer_outlives_graph{},
+         exp_ext::property::graph::assume_data_outlives_buffer{}}};
+
+    auto NodeA = add_node(Graph, Queue, [&](handler &CGH) {
+      auto Acc = Buffer.get_access(CGH);
+      CGH.fill(Acc, Pattern);
+    });
+
+    auto ExecGraph = Graph.finalize();
+
+    Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }).wait();
+  }
+  host_accessor HostData(Buffer);
+  for (int i = 0; i < N; i++)
+    assert(HostData[i] == Pattern);
+
+  return 0;
+}

sycl/test-e2e/Graph/Inputs/usm_memset.cpp

@@ -0,0 +1,30 @@
+// Tests adding a USM memset operation as a graph node.
+
+#include "../graph_common.hpp"
+
+int main() {
+
+  queue Queue;
+
+  exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
+
+  const size_t N = 10;
+  unsigned char *Arr = malloc_device<unsigned char>(N, Queue);
+
+  int Value = 77;
+  auto NodeA =
+      add_node(Graph, Queue, [&](handler &CGH) { CGH.memset(Arr, Value, N); });
+
+  auto ExecGraph = Graph.finalize();
+
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }).wait();
+
+  std::vector<unsigned char> Output(N);
+  Queue.memcpy(Output.data(), Arr, N).wait();
+  for (int i = 0; i < N; i++)
+    assert(Output[i] == Value);
+
+  sycl::free(Arr, Queue);
+
+  return 0;
+}

sycl/test-e2e/Graph/RecordReplay/buffer_fill.cpp

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_RECORD_REPLAY
+
+#include "../Inputs/buffer_fill.cpp"

sycl/test-e2e/Graph/RecordReplay/usm_memset.cpp

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_RECORD_REPLAY
+
+#include "../Inputs/usm_memset.cpp"