From 66867d4faf87e03b855e3dc3d004f6b39c7553cd Mon Sep 17 00:00:00 2001 From: Pablo Reble Date: Wed, 6 Nov 2024 12:54:20 -0600 Subject: [PATCH] [SYCL][Graph] Adding new graph enqueue function to spec (#15677) replacing https://github.com/intel/llvm/pull/15385 after offline discussion. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 28 +++++++++++++++++++ .../sycl_ext_oneapi_graph.asciidoc | 4 ++- .../oneapi/experimental/enqueue_functions.hpp | 12 ++++++++ .../ext_oneapi_enqueue_functions.cpp | 2 +- .../ext_oneapi_enqueue_functions_submit.cpp | 2 +- ...pi_enqueue_functions_submit_with_event.cpp | 5 ++-- 6 files changed, 48 insertions(+), 5 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index ad8d6a7f50194..933a6aabd2bd4 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/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 &g); + +void execute_graph(sycl::handler &h, command_graph &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 diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 56f09c04d3055..5dff0396f07fb 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/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 diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index b3c758aaa891d..7ecf5ce4c8b14 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -12,6 +12,7 @@ #include #include +#include #include #include #include @@ -383,6 +384,17 @@ inline void partial_barrier(queue Q, const std::vector &Events, submit(Q, [&](handler &CGH) { partial_barrier(CGH, Events); }, CodeLoc); } +inline void execute_graph(queue Q, command_graph &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 &G) { + CGH.ext_oneapi_graph(G); +} + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions.cpp b/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions.cpp index ef3b440790c6b..6c6fe20337dab 100644 --- a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions.cpp +++ b/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); diff --git a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit.cpp b/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit.cpp index 8eccce2ea8ef9..623d6fc817879 100644 --- a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit.cpp +++ b/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(); } diff --git a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit_with_event.cpp b/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit_with_event.cpp index f66731d745bd2..4b8294be7e989 100644 --- a/sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions_submit_with_event.cpp +++ b/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);