[SYCL][Graph] Adding new graph enqueue function to spec

sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc

@@ -709,6 +709,52 @@ 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 submit(sycl::queue q, 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`.
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
+sycl::event submit_with_event(sycl::queue q, 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`.
+
+_Returns_: A `sycl::event` associated with the submitted command.
+
+|====
 
 == Issues
 

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>
@@ -103,13 +104,22 @@ void submit(queue Q, CommandGroupFunc &&CGF,
       Q, std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
+__SYCL_EXPORT void submit(queue Q, command_graph<graph_state::executable> G,
+                          const sycl::detail::code_location &CodeLoc =
+                              sycl::detail::code_location::current());
+
 template <typename CommandGroupFunc>
 event submit_with_event(queue Q, CommandGroupFunc &&CGF,
                         const sycl::detail::code_location &CodeLoc =
                             sycl::detail::code_location::current()) {
   return Q.submit(std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
+__SYCL_EXPORT event
+submit_with_event(queue Q, command_graph<graph_state::executable> G,
+                  const sycl::detail::code_location &CodeLoc =
+                      sycl::detail::code_location::current());
+
 template <typename KernelName = sycl::detail::auto_name, typename KernelType>
 void single_task(handler &CGH, const KernelType &KernelObj) {
   CGH.single_task<KernelName>(KernelObj);

sycl/source/enqueue_functions.cpp

@@ -38,6 +38,17 @@ __SYCL_EXPORT void mem_advise(queue Q, void *Ptr, size_t NumBytes, int Advice,
                            /*CallerNeedsEvent=*/false);
 }
 
+__SYCL_EXPORT event
+submit_with_event(queue Q, command_graph<graph_state::executable> G,
+                  const sycl::detail::code_location &CodeLoc) {
+  return Q.ext_oneapi_graph(G, CodeLoc);
+}
+
+__SYCL_EXPORT void submit(queue Q, command_graph<graph_state::executable> G,
+                          const sycl::detail::code_location &CodeLoc) {
+  Q.ext_oneapi_graph(G, CodeLoc);
+}
+
 } // 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::submit(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::submit(Queue, GraphExec);
     Queue.wait_and_throw();
   }
 

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

@@ -52,8 +52,7 @@ 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, GraphExec).wait();
 
   free(PtrA, Queue);
   free(PtrB, Queue);

sycl/test/abi/sycl_symbols_linux.dump

@@ -3012,6 +3012,7 @@ _ZN4sycl3_V13ext6oneapi12experimental15get_image_rangeENS3_16image_mem_handleERK
 _ZN4sycl3_V13ext6oneapi12experimental15get_image_rangeENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE
 _ZN4sycl3_V13ext6oneapi12experimental15set_access_modeEPKvmNS3_19address_access_modeERKNS0_7contextE
 _ZN4sycl3_V13ext6oneapi12experimental16free_virtual_memEmmRKNS0_7contextE
+_ZN4sycl3_V13ext6oneapi12experimental17submit_with_eventENS0_5queueENS3_13command_graphILNS3_11graph_stateE1EEERKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_6deviceERKNS0_7contextENS3_16granularity_modeE
 _ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_7contextENS3_16granularity_modeE
 _ZN4sycl3_V13ext6oneapi12experimental19reserve_virtual_memEmmRKNS0_7contextE
@@ -3092,6 +3093,7 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKN
 _ZN4sycl3_V13ext6oneapi12experimental6detail33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE
 _ZN4sycl3_V13ext6oneapi12experimental6memcpyENS0_5queueEPvPKvmRKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental6memsetENS0_5queueEPvimRKNS0_6detail13code_locationE
+_ZN4sycl3_V13ext6oneapi12experimental6submitENS0_5queueENS3_13command_graphILNS3_11graph_stateE1EEERKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC2ERKNS3_16image_descriptorERKNS0_5queueE

sycl/test/abi/sycl_symbols_windows.dump

@@ -4254,6 +4254,7 @@
 ?storeRawArg@handler@_V1@sycl@@AEAAPEAXAEBVraw_kernel_arg@experimental@oneapi@ext@23@@Z
 ?storeRawArg@handler@_V1@sycl@@AEAAPEAXPEBX_K@Z
 ?stringifyErrorCode@detail@_V1@sycl@@YAPEBDH@Z
+?submit@experimental@oneapi@ext@_V1@sycl@@YAXVqueue@45@V?$command_graph@$00@12345@AEBUcode_location@detail@45@@Z
 ?submit_impl@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@@Z
 ?submit_impl@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@_N@Z
 ?submit_impl@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V123@AEBUcode_location@detail@23@@Z
@@ -4264,6 +4265,7 @@
 ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V123@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@_N@Z
 ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@@Z
 ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@_N@Z
+?submit_with_event@experimental@oneapi@ext@_V1@sycl@@YA?AVevent@45@Vqueue@45@V?$command_graph@$00@12345@AEBUcode_location@detail@45@@Z
 ?supportsUSMFill2D@handler@_V1@sycl@@AEAA_NXZ
 ?supportsUSMMemcpy2D@handler@_V1@sycl@@AEAA_NXZ
 ?supportsUSMMemset2D@handler@_V1@sycl@@AEAA_NXZ