Skip to content

Commit

Permalink
[SYCL][Graph] Add spec wording for dynamic events
Browse files Browse the repository at this point in the history
- Adds dynamic events which are updatable
- Removes limitations on events form outside a graph
- depends_on prop can now take events
- Added get_event to get event for node in graph
  • Loading branch information
Bensuo committed May 20, 2024
1 parent 0678c5c commit 2cb1e05
Showing 1 changed file with 123 additions and 29 deletions.
152 changes: 123 additions & 29 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc
Original file line number Diff line number Diff line change
Expand Up @@ -281,11 +281,14 @@ 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.
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.
`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.

|===

==== Sub-Graph
Expand Down Expand Up @@ -565,16 +568,24 @@ class depends_on {
public:
template<typename... NodeTN>
depends_on(NodeTN... nodes);
depends_on(const event& ev);
depends_on(const std::vector<event>& events);
};
}
----

The API for explicitly adding nodes to a `command_graph` includes a
`property_list` parameter. This extension defines the `depends_on` property to
be passed here. `depends_on` defines any `node` objects for the created node to
be dependent on, and therefore form an edge with. These nodes are in addition to
the dependent nodes identified from the command-group requisites of the created
node.
be passed here. `depends_on` may be used in two ways:

* Passing nodes from the same command_graph which will create dependencies and
graph edges between those nodes and the node being added.

* Passing SYCL events will create runtime dependencies for execution of the
graph node but will only create edges if those events are associated with other
nodes in the same graph.

==== Depends-On-All-Leaves Property
[source,c++]
Expand Down Expand Up @@ -646,6 +657,8 @@ public:
void update(node& node);
void update(const std::vector<node>& nodes);
event get_event(const node& n);
};
} // namespace sycl::ext::oneapi::experimental
Expand Down Expand Up @@ -761,6 +774,15 @@ dynamic parameter for the buffer can be registered with all the nodes which
use the buffer as a parameter. Then a single `dynamic_parameter::update()` call
will maintain the graphs data dependencies.

===== Node Event Dependency Update

Event dependencies for nodes can be updated using <<dynamic-events, Dynamic
Events>> in a similar usage to Dynamic Parameters.

Event updates are performed using a `dynamic_event` instance and calling
`dynamic_event::update()` to update all the associated event dependencies of
nodes which the `dynamic_event` is associated with.

==== Graph Properties [[graph-properties]]

===== No-Cycle-Check Property
Expand Down Expand Up @@ -1094,6 +1116,15 @@ std::vector<node> get_root_nodes() const;
----
|Returns a list of all nodes in the graph which have no dependencies.

|
[source,c++]
----
event get_event(const node& n);
----
|Returns a SYCL event which represents the completion of node `n` 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.

|===

Table {counter: tableNumber}. Member functions of the `command_graph` class for
Expand Down Expand Up @@ -1571,6 +1602,89 @@ 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.

|===

=== Dynamic Events [[dynamic-events]]

[source,c++]
----
namespace ext::oneapi::experimental{
class dynamic_event{
dynamic_event();
dynamic_event(const event& ev);
void update(const event& ev);
};
}
----

Dynamic events represent SYCL events from outside of a given `command_graph`
(either obtained from normal SYCL operations or from another `command_graph`)
that nodes in that graph may depend on. Dynamic events also allow for 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
`handler::depends_on()` inside the CGF which represents the node.

[source,c++]
----
// Obtain an event from a normal queue submission
event OutsideEvent = queue.submit(...);
// Create a dynamic event to wrap that event
ext::oneapi::experimental::dynamic_event DynEvent {OutsideEvent};
// Add a graph node which depends on that dynamic event
Graph.add([&](handler& CGH){
CGH.depends_on(DynEvent);
CGH.parallel_for(...);
});
----
Dynamic events created with an event from a `command_graph` cannot then be
associated with other nodes in that same graph as this would change the shape of
the graph. Attempting to call `handler::depends_on()` with such a
`dynamic_event` in that situation will result in an error.

Dynamic events can be created with no event but must be updated with a valid
event before any executable graph which depends on that event is executed.
Failing to do so will result in an error.

Table {counter: tableNumber}. Member functions of the `dynamic_event` class.
[cols="2a,a"]
|===
|Member function|Description

|
[source,c++]
----
dynamic_event();
----

| Constructs a default `dynamic_event` which is not associated with any SYCL event.

|
[source,c++]
----
dynamic_event(const event& ev);
----

| Constructs a `dynamic_event` which is associated with the SYCL event `ev`.


|
[source,c++]
----
void update(const event& ev);
----

| Updates the SYCL event associated with `dynamic_event`. This update will be
reflected immediately in the associated modifiable graph nodes. An executable
graph can then be updated to reflect these new event dependencies using
<<executable-graph-update, Executable Graph Update>>.



|===

=== Thread Safety
Expand Down Expand Up @@ -1646,26 +1760,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<info::event::command_execution_status>()` 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

A host-side wait on a queue in the recording state is an error and will
Expand Down

0 comments on commit 2cb1e05

Please sign in to comment.