[SYCL][Graph] Modified the adapters such that it is valid to call release on... (#18619)

Command Buffer, while it is still executing.

Author: konradkusiak97

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -765,10 +765,11 @@ _executable_ by the user invoking `command_graph::finalize()` to create a
 new executable instance of the graph. An executable graph cannot be converted
 to a modifiable graph. After finalizing a graph in the modifiable state, it is
 valid for a user to add additional nodes and finalize again to create subsequent
-executable graphs. The state of a `command_graph` object is made explicit by
-templating on state to make the class strongly typed, with the default template
-argument being `graph_state::modifiable` to reduce code verbosity on
-construction.
+executable graphs. When an executable graph is destroyed, the underlying
+resources will be freed only once any enqueued submissions of the graph have
+completed. The state of a `command_graph` object is made explicit by templating
+on state to make the class strongly typed, with the default template argument
+being `graph_state::modifiable` to reduce code verbosity on construction.
 
 .Graph State Diagram
 [source, mermaid]

sycl/source/detail/graph_impl.cpp

@@ -1002,11 +1002,6 @@ exec_graph_impl::~exec_graph_impl() {
     const sycl::detail::AdapterPtr &Adapter =
         sycl::detail::getSyclObjImpl(MContext)->getAdapter();
     MSchedule.clear();
-    // We need to wait on all command buffer executions before we can release
-    // them.
-    for (auto &Event : MExecutionEvents) {
-      Event->wait(Event);
-    }
 
     // Clean up any graph-owned allocations that were allocated
     MGraphImpl->getMemPool().deallocateAndUnmapAll();

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

@@ -0,0 +1,10 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// Extra run to check for immediate-command-list in Level Zero
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/release_while_executing.cpp"

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

@@ -0,0 +1,47 @@
+// Tests destroying finalized command_graph before it is finished executing,
+// relying on the backends to properly synchronize and wait for the submitted
+// work to finish.
+
+#include "../graph_common.hpp"
+
+int main() {
+  queue Queue;
+
+  using T = int;
+
+  std::vector<T> DataA(Size), ReferenceA(Size);
+
+  std::iota(DataA.begin(), DataA.end(), 1);
+  std::iota(ReferenceA.begin(), ReferenceA.end(), 2);
+
+  T *PtrA = malloc_device<T>(Size, Queue);
+
+  // Create the command_graph in a seperate scope so that it's destroyed before
+  // Queue.wait()
+  {
+    exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
+
+    Queue.copy(DataA.data(), PtrA, Size);
+    Queue.wait_and_throw();
+
+    auto Node = add_node(Graph, Queue, [&](handler &CGH) {
+      CGH.parallel_for(Size, [=](item<1> Item) { PtrA[Item.get_id()] += 1; });
+    });
+
+    auto GraphExec = Graph.finalize();
+    Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); });
+  }
+
+  Queue.wait_and_throw();
+
+  Queue.copy(PtrA, DataA.data(), Size);
+  Queue.wait_and_throw();
+
+  free(PtrA, Queue);
+
+  for (size_t i = 0; i < Size; i++) {
+    assert(check_value(i, ReferenceA[i], DataA[i], "DataA"));
+  }
+
+  return 0;
+}

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

@@ -0,0 +1,10 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// Extra run to check for immediate-command-list in Level Zero
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+#define GRAPH_E2E_RECORD_REPLAY
+
+#include "../Inputs/release_while_executing.cpp"

unified-runtime/include/ur_api.h

@@ -563,3 +563,4 @@ Contributors
 * Maxime France-Pillois `[email protected] <[email protected]>`_
 * Aaron Greig `[email protected] <[email protected]>`_
 * Fábio Mestre `[email protected] <[email protected]>`_
+* Konrad Kusiak `[email protected] <[email protected]>`_

unified-runtime/scripts/core/EXP-COMMAND-BUFFER.rst

@@ -310,7 +310,7 @@ returns:
     - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY
 --- #--------------------------------------------------------------------------
 type: function
-desc: "Decrement the command-buffer object's reference count and delete the command-buffer object if the reference count becomes zero."
+desc: "Decrement the command-buffer object's reference count and delete the command-buffer object if the reference count becomes zero. It is legal to call the entry-point while `hCommandBuffer` is still executing, which will block on completion if the reference count of `hCommandBuffer` becomes zero."
 class: $xCommandBuffer
 name: ReleaseExp
 params:

unified-runtime/scripts/core/exp-command-buffer.yml

@@ -671,6 +671,25 @@ ur_result_t createMainCommandList(ur_context_handle_t Context,
   return UR_RESULT_SUCCESS;
 }
 
+/**
+ * Waits for any ongoing executions of the command-buffer to finish
+ * @param CommandBuffer The command-buffer to wait for.
+ * @return UR_RESULT_SUCCESS or an error code on failure
+ */
+ur_result_t
+waitForOngoingExecution(ur_exp_command_buffer_handle_t CommandBuffer) {
+
+  if (ur_event_handle_t &CurrentSubmissionEvent =
+          CommandBuffer->CurrentSubmissionEvent) {
+    ZE2UR_CALL(zeEventHostSynchronize,
+               (CurrentSubmissionEvent->ZeEvent, UINT64_MAX));
+    UR_CALL(urEventReleaseInternal(CurrentSubmissionEvent));
+    CurrentSubmissionEvent = nullptr;
+  }
+
+  return UR_RESULT_SUCCESS;
+}
+
 /**
  * Checks whether the command-buffer can be constructed using in order
  * command-lists.
@@ -832,6 +851,7 @@ urCommandBufferReleaseExp(ur_exp_command_buffer_handle_t CommandBuffer) {
   if (!CommandBuffer->RefCount.decrementAndTest())
     return UR_RESULT_SUCCESS;
 
+  UR_CALL(waitForOngoingExecution(CommandBuffer));
   CommandBuffer->cleanupCommandBufferResources();
   delete CommandBuffer;
   return UR_RESULT_SUCCESS;
@@ -1453,25 +1473,6 @@ ur_result_t getZeCommandQueue(ur_queue_handle_t Queue, bool UseCopyEngine,
   return UR_RESULT_SUCCESS;
 }
 
-/**
- * Waits for any ongoing executions of the command-buffer to finish.
- * @param CommandBuffer The command-buffer to wait for.
- * @return UR_RESULT_SUCCESS or an error code on failure
- */
-ur_result_t
-waitForOngoingExecution(ur_exp_command_buffer_handle_t CommandBuffer) {
-
-  if (ur_event_handle_t &CurrentSubmissionEvent =
-          CommandBuffer->CurrentSubmissionEvent) {
-    ZE2UR_CALL(zeEventHostSynchronize,
-               (CurrentSubmissionEvent->ZeEvent, UINT64_MAX));
-    UR_CALL(urEventReleaseInternal(CurrentSubmissionEvent));
-    CurrentSubmissionEvent = nullptr;
-  }
-
-  return UR_RESULT_SUCCESS;
-}
-
 /**
  * Waits for the all the dependencies of the command-buffer
  * @param[in] CommandBuffer The command-buffer.

unified-runtime/source/adapters/level_zero/command_buffer.cpp

@@ -263,6 +263,10 @@ urCommandBufferReleaseExp(ur_exp_command_buffer_handle_t hCommandBuffer) try {
   if (!hCommandBuffer->RefCount.decrementAndTest())
     return UR_RESULT_SUCCESS;
 
+  if (auto executionEvent = hCommandBuffer->getExecutionEventUnlocked()) {
+    ZE2UR_CALL(zeEventHostSynchronize,
+               (executionEvent->getZeEvent(), UINT64_MAX));
+  }
   delete hCommandBuffer;
   return UR_RESULT_SUCCESS;
 } catch (...) {

unified-runtime/source/adapters/level_zero/v2/command_buffer.cpp

@@ -190,3 +190,14 @@ TEST_P(urEnqueueCommandBufferExpTest, SerializeInOrOutOfOrderQueue) {
     ASSERT_EQ(reference, Output[i]);
   }
 }
+
+// Tests releasing command-buffer while it is still executing relying
+// on synchronization during urCommandBufferReleaseExp call.
+TEST_P(urEnqueueCommandBufferExpTest, EnqueueAndRelease) {
+  ASSERT_SUCCESS(urEnqueueCommandBufferExp(
+      in_or_out_of_order_queue, cmd_buf_handle, 0, nullptr, nullptr));
+
+  // Release the command buffer without explicitly waiting beforehand
+  EXPECT_SUCCESS(urCommandBufferReleaseExp(cmd_buf_handle));
+  cmd_buf_handle = nullptr;
+}