Skip to content

Commit

Permalink
Add some specific examples and detail about event usage
Browse files Browse the repository at this point in the history
- Also includes basic API examples for both explicit and recording
  • Loading branch information
Bensuo committed Aug 19, 2024
1 parent 71116ea commit 8836107
Showing 1 changed file with 102 additions and 0 deletions.
102 changes: 102 additions & 0 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc
Original file line number Diff line number Diff line change
Expand Up @@ -290,6 +290,36 @@ SYCL events>>) will not create edges in the graph, but will create runtime
dependencies for a graph node on those other events.
|===

===== Explicit API Example

Simple example that shows using the explicit API to add two nodes to a graph
with the <<depends-on-property, depends_on property>> used to define
dependencies between them.

[source, c++]
----
namespace sycl_ext = sycl::ext::oneapi::experimental;
sycl_ext::command_graph Graph {SyclContext, SyclDevice};
sycl_ext::node NodeA = Graph.add(
[&](sycl::handler& CGH){
CGH.parallel_for(...);
}
);
Graph.add(
[&](sycl::handler& CGH){
CGH.parallel_for(...);
},
sycl_ext::property::node::depends_on{NodeA}
);
sycl_ext::command_graph<sycl_ext::graph_state::executable> ExecGraph = Graph.finalize();
SyclQueue.ext_oneapi_graph(ExecGraph);
----

==== Queue Recording API

When using the record & replay API to construct a graph by recording a queue,
Expand Down Expand Up @@ -328,6 +358,37 @@ the in-order queue.

|===

===== Queue Recording API Example

Simple example that shows using the Queue Recording API to add two nodes to a
graph with a `sycl::event` used to define the dependency between them.

[source, c++]
----
namespace sycl_ext = sycl::ext::oneapi::experimental;
sycl_ext::command_graph Graph {SyclContext, SyclDevice};
Graph.begin_recording(SyclQueue);
sycl::event EventA = SyclQueue.submit(
[&](sycl::handler& CGH){
CGH.parallel_for(...);
}
);
SyclQueue.submit(
[&](sycl::handler& CGH){
CGH.depends_on(EventA);
CGH.parallel_for(...);
}
);
sycl_ext::command_graph<sycl_ext::graph_state::executable> ExecGraph = Graph.finalize();
SyclQueue.ext_oneapi_graph(ExecGraph);
----

==== Sub-Graph

A node in a graph can take the form of a nested sub-graph. This occurs when
Expand Down Expand Up @@ -2057,6 +2118,17 @@ default constructed `dynamic_event` with no underlying SYCL event.

=== Events

Events can be used with graphs in the following ways:

- Defining dependencies between nodes in the same graph.
- Defining dependencies between <<inter-graph-dependencies, different graphs>>.
- Obtaining <<node-execution-events, execution events for invidual nodes>>
within an executable graph, which can be waited on or used as dependencies for
eager SYCL operations.
- Creating external event dependencies for a graph, either
<<external-event-dependencies, one-time>> or <<dynamic-events, updatable between
executions>>.

==== Node-Level Execution Events [[node-execution-events]]

Events representing the completion of an individual node inside a given
Expand Down Expand Up @@ -2092,6 +2164,36 @@ if they are to be performed for every graph execution.
These events cannot be used to define dependencies between graphs. These should
instead be defined as described in <<inter-graph-dependencies, this section>>.

==== Adding External Event Dependencies To Graphs [[external-event-dependencies]]

<<event-terminology, Regular SYCL events>> can be passed as dependencies to
graph nodes to create runtime dependencies at graph execution time on regular,
eager SYCL operations. This is done in the same way as creating dependencies
between graph nodes using events, for example:

[source, c++]
----
// Submit an eager SYCL operation
event ExternalEvent = SyclQueue.submit(...);
// Record a graph node which depends on this external event
Graph.begin_recording(SyclQueue);
SyclQueue.submit(
[&](handler& CGH){
CGH.depends_on(ExternalEvent);
CGH.parallel_for(...);
}
);
Graph.end_recording();
----

This can be useful for things such as one-time warmups which must be executed
before a given graph node executes. For external dependencies which need to be
updated between graph execution, <<dynamic-events, Dynamic Events>> should be
used instead.

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

[source,c++]
Expand Down

0 comments on commit 8836107

Please sign in to comment.