From d7f5d51f59467ae241b3b7ca7820984775c90391 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 24 Jan 2024 18:29:58 +0000 Subject: [PATCH 01/23] [SYCL][Graph] API for explicit update using indices - Adds an API to the spec for updating graph nodes using explicit indices - Also includes functionality for updating ND-range of kernel nodes - Note: Current design is only for kernel execution nodes --- .../sycl_ext_oneapi_graph.asciidoc | 382 +++++++++++++++++- 1 file changed, 374 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 4ed2abdf0e880..3dabf37871b41 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -246,9 +246,10 @@ 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 -buffer accessors to existing nodes in the graph are captured as an edge. Using +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. +returned from a queue submission captured by a queue recording to the same +graph. |=== ==== Queue Recording API @@ -328,6 +329,11 @@ class assume_buffer_outlives_graph { public: assume_buffer_outlives_graph() = default; }; + +class updatable { + public: + updatable() = default; +}; } // namespace graph namespace node { @@ -355,7 +361,23 @@ struct graphs_support; } // namespace device } // namespace info -class node {}; +template +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph); + + void register(int argIndex, const handler& cgh); + + void update_value(const ValueT& newValue); +}; + +class node { +public: + node() = delete; + + template + void update_nd_range(nd_range executionRange); +}; // State of a graph enum class graph_state { @@ -399,7 +421,10 @@ public: template<> class command_graph { public: - command_graph() = delete; + command_graph() = delete; + + void update(node& node); + void update(const std::vector& nodes); }; } // namespace ext::oneapi::experimental @@ -467,10 +492,135 @@ The `node` class provides the {crs}[common reference semantics]. [source,c++] ---- namespace sycl::ext::oneapi::experimental { - class node {}; + class node { + template + void update_nd_range(nd_range executionRange); + }; } ---- +Table {counter: tableNumber}. Member functions of the `node` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +template +void update_nd_range(nd_range executionRange); +---- +| Updates the ND-Range for this node with a new value. This new value will not +affect any executable graphs this node is part of until it is passed to the +executable graph's update function. +See <> for more information +about updating node parameters. + +Parameters: + +* `executionRange` - The new value for the ND-Range. + +Exceptions: + +* Throws with error code `nd_range` if `executionRange` is an invalid ND-Range + for this node. + +* Throws with error code `invalid` if the type of the node is not a kernel + execution. + +|=== + +==== Dynamic Parameters + +Dynamic parameters are arguments to a node's command-group which can be updated +by the user after the node has been added to a graph. Updating the value of a +dynamic parameter will be reflected in the modifiable graph which contains this +node. These updated nodes can then be passed to an executable graph to update +it with new values. + +The type of the underlying object a dynamic parameter represents is set at +compile time using a template parameter. This underlying type can be an +accessor, a pointer to a USM allocation, scalar passed by value, or a raw byte +representation of the argument. The raw byte representation is intended to +enable updating arguments set using +link:../proposed/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg]. + +Dynamic parameters are registered with nodes in a modifiable graph, with each +registration associating one or more node arguments to the dynamic parameter +instance. Registration happens inside the command-group that the node represents +, and is done using the argument index (known when argument originally set using +`set_arg()`/`set_args()`). It is valid for a node argument to be registered with +more than one dynamic parameter instance. + +Registration is done by using the argument index (known when argument +originally set using `set_arg()`/`set_args()`). + +See <> for more information +about updating node parameters. + +The `dynamic_parameter` class provides the {crs}[common reference semantics]. + +Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +dynamic_parameter(command_graph graph); +---- +|Constructs a dynamic parameter object that can be registered with node +arguments. + +Parameters: + +* `graph` - Graph containing the nodes which will have dynamic parameters. + +| +[source,c++] +---- +void register(int argIndex, const handler &cgh); +---- +|Callable only from command-group scope when adding a node to a graph +using the explicit API. Associate the dynamic parameter +with the kernel parameter of the node that encapsulates this command-group. +Parameters are associated by their index which was given when they were +registered with `handler`. + +Parameters: + +* `argIndex` - Kernel argument index starting from zero to associate with + the dynamic parameter. +* `handler` - SYCL handler object. + +Exceptions: + +* Throws with error code `invalid` if `argIndex` is not a valid argument index + for the resulting node. +* Throws with error code `invalid` if the type of the command-group is not a + kernel execution. +* Throws with error code `invalid` if the resulting node is not associated with + the graph that this dynamic parameter was created with. + +| +[source,c++] +---- +void update(const ValueT& newValue); +---- + +|Updates parameters in all nodes registered with this dynamic +parameter to `newValue`. This new value will be reflect immediately in the +modifiable graph which contains the registered nodes. It is not an +error if `newValue` is set to the current parameter value in any registered +nodes. + +Parameters: + +* `newValue` - Value to update the registered node parameters to. + +|=== + ==== Depends-On Property The API for explicitly adding nodes to a `command_graph` includes a @@ -548,6 +698,64 @@ 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. Updates to a graph take effect from the next +submission of that graph and will not affect any previous submissions or +in-flight executions of the same graph. + +The only type of nodes that are currently supported for updating in a graph are +kernel execution nodes. + +The aspects of a kernel execution node that can be configured during update are: + +* Parameters to the kernel. +* Execution ND-Range of the kernel. + +To update an executable graph, the `property::graph::updatable` 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. + +===== Individual Node Update + +Memory parameters to individual nodes in a graph in the `executable` state +can be updated between graph executions using dynamic parameters. When a +`dynamic_parameter` instance is created with a modifiable state graph it is +empty, with no associated kernel arguments. The `dynamic_parameter` +can be used to update that node argument in the future. + +Parameter updates are performed using a `dynamic_parameter` instance by +calling `dynamic_parameter::update()` to update all the parameters of nodes to +which the `dynamic_parameter` is registered. + +The other node configuration that can be updated is the ND-Range, this can +be set through `node::update_nd_range()` but does not require any prior +registration. + +These updated nodes can then be passed to +`command_graph::update()` which will update the +executable graph with the current state of the nodes. + +Since the structure of the graph became fixed when finalizing, updating +parameters on a node will not change the already defined dependencies between +nodes. This is important to note when updating buffer parameters to a node, +since no edges will be automatically created or removed based on this change. +Care should be taken that updates of buffer parameters do not change the +behavior of a graph when executed. + +For example, if there are two nodes (NodeA and NodeB) which are connected +by an edge due to a dependency on the same buffer, both nodes must have +this buffer parameter updated to the new value. This maintains the correct +data dependency and prevents unexpected behavior. To achieve this, one +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. + ==== Graph Properties [[graph-properties]] ===== No-Cycle-Check Property @@ -775,8 +983,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`, to enable the + returned executable graph to be <>. Returns: A new executable graph object which can be submitted to a queue. @@ -808,7 +1017,70 @@ the file extension is not supported or if the write operation failed. |=== -Table {counter: tableNumber}. Member functions of the `command_graph` class for queue recording. +Table {counter: tableNumber}. Member functions of the `command_graph` class for +graph update. +[cols="2a,a"] +|=== +|Member function|Description + +| +[source,c++] +---- +void update(node& node); +---- + +| Updates an executable graph node that corresponds to `node`. `node` must be a +kernel execution node. Argument values which have associated dynamic parameters +will be updated, and the ND-range of the node will be updated inside the +executable graph to reflect the current values in `node`. + +Preconditions: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `node` - The node in this graph which will be updated. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when the executable graph was + created. +* Throws with error code `invalid` if `node` is not part of the + graph. + +| +[source,c++] +---- +void update(const std::vector& nodes); +---- + +| Updates all executable graph nodes that corresponds to the nodes contained in +`nodes`. All nodes must be kernel nodes. Argument values which have associated +dynamic parameters will be updated, and the ND-range of the nodes will be +updated inside the executable graph to reflect the current values in each node. + +Preconditions: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `nodes` - The nodes in this graph which will be updated. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when the executable graph was created. +* Throws with error code `invalid` if any node in `nodes` is not part of the + graph. +|=== + +Table {counter: tableNumber}. Member functions of the `command_graph` class for +queue recording. [cols="2a,a"] |=== |Member function|Description @@ -844,6 +1116,7 @@ Exceptions: * Throws synchronously with error code `invalid` if `recordingQueue` is associated with a device or context that is different from the device and context used on creation of the graph. + | [source, c++] ---- @@ -1548,6 +1821,82 @@ submitted in its entirety for execution via ... ---- +=== Dynamic Parameter Update + +Example showing a graph with a single kernel node that is created using a kernel +bundle with `handler::set_args()` and having its node arguments updated. + +[source,c++] +---- +... + +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + +queue myQueue; +auto myContext = myQueue.get_context(); +auto myDevice = myQueue.get_device(); + +// USM allocations for kernel input/output +const size_t n = 1024; +int *ptrX = malloc_shared(n, myQueue); +int *ptrY = malloc_device(n, myQueue); + +int *ptrZ = malloc_shared(n, myQueue); +int *ptrQ = malloc_device(n, myQueue); + +// Kernel loaded from kernel bundle +const std::vector builtinKernelIds = + myDevice.get_info(); +kernel_bundle myBundle = + get_kernel_bundle(myContext, { myDevice }, builtinKernelIds); +kernel builtinKernel = myBundle.get_kernel(builtinKernelIds[0]); + +// Graph containing a two kernels node +sycl_ext::command_graph myGraph(myContext, myDevice); + +// Create graph dynamic parameters +dynamic_parameter dynParamInput(myGraph); +dynamic_parameter dynParamScalar(myGraph); + +// First node uses ptrX as an input & output parameter, with operand +// mySclar as another argument. +int myScalar = 42; +node nodeA = myGraph.add([&](handler& cgh) { + cgh.set_args(ptrX, ptrY, myScalar); + // Register nodeA dynamic parameters + dynParamInput.register(0, cgh); // Argument index 0 is ptrX + dynParamScalar.register(2, cgh); // Argument index 2 is myScalar + cgh.parallel_for(range {n}, builtinKernel); +}); + + +// Create an executable graph with the updatable property. +auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); + +// Execute graph, then update without needing to wait for it to complete +myQueue.ext_oneapi_graph(execGraph); + +// Change ptrX argument to node A to ptrZ +dynParamInput.update(ptrZ); + +// Change myScalar argument to node A to newScalar +int newScalar = 12; +dynParamScalar.update(newScalar); + +// Update nodeA in the executable graph with the new parameters +execGraph.update(nodeA); +// Execute graph again +myQueue.ext_oneapi_graph(execGraph); +myQueue.wait(); + +sycl::free(ptrX, myQueue); +sycl::free(ptrY, myQueue); +sycl::free(ptrZ, myQueue); +sycl::free(ptrQ, myQueue); + +---- + == Future Direction [[future-direction]] This section contains both features of the specification which have been @@ -1752,6 +2101,23 @@ runtime. == Issues +=== Update More Command Types + +Support updating arguments to types of nodes other that kernel execution +commands. + +**UNRESOLVED** Should be added for at least memory copy nodes, however +full scope of support needs to be designed and implemented. + +=== Updateable Property Graph Resubmission + +It has been suggested that updateable graphs could remove the dependencies +generated between graphs upon resubmission while a previous submission of the +same graph is still executing. However, this requires further design discussion +to ensure this is desired and makes sense to users. + +**UNRESOLVED** Needs more discussion + === Multi Device Graph Allow an executable graph to contain nodes targeting different devices. From bfdf7508d399e3332a1f5ded4a982c1e05f843a9 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 29 Jan 2024 10:11:12 +0000 Subject: [PATCH 02/23] Update sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc Fix typo Co-authored-by: Pablo Reble --- sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 3dabf37871b41..e852d09f4600c 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -610,7 +610,7 @@ void update(const ValueT& newValue); ---- |Updates parameters in all nodes registered with this dynamic -parameter to `newValue`. This new value will be reflect immediately in the +parameter to `newValue`. This new value will be reflected immediately in the modifiable graph which contains the registered nodes. It is not an error if `newValue` is set to the current parameter value in any registered nodes. From ac29d0a33cb9dc567409cb5b4e6051ab392e9562 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 14 Feb 2024 16:29:29 +0000 Subject: [PATCH 03/23] Fix update function name, add initial value constructor for dynamic parameter --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 17 ++++++++++++++++- 1 file changed, 16 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e852d09f4600c..57ededd0d5aff 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -366,9 +366,11 @@ class dynamic_parameter { public: dynamic_parameter(command_graph graph); + dynamic_parameter(const ValueT &initialValue, command_graph graph); + void register(int argIndex, const handler& cgh); - void update_value(const ValueT& newValue); + void update(const ValueT& newValue); }; class node { @@ -577,6 +579,19 @@ Parameters: * `graph` - Graph containing the nodes which will have dynamic parameters. +| +[source,c++] +---- +dynamic_parameter(const ValueT &initialValue, command_graph graph); +---- +|Constructs a dynamic parameter object that can be registered with node +arguments, while providing an intiial value. + +Parameters: + +* `initialValue` - Initial value for the dynamic parameter. +* `graph` - Graph containing the nodes which will have dynamic parameters. + | [source,c++] ---- From 5431de43133fe6a81bfdc3baeab2f9657aea471a Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 20 Feb 2024 16:48:58 +0000 Subject: [PATCH 04/23] Addressing review comments - Add update_range node member function for updating using a sycl::range - Swap parameter order when registering - Clarify some wording --- .../sycl_ext_oneapi_graph.asciidoc | 32 ++++++++++++++++--- 1 file changed, 27 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 57ededd0d5aff..482852e310a24 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -527,6 +527,27 @@ Exceptions: * Throws with error code `nd_range` if `executionRange` is an invalid ND-Range for this node. +* Throws with error code `invalid` if the type of the node is not a kernel + execution. + +| +[source,c++] +---- +template +void update_range(range executionRange); +---- +| Updates the execution Range for this node with a new value. This new value +will not affect any executable graphs this node is part of until it is +passed to the executable graph's update function. +See <> for more information +about updating node parameters. + +Parameters: + +* `executionRange` - The new value for the Range. + +Exceptions: + * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -585,17 +606,18 @@ Parameters: dynamic_parameter(const ValueT &initialValue, command_graph graph); ---- |Constructs a dynamic parameter object that can be registered with node -arguments, while providing an intiial value. +arguments, while providing an intiial value which is used solely for +deducing ValueT. Parameters: -* `initialValue` - Initial value for the dynamic parameter. +* `initialValue` - Initial value used for CTAD. * `graph` - Graph containing the nodes which will have dynamic parameters. | [source,c++] ---- -void register(int argIndex, const handler &cgh); +void register_with_node(const handler &cgh, int argIndex); ---- |Callable only from command-group scope when adding a node to a graph using the explicit API. Associate the dynamic parameter @@ -605,9 +627,9 @@ registered with `handler`. Parameters: +* `handler` - SYCL handler object. * `argIndex` - Kernel argument index starting from zero to associate with the dynamic parameter. -* `handler` - SYCL handler object. Exceptions: @@ -738,7 +760,7 @@ possible. ===== Individual Node Update -Memory parameters to individual nodes in a graph in the `executable` state +Parameters to individual nodes in a graph in the `executable` state can be updated between graph executions using dynamic parameters. When a `dynamic_parameter` instance is created with a modifiable state graph it is empty, with no associated kernel arguments. The `dynamic_parameter` From 961ae13d7530f38807d0c8567c7b96a48b7f13b4 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 21 Feb 2024 13:05:52 +0000 Subject: [PATCH 05/23] Remove dynamic_param ctor, fix example based on changes, minor wording updates --- .../sycl_ext_oneapi_graph.asciidoc | 38 ++++++------------- 1 file changed, 12 insertions(+), 26 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 482852e310a24..167456ba63f0b 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -364,11 +364,9 @@ struct graphs_support; template class dynamic_parameter { public: - dynamic_parameter(command_graph graph); + dynamic_parameter(command_graph graph, const ValueT ¶m); - dynamic_parameter(const ValueT &initialValue, command_graph graph); - - void register(int argIndex, const handler& cgh); + void register_with_node(int argIndex, const handler& cgh); void update(const ValueT& newValue); }; @@ -591,28 +589,16 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. | [source,c++] ---- -dynamic_parameter(command_graph graph); ----- -|Constructs a dynamic parameter object that can be registered with node -arguments. - -Parameters: - -* `graph` - Graph containing the nodes which will have dynamic parameters. - -| -[source,c++] ----- -dynamic_parameter(const ValueT &initialValue, command_graph graph); +dynamic_parameter(command_graph graph, const ValueT ¶m); ---- -|Constructs a dynamic parameter object that can be registered with node -arguments, while providing an intiial value which is used solely for -deducing ValueT. +|Constructs a dynamic parameter object that can be registered with command graph +nodes. The param argument is used only to deduce the template parameter ValueT, +which is the type of the kernel argument this `dynamic_parameter` represents. Parameters: -* `initialValue` - Initial value used for CTAD. * `graph` - Graph containing the nodes which will have dynamic parameters. +* `param` - Value used for CTAD. | [source,c++] @@ -1892,18 +1878,18 @@ kernel builtinKernel = myBundle.get_kernel(builtinKernelIds[0]); // Graph containing a two kernels node sycl_ext::command_graph myGraph(myContext, myDevice); +int myScalar = 42; // Create graph dynamic parameters -dynamic_parameter dynParamInput(myGraph); -dynamic_parameter dynParamScalar(myGraph); +dynamic_parameter dynParamInput(myGraph, ptrX); +dynamic_parameter dynParamScalar(myGraph, myScalar); // First node uses ptrX as an input & output parameter, with operand // mySclar as another argument. -int myScalar = 42; node nodeA = myGraph.add([&](handler& cgh) { cgh.set_args(ptrX, ptrY, myScalar); // Register nodeA dynamic parameters - dynParamInput.register(0, cgh); // Argument index 0 is ptrX - dynParamScalar.register(2, cgh); // Argument index 2 is myScalar + dynParamInput.register_with_node(cgh, 0); // Argument index 0 is ptrX + dynParamScalar.register_with_node(cgh, 2); // Argument index 2 is myScalar cgh.parallel_for(range {n}, builtinKernel); }); From ebc2adb74e8b27cb1812d01490951d4cfa3d5827 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 21 Feb 2024 13:16:22 +0000 Subject: [PATCH 06/23] Rework command_graph::update() description --- .../sycl_ext_oneapi_graph.asciidoc | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 167456ba63f0b..49c173f5d64ed 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -1053,9 +1053,10 @@ void update(node& node); ---- | Updates an executable graph node that corresponds to `node`. `node` must be a -kernel execution node. Argument values which have associated dynamic parameters -will be updated, and the ND-range of the node will be updated inside the -executable graph to reflect the current values in `node`. +kernel execution node. Kernel arguments and the ND-range of the node will be +updated inside the executable graph to reflect the current values in `node`. + +Updating these values will not change the structure of the graph. Preconditions: @@ -1064,7 +1065,8 @@ Preconditions: Parameters: -* `node` - The node in this graph which will be updated. +* `node` - The node with which the equivalent node in this graph will be +updated. Exceptions: @@ -1081,9 +1083,9 @@ void update(const std::vector& nodes); ---- | Updates all executable graph nodes that corresponds to the nodes contained in -`nodes`. All nodes must be kernel nodes. Argument values which have associated -dynamic parameters will be updated, and the ND-range of the nodes will be -updated inside the executable graph to reflect the current values in each node. +`nodes`. All nodes must be kernel nodes. Kernel arguments and the ND-range of +each node will be updated inside the executable graph to reflect the current +values in each node in `nodes`. Preconditions: @@ -1092,7 +1094,8 @@ Preconditions: Parameters: -* `nodes` - The nodes in this graph which will be updated. +* `nodes` - The nodes with which the equivalent nodes in this graph will be +updated. Exceptions: From e7fd524d3f13fc6f79c80482f9fd91cff9391dc5 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 21 Feb 2024 18:13:18 +0000 Subject: [PATCH 07/23] Remove duplicate node synopsis --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 10 ---------- 1 file changed, 10 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 49c173f5d64ed..c3523568f2323 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -489,16 +489,6 @@ edges. The `node` class provides the {crs}[common reference semantics]. -[source,c++] ----- -namespace sycl::ext::oneapi::experimental { - class node { - template - void update_nd_range(nd_range executionRange); - }; -} ----- - Table {counter: tableNumber}. Member functions of the `node` class. [cols="2a,a"] |=== From ae5d62dbd09f2935fe407994151962182d2f3193 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 21 Feb 2024 18:30:17 +0000 Subject: [PATCH 08/23] Improve update_range/nd_range errors, clarify register usage --- .../sycl_ext_oneapi_graph.asciidoc | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index c3523568f2323..e13f2087017e1 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -377,6 +377,9 @@ public: template void update_nd_range(nd_range executionRange); + + template + void update_range(range executionRange); }; // State of a graph @@ -512,8 +515,11 @@ Parameters: Exceptions: -* Throws with error code `nd_range` if `executionRange` is an invalid ND-Range - for this node. +* Throws with error code `invalid` if `Dimensions` does not match the dimensions + of the nd_range the kernel was originally created with. + +* Throws with error code `invalid` if the kernel node was originally created + with a `sycl::range` * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -536,6 +542,12 @@ Parameters: Exceptions: +* Throws with error code `invalid` if `Dimensions` does not match the dimensions + of the range the kernel was originally created with. + +* Throws with error code `invalid` if the kernel node was originally created + with a `sycl::nd_range` + * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -601,6 +613,9 @@ with the kernel parameter of the node that encapsulates this command-group. Parameters are associated by their index which was given when they were registered with `handler`. +This function should only be used when kernel arguments are being set explicitly +using indices, such as when using `handler::set_arg()`. + Parameters: * `handler` - SYCL handler object. From 7b4001912bdcd5d19444885ec9a98e6edaaf1723 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 26 Feb 2024 17:22:27 +0000 Subject: [PATCH 09/23] Update wording - Update wording of dynamic_parameter::update - Move error on non-kernel node registration to add() since it is unimplementable in register --- .../sycl_ext_oneapi_graph.asciidoc | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e13f2087017e1..ce6d1df73ab9e 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -626,8 +626,6 @@ Exceptions: * Throws with error code `invalid` if `argIndex` is not a valid argument index for the resulting node. -* Throws with error code `invalid` if the type of the command-group is not a - kernel execution. * Throws with error code `invalid` if the resulting node is not associated with the graph that this dynamic parameter was created with. @@ -637,11 +635,15 @@ Exceptions: void update(const ValueT& newValue); ---- -|Updates parameters in all nodes registered with this dynamic -parameter to `newValue`. This new value will be reflected immediately in the -modifiable graph which contains the registered nodes. It is not an -error if `newValue` is set to the current parameter value in any registered -nodes. +|Updates parameters in all nodes registered with this dynamic parameter to +`newValue`. This new value will be reflected immediately in the modifiable graph +which contains the registered nodes. The new value will not be reflected in any +executable graphs created from that modifiable graph until +`command_graph::update()` is called passing the modified nodes, or a new +executable graph is finalized from the modifiable graph. + +It is not an error if `newValue` is set to the current parameter value in any +registered nodes. Parameters: @@ -950,12 +952,13 @@ Exceptions: * Throws synchronously with error code `invalid` if a queue is recording commands to the graph. - * Throws synchronously with error code `invalid` if the graph wasn't created with the `property::graph::assume_buffer_outlives_graph` property and this command uses a buffer. See the <> property for more information. +* Throws with error code `invalid` if the type of the command-group is not a + kernel execution and a `dynamic_parameter` was registered inside `cgf`. | [source,c++] From 0186285f4e75b04c618c1aee791fcb8f84883212 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 27 Feb 2024 13:04:45 +0000 Subject: [PATCH 10/23] Fix minor formatting issues --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index ce6d1df73ab9e..3c3c74435ae6b 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -519,7 +519,7 @@ Exceptions: of the nd_range the kernel was originally created with. * Throws with error code `invalid` if the kernel node was originally created - with a `sycl::range` + with a `sycl::range`. * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -546,7 +546,7 @@ Exceptions: of the range the kernel was originally created with. * Throws with error code `invalid` if the kernel node was originally created - with a `sycl::nd_range` + with a `sycl::nd_range`. * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -594,8 +594,9 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. dynamic_parameter(command_graph graph, const ValueT ¶m); ---- |Constructs a dynamic parameter object that can be registered with command graph -nodes. The param argument is used only to deduce the template parameter ValueT, -which is the type of the kernel argument this `dynamic_parameter` represents. +nodes. The param argument is used only to deduce the template parameter +`ValueT`, which is the type of the kernel argument this `dynamic_parameter` +represents. Parameters: From 8cb86eeb2dfff1cbf1aebbb0dc7efbc7913f917c Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 27 Feb 2024 13:17:43 +0000 Subject: [PATCH 11/23] Remove extra spaces --- .../extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 81c823df53515..21a0f9da38eeb 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -387,7 +387,7 @@ public: class node { public: node() = delete; - + node_type get_type() const; std::vector get_predecessors() const; @@ -395,7 +395,7 @@ public: std::vector get_successors() const; static node get_node_from_event(event nodeEvent); - + template void update_nd_range(nd_range executionRange); @@ -525,7 +525,7 @@ Table {counter: tableNumber}. Member functions of the `node` class. [cols="2a,a"] |=== |Member Function|Description - + | [source,c++] ---- From 48f97704c051d8faf797dd9c43dc4b5b56a5f3ba Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 5 Mar 2024 18:22:42 +0000 Subject: [PATCH 12/23] Add aspect for querying update support --- .../sycl_ext_oneapi_graph.asciidoc | 20 ++++++++++++++++++- 1 file changed, 19 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 21a0f9da38eeb..97673908da991 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -484,7 +484,7 @@ public: } // namespace sycl ---- -=== Device Info Query +=== Querying Graph Supoport Due to the experimental nature of the extension, support is not available across all devices. The following device support query is added to the @@ -507,6 +507,24 @@ support using this graph extension. |=== +Command-graph update support may also not be available across all devices. The +following extension to the `aspect` enum class can be used to query for graph +update support: + +[source] +---- +namespace sycl { +enum class aspect { + ... + sycl_ext_oneapi_graph_update +} +} +---- + +If a SYCL device has the `sycl_ext_oneapi_graph_update` aspact, then it supports +the functionality described in the <> section. + === Node From b63d6c31d831117dd88870e52c6ebd7a132f9da1 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 6 Mar 2024 13:25:25 +0000 Subject: [PATCH 13/23] Rename update aspect --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 97673908da991..cf39d5aa760a3 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -516,14 +516,14 @@ update support: namespace sycl { enum class aspect { ... - sycl_ext_oneapi_graph_update + sycl_ext_oneapi_limited_graph } } ---- -If a SYCL device has the `sycl_ext_oneapi_graph_update` aspact, then it supports -the functionality described in the <> section. +If a SYCL device has the `sycl_ext_oneapi_limited_graph` aspect, then it +supports all graph functionality except that which is described in the +<> section. === Node From e4ec0d429e9b00e1fd82fe4ee6e36fb0bedb95bb Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 6 Mar 2024 13:28:29 +0000 Subject: [PATCH 14/23] Update aspect description --- .../extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index cf39d5aa760a3..f629cb85ffa7f 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -523,7 +523,9 @@ enum class aspect { If a SYCL device has the `sycl_ext_oneapi_limited_graph` aspect, then it supports all graph functionality except that which is described in the -<> section. +<> section. This is a +temporary aspect that we intend to remove once devices with full graph support +are more prevalent. === Node From 50db9ad8e9b8c7734c9299d9c70f1cc033ffc490 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 6 Mar 2024 13:51:43 +0000 Subject: [PATCH 15/23] updateable -> updatable --- .../extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index f629cb85ffa7f..1169cd66409d0 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -2266,9 +2266,9 @@ commands. **UNRESOLVED** Should be added for at least memory copy nodes, however full scope of support needs to be designed and implemented. -=== Updateable Property Graph Resubmission +=== Updatable Property Graph Resubmission -It has been suggested that updateable graphs could remove the dependencies +It has been suggested that updatable graphs could remove the dependencies generated between graphs upon resubmission while a previous submission of the same graph is still executing. However, this requires further design discussion to ensure this is desired and makes sense to users. From 605f56fed24214e279bf5d5db180387c3806393c Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 6 Mar 2024 16:45:28 +0000 Subject: [PATCH 16/23] Fix aspect name --- .../extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 1169cd66409d0..e4908c68e2e87 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -516,12 +516,12 @@ update support: namespace sycl { enum class aspect { ... - sycl_ext_oneapi_limited_graph + ext_oneapi_limited_graph } } ---- -If a SYCL device has the `sycl_ext_oneapi_limited_graph` aspect, then it +If a SYCL device has the `ext_oneapi_limited_graph` aspect, then it supports all graph functionality except that which is described in the <> section. This is a temporary aspect that we intend to remove once devices with full graph support From 8bcf1256ca1a436fa7b2ce2adc15f3db7a0d94a5 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 6 Mar 2024 18:23:48 +0000 Subject: [PATCH 17/23] Add example of accessor usage with dynamic parameters --- .../sycl_ext_oneapi_graph.asciidoc | 24 +++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e4908c68e2e87..17559a2bbdac7 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -2054,6 +2054,30 @@ sycl::free(ptrQ, myQueue); ---- +Example snippet showing how to use accessors with `dynamic_parameter` update: +[source,c++] +---- +sycl::buffer bufferA{...}; +sycl::buffer bufferB{...}; + +// Create graph dynamic parameter using a placeholder accessor, since the +// sycl::handler is not available here outside of the command-group scope. +dynamic_parameter dynParamAccessor(myGraph, bufferA.get_access()); + +node nodeA = myGraph.add([&](handler& cgh) { + // Create a non-placeholder accessor for the kernel + auto accA = bufferA.get_access(cgh); + cgh.set_args(accA); + // Register nodeA dynamic parameters + dynParamAccessor.register_with_node(cgh, 0); // Argument index 0 is accA + cgh.parallel_for(range {n}, builtinKernel); +}); + +... +// Update the dynamic parameter with a placeholder accessor from bufferB instead +dynParamAccessor.update(bufferB.get_access()); +---- + == Future Direction [[future-direction]] This section contains both features of the specification which have been From 8ede6b73363d2b50c82fdfae962b6c2561be3d53 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 13 Mar 2024 14:39:20 +0000 Subject: [PATCH 18/23] Remove register_with_node --- .../sycl_ext_oneapi_graph.asciidoc | 123 +++++++++++------- 1 file changed, 73 insertions(+), 50 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 3af113039e7b3..973c6db1c238b 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -481,9 +481,7 @@ namespace ext::oneapi::experimental{ template class dynamic_parameter { public: - dynamic_parameter(command_graph graph, const ValueT ¶m); - - void register_with_node(int argIndex, const handler& cgh); + dynamic_parameter(command_graph graph, const ValueT &initialValue); void update(const ValueT& newValue); }; @@ -506,12 +504,9 @@ link:../proposed/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_ker Dynamic parameters are registered with nodes in a modifiable graph, with each registration associating one or more node arguments to the dynamic parameter instance. Registration happens inside the command-group that the node represents -, and is done using the argument index (known when argument originally set using -`set_arg()`/`set_args()`). It is valid for a node argument to be registered with -more than one dynamic parameter instance. - -Registration is done by using the argument index (known when argument -originally set using `set_arg()`/`set_args()`). +, and is done when dynamic parameters are set as parameters to the kernel using +`handler::set_arg()`/`handler::set_args()`). It is valid for a node argument to +be registered with more than one dynamic parameter instance. See <> for more information about updating node parameters. @@ -526,44 +521,16 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. | [source,c++] ---- -dynamic_parameter(command_graph graph, const ValueT ¶m); +dynamic_parameter(command_graph graph, const ValueT +&initialValue); ---- |Constructs a dynamic parameter object that can be registered with command graph -nodes. The param argument is used only to deduce the template parameter -`ValueT`, which is the type of the kernel argument this `dynamic_parameter` -represents. +nodes with an initial value. Parameters: * `graph` - Graph containing the nodes which will have dynamic parameters. -* `param` - Value used for CTAD. - -| -[source,c++] ----- -void register_with_node(const handler &cgh, int argIndex); ----- -|Callable only from command-group scope when adding a node to a graph -using the explicit API. Associate the dynamic parameter -with the kernel parameter of the node that encapsulates this command-group. -Parameters are associated by their index which was given when they were -registered with `handler`. - -This function should only be used when kernel arguments are being set explicitly -using indices, such as when using `handler::set_arg()`. - -Parameters: - -* `handler` - SYCL handler object. -* `argIndex` - Kernel argument index starting from zero to associate with - the dynamic parameter. - -Exceptions: - -* Throws with error code `invalid` if `argIndex` is not a valid argument index - for the resulting node. -* Throws with error code `invalid` if the resulting node is not associated with - the graph that this dynamic parameter was created with. +* `initialValue` - Initial value of this parameter. | [source,c++] @@ -1507,6 +1474,66 @@ Exceptions: to a queue which is associated with a device or context that is different from the device and context used on creation of the graph. +| +[source,c++] +---- +template void +handler::require(ext::oneapi::experimental::dynamic_parameter< + accessor> + DynamicParamAcc) +---- + +|Requires access to a memory object associated with an accessor contained in a +dynamic parameter. + +Parameters: + +* `DynamicParamAcc` - The dynamic parameter which contains the accessor that is +required. + +Exceptions: + +* Throws synchronously with error code `invalid` if this function is called from +a command-group submitted to a queue with is currently recording to a graph. + +* Throws synchronously with error code `invalid` if this function is called from +a normal SYCL command-group submission. + +* Throws synchronously with error code `invalid` if the graph which will be +associated with the graph node resulting from this command-group submission is +different from the one with which the dynamic_parameter was created. + +| +[source,c++] +---- +template +void handler::set_arg(int ArgIndex, + ext::oneapi::experimental::dynamic_parameter &DynamicParam); +---- + +|Sets an argument to a kernel based on the value inside a dynamic parameter, and +registers that dynamic parameter with the graph node encapsulating the +submission of the command-group that calls this function. + +Parameters: + +* `ArgIndex` - The index of the kernel argument. + +* `DynamicParam` - The dynamic parameter which contains the argument. + +Exceptions: + +* Throws synchronously with error code `invalid` if this function is called from +a command-group submitted to a queue with is currently recording to a graph. + +* Throws synchronously with error code `invalid` if this function is called from +a normal SYCL command-group submission. + +* Throws synchronously with error code `invalid` if the graph which will be +associated with the graph node resulting from this command-group submission is +different from the one with which the dynamic_parameter was created. + |=== === Thread Safety @@ -1992,10 +2019,7 @@ dynamic_parameter dynParamScalar(myGraph, myScalar); // First node uses ptrX as an input & output parameter, with operand // mySclar as another argument. node nodeA = myGraph.add([&](handler& cgh) { - cgh.set_args(ptrX, ptrY, myScalar); - // Register nodeA dynamic parameters - dynParamInput.register_with_node(cgh, 0); // Argument index 0 is ptrX - dynParamScalar.register_with_node(cgh, 2); // Argument index 2 is myScalar + cgh.set_args(dynParamInput, ptrY, dynParamScalar); cgh.parallel_for(range {n}, builtinKernel); }); @@ -2037,11 +2061,10 @@ sycl::buffer bufferB{...}; dynamic_parameter dynParamAccessor(myGraph, bufferA.get_access()); node nodeA = myGraph.add([&](handler& cgh) { - // Create a non-placeholder accessor for the kernel - auto accA = bufferA.get_access(cgh); - cgh.set_args(accA); - // Register nodeA dynamic parameters - dynParamAccessor.register_with_node(cgh, 0); // Argument index 0 is accA + // Require the accessor contained in the dynamic paramter + cgh.require(dynParamAccessor); + // Set the arg on the kernel using the dynamic parameter directly + cgh.set_args(dynParamAccessor); cgh.parallel_for(range {n}, builtinKernel); }); From 6a6bf5e6a342cd99deef0a5bf264a53786854dec Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 14 Mar 2024 15:30:58 +0000 Subject: [PATCH 19/23] Clarify update blocking/wait behaviour - User is not required to wait on previous graph submissions - update() may perform a blocking wait if required. --- .../sycl_ext_oneapi_graph.asciidoc | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 973c6db1c238b..f3bf3a492765b 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -693,10 +693,12 @@ graph LR 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. Updates to a graph take effect from the next -submission of that graph and will not affect any previous submissions or -in-flight executions of the same graph. +finalize a new executable graph when only the inputs & outputs to a graph change +between submissions. + +Updates to a graph will be scheduled after any in-flight executions of the same +graph and will not affect previous submissions of the same graph. The user is +not required to wait on any previous submissions of a graph before updating it. The only type of nodes that are currently supported for updating in a graph are kernel execution nodes. @@ -1079,6 +1081,9 @@ updated inside the executable graph to reflect the current values in `node`. Updating these values will not change the structure of the graph. +Note: The implementation may perform a blocking wait during this call on any +in-flight executions of that same graph if required by the backend. + Constraints: * This member function is only available when the `command_graph` state is @@ -1108,6 +1113,11 @@ void update(const std::vector& nodes); each node will be updated inside the executable graph to reflect the current values in each node in `nodes`. +Updating these values will not change the structure of the graph. + +Note: The implementation may perform a blocking wait during this call on any +in-flight executions of that same graph if required by the backend. + Constraints: * This member function is only available when the `command_graph` state is From 40632a1320eec0483cbf44347b384203ffe25c21 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 18 Mar 2024 14:50:45 +0000 Subject: [PATCH 20/23] Addressing MR comments - Minor formatting and wording changes --- .../sycl_ext_oneapi_graph.asciidoc | 50 ++++++++++--------- 1 file changed, 26 insertions(+), 24 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index f3bf3a492765b..4f6e861855131 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -503,10 +503,10 @@ link:../proposed/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_ker Dynamic parameters are registered with nodes in a modifiable graph, with each registration associating one or more node arguments to the dynamic parameter -instance. Registration happens inside the command-group that the node represents -, and is done when dynamic parameters are set as parameters to the kernel using -`handler::set_arg()`/`handler::set_args()`). It is valid for a node argument to -be registered with more than one dynamic parameter instance. +instance. Registration happens inside the command-group that the node +represents, and is done when dynamic parameters are set as parameters to the +kernel using `handler::set_arg()`/`handler::set_args()`. It is valid for a node +argument to be registered with more than one dynamic parameter instance. See <> for more information about updating node parameters. @@ -521,8 +521,8 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. | [source,c++] ---- -dynamic_parameter(command_graph graph, const ValueT -&initialValue); +dynamic_parameter(command_graph graph, + const ValueT &initialValue); ---- |Constructs a dynamic parameter object that can be registered with command graph nodes with an initial value. @@ -716,19 +716,21 @@ possible. ===== Individual Node Update -Parameters to individual nodes in a graph in the `executable` state -can be updated between graph executions using dynamic parameters. When a -`dynamic_parameter` instance is created with a modifiable state graph it is -empty, with no associated kernel arguments. The `dynamic_parameter` -can be used to update that node argument in the future. +Parameters to individual nodes in a graph in the `executable` state can be +updated between graph executions using dynamic parameters. A `dynamic_parameter` +object is created with a modifiable state graph and an initial value for the +parameter. Dynamic parameters can then be registered with nodes in that graph +when passed to calls to `set_arg()/set_args()`. -Parameter updates are performed using a `dynamic_parameter` instance by -calling `dynamic_parameter::update()` to update all the parameters of nodes to -which the `dynamic_parameter` is registered. +Parameter updates are performed using a `dynamic_parameter` instance by calling +`dynamic_parameter::update()` to update all the parameters of nodes to which the +`dynamic_parameter` is registered. Updates will not affect any nodes which were +not registered, even if they use the same parameter value as a +`dynamic_parameter`. -The other node configuration that can be updated is the ND-Range, this can -be set through `node::update_nd_range()` but does not require any prior -registration. +The other node configuration that can be updated is the execution range of the +kernel, this can be set through `node::update_nd_range()` or +`node::update_range()` but does not require any prior registration. These updated nodes can then be passed to `command_graph::update()` which will update the @@ -1081,7 +1083,7 @@ updated inside the executable graph to reflect the current values in `node`. Updating these values will not change the structure of the graph. -Note: The implementation may perform a blocking wait during this call on any +The implementation may perform a blocking wait during this call on any in-flight executions of that same graph if required by the backend. Constraints: @@ -1491,7 +1493,7 @@ template void handler::require(ext::oneapi::experimental::dynamic_parameter< accessor> - DynamicParamAcc) + dynamicParamAcc) ---- |Requires access to a memory object associated with an accessor contained in a @@ -1499,7 +1501,7 @@ dynamic parameter. Parameters: -* `DynamicParamAcc` - The dynamic parameter which contains the accessor that is +* `dynamicParamAcc` - The dynamic parameter which contains the accessor that is required. Exceptions: @@ -1518,8 +1520,8 @@ different from the one with which the dynamic_parameter was created. [source,c++] ---- template -void handler::set_arg(int ArgIndex, - ext::oneapi::experimental::dynamic_parameter &DynamicParam); +void handler::set_arg(int argIndex, + ext::oneapi::experimental::dynamic_parameter &dynamicParam); ---- |Sets an argument to a kernel based on the value inside a dynamic parameter, and @@ -1528,9 +1530,9 @@ submission of the command-group that calls this function. Parameters: -* `ArgIndex` - The index of the kernel argument. +* `argIndex` - The index of the kernel argument. -* `DynamicParam` - The dynamic parameter which contains the argument. +* `dynamicParam` - The dynamic parameter which contains the argument. Exceptions: From 6ec17e3d5ca45ec70e6a430de595e8bb23285ead Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 25 Mar 2024 08:21:32 +0000 Subject: [PATCH 21/23] Minor Tweaks - Remove redundant "Note: " annotation. - Avoid duplicating "Executable Graph Update" section name. - Reference parameter name in entry-point error description --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 4f6e861855131..c77181049bd89 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -694,7 +694,7 @@ graph LR 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. +between submissions. Updates to a graph will be scheduled after any in-flight executions of the same graph and will not affect previous submissions of the same graph. The user is @@ -1117,7 +1117,7 @@ values in each node in `nodes`. Updating these values will not change the structure of the graph. -Note: The implementation may perform a blocking wait during this call on any +The implementation may perform a blocking wait during this call on any in-flight executions of that same graph if required by the backend. Constraints: @@ -1514,7 +1514,7 @@ a normal SYCL command-group submission. * Throws synchronously with error code `invalid` if the graph which will be associated with the graph node resulting from this command-group submission is -different from the one with which the dynamic_parameter was created. +different from the one with which `dynamicParameterAcc` was created. | [source,c++] @@ -2145,7 +2145,7 @@ if all the commands accessing this buffer use `access_mode::write` or the Note, however, that these cases require the application to disable copy-back as described in <>. -==== Executable Graph Update +==== Whole 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 From d981ebfdcd9e290d14347661f52d38fbadd50556 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 25 Mar 2024 08:24:03 +0000 Subject: [PATCH 22/23] Fix typo in section tag --- sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index c77181049bd89..1f3bc65120cfa 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -689,7 +689,7 @@ graph LR Modifiable -->|Finalize| Executable .... -==== Executable Graph Update [[executable-graph-upate]] +==== Executable Graph Update [[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 From 6fd83ff1d289a6fda2c73f204f14351495217564 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Tue, 26 Mar 2024 10:57:10 +0000 Subject: [PATCH 23/23] Minor editorial touchups --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 1f3bc65120cfa..a6bd9aeae1da0 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -427,7 +427,7 @@ void update_nd_range(nd_range executionRange); affect any executable graphs this node is part of until it is passed to the executable graph's update function. See <> for more information -about updating node parameters. +about updating kernel nodes. Parameters: @@ -454,7 +454,7 @@ void update_range(range executionRange); will not affect any executable graphs this node is part of until it is passed to the executable graph's update function. See <> for more information -about updating node parameters. +about updating kernel nodes. Parameters: @@ -1490,9 +1490,9 @@ Exceptions: [source,c++] ---- template void +AccTarget, access::placeholder IsPlaceholder> void handler::require(ext::oneapi::experimental::dynamic_parameter< - accessor> + accessor> dynamicParamAcc) ---- @@ -2035,7 +2035,6 @@ node nodeA = myGraph.add([&](handler& cgh) { cgh.parallel_for(range {n}, builtinKernel); }); - // Create an executable graph with the updatable property. auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); @@ -2301,10 +2300,10 @@ full scope of support needs to be designed and implemented. It has been suggested that updatable graphs could remove the dependencies generated between graphs upon resubmission while a previous submission of the -same graph is still executing. However, this requires further design discussion +same graph is still executing. However, this requires further design discussion to ensure this is desired and makes sense to users. -**UNRESOLVED** Needs more discussion +**UNRESOLVED** Needs more discussion. === Multi Device Graph