From 043a8971f59b14c92dcb7c3d21aa58e629e8cb93 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 8 Aug 2024 13:16:31 +0100 Subject: [PATCH] Addressing review comments and action items --- .../sycl_ext_oneapi_graph.asciidoc | 218 +++++++++++++----- 1 file changed, 160 insertions(+), 58 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e7786d6bbf9e..58f992f7cc05 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -225,6 +225,32 @@ Table {counter: tableNumber}. Terminology. |=== +==== Event Terminology [[event-terminology]] + +:events-spec: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:interface.event + +For the purposes of clarity when talking about events in this specification we +will split events into two categories: + +- *Limited graph events*: These are events returned from a queue submission +which is recorded to a `command_graph`. These events are only valid for use with +other queue submissions recorded to the same `command_graph`. These events +cannot be waited on or used as dependencies for normal SYCL operations, or used +as dependencies for queue submissions recorded to a `command_graph` other than +the one they originated from. See the section on <> for a more detailed overview of the limitations of these events. + +- *Regular SYCL events*: These are normal SYCL events as defined in the SYCL +specification. See {events-spec}[the SYCL specification] for reference. These +include normal submissions to SYCL queue, events returned from submitting an +executable `command_graph` for execution and events obtained via +`command_graph::get_event()`. + +Please note that these definitions are only for clarity within this +specification. There are no distinct event object types, and all events +referenced in this specification are of the type `sycl::event`. Errors will be +thrown on invalid usage of limited graph events. + ==== Explicit Graph Building API When using the explicit graph building API to construct a graph, nodes and @@ -248,10 +274,15 @@ 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. +accessors to existing nodes in the graph are captured as an edge. + +Using `handler::depends_on()` inside the node's command-group function can also +be used for defining graph edges. 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 (a <>). Passing events from other sources (<>) will not create edges in the graph, but will create runtime +dependencies for a graph node on those other events. |=== ==== Queue Recording API @@ -281,11 +312,12 @@ 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. Passing events from other sources (other graph submissions, -regular SYCL submissions) will not create edges in the graph, but will create -runtime dependencies for a graph node on those other events. 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. +the same graph (a <>). Passing events +from other sources (<>) will not create +edges in the graph, but will create runtime dependencies for a graph node on +those other events. 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. |=== @@ -314,10 +346,15 @@ Table {counter: tableNumber}. Device Support Aspect. |`aspect::ext_oneapi_graph` | Indicates that the device supports all the APIs described in this extension. |`aspect::ext_oneapi_limited_graph` -| Indicates that the device supports all the APIs described in this extension -except for those described in the <> section. This is a temporary aspect that we intend to remove once -devices with full graph support are more prevalent. +a| Indicates that the device supports all the APIs described in this extension +except for the following: + + * <> + * <> + +This is a temporary aspect that we intend to remove once devices with full graph +support are more prevalent. |=== @@ -404,13 +441,12 @@ std::vector get_successors() const; ---- static node get_node_from_event(event nodeEvent); ---- -|Finds the node associated with an event created from a submission to a queue - in the recording state. +|Finds the node associated with a <> +created from a submission to a queue in the recording state. Parameters: -* `nodeEvent` - Event returned from a submission to a queue in the recording - state. +* `nodeEvent` - A limited graph event from a recorded submission to this graph. Returns: Graph node that was created when the command that returned `nodeEvent` was submitted. @@ -581,13 +617,13 @@ be passed here. `depends_on` may be used in two ways: graph edges between those nodes and the node being added. * Passing SYCL events, including <>. If an event -represents a recorded node in the same graph, then a graph edge is created -between this node and the other node. Passing an event from a recorded -submission to another graph is an error (see <> for more information). For dynamic events, events obtained from -another graph via `command_graph::get_event()` or -normal SYCL operations, a runtime dependency is created between this node and -the command that is associated with the event. +is a <>, then a graph edge is created +between this node and the other node. Passing a limited graph event associated +with another graph is an error (see <> for +more information). For dynamic events, or <>, a runtime dependency is created between this node and the command that +is associated with the event. Passing a default constructed `dynamic_event` with +no associated SYCL event will result in a synchronous error being thrown. The only permitted types for `NodeTN` and `EventTN` are `node` and `event`/`dynamic_event` respectively. @@ -919,6 +955,32 @@ graph finalized with profiling enabled is longer than that of a graph without profiling capability. An error will be thrown when attempting to profile an event from a graph submission that was created without this property. +==== Requires-Execution-Event Property [[requires-execution-event]] + +[source,c++] +---- +namespace sycl::ext::oneapi::experimental::property::graph { +class requires_execution_event { + public: + requires_execution_event() = default; +}; +} +---- + +The `property::graph::requires_execution_event` property is used to indicate +that the user intends to obtain events for the execution of specific nodes in an +executable-state graph using `command_graph::get_event()`. + +This property can be used with the following functions: + +* All overloads of `command_graph::add()` - this will +allow obtaining an execution event for the specific node added by this +function call. + +* All overloads of `command_graph::begin_recording()` - +this will allowing obtaining an execution event for every node added to this +queue before `end_recording()` is called. + ==== Graph Member Functions Table {counter: tableNumber}. Constructor of the `command_graph` class. @@ -1083,12 +1145,12 @@ Exceptions: kernel execution and a `dynamic_parameter` was registered inside `cgf`. * Throws synchronously with error code `invalid` if an `event` dependency is - passed via the `depends_on` property and that dependency comes from a recorded - submission to a different graph. + passed via the `depends_on` property and that dependency is a + <>. * Throws synchronously with error code `invalid` if a `node` dependency is - passed via the `depends_on` property and that dependency comes from a different - graph. + passed via the `depends_on` property and that dependency comes from a + different graph. | [source,c++] @@ -1198,9 +1260,11 @@ std::vector get_root_nodes() const; ---- event get_event(const node& node); ---- -|Returns a SYCL event which represents the completion of node `node` which is -valid only for the next execution of the graph. This event can be used as a -dependency in the same way as normal SYCL events. +|Returns a <> which represents the +completion of node `node` which is valid only for the most recent execution of +the graph. This event can be used as a dependency in the same way as normal SYCL +events. Nodes must have been created using the <> property to allow obtaining an event here. Constraints: @@ -1216,6 +1280,9 @@ Exceptions: * Throws synchronously with error code `invalid` if `node` is not a node within the graph. +* Throws synchronously with error code `invalid` if `node` was not created with +`property::graph::requires_execution_event`. + |=== Table {counter: tableNumber}. Member functions of the `command_graph` class for @@ -1358,8 +1425,11 @@ begin_recording(queue& recordingQueue, ---- |Synchronously changes the state of `recordingQueue` to the -`queue_state::recording` state. This operation is a no-op if `recordingQueue` -is already in the `queue_state::recording` state. +`queue_state::recording` state. If `recordingQueue` is already in the +`queue_state::recording` state calling this function will not change the state, +but will reflect any changes in the properties passed via `propList`. Queues +which are in the recording state will return <> from submissions to that queue. Parameters: @@ -1368,7 +1438,10 @@ Parameters: instance. * `propList` - Optional parameter for passing properties. Properties for - the `command_graph` class are defined in <>. + the `command_graph` class are defined in <>. When `begin_recording()` has been called multiple times for the + same queue, only the most recently passed property list will apply to + subsequent queue operations. Exceptions: @@ -1388,8 +1461,11 @@ begin_recording(const std::vector& recordingQueues, ---- |Synchronously changes the state of each queue in `recordingQueues` to the -`queue_state::recording` state. This operation is a no-op for any queue in -`recordingQueues` that is already in the `queue_state::recording` state. +`queue_state::recording` state. If any of `recordingQueues` is already in the +`queue_state::recording` state calling this function will not change the state, +but will reflect any changes in the properties passed via `propList`. Queues +which are in the recording state will return <> from submissions to that queue. Parameters: @@ -1398,7 +1474,10 @@ Parameters: instance. * `propList` - Optional parameter for passing properties. Properties for - the `command_graph` class are defined in <>. + the `command_graph` class are defined in <>. When `begin_recording()` has been called multiple times for the + same queue, only the most recently passed property list will apply to + subsequent queue operations. Exceptions: @@ -1518,10 +1597,12 @@ submitted command-groups being immediately scheduled for asynchronous execution. The alternative `queue_state::recording` state is used for graph construction. Instead of being scheduled for execution, command-groups submitted to the queue -are recorded to a graph object as new nodes for each submission. After recording -has finished and the queue returns to the executing state, the recorded commands are -not then executed, they are transparent to any following queue operations. The state -of a queue can be queried with `queue::ext_oneapi_get_state()`. +are recorded to a graph object as new nodes for each submission. Queues which +are in the recording state will return <> from submissions to that queue. After recording has finished and the +queue returns to the executing state, the recorded commands are not then +executed, they are transparent to any following queue operations. The state of a +queue can be queried with `queue::ext_oneapi_get_state()`. .Queue State Diagram [source, mermaid] @@ -1671,6 +1752,11 @@ queue::ext_oneapi_barrier(const std::vector& waitList); This function has the same semantics as `ext_oneapi_barrier(const std::vector&)`. + +Exceptions: + +* Throws synchronously with error code `invalid` if any of `waitList` is a +default constructed `dynamic_event` with no associated SYCL event. |=== ==== New Handler Member Functions @@ -1787,6 +1873,9 @@ a normal SYCL command-group submission. associated with `waitList` came from the same graph that the graph node resulting from this command-group submission is associated with. +* Throws synchronously with error code `invalid` if any of `waitList` is a +default constructed `dynamic_event` with no underlying SYCL event. + | [source,c++] ---- @@ -1810,6 +1899,9 @@ a normal SYCL command-group submission. with `depEvent` came from the same graph that the graph node resulting from this command-group submission is associated with. +* Throws synchronously with error code `invalid` if `depEvent` is a default +constructed `dynamic_event` with no underlying SYCL event. + | [source,c++] ---- @@ -1833,9 +1925,14 @@ a normal SYCL command-group submission. objects associated with `depEvents` came from the same graph that the graph node resulting from this command-group submission is associated with. +* Throws synchronously with error code `invalid` if any of `depEvents` is a +default constructed `dynamic_event` with no underlying SYCL event. |=== -=== Dynamic Events [[dynamic-events]] +=== Events + + +==== Dynamic Events [[dynamic-events]] [source,c++] ---- @@ -1849,11 +1946,11 @@ namespace ext::oneapi::experimental { } ---- -Dynamic events represent SYCL events from outside of a given `command_graph` -which nodes in that graph may depend on. These events are either obtained from -normal SYCL operations or from another `command_graph` via `get_event()`. The -`dynamic_event` object enables these dependent events to be updated between -graph executions. +Dynamic events represent <> from +outside of a given `command_graph` which nodes in that graph may depend on. +These events are either obtained from normal SYCL operations or from another +`command_graph` via `get_event()`. The `dynamic_event` object enables these +dependent events to be updated between graph executions. Dynamic events can be used to add dependencies to a graph node in the same way that regular SYCL events can, by passing them as parameters to @@ -1873,7 +1970,7 @@ Graph.add([&](handler& CGH){ CGH.parallel_for(...); }); ---- -Dynamic events created with an event from a `command_graph` cannot then be +Dynamic events created with a regular SYCL event from a `command_graph` cannot then be associated with other nodes in that same graph as this could be used change the topology of the graph. Attempting to call `handler::depends_on()` with such a `dynamic_event` in that situation will result in an error. @@ -1895,7 +1992,8 @@ Table {counter: tableNumber}. Member functions of the `dynamic_event` class. dynamic_event(); ---- -| Constructs a default `dynamic_event` which is not associated with any SYCL event. +| Constructs a default `dynamic_event` which is not associated with any SYCL +event. | [source,c++] @@ -1903,7 +2001,8 @@ dynamic_event(); dynamic_event(const event& syclEvent); ---- -| Constructs a `dynamic_event` which is associated with the SYCL event `syclEvent`. +| Constructs a `dynamic_event` which is associated with the SYCL event +`syclEvent`. Parameters: @@ -1932,12 +2031,15 @@ Parameters: Exceptions: -* Throws synchronously with error code `invalid` if `syclEvent` is an event -obtained from a queue submission which was recorded to the same graph any of the -`node` objects associated with this `dynamic_event` are from. +* Throws synchronously with error code `invalid` if `syclEvent` is a +<> obtained from the same executable +graph any of the `node` objects associated with this `dynamic_event` are from. * Throws synchronously with error code `invalid` if `syclEvent` is an event -returned from enqueuing a `host_task`. +returned from enqueuing a `host_task`. + +* Throws synchronously with error code `invalid` if `syclEvent` is a +<>. |=== @@ -2017,8 +2119,8 @@ they do in non-recording mode. ==== Event Limitations [[event-limitations]] -Other limitations on the events returned from a submission to a queue in the -recording state are: +The limitations on the <> 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 @@ -2027,8 +2129,8 @@ recording state are: - 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`. +- Using the event as a dependency outside of the recording scope will throw +synchronously with error code `invalid`. ==== Queue Limitations