From 80a40008c2c47ff3ab87a45a32007625233b7209 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Tue, 13 Feb 2024 16:25:10 +0000 Subject: [PATCH 01/10] [SYCL][Graph] Add External Events APIs Updates specification with new APIs for handling external events in Record&Replay and Explicit mode. Adds 3 functions to `graph`: - add_barrier() - make_external_event() - update_external_event() Adds 1 function to `queue`: - ext_oneapi_external_event() Adds examples to show how to use these functions. update spec Fix issue in the second code example. --- .../sycl_ext_oneapi_graph.asciidoc | 264 ++++++++++++++++++ 1 file changed, 264 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index a9690ab73764c..db6a6dc764606 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -632,6 +632,10 @@ public: void make_edge(node& src, node& dest); + node add_barrier(const std::vector eventList = {}); + + node make_external_event(event& recordedEvent, event& externalEvent, bool updatable = false); + void print_graph(std::string path, bool verbose = false) const; std::vector get_nodes() const; @@ -645,6 +649,8 @@ public: void update(node& node); void update(const std::vector& nodes); + + void update_external_event(node externalEventNode, event& newExternalEvent); }; } // namespace sycl::ext::oneapi::experimental @@ -996,6 +1002,93 @@ Exceptions: * Throws synchronously with error code `invalid` if the resulting dependency would lead to a cycle. This error is omitted when `property::graph::no_cycle_check` is set. +| +[source,c++] +---- +node add_barrier(const std::vector eventList = {}); +---- + +|Adds a barrier to the graph. + +Constraints: + +* This member function is only available when the `command_graph` state is + `graph_state::modifiable`. + +Parameters: + +* `eventList` - Zero or more events to wait for. + These events can be graph-limited events (from graph recording) + or general sycl events (such as events external to the graph). + If no event are provided, the barrier waits for all the previous nodes + of the graph to complete. + +Returns: The command-group function object node which has been added to the graph. + +| +[source,c++] +---- +node make_external_event(event& recordedEvent, event& externalEvent, bool updatable = false); +---- + +|Turns a graph-limited event into an external event on which other operations +can wait (including other graphs). + +Constraints: + +* This member function is only available when the `command_graph` state is + `graph_state::modifiable`. + +Parameters: + +* `recordedEvent` - Graph-limited event that will be linked to + the external event. + +* `externalEvent` - Sycl event to be used as external event. + +* `updatable` - Flag indicating that users will update the external event + after the graph finalization. Please, note that setting this flag may degrade + graph execution performance for level-zero backend, as updatable events + are handled by the host. + +Returns: The command-group function object node which has been added to the graph. + +Exceptions: + +* Throws synchronously with error code `invalid` if `recordedEvent` + is not a valid event assigned to the graph object. + +| +[source,c++] +---- +void update_external_event(node externalEventNode, event& newExternalEvent) +---- + +|Updates the external event associated to a graph-limited event. + +Constraints: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `externalEventNode` - Node that has been added to signal or wait for + the external event. + +* `newExternalEvent` - Sycl event that must be used as external event. + +Exceptions: + +* Throws synchronously with error code `invalid` if `externalEventNode` + is not a valid node of the graph. + +* Throws synchronously with error code `invalid` if the external event + associated with `externalEventNode` has not been created with + the `updatable` flag set. + +* Throws synchronously with error code `invalid` if the graph is executing. + | [source,c++] ---- @@ -1297,6 +1390,10 @@ public: event depEvent); event ext_oneapi_graph(command_graph& graph, const std::vector& depEvents); + + /* -- External events management -- */ + void ext_oneapi_external_event(event& recordedEvent, event& externalEvent, + bool updatable = false); }; } // namespace sycl ---- @@ -1456,6 +1553,7 @@ all the nodes have finished execution. The queue should be associated with a device and context that are the same as the device and context used on creation of the graph. + |=== ==== New Handler Member Functions @@ -1547,6 +1645,41 @@ a normal SYCL command-group submission. associated with the graph node resulting from this command-group submission is different from the one with which the dynamic_parameter was created. +| +[source,c++] +---- +void handler::ext_oneapi_external_event(event& recordedEvent, + event& externalEvent, + bool updatable = false); +---- + +|Turns a graph-limited event into an external event on which other operations +can wait (including other graphs). + +Parameters: + +* `recordedEvent` - Graph-limited event that will be linked to + the external event. + +* `externalEvent` - Sycl event to be used as external event. + +* `updatable` - Flag indicating that users will update the external event + after the graph finalization. Please, note that setting this flag may degrade + graph execution performance for level-zero backend, as updatable events + are handled by the host. + +Returns: The graph-limited event referencing the command-group function +object node which has been added to the graph. + +Exceptions: + +* Throws synchronously with error code `invalid` if the queue is not + in recording mode. + +* Throws synchronously with error code `invalid` if `recordedEvent` + is not a valid event assigned to the graph object. + + |=== === Thread Safety @@ -1684,6 +1817,137 @@ no queue state is changed. This design is because the queues are already in the state the user desires, so if the function threw an exception in this case, the application would likely swallow it and then proceed. +=== External events + +The default behavior is that the event associated with a graph node cannot be +used to synchronize workload outside the graph. +To overcome this limitation, users can define and use graph exernal events. +These events can be used as regular sycl events. Users can therefore use them +to manage the synchronizations between different graphs running concurrently +or between the graph's kernels and regular operations submit to a sycl queue. +If a graph containing one or more external events is run multiple times, +the external events are automatically reset when the graph is resubmitted to +the queue. + +[source,c++] +---- +/* Basic usage example using the Record&Reply API. */ + +// Create the graph `A`. +sycl::queue queueA; +sycl_ext::command_graph graphA(queueA.get_context(), queueA.get_device()); + +// Start recording the queue. +graphA.begin_recording(queueA); + +// Submit kernel on the recorded queue. +// EventA is a graph-limited event that can only be used for managing +// graph internal dependencies. +EventA = queueA.submit([&](handler& cgh) {...}); + +event externalEventA; +queueA.submit([&](handler& cgh) {ext_oneapi_external_event(EventA, + externalEventA, + false /* non-updatable */)});; + +// Stop recording the queue. +graphA.end_recording(); + +// Create another graph `B` +sycl::queue queueB; +sycl_ext::command_graph graphB(queueB.get_context(), queueB.get_device()); + +// Start recording the queue. +graphB.begin_recording(queueB); + +// Submit kernel on the recorded queue. +queueB.submit([&](handler& cgh) {...}); + +queueB.ext_oneapi_submit_barrier({externalEventA}); + +// Stop recording the queue. +graphB.end_recording(); + +// Finalize the modifiable graphs to create an executable graphs that can be +// submitted for execution. +auto execGraphA = graphA.finalize(); +auto execGraphB = graphB.finalize(); + +// Execute graph +for(int i = 0 ; i < Iterations, i++) { + queueA.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphA); + }) + + queueB.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphB); + }) +} +---- + +[source,c++] +---- +/* Example: Event update. */ + +// Create the graph `A`. +sycl::queue queueA; +sycl_ext::command_graph graphA(queueA.get_context(), queueA.get_device()); + +// Start recording the queue. +graphA.begin_recording(queueA); + +// Submit kernel on the recorded queue. +// EventA is a graph-limited event that can only be used for managing +// graph internal dependencies. +EventA = queueA.submit([&](handler& cgh) {...}); + +// Stop recording the queue. +graphA.end_recording(); + +event externalEventA; +auto nodeSignalExternalEvent = graphA.make_external_event(EventA, + externalEventA, + true /* updatable */); + +// Create another graph `B` +sycl::queue queueB; +sycl_ext::command_graph graphB(queueB.get_context(), queueB.get_device()); + +// Submit kernel on the recorded queue. +graphB.add([&](handler& cgh) {...}); + +auto nodeWaitExternalEvent = graphB.add_barrier({externalEventA}); + +// Finalize the modifiable graphs to create an executable graphs that can be +// submitted for execution. +auto execGraphA = graphA.finalize(); +auto execGraphB = graphB.finalize(); + +// Execute graph +// first execution +queueA.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphA); +}) +queueB.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphB); +}) + +event newExternalEventA; +execGraphA.update_external_event(nodeSignalExternalEvent, newExternalEventA); + +event newExternalEventB; +execGraphB.update_external_event(nodeWaitExternalEvent, newExternalEventB); + +// second execution +queueA.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphA); +}) +queueB.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphB); +}) +---- + + === Interaction With Other Extensions [[extension-interaction]] This section defines the interaction of `sycl_ext_oneapi_graph` with other From 00fae5ba0cba4dfd930ac6f436f2e119c9e106b1 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 4 Apr 2024 17:00:35 +0100 Subject: [PATCH 02/10] [SYCL][Graph] External event changes - Remove updating and profiling capabilities - Simplify and align interfaces for getting external events - Restrict external event use within the graph (prevents needing graph-limited event from queue shortcut) --- .../sycl_ext_oneapi_graph.asciidoc | 430 +++++++++--------- 1 file changed, 203 insertions(+), 227 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index db6a6dc764606..750167d76d330 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -336,6 +336,7 @@ enum class node_type { memadvise, ext_oneapi_barrier, host_task, + external_event, }; class node { @@ -350,6 +351,8 @@ public: static node get_node_from_event(event nodeEvent); + event get_external_event() const; + template void update_nd_range(nd_range executionRange); @@ -418,6 +421,24 @@ Exceptions: * Throws with error code `invalid` if `nodeEvent` is not associated with a graph node. +| +[source,c++] +---- +event get_external_event(); +---- +|Gets the external event associated with this node. Can only be called for nodes +with type `node_type::external_event`. + +For more information on external events see <>. + +Returns: The external event associated with this node. + +Exceptions: + +* Throws with error code `invalid` if the type of this node is not +`node_type::external_event`. + | [source,c++] ---- @@ -634,7 +655,7 @@ public: node add_barrier(const std::vector eventList = {}); - node make_external_event(event& recordedEvent, event& externalEvent, bool updatable = false); + node add_external_event(event& externalEvent, const property_list& propList = {}); void print_graph(std::string path, bool verbose = false) const; @@ -649,8 +670,6 @@ public: void update(node& node); void update(const std::vector& nodes); - - void update_external_event(node externalEventNode, event& newExternalEvent); }; } // namespace sycl::ext::oneapi::experimental @@ -1028,66 +1047,27 @@ Returns: The command-group function object node which has been added to the grap | [source,c++] ---- -node make_external_event(event& recordedEvent, event& externalEvent, bool updatable = false); +node add_external_event(const property_list& propList = {}); ---- -|Turns a graph-limited event into an external event on which other operations -can wait (including other graphs). - -Constraints: - -* This member function is only available when the `command_graph` state is - `graph_state::modifiable`. - -Parameters: - -* `recordedEvent` - Graph-limited event that will be linked to - the external event. - -* `externalEvent` - Sycl event to be used as external event. - -* `updatable` - Flag indicating that users will update the external event - after the graph finalization. Please, note that setting this flag may degrade - graph execution performance for level-zero backend, as updatable events - are handled by the host. - -Returns: The command-group function object node which has been added to the graph. - -Exceptions: - -* Throws synchronously with error code `invalid` if `recordedEvent` - is not a valid event assigned to the graph object. - -| -[source,c++] ----- -void update_external_event(node externalEventNode, event& newExternalEvent) ----- +|Adds an external event node to the graph, this node contains an external graph +event which represents the execution at this point in the graph. -|Updates the external event associated to a graph-limited event. +For more information on external events see <>. Constraints: * This member function is only available when the `command_graph` state is - `graph_state::executable`. + `graph_state::modifiable`. Parameters: -* `externalEventNode` - Node that has been added to signal or wait for - the external event. - -* `newExternalEvent` - Sycl event that must be used as external event. - -Exceptions: - -* Throws synchronously with error code `invalid` if `externalEventNode` - is not a valid node of the graph. - -* Throws synchronously with error code `invalid` if the external event - associated with `externalEventNode` has not been created with - the `updatable` flag set. +* `propList` - Zero or more properties can be provided to the constructed node + via an instance of `property_list`. The `property::node::depends_on` property + can be passed here with a list of nodes to create dependency edges on. -* Throws synchronously with error code `invalid` if the graph is executing. +Returns: The external event node which has been added to the graph. | [source,c++] @@ -1392,8 +1372,8 @@ public: const std::vector& depEvents); /* -- External events management -- */ - void ext_oneapi_external_event(event& recordedEvent, event& externalEvent, - bool updatable = false); + event ext_oneapi_external_event(const event& depEvent); + event ext_oneapi_external_event(const std::vector& depEvents); }; } // namespace sycl ---- @@ -1554,6 +1534,58 @@ all the nodes have finished execution. The queue should be associated with a device and context that are the same as the device and context used on creation of the graph. +| +[source,c++] +---- +event queue::ext_oneapi_external_event(const event& depEvent); +---- + +|Create and return a graph external event representing the execution of +`depEvent`. This function should only be called on a queue which is in the +recording state. + +For more information on external events see <>. + +Parameters: + +* `depEvent` - The event which this external event depends on. The +external event will be considered complete when all of its dependencies are +complete. + +Returns: The external event representing this point of execution in the graph. + +Exceptions: + +* Throws synchronously with error code `invalid` if the queue is not + in recording mode. + +| +[source,c++] +---- +event queue::ext_oneapi_external_event(const std::vector& depEvents); +---- + +|Create and return a graph external event representing the execution of +`depEvents`. This function should only be called on a queue which is in the +recording state. + +For more information on external events see <>. + +Parameters: + +* `depEvents` - List of events which this external event depends on. The +external event will be considered complete when all of its dependencies are +complete. + +Returns: The external event representing this point of execution in the graph. + +Exceptions: + +* Throws synchronously with error code `invalid` if the queue is not + in recording mode. + |=== ==== New Handler Member Functions @@ -1645,42 +1677,136 @@ a normal SYCL command-group submission. associated with the graph node resulting from this command-group submission is different from the one with which the dynamic_parameter was created. -| -[source,c++] +|=== + +=== Event Class Modifications [[event-class-modifications]] + +[source, c++] ---- -void handler::ext_oneapi_external_event(event& recordedEvent, - event& externalEvent, - bool updatable = false); +namespace sycl { +// New methods added to the sycl::event class +using namespace ext::oneapi::experimental; +class event { +public: + bool is_graph_limited() const; +}; +} // namespace sycl ---- -|Turns a graph-limited event into an external event on which other operations -can wait (including other graphs). +:event-class: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:interface.event -Parameters: +==== Graph-Limited Events -* `recordedEvent` - Graph-limited event that will be linked to - the external event. +Events associated with a graph node (those returned from recorded queue +submissions) cannot be used as normal SYCL events outside of their associated +graph, these are known as "graph-limited" events. They can only be used as +parameters to `handler::depends_on()`, or as dependent events for queue +shortcuts like `queue::parallel_for()` for other submissions recorded to +the same modifiable `command_graph`. -* `externalEvent` - Sycl event to be used as external event. +The following limitations apply to graph-limited events: -* `updatable` - Flag indicating that users will update the external event - after the graph finalization. Please, note that setting this flag may degrade - graph execution performance for level-zero backend, as updatable events - are handled by the host. +- Calling `event::get_info()` or +`event::get_profiling_info()` will throw synchronously with error code +`invalid`. -Returns: The graph-limited event referencing the command-group function -object node which has been added to the graph. +- A host-side wait on the event will throw synchronously with error +code `invalid`. -Exceptions: +- Using the event outside of the recording scope will throw synchronously with +error code `invalid`. -* Throws synchronously with error code `invalid` if the queue is not - in recording mode. +==== External Events -* Throws synchronously with error code `invalid` if `recordedEvent` - is not a valid event assigned to the graph object. +To overcome the limitations of "graph-limited" events, users can add external +event nodes to a graph. These represent a point of execution within the graph +and provide a SYCL event object which can be used to both manage synchronization +between different graphs as well as between the graph and regular operations +submitted to a SYCL queue. External graph events are associated with a single +external event node inside a graph. External events are functionally normal SYCL +events with some restrictions (see below). +External events can be obtained by either: -|=== +* Calling `command_graph::add_external_event()` on a modifiable command_graph to +add an external event node, then obtain the SYCL event object by calling +`node::get_external_event()`. + +* Calling `queue::ext_oneapi_external_event()` on a queue in the recording +state, which will return an external SYCL event. + +External events may be used as parameters to `handler::depends_on()` or as +dependent events for queue submissions outside of the graph they +are associated with. For example, this allows enqueing work outside of the graph +part-way through the graph execution, rather than having to wait for the entire +graph to finish. + +External events are considered complete when all their dependent graph nodes +have finished execution, and their execution status is automatically reset when +the executable graph which contains them is submitted to a queue for execution. + +External events have the following restrictions: + +* Profiling is not currently supported for external graph events and calling +`event::get_profiling_info()` will throw synchronously with error code +`invalid`. + +* External events cannot be passed as dependent events for submissions within +the same graph, and doing so will throw synchronously with error code `invalid`. +Graph-limited events should be used for this purpose. + +[source,c++] +---- +/* Basic usage example using the Record&Replay API. */ + +// Create the graph `A`. +sycl::queue queueA; +sycl_ext::command_graph graphA(queueA.get_context(), queueA.get_device()); + +// Start recording the queue. +graphA.begin_recording(queueA); + +// Submit kernel on the recorded queue. +// EventA is a graph-limited event that can only be used for managing +// graph internal dependencies. +EventA = queueA.submit([&](handler& cgh) {...}); + +event externalEventA = queueA.ext_oneapi_external_event({EventA}); + +// Stop recording the queue. +graphA.end_recording(); + +// Create another graph `B` +sycl::queue queueB; +sycl_ext::command_graph graphB(queueB.get_context(), queueB.get_device()); + +// Start recording the queue. +graphB.begin_recording(queueB); + +// Submit kernel on the recorded queue. +queueB.submit([&](handler& cgh) {...}); + +queueB.ext_oneapi_submit_barrier({externalEventA}); + +// Stop recording the queue. +graphB.end_recording(); + +// Finalize the modifiable graphs to create an executable graphs that can be +// submitted for execution. +auto execGraphA = graphA.finalize(); +auto execGraphB = graphB.finalize(); + +// Execute graph +for(int i = 0 ; i < Iterations, i++) { + queueA.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphA); + }) + + queueB.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(execGraphB); + }) +} +---- === Thread Safety @@ -1742,25 +1868,6 @@ of failure. The following list describes the behavior that changes during recording mode. Features not listed below behave the same in recording mode as they do in non-recording mode. -==== Event Limitations - -For queue submissions that are being recorded to a modifiable `command_graph`, -the only events that can be used as parameters to `handler::depends_on()`, or -as dependent events for queue shortcuts like `queue::parallel_for()`, are events -that have been returned from queue submissions recorded to the same modifiable -`command_graph`. - -Other limitations on the events returned from a submission to a queue in the -recording state are: - -- Calling `event::get_info()` or -`event::get_profiling_info()` will throw synchronously with error code `invalid`. - -- A host-side wait on the event will throw synchronously with error -code `invalid`. - -- Using the event outside of the recording scope will throw synchronously with error code -`invalid`. ==== Queue Limitations @@ -1817,137 +1924,6 @@ no queue state is changed. This design is because the queues are already in the state the user desires, so if the function threw an exception in this case, the application would likely swallow it and then proceed. -=== External events - -The default behavior is that the event associated with a graph node cannot be -used to synchronize workload outside the graph. -To overcome this limitation, users can define and use graph exernal events. -These events can be used as regular sycl events. Users can therefore use them -to manage the synchronizations between different graphs running concurrently -or between the graph's kernels and regular operations submit to a sycl queue. -If a graph containing one or more external events is run multiple times, -the external events are automatically reset when the graph is resubmitted to -the queue. - -[source,c++] ----- -/* Basic usage example using the Record&Reply API. */ - -// Create the graph `A`. -sycl::queue queueA; -sycl_ext::command_graph graphA(queueA.get_context(), queueA.get_device()); - -// Start recording the queue. -graphA.begin_recording(queueA); - -// Submit kernel on the recorded queue. -// EventA is a graph-limited event that can only be used for managing -// graph internal dependencies. -EventA = queueA.submit([&](handler& cgh) {...}); - -event externalEventA; -queueA.submit([&](handler& cgh) {ext_oneapi_external_event(EventA, - externalEventA, - false /* non-updatable */)});; - -// Stop recording the queue. -graphA.end_recording(); - -// Create another graph `B` -sycl::queue queueB; -sycl_ext::command_graph graphB(queueB.get_context(), queueB.get_device()); - -// Start recording the queue. -graphB.begin_recording(queueB); - -// Submit kernel on the recorded queue. -queueB.submit([&](handler& cgh) {...}); - -queueB.ext_oneapi_submit_barrier({externalEventA}); - -// Stop recording the queue. -graphB.end_recording(); - -// Finalize the modifiable graphs to create an executable graphs that can be -// submitted for execution. -auto execGraphA = graphA.finalize(); -auto execGraphB = graphB.finalize(); - -// Execute graph -for(int i = 0 ; i < Iterations, i++) { - queueA.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(execGraphA); - }) - - queueB.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(execGraphB); - }) -} ----- - -[source,c++] ----- -/* Example: Event update. */ - -// Create the graph `A`. -sycl::queue queueA; -sycl_ext::command_graph graphA(queueA.get_context(), queueA.get_device()); - -// Start recording the queue. -graphA.begin_recording(queueA); - -// Submit kernel on the recorded queue. -// EventA is a graph-limited event that can only be used for managing -// graph internal dependencies. -EventA = queueA.submit([&](handler& cgh) {...}); - -// Stop recording the queue. -graphA.end_recording(); - -event externalEventA; -auto nodeSignalExternalEvent = graphA.make_external_event(EventA, - externalEventA, - true /* updatable */); - -// Create another graph `B` -sycl::queue queueB; -sycl_ext::command_graph graphB(queueB.get_context(), queueB.get_device()); - -// Submit kernel on the recorded queue. -graphB.add([&](handler& cgh) {...}); - -auto nodeWaitExternalEvent = graphB.add_barrier({externalEventA}); - -// Finalize the modifiable graphs to create an executable graphs that can be -// submitted for execution. -auto execGraphA = graphA.finalize(); -auto execGraphB = graphB.finalize(); - -// Execute graph -// first execution -queueA.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(execGraphA); -}) -queueB.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(execGraphB); -}) - -event newExternalEventA; -execGraphA.update_external_event(nodeSignalExternalEvent, newExternalEventA); - -event newExternalEventB; -execGraphB.update_external_event(nodeWaitExternalEvent, newExternalEventB); - -// second execution -queueA.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(execGraphA); -}) -queueB.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(execGraphB); -}) ----- - - === Interaction With Other Extensions [[extension-interaction]] This section defines the interaction of `sycl_ext_oneapi_graph` with other From e0dc8d1912ee772fbd70e81da9761f4ad19ea87d Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 8 Apr 2024 11:00:27 +0100 Subject: [PATCH 03/10] Fix typos --- .../extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 750167d76d330..8ac2cd2c01b33 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -1039,7 +1039,7 @@ Parameters: * `eventList` - Zero or more events to wait for. These events can be graph-limited events (from graph recording) or general sycl events (such as events external to the graph). - If no event are provided, the barrier waits for all the previous nodes + If no events are provided, the barrier waits for all the previous nodes of the graph to complete. Returns: The command-group function object node which has been added to the graph. @@ -1728,8 +1728,8 @@ events with some restrictions (see below). External events can be obtained by either: -* Calling `command_graph::add_external_event()` on a modifiable command_graph to -add an external event node, then obtain the SYCL event object by calling +* Calling `command_graph::add_external_event()` on a modifiable `command_graph` +to add an external event node, then obtain the SYCL event object by calling `node::get_external_event()`. * Calling `queue::ext_oneapi_external_event()` on a queue in the recording From 69b1d047f4cc8c3a1120c0d6f73b8706df40f97d Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 8 Apr 2024 16:46:35 +0100 Subject: [PATCH 04/10] Restore event limitations - Restore forbidding outside graph deps that are not external events --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 8ac2cd2c01b33..366c3cc1b4696 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -1808,6 +1808,15 @@ for(int i = 0 ; i < Iterations, i++) { } ---- +==== Event Limitations + +For queue submissions that are being recorded to a modifiable `command_graph`, +the only events that can be used as parameters to `handler::depends_on()`, or as +dependent events for queue shortcuts like `queue::parallel_for()`, are +graph-limited events that have been returned from queue submissions recorded to +the same modifiable `command_graph` and external events from graphs other than +the one currently being recorded to. + === Thread Safety The new functions in this extension are thread-safe, the same as member From 898e58de76e8676713426404e6752a40264d564e Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 8 Apr 2024 17:08:42 +0100 Subject: [PATCH 05/10] Replace is_graph_limited with event info query --- .../sycl_ext_oneapi_graph.asciidoc | 38 +++++++++++++++---- 1 file changed, 30 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 366c3cc1b4696..9b1250d5da121 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -1683,18 +1683,40 @@ different from the one with which the dynamic_parameter was created. [source, c++] ---- -namespace sycl { -// New methods added to the sycl::event class -using namespace ext::oneapi::experimental; -class event { -public: - bool is_graph_limited() const; -}; -} // namespace sycl +namespace sycl::info { + enum class ext_oneapi_event_t{ + normal, + graph_limited, + graph_external + }; +} // namespace sycl::info ---- :event-class: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:interface.event +==== Event Type Query + +This extension adds the following new information descriptor to the event class +for use with `event::get_info<>()` which can be used to query the type of a SYCL +event. + +Table {counter: tableNumber}. Additions to the `sycl::event` class information descriptors. +[%header,cols="2a,a,a"] +|=== +|Event Descriptor +|Return Type +|Description + +|info::event::ext_oneapi_event_type +|info::ext_oneapi_event_t +|Returns the type of this SYCL event +|=== + +Events of type `ext_oneapi_event_t::normal` are normal SYCL events as described +in the SYCL specification. `ext_oneapi_event_t::graph_limited` and +`ext_oneapi_event_t::graph_external` events are new as part of this extension, +and are differentiated from normal SYCL events in how they can be used. + ==== Graph-Limited Events Events associated with a graph node (those returned from recorded queue From 3714f8ffc7e8f5d6fd93554bf0d42fe3f2ce8a13 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 9 Apr 2024 14:30:30 +0100 Subject: [PATCH 06/10] Addressing review comments - Rename add_barrier to add_wait_external_event - Clarify wording about edges and ext events - Minor formatting and wording changes --- .../sycl_ext_oneapi_graph.asciidoc | 61 ++++++++++++------- 1 file changed, 38 insertions(+), 23 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 9b1250d5da121..35e36adf1bd63 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -248,9 +248,9 @@ to define an edge between existing nodes, or using a Edges can also be created when explicitly adding nodes to the graph through existing SYCL mechanisms for expressing dependencies. Data dependencies from accessors to existing nodes in the graph are captured as an edge. Using -`handler::depends_on()` will also create a graph edge when passed an event -returned from a queue submission captured by a queue recording to the same -graph. +`handler::depends_on()` will also create a graph edge when passed certain kinds +of events, for more information see <>. |=== ==== Queue Recording API @@ -279,12 +279,14 @@ dependencies in one of three ways. Firstly, through buffer accessors that represent data dependencies between two command groups captured as nodes. Secondly, by using the `handler::depends_on()` mechanism inside a command group captured as a node. However, for an event passed to `handler::depends_on()` to -create an edge, it must be an event returned from a queue -submission captured by the same graph. Otherwise, a synchronous error will be -thrown with error code `invalid`. `handler::depends_on()` can be -used to express edges when a user is working with USM memory rather than SYCL -buffers. Thirdly, for a graph recorded with an in-order queue, an edge is added -automatically between two sequential command groups submitted to the in-order queue. +create an edge, it must be of a specific type (see <> for more information). Using normal +SYCL events from regular queue submissions is not allowed. +`handler::depends_on()` can be used to express edges when a user is working with +USM memory rather than SYCL buffers. For more information about event usage see +<>. For a +graph recorded with an in-order queue, an edge is added automatically between +two sequential command groups submitted to the in-order queue. |=== ==== Sub-Graph @@ -337,6 +339,7 @@ enum class node_type { ext_oneapi_barrier, host_task, external_event, + wait_external_event }; class node { @@ -653,9 +656,9 @@ public: void make_edge(node& src, node& dest); - node add_barrier(const std::vector eventList = {}); + node add_wait_external_event(const std::vector eventList = {}, const property_list& propList = {}); - node add_external_event(event& externalEvent, const property_list& propList = {}); + node add_external_event(const property_list& propList = {}); void print_graph(std::string path, bool verbose = false) const; @@ -1024,10 +1027,13 @@ Exceptions: | [source,c++] ---- -node add_barrier(const std::vector eventList = {}); +node add_wait_external_event(const std::vector eventList, const property_list& propList = {}); ---- -|Adds a barrier to the graph. +|Adds a node to the graph which waits on a number of external events. + +For more information on external events see <>. Constraints: @@ -1036,13 +1042,21 @@ Constraints: Parameters: -* `eventList` - Zero or more events to wait for. - These events can be graph-limited events (from graph recording) - or general sycl events (such as events external to the graph). - If no events are provided, the barrier waits for all the previous nodes - of the graph to complete. +* `eventList` - The events to wait for. + These events can only be external events from other graph's external event + nodes. -Returns: The command-group function object node which has been added to the graph. +* `propList` - Zero or more properties can be provided to the constructed node + via an instance of `property_list`. The `property::node::depends_on` property + can be passed here with a list of nodes to create dependency edges on. + +Returns: The external event wait node which has been added to the +graph. + +Exceptions: + +* Throws synchronously with error code `invalid` if any event in `eventList` is + not an external event. | [source,c++] @@ -1707,9 +1721,9 @@ Table {counter: tableNumber}. Additions to the `sycl::event` class information d |Return Type |Description -|info::event::ext_oneapi_event_type -|info::ext_oneapi_event_t -|Returns the type of this SYCL event +|`info::event::ext_oneapi_event_type` +|`info::ext_oneapi_event_t` +|Returns the type of this SYCL event. |=== Events of type `ext_oneapi_event_t::normal` are normal SYCL events as described @@ -1837,7 +1851,8 @@ the only events that can be used as parameters to `handler::depends_on()`, or as dependent events for queue shortcuts like `queue::parallel_for()`, are graph-limited events that have been returned from queue submissions recorded to the same modifiable `command_graph` and external events from graphs other than -the one currently being recorded to. +the one currently being recorded to. Normal SYCL events returned from regular +queue submissions outside of a graph are not valid. === Thread Safety From 9c27142c4de2a78be0f7a8d8f6d7accca021dd07 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 10 Apr 2024 11:00:35 +0100 Subject: [PATCH 07/10] Remove duplicate wording --- .../extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 35e36adf1bd63..ace4b1563dfa5 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -283,10 +283,9 @@ create an edge, it must be of a specific type (see <> for more information). Using normal SYCL events from regular queue submissions is not allowed. `handler::depends_on()` can be used to express edges when a user is working with -USM memory rather than SYCL buffers. For more information about event usage see -<>. For a -graph recorded with an in-order queue, an edge is added automatically between -two sequential command groups submitted to the in-order queue. +USM memory rather than SYCL buffers. For a graph recorded with an in-order +queue, an edge is added automatically between two sequential command groups +submitted to the in-order queue. |=== ==== Sub-Graph From df48b8a3bbb9565c84a0cd71ad1b1fd29ea802e9 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 16 Apr 2024 18:59:12 +0100 Subject: [PATCH 08/10] Minor wording update --- sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index ace4b1563dfa5..cad7618cf0c65 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -281,7 +281,7 @@ Secondly, by using the `handler::depends_on()` mechanism inside a command group captured as a node. However, for an event passed to `handler::depends_on()` to create an edge, it must be of a specific type (see <> for more information). Using normal -SYCL events from regular queue submissions is not allowed. +SYCL events from regular queue submissions outside of the graph is not allowed. `handler::depends_on()` can be used to express edges when a user is working with USM memory rather than SYCL buffers. For a graph recorded with an in-order queue, an edge is added automatically between two sequential command groups From 75f8c1427ca756dd617110c729fd7a3c36edf1af Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 22 Apr 2024 16:22:43 +0100 Subject: [PATCH 09/10] Restore using events from outside graph - Useful for warmups etc. --- .../sycl_ext_oneapi_graph.asciidoc | 30 +++++++++---------- 1 file changed, 14 insertions(+), 16 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index cad7618cf0c65..f51d3b503d78a 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -280,8 +280,7 @@ represent data dependencies between two command groups captured as nodes. Secondly, by using the `handler::depends_on()` mechanism inside a command group captured as a node. However, for an event passed to `handler::depends_on()` to create an edge, it must be of a specific type (see <> for more information). Using normal -SYCL events from regular queue submissions outside of the graph is not allowed. +the "Event Class Modifications" section>> for more information). `handler::depends_on()` can be used to express edges when a user is working with USM memory rather than SYCL buffers. For a graph recorded with an in-order queue, an edge is added automatically between two sequential command groups @@ -1770,13 +1769,13 @@ to add an external event node, then obtain the SYCL event object by calling * Calling `queue::ext_oneapi_external_event()` on a queue in the recording state, which will return an external SYCL event. -External events may be used as parameters to `handler::depends_on()` or as -dependent events for queue submissions outside of the graph they -are associated with. For example, this allows enqueing work outside of the graph -part-way through the graph execution, rather than having to wait for the entire -graph to finish. +External events may be used as parameters to `handler::depends_on()`, as +dependent events for queue submissions outside of the graph they are associated +with or as parameters to `command_graph::add_wait_external_event()`. For +example, this allows enqueing work outside of the graph part-way through the +graph execution, rather than having to wait for the entire graph to finish. -External events are considered complete when all their dependent graph nodes +External events are considered complete when all of their dependent graph nodes have finished execution, and their execution status is automatically reset when the executable graph which contains them is submitted to a queue for execution. @@ -1843,15 +1842,14 @@ for(int i = 0 ; i < Iterations, i++) { } ---- -==== Event Limitations +==== Normal SYCL Events -For queue submissions that are being recorded to a modifiable `command_graph`, -the only events that can be used as parameters to `handler::depends_on()`, or as -dependent events for queue shortcuts like `queue::parallel_for()`, are -graph-limited events that have been returned from queue submissions recorded to -the same modifiable `command_graph` and external events from graphs other than -the one currently being recorded to. Normal SYCL events returned from regular -queue submissions outside of a graph are not valid. +Normal SYCL events from queue submissions outside of a `command_graph` may be +used as dependent events for queue submissions recorded to a `command_graph` or +passed to `handler::depends_on()` for nodes added explicitly to a +`command_graph`. These types of dependencies may be useful for one-off +operations that do not need to run on every graph execution but must be executed +before the main graph commands, such as warm-ups or initialization. === Thread Safety From 35788b5fa1e95835eba54fe83602da6bccdf0392 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 22 Apr 2024 16:59:19 +0100 Subject: [PATCH 10/10] Add some more external event examples --- .../sycl_ext_oneapi_graph.asciidoc | 36 +++++++++++++++++++ 1 file changed, 36 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index f51d3b503d78a..b9bff2d226c09 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -1842,6 +1842,19 @@ for(int i = 0 ; i < Iterations, i++) { } ---- +[source,c++] +---- +// Explicit API usage snippet + +// Add an external event node to Graph A +auto ExtEventNode = GraphA.add_external_event(); + +... + +// Add a wait external event node in GraphB +auto ExtWaitNode = GraphB.add_wait_external_event({ExtEventNode.get_external_event()}); +---- + ==== Normal SYCL Events Normal SYCL events from queue submissions outside of a `command_graph` may be @@ -1851,6 +1864,29 @@ passed to `handler::depends_on()` for nodes added explicitly to a operations that do not need to run on every graph execution but must be executed before the main graph commands, such as warm-ups or initialization. +[source,c++] +---- + +// Normal submission to a queue +auto NormalEvent = Queue.submit(...); + +Graph.begin_recording(Queue); + +auto GraphEvent = Queue.submit([&](handler& CGH){ + // OK! This node will depend on NormalEvent + CGH.depends_on(NormalEvent); + CGH.parallel_for(...); +}); + +Graph.end_recording(Queue); + +auto ExecGraph = Graph.finalize(); + +// The node associated with GraphEvent will be correctly ordered with respect to +// NormalEvent +Queue.ext_oneapi_graph(ExecGraph); +---- + === Thread Safety The new functions in this extension are thread-safe, the same as member