[SYCL] Use handler to execute graph

sycl/include/sycl/handler.hpp

@@ -30,6 +30,8 @@
 #include <sycl/sampler.hpp>
 #include <sycl/stl.hpp>
 
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+
 #include <functional>
 #include <limits>
 #include <memory>
@@ -2516,6 +2518,13 @@ class __SYCL_EXPORT handler {
   /// \param Advice is a device-defined advice for the specified allocation.
   void mem_advise(const void *Ptr, size_t Length, int Advice);
 
+  /// Executes a command_graph.
+  ///
+  /// \param Graph Executable command_graph to run
+  void exec_graph(ext::oneapi::experimental::command_graph<
+                  ext::oneapi::experimental::graph_state::executable>
+                      Graph);
+
 private:
   std::shared_ptr<detail::handler_impl> MImpl;
   std::shared_ptr<detail::queue_impl> MQueue;

sycl/include/sycl/queue.hpp

@@ -23,7 +23,6 @@
 #include <sycl/property_list.hpp>
 #include <sycl/stl.hpp>
 
-#include <sycl/ext/oneapi/experimental/graph.hpp>
 
 // Explicitly request format macros
 #ifndef __STDC_FORMAT_MACROS
@@ -1060,6 +1059,55 @@ class __SYCL_EXPORT queue {
 // Clean KERNELFUNC macros.
 #undef _KERNELFUNCPARAM
 
+  /// Shortcut for executing a graph of commands.
+  ///
+  /// \param Graph the graph of commands to execute
+  /// \return an event representing graph execution operation.
+  event exec_graph(ext::oneapi::experimental::command_graph<
+                   ext::oneapi::experimental::graph_state::executable>
+                       Graph) {
+    const detail::code_location CodeLoc = {};
+    return submit([&](handler &CGH) { CGH.exec_graph(Graph); }, CodeLoc);
+  }
+
+  /// Shortcut for executing a graph of commands.
+  ///
+  /// \param Graph the graph of commands to execute
+  /// \param DepEvent is an event that specifies the graph execution
+  /// dependencies.
+  /// \return an event representing graph execution operation.
+  event exec_graph(ext::oneapi::experimental::command_graph<
+                       ext::oneapi::experimental::graph_state::executable>
+                       Graph,
+                   event DepEvent) {
+    const detail::code_location CodeLoc = {};
+    return submit(
+        [&](handler &CGH) {
+          CGH.depends_on(DepEvent);
+          CGH.exec_graph(Graph);
+        },
+        CodeLoc);
+  }
+
+  /// Shortcut for executing a graph of commands.
+  ///
+  /// \param Graph the graph of commands to execute
+  /// \param DepEvents is a vector of events that specifies the graph
+  /// execution dependencies.
+  /// \return an event representing graph execution operation.
+  event exec_graph(ext::oneapi::experimental::command_graph<
+                       ext::oneapi::experimental::graph_state::executable>
+                       Graph,
+                   const std::vector<event> &DepEvents) {
+    const detail::code_location CodeLoc = {};
+    return submit(
+        [&](handler &CGH) {
+          CGH.depends_on(DepEvents);
+          CGH.exec_graph(Graph);
+        },
+        CodeLoc);
+  }
+
   /// Returns whether the queue is in order or OoO
   ///
   /// Equivalent to has_property<property::queue::in_order>()
@@ -1070,14 +1118,6 @@ class __SYCL_EXPORT queue {
   /// \return the backend associated with this queue.
   backend get_backend() const noexcept;
 
-public:
-  /// Submits an executable command_graph for execution on this queue
-  ///
-  /// \return an event representing the execution of the command_graph
-  event submit(ext::oneapi::experimental::command_graph<
-               ext::oneapi::experimental::graph_state::executable>
-                   graph);
-
 private:
   pi_native_handle getNative() const;
 

sycl/source/handler.cpp

@@ -10,6 +10,7 @@
 
 #include <detail/config.hpp>
 #include <detail/global_handler.hpp>
+#include <detail/graph_impl.hpp>
 #include <detail/handler_impl.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
@@ -698,5 +699,12 @@ void handler::depends_on(const std::vector<event> &Events) {
   }
 }
 
+void handler::exec_graph(ext::oneapi::experimental::command_graph<
+                         ext::oneapi::experimental::graph_state::executable>
+                             Graph) {
+  auto GraphImpl = detail::getSyclObjImpl(Graph);
+  GraphImpl->exec_and_wait(MQueue);
+}
+
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/source/queue.cpp

@@ -8,7 +8,6 @@
 
 #include <detail/backend_impl.hpp>
 #include <detail/event_impl.hpp>
-#include <detail/graph_impl.hpp>
 #include <detail/queue_impl.hpp>
 #include <sycl/event.hpp>
 #include <sycl/exception_list.hpp>
@@ -213,13 +212,5 @@ bool queue::device_has(aspect Aspect) const {
   // avoid creating sycl object from impl
   return impl->getDeviceImplPtr()->has(Aspect);
 }
-
-event queue::submit(ext::oneapi::experimental::command_graph<
-                    ext::oneapi::experimental::graph_state::executable>
-                        Graph) {
-  auto GraphImpl = detail::getSyclObjImpl(Graph);
-  GraphImpl->exec_and_wait(this->impl);
-  return {};
-}
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/test/graph/graph-explicit-dotp.cpp

@@ -78,9 +78,10 @@ int main() {
       },
       {node_a, node_b});
 
-  auto exec_graph = g.finalize(q.get_context());
+  auto executable_graph = g.finalize(q.get_context());
 
-  q.submit(exec_graph).wait();
+  // Using shortcut for executing a graph of commands
+  q.exec_graph(executable_graph).wait();
 
   if (*dotp != host_gold_result()) {
     std::cout << "Error unexpected result!\n";

sycl/test/graph/graph-explicit-queue-shortcuts.cpp

@@ -0,0 +1,39 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+#include <CL/sycl.hpp>
+#include <iostream>
+
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+
+int main() {
+
+  sycl::property_list properties{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::lazy_execution{}};
+
+  sycl::queue q{sycl::gpu_selector_v, properties};
+
+  sycl::ext::oneapi::experimental::command_graph g;
+
+  const size_t n = 10;
+  float *arr = sycl::malloc_shared<float>(n, q);
+
+  g.add([&](sycl::handler &h) {
+    h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx) {
+      size_t i = idx;
+      arr[i] = 1;
+    });
+  });
+
+  auto executable_graph = g.finalize(q.get_context());
+
+  auto e1 = q.exec_graph(executable_graph);
+  auto e2 = q.exec_graph(executable_graph, e1);
+  auto e3 = q.exec_graph(executable_graph, e1);
+  q.exec_graph(executable_graph, {e2, e3}).wait();
+
+  sycl::free(arr, q);
+
+  std::cout << "done " << arr[0] << std::endl;
+
+  return 0;
+}

sycl/test/graph/graph-explicit-simple.cpp

@@ -26,11 +26,11 @@ int main() {
 
   auto result_before_exec1 = arr[0];
 
-  auto exec_graph = g.finalize(q.get_context());
+  auto executable_graph = g.finalize(q.get_context());
 
   auto result_before_exec2 = arr[0];
 
-  q.submit(exec_graph).wait();
+  q.submit([&](sycl::handler &h) { h.exec_graph(executable_graph); });
 
   auto result = arr[0];