Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[SYCL][Graph][Doc] Reinstate whole graph update #343

Closed
wants to merge 2 commits into from
Closed

[SYCL][Graph][Doc] Reinstate whole graph update #343

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
b781a40
[SYCL][Graph][Doc] Reinstate whole graph update
EwanC Nov 15, 2023
398bb8f
Add finalization property to allow a graph to be updated.
EwanC Nov 29, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
212 changes: 128 additions & 84 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -326,6 +326,11 @@ class assume_buffer_outlives_graph {
public:
assume_buffer_outlives_graph() = default;
};

class updatable_graph {
public:
updatable_graph() = default;
};
} // namespace graph

namespace node {
Expand Down Expand Up @@ -394,7 +399,9 @@ public:
template<>
class command_graph<graph_state::executable> {
public:
command_graph() = delete;
command_graph() = delete;

void update(const command_graph<graph_state::modifiable>& graph);
};
} // namespace ext::oneapi::experimental

Expand Down Expand Up @@ -543,6 +550,49 @@ graph LR
Modifiable -->|Finalize| Executable
....

==== Executable Graph Update
EwanC marked this conversation as resolved.
Show resolved Hide resolved

A graph in the executable state can have the configuration of its nodes modified
using a concept called graph _update_. This avoids a user having to rebuild and
finalize a new executable graph when only the inputs & outputs to a graph
change between submissions.

The only type of node that is currently supported for updating in a graph is
kernel execution nodes.

The aspects of a kernel execution node that can be configured during update are:

* Parameters to the kernel.
* ND-Range of the kernel.

Two methods are provided by the API to the user for performing this update.
An API for updating the complete graph object, useful when the graph was
created from a library, and an individual node update API.

To update an executable graph, the `property::graph::updatable_graph` property
must have been set when the graph was created during finalization. Otherwise, an
exception will be thrown if a user tries to update an executable graph. This
guarantee allows the backend to provide a more optimized implementation, if
possible.

The `property::graph::updatable_graph` property also allows an executable graph
to be submitted for execution while a previous submission of the same
executable graph instance is still executing. This is because the ability to
change the graph inputs/outputs can remove the data race conditions that could
otherwise exist if the same executable graph was executed concurrently.

===== Whole Graph Update

The `command_graph::update()` method takes a graph in the modifiable state and
updates the executable graph to the configuration of the nodes in the
modifiable graph, a technique called _Whole Graph Update_.

The modifiable graph must have the same topology as the graph originally used
to create the executable graphs, with the nodes targeting the same devices and
added in the same order.

===== Individual Node Update

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

===== No-Cycle-Check Property
Expand Down Expand Up @@ -615,7 +665,8 @@ with `syclDevice` is not supported.

|===

Table {counter: tableNumber}. Member functions of the `command_graph` class.
Table {counter: tableNumber}. Member functions of the modifiable `command_graph`
class.
[cols="2a,a"]
|===
|Member function|Description
Expand Down Expand Up @@ -744,8 +795,9 @@ Preconditions:

Parameters:

* `propList` - Optional parameter for passing properties. No finalization
properties are defined by this extension.
* `propList` - Optional parameter for passing properties. The only property
that is valid to pass here is `property::graph::updatable_graph`, to enable
the returned executable graph to be <<executable-graph-update, updated>>.

Returns: A new executable graph object which can be submitted to a queue.

Expand Down Expand Up @@ -902,6 +954,74 @@ Exceptions:

|===

:sycl-kernel-function: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sycl-kernel-function

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

|
[source, c++]
----
void
update(const command_graph<graph_state::modifiable>& graph);
----

|Updates the executable graph node inputs & outputs from a topologically
identical modifiable graph. A topologically identical graph is one with the
same structure of nodes and edges, and the nodes added in the same order to
both graphs. Equivalent nodes in topologically identical graphs each have the
same command, targeting the same device. There is the additional limitation that
to update an executable graph every node in the graph must be either a kernel
command.

The only characteristic that can differ between two topologically identical
graphs during an update are the arguments to kernel nodes. For example,
the graph may capture different values for the USM pointers or accessors used
in the graph. It is these kernels arguments in `graph` that constitute the
inputs & outputs to update to.

Modifying the {sycl-kernel-function}[kernel function] of a kernel node between
two graphs during an update results in undefined behavior.

The effects of the update will be visible on the next submission of the
executable graph without the need for additional user synchronization.

Preconditions:

* This member function is only available when the `command_graph` state is
`graph_state::executable`.

Parameters:

* `graph` - Modifiable graph object to update graph node inputs & outputs with.
This graph must have the same topology as the original graph used on
executable graph creation.

Exceptions:

* Throws synchronously with error code `invalid` if the topology of `graph` is
not the same as the existing graph topology, or if the nodes were not added in
the same order.

* Throws synchronously with error code `invalid` if
`property::graph::updatable_graph` was not set when the executable graph
was created.

:handler-copy-functions: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#table.members.handler.copy

* Throws synchronously with error code `invalid` if `graph` contains any node
which is not a kernel command. For example, a host-task or
{handler-copy-functions}[memory operations].

* Throws synchronously with error code `invalid` if the context or device
associated with `graph` does not match that of the `command_graph` being
updated.

|===

=== Queue Class Modifications

:queue-class: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:interface.queue.class
Expand Down Expand Up @@ -1060,8 +1180,8 @@ handler::ext_oneapi_graph(command_graph<graph_state::executable>& graph)
----

|Invokes the execution of a graph. Only one instance of `graph` may be executing,
or pending execution, at any time. Concurrent graph execution can be achieved by
finalizing a graph in modifiable state into multiple graphs in executable state.
or pending execution, at any time unless `property::graph::updatable_graph` is
set.

Parameters:

Expand All @@ -1074,7 +1194,8 @@ Exceptions:
from the device and context used on creation of the graph.

* Throws synchronously with error code `invalid` if a previous submission of
`graph` has yet to complete execution.
`graph` has yet to complete execution and `property::graph::updatable_graph`
is not set on `graph`.
|===

=== Thread Safety
Expand Down Expand Up @@ -1591,83 +1712,6 @@ auto node = graph.add([&](sycl::handler& cgh){
});
----

==== Executable Graph Update

A graph in the executable state can have each nodes inputs & outputs updated
using the `command_graph::update()` method. This takes a graph in the
modifiable state and updates the executable graph to use the node input &
outputs of the modifiable graph, a technique called _Whole Graph Update_. The
modifiable graph must have the same topology as the graph originally used to
create the executable graphs, with the nodes targeting the same devices and
added in the same order.

:sycl-kernel-function: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sycl-kernel-function

Table {counter: tableNumber}. Member functions of the `command_graph` class (executable graph update).
[cols="2a,a"]
|===
|Member function|Description

|
[source, c++]
----
void
update(const command_graph<graph_state::modifiable>& graph);
----


|Updates the executable graph node inputs & outputs from a topologically
identical modifiable graph. A topologically identical graph is one with the
same structure of nodes and edges, and the nodes added in the same order to
both graphs. Equivalent nodes in topologically identical graphs each have the
same command, targeting the same device. There is the additional limitation that
to update an executable graph, every node in the graph must be either a kernel
command or a host task.

The only characteristic that can differ between two topologically identical
graphs during an update are the arguments to kernel nodes. For example,
the graph may capture different values for the USM pointers or accessors used
in the graph. It is these kernels arguments in `graph` that constitute the
inputs & outputs to update to.

Differences in the following characteristics between two graphs during an
update results in undefined behavior:

* Modifying the native C++ callable of a `host task` node.
* Modifying the {sycl-kernel-function}[kernel function] of a kernel node.

The effects of the update will be visible on the next submission of the
executable graph without the need for additional user synchronization.

Preconditions:

* This member function is only available when the `command_graph` state is
`graph_state::executable`.

Parameters:

* `graph` - Modifiable graph object to update graph node inputs & outputs with.
This graph must have the same topology as the original graph used on
executable graph creation.

Exceptions:

* Throws synchronously with error code `invalid` if the topology of `graph` is
not the same as the existing graph topology, or if the nodes were not added in
the same order.

:handler-copy-functions: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#table.members.handler.copy

* Throws synchronously with error code `invalid` if `graph` contains any node
which is not a kernel command or host task, e.g.
{handler-copy-functions}[memory operations].

* Throws synchronously with error code `invalid` if the context or device
associated with `graph` does not match that of the `command_graph` being
updated.

|===

=== Features Still in Development

==== Memory Allocation Nodes
Expand Down