[SYCL][Graph] Adding new graph enqueue function to spec (#15677)

replacing #15385 after offline
discussion.

sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc

@@ -716,6 +716,34 @@ optimize such partial barriers.
 _{endnote}_]
 |====
 
+==== Command Graph
+
+The functions in this section are only available if the
+link:./sycl_ext_oneapi_graph.asciidoc[
+  sycl_ext_oneapi_graph] extension is supported.
+
+|====
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
+void execute_graph(sycl::queue q, command_graph<graph_state::executable> &g);
+
+void execute_graph(sycl::handler &h, command_graph<graph_state::executable> &g);
+
+}
+----
+!====
+_Constraints_: Device and context associated with queue need to be identical 
+to device and context provided at command graph creation.
+
+_Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::handler`.
+
+|====
 
 == Issues
 

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -1975,7 +1975,9 @@ Removing this restriction is something we may look at for future revisions of
 
 The command submission functions defined in
 link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions]
-can be used to add nodes to a graph when creating a graph from queue recording.
+can be used adding nodes to a graph when creating a graph from queue recording.
+New methods are also defined that enable submitting an executable graph, 
+e.g. directly to a queue without returning an event.
 
 == Examples and Usage Guide
 

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -12,6 +12,7 @@
 
 #include <sycl/detail/common.hpp>
 #include <sycl/event.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/nd_range.hpp>
@@ -383,6 +384,17 @@ inline void partial_barrier(queue Q, const std::vector<event> &Events,
   submit(Q, [&](handler &CGH) { partial_barrier(CGH, Events); }, CodeLoc);
 }
 
+inline void execute_graph(queue Q, command_graph<graph_state::executable> &G,
+                          const sycl::detail::code_location &CodeLoc =
+                              sycl::detail::code_location::current()) {
+  Q.ext_oneapi_graph(G, CodeLoc);
+}
+
+inline void execute_graph(handler &CGH,
+                          command_graph<graph_state::executable> &G) {
+  CGH.ext_oneapi_graph(G);
+}
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

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

@@ -43,7 +43,7 @@ int main() {
 
   auto GraphExec = Graph.finalize();
 
-  InOrderQueue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); });
+  exp_ext::execute_graph(InOrderQueue, GraphExec);
   InOrderQueue.wait_and_throw();
 
   free(PtrA, InOrderQueue);

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

@@ -60,7 +60,7 @@ int main() {
 
     auto GraphExec = Graph.finalize();
 
-    Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); });
+    exp_ext::execute_graph(Queue, GraphExec);
     Queue.wait_and_throw();
   }
 

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

@@ -52,8 +52,9 @@ int main() {
 
   auto GraphExec = Graph.finalize();
 
-  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); });
-  Queue.wait_and_throw();
+  exp_ext::submit_with_event(Queue, [&](handler &CGH) {
+    exp_ext::execute_graph(CGH, GraphExec);
+  }).wait();
 
   free(PtrA, Queue);
   free(PtrB, Queue);