From dbc74e8b06cce2563ca4361f591cf4a3f3a1e208 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 14 Nov 2022 10:46:08 +0000 Subject: [PATCH] Execute graph using handler Change the `queue::submit(command_graph)` API for launching an executable graph to `handler::exec_graph(command_graph)`. See Issue https://github.com/reble/llvm/issues/21 Using the handler is more in-keeping with the existing SYCL API and allows the execution of graph to depend on an arbitrary event through `handler::depends_on`. This design makes it easier for users to write the following code without have to block in host for waits. ```cpp auto ev1 = q.submit([&](handler& cgh){ cgh.exec_graph(g); }); // dest is some input to graph `g` auto ev2 = q.memcpy(dest, src, numBytes, ev1); auto ev3 = q.submit([&](handler& cgh){ cgh.depends_on(ev2); cgh.exec_graph(g); }); ``` Queue shortcut functions are also included, as is the case in the core SYCL spec for other handler functionality. This change should also enable the explicit API to capture nested sub-graph executions, which is not currently possible in the explicit API but is possible in the record & replay API. See issue https://github.com/reble/llvm/issues/23 For example, a user can now do: ```cpp command_graph executable_graph; auto node = recordable_graph.add([&](handler& cgh){ cgh.exec_graph(executable_graph); }); ``` --- .../proposed/sycl_ext_oneapi_graph.asciidoc | 91 +++++++++++++++---- 1 file changed, 71 insertions(+), 20 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc index a65273a766978..7f88daac17766 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc @@ -253,6 +253,12 @@ accessors. USM pointers also convey data dependencies, however offsets into system allocations (`malloc`/`new`) are not supported. |=== +==== Sub-Graph + +A node in a graph can take the form of a nested sub-graph. This occurs when +a command-group submission that invokes `handler::exec_graph()` with an +executable graph object is added to the graph as a node. + === API Modifications [source, c++] @@ -321,8 +327,22 @@ class queue { public: bool begin_recording(command_graph &graph); bool end_recording(); - event submit(command_graph graph); + + /* -- graph convenience shortcuts -- */ + + event exec_graph(command_graph graph); + event exec_graph(command_graph graph, + event depEvent); + event exec_graph(command_graph graph, + const std::vector &depEvents); }; + +// New methods added to the sycl::handler class +class handler { +public: + void exec_graph(command_graph graph); +} + } // namespace sycl ---- @@ -788,22 +808,55 @@ state, `false` otherwise. [source,c++] ---- using namespace ext::oneapi::experimental; -event queue::submit(command_graph graph) +event queue::exec_graph(command_graph graph) +---- + +|Queue shortcut function that is equivalent to submitting a command-group +containing `handler::exec_graph(graph)`. + +| +[source,c++] +---- +using namespace ext::oneapi::experimental; +event queue::exec_graph(command_graph graph, + event depEvent); +---- + +|Queue shortcut function that is equivalent to submitting a command-group +containing `handler::depends_on(depEvent)` and `handler::exec_graph(graph)`. + +| +[source,c++] +---- +using namespace ext::oneapi::experimental; +event queue::exec_graph(command_graph graph, + const std::vector &depEvents); ---- -|When invoked with the queue in the `queue_state::recording` state, a graph is -added as a subgraph node. When invoked with the queue in the default -`queue_state::executing` state, the graph is submitted for execution. Support -for submitting a graph for execution, before a previous execution has been -completed is backend specific. The runtime may throw an error. +|Queue shortcut function that is equivalent to submitting a command-group +containing `handler::depends_on(depEvents)` and `handler::exec_graph(graph)`. +|=== + +==== New Handler Member Functions + +Table 10. Additional member functions of the `sycl::handler` class. +[cols="2a,a"] +|=== +|Member function|Description +[source,c++] +---- +using namespace ext::oneapi::experimental; +void handler::exec_graph(command_graph graph) +---- + +|Invokes the execution of a graph. Support for invoking an executable graph, +before a previous execution of the same graph has been completed is backend +specific. The runtime may throw an error. Parameters: * `graph` - Graph object to execute. -When the queue is in the execution state, an `event` object used to synchronize -graph submission with other command-groups is returned. Otherwise the queue is -in the recording state, and a default event is returned. |=== === Thread Safety @@ -966,7 +1019,8 @@ int main() { auto exec = g.finalize(q.get_context()); - q.submit(exec).wait(); + // use queue shortcut for graph submission + q.exec_graph(exec).wait(); // memory can be freed inside or outside the graph sycl::free(z, q.get_context()); @@ -987,7 +1041,7 @@ command-groups submitted to the queue. Once the graph is complete, recording finishes on the queue to put it back into the default executing state. The graph is then finalized so that no more nodes can be added. Lastly, the graph is submitted in its entirety for execution via -`queue::submit(command_graph)`. +`handler::exec_graph(command_graph)`. [source, c++] ---- @@ -1055,7 +1109,10 @@ submitted in its entirety for execution via auto exec_graph = graph.finalize(q.get_context()); // Execute graph - q.submit(exec_graph); + q.submit([&](handler &cgh) { + cgh.exec_graph(exec_graph); + }); + ---- == Issues @@ -1082,13 +1139,6 @@ this feature in the extension. **Outcome:** Unresolved -=== Graph Submission Synchronization - -Should we provide a mechanism for a graph submission to depend on other graph -submission events or any arbitrary sycl event? - -**Outcome:** Unresolved - == Revision History [cols="5,15,15,70"] @@ -1101,4 +1151,5 @@ submission events or any arbitrary sycl event? |3|2022-05-25|Pablo Reble|Extend API and Example |4|2022-08-10|Pablo Reble|Adding USM shortcuts |5|2022-10-21|Ewan Crawford|Merge in Codeplay vendor extension +|6|2022-11-14|Ewan Crawford|Change graph execution to be a function on the handler |========================================