[SYCL][Graphs] Remove blocking wait from graph enqueue.

The current approach uses the same LO PI event for all submissions
of the same graph, and also doesn't use the wait events to
enforce dependencies on the command-list submission. By doing these
in the L0 adapter, we can remove the blocking queue wait from
our graphs submission code in the runtime.

Closes issue #139

Author: Ewan Crawford

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -8829,7 +8829,64 @@ pi_result _pi_buffer::free() {
   return PI_SUCCESS;
 }
 
-/// command-buffer Extension
+/* Command-buffer Extension
+
+   The PI interface for submitting a PI command-buffer takes a list
+   of events to wait on, and an event representing the completion of
+   that particular submission of the command-buffer.
+
+   However, in `zeCommandQueueExecuteCommandLists` there are no parameters to
+   take a waitlist and also the only sync primitive returned is to block on
+   host.
+
+   In order to get the PI command-buffer enqueue semantics we want with L0
+   this adapter adds extra commands to the L0 command-list representing a
+   PI command-buffer.
+
+   Prefix - Commands added to the start of the L0 command-list by L0 adapter.
+   Suffix - Commands added to the end of the L0 command-list by L0 adapter.
+
+   These extra commands operate on L0 event synchronisation primitives used by
+   the command-list to interact with the external PI wait-list and PI return
+   event required for the enqueue interface.
+
+   The `pi_ext_command_buffer` class for this adapter contains a SignalEvent
+   which signals the completion of the command-list in the suffix, and
+   is reset in the prefix. This signal is detected by a new PI return
+   event created on PI command-buffer enqueue.
+
+   There is also a WaitEvent used by the `pi_ext_command_buffer` class
+   in the prefix to wait on any dependencies passed in the enqueue wait-list.
+
+  ┌──────────┬────────────────────────────────────────────────┬─────────┐
+  │  Prefix  │ Commands added to PI command-buffer by PI user │ Suffix  │
+  └──────────┴────────────────────────────────────────────────┴─────────┘
+
+            ┌───────────────────┬──────────────────────────────┐
+  Prefix    │Reset signal event │ Barrier waiting on wait event│
+            └───────────────────┴──────────────────────────────┘
+
+            ┌─────────────────────────────────────────┐
+  Suffix    │Signal the PI command-buffer signal event│
+            └─────────────────────────────────────────┘
+
+
+  For a call to `piextEnqueueCommandBuffer` with an event_list `EL`,
+  command-buffer `CB`, and return event `RE` our implementation has to create
+  and submit two new command-lists for the above approach to work. One before
+  the command-list with extra commands associated with `CB`, and the other
+  after `CB`.
+
+  Command-list created on `piextEnqueueCommandBuffer` to execution before `CB`:
+  ┌───────────────────────────────────────────────────────────┐
+  │Barrier on `EL` than signals `CB` WaitEvent when completed │
+  └───────────────────────────────────────────────────────────┘
+
+  Command-list created on `piextEnqueueCommandBuffer` to execution after `CB`:
+  ┌─────────────────────────────────────────────────────────────┐
+  │Barrier on `CB` SignalEvent that signals `RE` when completed │
+  └─────────────────────────────────────────────────────────────┘
+*/
 
 /// Helper function to take a list of pi_ext_sync_points and fill the provided
 /// vector with the associated ZeEvents
@@ -8872,6 +8929,19 @@ pi_result piextCommandBufferCreate(pi_context Context, pi_device Device,
   } catch (...) {
     return PI_ERROR_UNKNOWN;
   }
+
+  // Create signal & wait events to be used in the command-list for sync
+  // on command-buffer enqueue.
+  auto CommandBuffer = *RetCommandBuffer;
+  PI_CALL(EventCreate(Context, nullptr, true, &CommandBuffer->SignalEvent));
+  PI_CALL(EventCreate(Context, nullptr, false, &CommandBuffer->WaitEvent));
+
+  // Add prefix commands
+  ZE_CALL(zeCommandListAppendEventReset,
+          (ZeCommandList, CommandBuffer->SignalEvent->ZeEvent));
+  ZE_CALL(zeCommandListAppendBarrier,
+          (ZeCommandList, nullptr, 1, &CommandBuffer->WaitEvent->ZeEvent));
+
   return PI_SUCCESS;
 }
 
@@ -8891,13 +8961,10 @@ pi_result piextCommandBufferRelease(pi_ext_command_buffer CommandBuffer) {
 }
 
 pi_result piextCommandBufferFinalize(pi_ext_command_buffer CommandBuffer) {
-  // We need to append some signal that will indicate that command-buffer has
+  // We need to append signal that will indicate that command-buffer has
   // finished executing.
-  EventCreate(CommandBuffer->Context, nullptr, true,
-              &CommandBuffer->ExecutionEvent);
-  ZE_CALL(
-      zeCommandListAppendSignalEvent,
-      (CommandBuffer->ZeCommandList, CommandBuffer->ExecutionEvent->ZeEvent));
+  ZE_CALL(zeCommandListAppendSignalEvent,
+          (CommandBuffer->ZeCommandList, CommandBuffer->SignalEvent->ZeEvent));
   // Close the command list and have it ready for dispatch.
   ZE_CALL(zeCommandListClose, (CommandBuffer->ZeCommandList));
   return PI_SUCCESS;
@@ -9026,17 +9093,11 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
                                     pi_uint32 NumEventsInWaitList,
                                     const pi_event *EventWaitList,
                                     pi_event *Event) {
-
-  // Execute command list asynchronously, as the event will be used
-  // to track down its completion.
-
-  uint32_t QueueGroupOrdinal;
-  // TODO: Revisit forcing compute engine
-  auto UseCopyEngine = false;
+  // Use compute engine rather than copy engine
+  const auto UseCopyEngine = false;
   auto &QGroup = Queue->getQueueGroup(UseCopyEngine);
-  auto &ZeCommandQueue =
-      // ForcedCmdQueue ? *ForcedCmdQueue :
-      QGroup.getZeQueue(&QueueGroupOrdinal);
+  uint32_t QueueGroupOrdinal;
+  auto &ZeCommandQueue = QGroup.getZeQueue(&QueueGroupOrdinal);
 
   ze_fence_handle_t ZeFence;
   ZeStruct<ze_fence_desc_t> ZeFenceDesc;
@@ -9050,25 +9111,69 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
           CommandBuffer->ZeCommandList,
           {ZeFence, false, false, ZeCommandQueue, QueueGroupOrdinal}));
 
-  Queue->insertActiveBarriers(CommandListPtr, UseCopyEngine);
-
+  // Previous execution will have closed the command list, we need to reopen
+  // it otherwise calling `executeCommandList` will return early.
+  CommandListPtr->second.IsClosed = false;
   CommandListPtr->second.ZeFenceInUse = true;
 
-  // Return the command-buffer's execution event as the user visible pi_event
-  *Event = CommandBuffer->ExecutionEvent;
-  (*Event)->Queue = Queue;
-  (*Event)->RefCount.increment();
-  Queue->RefCount.increment();
+  // Create command-list to execute before `CommandListPtr` and will signal
+  // when `EventWaitList` dependencies are complete.
+  pi_command_list_ptr_t WaitCommandList{};
+  if (NumEventsInWaitList) {
+    _pi_ze_event_list_t TmpWaitList;
+    if (auto Res = TmpWaitList.createAndRetainPiZeEventList(
+            NumEventsInWaitList, EventWaitList, Queue, UseCopyEngine))
+      return Res;
 
-  PI_CALL(piEventRetain(*Event));
+    if (auto Res = Queue->Context->getAvailableCommandList(
+            Queue, WaitCommandList, false, false))
+      return Res;
 
-  // Previous execution will have closed the command list so we need to reopen
-  // it.
-  CommandListPtr->second.IsClosed = false;
+    ZE_CALL(zeCommandListAppendBarrier,
+            (WaitCommandList->first, CommandBuffer->WaitEvent->ZeEvent,
+             NumEventsInWaitList, TmpWaitList.ZeEventList));
+  } else {
+    if (auto Res = Queue->Context->getAvailableCommandList(
+            Queue, WaitCommandList, false, false))
+      return Res;
+
+    ZE_CALL(zeCommandListAppendSignalEvent,
+            (WaitCommandList->first, CommandBuffer->WaitEvent->ZeEvent));
+  }
+
+  // Execution event for this enqueue of the PI command-buffer
+  pi_event RetEvent{};
+  // Create a command-list to signal RetEvent on completion
+  pi_command_list_ptr_t SignalCommandList{};
+  if (Event) {
+    if (auto Res = Queue->Context->getAvailableCommandList(
+            Queue, SignalCommandList, false, false))
+      return Res;
+
+    if (auto Res = createEventAndAssociateQueue(
+            Queue, &RetEvent, PI_COMMAND_TYPE_EXT_COMMAND_BUFFER,
+            SignalCommandList, false))
+      return Res;
+
+    ZE_CALL(zeCommandListAppendBarrier,
+            (SignalCommandList->first, RetEvent->ZeEvent, 1,
+             &(CommandBuffer->SignalEvent->ZeEvent)));
+  }
+
+  // Execution our command-lists asynchronously
+  if (auto Res = Queue->executeCommandList(WaitCommandList, false, false))
+    return Res;
 
   if (auto Res = Queue->executeCommandList(CommandListPtr, false, false))
     return Res;
 
+  if (auto Res = Queue->executeCommandList(SignalCommandList, false, false))
+    return Res;
+
+  if (Event) {
+    *Event = RetEvent;
+  }
+
   return PI_SUCCESS;
 }
 
@@ -9088,8 +9193,11 @@ _pi_ext_command_buffer::~_pi_ext_command_buffer() {
   if (ZeCommandList) {
     ZE_CALL_NOCHECK(zeCommandListDestroy, (ZeCommandList));
   }
-  if (ExecutionEvent) {
-    ExecutionEvent->RefCount.decrementAndTest();
+  if (SignalEvent) {
+    SignalEvent->RefCount.decrementAndTest();
+  }
+  if (WaitEvent) {
+    WaitEvent->RefCount.decrementAndTest();
   }
   Context->RefCount.decrementAndTest();
 }

sycl/plugins/level_zero/pi_level_zero.hpp

@@ -1374,8 +1374,12 @@ struct _pi_ext_command_buffer : _ur_object {
   // Command list map so we can use queue::executeCommandList, TODO: Remove in
   // future if possible
   pi_command_list_map_t CommandListMap;
-  // Event which will signal the execution of the command-buffer has finished
-  pi_event ExecutionEvent;
+  // Event which will signals the most recent execution of the command-buffer
+  // has finished
+  pi_event SignalEvent = nullptr;
+  // Event which a command-buffer waits on until the wait-list dependencies
+  // passed to a command-buffer enqueue have been satisfied.
+  pi_event WaitEvent = nullptr;
 };
 
 #endif // PI_LEVEL_ZERO_HPP

sycl/source/detail/graph_impl.cpp

@@ -104,16 +104,6 @@ std::shared_ptr<node_impl> graph_impl::add_subgraph_nodes(
   return this->add(Outputs);
 }
 
-sycl::event
-exec_graph_impl::exec(const std::shared_ptr<sycl::detail::queue_impl> &Queue) {
-  sycl::event RetEvent = enqueue(Queue);
-  // TODO: Remove this queue wait. Currently waiting on the event returned from
-  // graph execution does not work.
-  Queue->wait();
-
-  return RetEvent;
-}
-
 void graph_impl::add_root(const std::shared_ptr<node_impl> &Root) {
   MRoots.insert(Root);
 }

sycl/source/detail/graph_impl.hpp

@@ -325,16 +325,11 @@ class exec_graph_impl {
   /// Add nodes to MSchedule.
   void schedule();
 
-  /// Enqueues the backend objects for the graph to the parametrized queue.
-  /// @param Queue Command-queue to submit backend objects to.
-  /// @return Event associated with enqueued object.
-  sycl::event enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue);
-
   /// Called by handler::ext_oneapi_command_graph() to schedule graph for
   /// execution.
   /// @param Queue Command-queue to schedule execution on.
-  /// @return Event associated with the execution of the graph
-  sycl::event exec(const std::shared_ptr<sycl::detail::queue_impl> &Queue);
+  /// @return Event associated with the execution of the graph.
+  sycl::event enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue);
 
   /// Turns our internal graph representation into PI command-buffers for a
   /// device.

sycl/source/handler.cpp

@@ -399,7 +399,7 @@ event handler::finalize() {
     // If we have a subgraph node we don't want to actually execute this command
     // graph submission.
     if (!MSubgraphNode) {
-      event GraphCompletionEvent = MExecGraph->exec(MQueue);
+      event GraphCompletionEvent = MExecGraph->enqueue(MQueue);
       MLastEvent = GraphCompletionEvent;
       return MLastEvent;
     }

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

@@ -0,0 +1,100 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests that buffer accessors exhibit the correct behaviour when:
+// * A node is added to the graph between two queue submissions which
+//   use the same buffer, but are not added to the graph.
+//
+// * A queue submission using the same buffer is made after finalization
+//   of the graph, but before graph execution.
+//
+// * The graph is submitted for execution twice separated by a queue
+//   submission using the same buffer, this should respect dependencies and
+//   create the correct ordering.
+
+#include "../graph_common.hpp"
+int main() {
+
+  queue Queue{gpu_selector_v};
+
+  exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
+
+  const size_t N = 10;
+  std::vector<float> Arr(N, 0.0f);
+
+  buffer<float> Buf{N};
+  Buf.set_write_back(false);
+
+  // Buffer elements set to 0.5
+  Queue.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access(CGH);
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+      size_t i = idx;
+      Acc[i] = 0.5f;
+    });
+  });
+
+  Graph.add([&](handler &CGH) {
+    auto Acc = Buf.get_access(CGH);
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+      size_t i = idx;
+      Acc[i] += 0.25f;
+    });
+  });
+
+  for (size_t i = 0; i < N; i++) {
+    assert(Arr[i] == 0.0f);
+  }
+
+  // Buffer elements set to 1.5
+  Queue.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access(CGH);
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+      size_t i = idx;
+      Acc[i] += 1.0f;
+    });
+  });
+
+  auto ExecGraph = Graph.finalize();
+
+  for (size_t i = 0; i < N; i++) {
+    assert(Arr[i] == 0.0f);
+  }
+
+  // Buffer elements set to 3.0
+  Queue.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access(CGH);
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+      size_t i = idx;
+      Acc[i] *= 2.0f;
+    });
+  });
+
+  // Buffer elements set to 3.25
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); });
+
+  // Buffer elements set to 6.5
+  Queue.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access(CGH);
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+      size_t i = idx;
+      Acc[i] *= 2.0f;
+    });
+  });
+
+  // Buffer elements set to 6.75
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); });
+
+  Queue.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access(CGH);
+    CGH.copy(Acc, Arr.data());
+  });
+  Queue.wait();
+
+  for (size_t i = 0; i < N; i++) {
+    assert(Arr[i] == 6.75f);
+  }
+
+  return 0;
+}