From b781a407e14e1286b217b88df3cd0201397ce5ee Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 15 Nov 2023 09:36:00 +0000 Subject: [PATCH 1/2] [SYCL][Graph][Doc] Reinstate whole graph update Take the whole graph update API out of the future work section to be a supported part of the spec. --- .../sycl_ext_oneapi_graph.asciidoc | 179 ++++++++++-------- 1 file changed, 100 insertions(+), 79 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 1dd5537a496f8..5cc381f6b7ae4 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -394,7 +394,9 @@ public: template<> class command_graph { public: - command_graph() = delete; + command_graph() = delete; + + void update(const command_graph& graph); }; } // namespace ext::oneapi::experimental @@ -543,6 +545,37 @@ graph LR Modifiable -->|Finalize| Executable .... +==== Executable Graph Update + +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. + +===== 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 @@ -615,7 +648,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 @@ -902,6 +936,70 @@ 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); +---- + +|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. + +: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 @@ -1591,83 +1689,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); ----- - - -|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 From 398bb8f63414ab34e8c71db4ff3efe41d92de665 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 29 Nov 2023 08:50:57 +0000 Subject: [PATCH 2/2] Add finalization property to allow a graph to be updated. --- .../sycl_ext_oneapi_graph.asciidoc | 33 ++++++++++++++++--- 1 file changed, 28 insertions(+), 5 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 5cc381f6b7ae4..060d96f4c28da 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -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 { @@ -564,6 +569,18 @@ 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 @@ -778,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 <>. Returns: A new executable graph object which can be submitted to a queue. @@ -988,6 +1006,10 @@ Exceptions: 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 @@ -1158,8 +1180,8 @@ handler::ext_oneapi_graph(command_graph& 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: @@ -1172,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