From 7da187afb1df89f5620ccf9cd9b2ae0d0d745b3a Mon Sep 17 00:00:00 2001
From: Ewan Crawford <ewan@codeplay.com>
Date: Wed, 10 May 2023 15:27:36 +0100
Subject: [PATCH] [SYCL][Graph] Create handler::depends_on edges in Explicit
 API

An event returned from a queue submission captured by
Record & Replay should be able to create an edge to
a node created by the Explicit API.

This edge is defined by passing the event to
`handler::depends_on` inside the command-group added
explicitly to the graph.

Closes https://github.com/reble/llvm/issues/89
---
 sycl/source/detail/graph_impl.cpp           |  4 +-
 sycl/test-e2e/Graph/Explicit/depends_on.cpp | 45 +++++++++++++++++++++
 2 files changed, 47 insertions(+), 2 deletions(-)
 create mode 100644 sycl/test-e2e/Graph/Explicit/depends_on.cpp

diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp
index e7ed46c457e61..8b590eeb41737 100644
--- a/sycl/source/detail/graph_impl.cpp
+++ b/sycl/source/detail/graph_impl.cpp
@@ -150,11 +150,11 @@ graph_impl::add(const std::shared_ptr<graph_impl> &Impl,
   if (Handler.MSubgraphNode) {
     return Handler.MSubgraphNode;
   }
-  // TODO: Do we need to pass event dependencies here for the explicit API?
+
   return this->add(Handler.MKernel, Handler.MNDRDesc, Handler.MOSModuleHandle,
                    Handler.MKernelName, Handler.MAccStorage,
                    Handler.MLocalAccStorage, Handler.MCGType, Handler.MArgs,
-                   Handler.MImpl->MAuxiliaryResources, Dep);
+                   Handler.MImpl->MAuxiliaryResources, Dep, Handler.MEvents);
 }
 
 std::shared_ptr<node_impl> graph_impl::add(
diff --git a/sycl/test-e2e/Graph/Explicit/depends_on.cpp b/sycl/test-e2e/Graph/Explicit/depends_on.cpp
new file mode 100644
index 0000000000000..58c3983cd411a
--- /dev/null
+++ b/sycl/test-e2e/Graph/Explicit/depends_on.cpp
@@ -0,0 +1,45 @@
+// REQUIRES: level_zero, gpu
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// Tests that an event returned from adding a graph node using the queue
+// recording API can be passed to `handler::depends_on` inside a node
+// added using the explicit API. This should create a graph edge.
+
+#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;
+  float *Arr = malloc_device<float>(N, Queue);
+
+  Graph.begin_recording(Queue);
+  // `Event` corresponds to a graph node
+  event Event = Queue.submit([&](handler &CGH) {
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) { Arr[idx] = 42.0f; });
+  });
+  Graph.end_recording(Queue);
+
+  Graph.add([&](handler &CGH) {
+    CGH.depends_on(Event); // creates edge to recorded graph node
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) { Arr[idx] *= 2.0f; });
+  });
+
+  auto ExecGraph = Graph.finalize();
+
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }).wait();
+
+  constexpr float ref = 42.0f * 2.0f;
+  std::vector<float> Output(N);
+  Queue.memcpy(Output.data(), Arr, N * sizeof(float)).wait();
+  for (int i = 0; i < N; i++)
+    assert(Output[i] == ref);
+
+  sycl::free(Arr, Queue);
+
+  return 0;
+}