From 0c7012da4741a919b66c6bfd65cad88980cddc1a Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Wed, 1 Nov 2023 17:34:55 +0000 Subject: [PATCH 01/22] [SYCL][Graph] Add APIs for updating single node parameters - New API and wording for updating single parameters - Add some queries to node and modifiable graphs to support more mixed use cases - Add deleted constructor to node which was missing --- .../sycl_ext_oneapi_graph.asciidoc | 144 +++++++++++++++++- 1 file changed, 139 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 4204515123ca4..5ad873f9fc810 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -313,6 +313,14 @@ enum class graph_support_level { emulated }; +enum class node_type { + empty = 0, + kernel, + memcpy, + memset, + memfill, +}; + namespace property { namespace graph { @@ -353,7 +361,21 @@ struct graphs_support; } // namespace device } // namespace info -class node {}; + +class node { + public: + class dynamic_parameter{ + dynamic_parameter() = delete; + const std::string& getLabel() const; + }; + + node() = delete; + + node_type getType() const; + std::vector getDependencies() const; + template + dynamic_parameter registerDynamicParameter(T Param, std::string label = "") const; +}; // State of a graph enum class graph_state { @@ -389,12 +411,17 @@ public: void make_edge(node& src, node& dest); void print_graph(std::string path, bool verbose = false) const; + + std::vector getNodes() const; }; template<> class command_graph { public: - command_graph() = delete; + command_graph() = delete; + + template + void updateParameter(node Node, dynamic_parameter DynamicParam, T NewValue); }; } // namespace ext::oneapi::experimental @@ -459,12 +486,75 @@ edges. The `node` class provides the {crs}[common reference semantics]. +==== Dynamic Parameters + +Dynamic parameters are inputs to a node's command-group which can be updated by +the user after the node has been added to a graph. They are updated through an +executable graph which contains the node associated with a dynamic parameter +object. + +Dynamic parameters may represent either a `sycl::buffer` or a pointer to a USM +allocation. When dynamic parameters are updated the new value must be +compatible with the existing value. For USM allocations this must be a pointer +to a valid USM allocation in the same context, and for buffers the new value +must have a total memory size that is greater than or equal to the original +value. + +Dynamic parameters can optionally have a user-specified string label associated +with them. This can provide meaningful names to help track what the parameters +are when updating them. These labels are not used by the SYCL runtime in any +way. + + +==== Node Member Functions + +Table {counter: tableNumber}. Member functions of the `node` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| [source,c++] ---- -namespace sycl::ext::oneapi::experimental { - class node {}; -} +node_type getType() const; +---- +|Returns a value representing the type of command this node represents. + +| +[source,c++] ---- +std::vector getDependencies() const; +---- +|Returns a list of the nodes which this node depends on. + +| +[source,c++] +---- +template +dynamic_parameter registerDynamicParameter(T param, std::string label = "") const; +---- +|Registers `param` as a dynamic parameter which can be updated in the future. +See for more information about updating node parameters. + +Parameters: + +* `param`` - The parameter used in the node which should be registered for +updating. It can be either a `sycl::buffer` or a pointer to a USM allocation. + +* `label` - An optional string label which can be used to help identify the +parameter which has beeen registered as dynamic. + +Returns: A `dynamic_parameter` object which can be used to reference this +parameter in future when updating it. + +Exceptions: + +* Throws with error code `invalid` if `param` has not been used in the node. + +* Throws with error code `invalid` if `param` is not a `sycl::buffer` or a +pointer to a USM allocation. +|=== + ==== Depends-On Property @@ -536,6 +626,20 @@ templating on state to make the class strongly typed, with the default template argument being `graph_state::modifiable` to reduce code verbosity on construction. +==== Executable Graph Update + +Memory parameters to individual nodes in a graph in the `executable` state +can be updated between graph executions using dynamic parameters. Updates +to these parameters will take effect from the next submission of a graph +and will not affect any previous submissions or in-flight executions of the +same graph. + +Parameters are associated with a `dynamic_parameter` object which is later +used to reference it when calling +`command_graph::updateDynamicParameter()`. This +prevents the user from having to keep track of the old parameter itself +until the update is performed. + .Graph State Diagram [source, mermaid] .... @@ -775,6 +879,36 @@ Exceptions: * Throws synchronously with error code `invalid` if the path is invalid or the file extension is not supported or if the write operation failed. +| +[source,c++] +---- +template +void updateParameter(node target, dynamic_parameter dynamicParam, T newValue); +---- + +| Updates a parameter in a target node with a new value. This new value will +not affected any prior submissions of this graph and will take affect only +for future submissions. + +Preconditions: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `target` - The node in this graph which will have its parameter value updated. + +* `dynamicParam` - The object which represents the parameter to be updated. + +* `newValue` - The new value for the parameter being updated. + +Exceptions: + +* Throws with error code `invalid` if `target` is not part of the graph. +* Throws with error code `invalid` if `dynamicParam` was not registered with `target`. +* Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. + |=== Table {counter: tableNumber}. Member functions of the `command_graph` class for queue recording. From c9a8b7754702b0ab0a6865a9cba43e3eff3aebc3 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 2 Nov 2023 13:14:15 +0000 Subject: [PATCH 02/22] Expand node_type enum --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 14 +++++++++----- 1 file changed, 9 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 5ad873f9fc810..e45ff81ba5cbe 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -314,11 +314,15 @@ enum class graph_support_level { }; enum class node_type { - empty = 0, - kernel, - memcpy, - memset, - memfill, + empty = 0, + subgraph, + kernel, + memcpy, + memset, + memfill, + prefetch, + memadvise, + ext_oneapi_barrier, }; namespace property { From 0e9b32198364999e1c49fa90f1894d451e7647de Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 2 Nov 2023 13:30:33 +0000 Subject: [PATCH 03/22] Detail compatible parameters in function defintion --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 8 +++++++- 1 file changed, 7 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 e45ff81ba5cbe..4cbe3c37e9093 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -911,7 +911,13 @@ Exceptions: * Throws with error code `invalid` if `target` is not part of the graph. * Throws with error code `invalid` if `dynamicParam` was not registered with `target`. -* Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. +* Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. Parameters are considered incompatible if: + +** For pointers to USM allocations `newValue` must be a pointer type. + +** For buffers `newValue` must have a total memory size that is greater +than or equal to the parameter represented by `dynamicParam`. + |=== From a93290ce7a35bd681b9510cff035a3277ea8a737 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Fri, 3 Nov 2023 15:36:11 +0000 Subject: [PATCH 04/22] Add wording about dynamic params not changing graph structure --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 4cbe3c37e9093..f078f299c88c5 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -644,6 +644,18 @@ used to reference it when calling prevents the user from having to keep track of the old parameter itself until the update is performed. +Since the structure of the graph became fixed when finalizing, updating +parameters on a node in a graph in the `executable` state 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 behaviour 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 behaviour. + .Graph State Diagram [source, mermaid] .... From f3408b02f66204ffad51d3b16fc09fe3b1a31fcc Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Mon, 6 Nov 2023 13:57:11 +0000 Subject: [PATCH 05/22] Fix typos and add new queries --- .../sycl_ext_oneapi_graph.asciidoc | 31 +++++++++++++++++-- 1 file changed, 28 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 f078f299c88c5..e3d3fa4b3dd70 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -377,6 +377,8 @@ class node { node_type getType() const; std::vector getDependencies() const; + std::vector getDependentNodes() const; + template dynamic_parameter registerDynamicParameter(T Param, std::string label = "") const; }; @@ -417,6 +419,7 @@ public: void print_graph(std::string path, bool verbose = false) const; std::vector getNodes() const; + std::vector getRootNodes() const; }; template<> @@ -498,7 +501,7 @@ executable graph which contains the node associated with a dynamic parameter object. Dynamic parameters may represent either a `sycl::buffer` or a pointer to a USM -allocation. When dynamic parameters are updated the new value must be +allocation. When dynamic parameters are updated, the new value must be compatible with the existing value. For USM allocations this must be a pointer to a valid USM allocation in the same context, and for buffers the new value must have a total memory size that is greater than or equal to the original @@ -531,6 +534,13 @@ std::vector getDependencies() const; ---- |Returns a list of the nodes which this node depends on. +| +[source,c++] +---- +std::vector getDependentNodes() const; +---- +|Returns a list of the nodes which depend on this node. + | [source,c++] ---- @@ -542,7 +552,7 @@ See for more information about updating node parameters. Parameters: -* `param`` - The parameter used in the node which should be registered for +* `param` - The parameter used in the node which should be registered for updating. It can be either a `sycl::buffer` or a pointer to a USM allocation. * `label` - An optional string label which can be used to help identify the @@ -903,7 +913,7 @@ void updateParameter(node target, dynamic_parameter dynamicParam, T newValue); ---- | Updates a parameter in a target node with a new value. This new value will -not affected any prior submissions of this graph and will take affect only +not affect any prior submissions of this graph and will take affect only for future submissions. Preconditions: @@ -1034,6 +1044,21 @@ Exceptions: * Throws synchronously with error code `invalid` if `recordingQueue` is recording to a different graph. +| +[source,c++] +---- +std::vector getNodes() const; +---- +|Returns a list of all the nodes present in the graph. + +| +[source,c++] +---- +std::vector getRootNodes() const; +---- +|Returns a list of all nodes in the graph which have no dependencies. + + | [source, c++] ---- From 9859ac11851e8c98b865e479909181e627c379e6 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 15 Nov 2023 10:18:18 +0000 Subject: [PATCH 06/22] Minor superficial tweaks --- .../sycl_ext_oneapi_graph.asciidoc | 50 +++++++++---------- 1 file changed, 24 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 e3d3fa4b3dd70..34c024b6053b3 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -367,7 +367,7 @@ struct graphs_support; class node { - public: +public: class dynamic_parameter{ dynamic_parameter() = delete; const std::string& getLabel() const; @@ -375,9 +375,9 @@ class node { node() = delete; - node_type getType() const; - std::vector getDependencies() const; - std::vector getDependentNodes() const; + node_type getType() const; + std::vector getDependencies() const; + std::vector getDependentNodes() const; template dynamic_parameter registerDynamicParameter(T Param, std::string label = "") const; @@ -512,7 +512,6 @@ with them. This can provide meaningful names to help track what the parameters are when updating them. These labels are not used by the SYCL runtime in any way. - ==== Node Member Functions Table {counter: tableNumber}. Member functions of the `node` class. @@ -530,14 +529,14 @@ node_type getType() const; | [source,c++] ---- -std::vector getDependencies() const; +std::vector getDependencies() const; ---- |Returns a list of the nodes which this node depends on. | [source,c++] ---- -std::vector getDependentNodes() const; +std::vector getDependentNodes() const; ---- |Returns a list of the nodes which depend on this node. @@ -548,7 +547,9 @@ template dynamic_parameter registerDynamicParameter(T param, std::string label = "") const; ---- |Registers `param` as a dynamic parameter which can be updated in the future. -See for more information about updating node parameters. + +See <> for more information +about updating node parameters. Parameters: @@ -556,20 +557,19 @@ Parameters: updating. It can be either a `sycl::buffer` or a pointer to a USM allocation. * `label` - An optional string label which can be used to help identify the -parameter which has beeen registered as dynamic. +parameter which has been registered as dynamic. -Returns: A `dynamic_parameter` object which can be used to reference this +Returns: A `dynamic_parameter` object which can be used to reference this parameter in future when updating it. Exceptions: * Throws with error code `invalid` if `param` has not been used in the node. -* Throws with error code `invalid` if `param` is not a `sycl::buffer` or a -pointer to a USM allocation. +* Throws with error code `invalid` if `param` is not a `sycl::buffer` or a + pointer to a USM allocation. |=== - ==== Depends-On Property The API for explicitly adding nodes to a `command_graph` includes a @@ -640,6 +640,13 @@ templating on state to make the class strongly typed, with the default template argument being `graph_state::modifiable` to reduce code verbosity on construction. +.Graph State Diagram +[source, mermaid] +.... +graph LR + Modifiable -->|Finalize| Executable +.... + ==== Executable Graph Update Memory parameters to individual nodes in a graph in the `executable` state @@ -649,7 +656,7 @@ and will not affect any previous submissions or in-flight executions of the same graph. Parameters are associated with a `dynamic_parameter` object which is later -used to reference it when calling +used to reference it when calling `command_graph::updateDynamicParameter()`. This prevents the user from having to keep track of the old parameter itself until the update is performed. @@ -666,13 +673,6 @@ 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 behaviour. -.Graph State Diagram -[source, mermaid] -.... -graph LR - Modifiable -->|Finalize| Executable -.... - ==== Graph Properties [[graph-properties]] ===== No-Cycle-Check Property @@ -933,14 +933,13 @@ Exceptions: * Throws with error code `invalid` if `target` is not part of the graph. * Throws with error code `invalid` if `dynamicParam` was not registered with `target`. -* Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. Parameters are considered incompatible if: +* Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. + Parameters are considered incompatible if: ** For pointers to USM allocations `newValue` must be a pointer type. ** For buffers `newValue` must have a total memory size that is greater -than or equal to the parameter represented by `dynamicParam`. - - +than or equal to the parameter represented by `dynamicParam`. |=== Table {counter: tableNumber}. Member functions of the `command_graph` class for queue recording. @@ -1058,7 +1057,6 @@ std::vector getRootNodes() const; ---- |Returns a list of all nodes in the graph which have no dependencies. - | [source, c++] ---- From 6046db08d753646d503063575ab0931931738abe Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 20 Nov 2023 12:01:29 +0000 Subject: [PATCH 07/22] Give `dynamic_parameter` its own constructor --- .../sycl_ext_oneapi_graph.asciidoc | 239 ++++++++++++++---- 1 file changed, 189 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 34c024b6053b3..0097dc36283db 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -365,22 +365,25 @@ struct graphs_support; } // namespace device } // namespace info +class dynamic_parameter { +public: + template + explicit dynamic_parameter(T param, std::string label = ""); + + const std::string& get_label() const; +}; class node { public: - class dynamic_parameter{ - dynamic_parameter() = delete; - const std::string& getLabel() const; - }; - node() = delete; - node_type getType() const; - std::vector getDependencies() const; - std::vector getDependentNodes() const; + node_type get_type() const; + std::vector get_dependencies() const; + std::vector get_dependent_nodes() const; - template - dynamic_parameter registerDynamicParameter(T Param, std::string label = "") const; + void register_dynamic_parameter(dynamic_parameter dynamicParam) const; + + static node get_node_from_event(event nodeEvent); }; // State of a graph @@ -418,8 +421,8 @@ public: void print_graph(std::string path, bool verbose = false) const; - std::vector getNodes() const; - std::vector getRootNodes() const; + std::vector get_nodes() const; + std::vector get_root_nodes() const; }; template<> @@ -427,8 +430,12 @@ class command_graph { public: command_graph() = delete; - template - void updateParameter(node Node, dynamic_parameter DynamicParam, T NewValue); + void update_dynamic_parameter(const std::vector& nodes, + dynamic_parameter dynamicParam, + T newValue); + template + void update_nd_range(const std::vector& nodes, + nd_range executionRange); }; } // namespace ext::oneapi::experimental @@ -495,9 +502,9 @@ The `node` class provides the {crs}[common reference semantics]. ==== Dynamic Parameters -Dynamic parameters are inputs to a node's command-group which can be updated by -the user after the node has been added to a graph. They are updated through an -executable graph which contains the node associated with a dynamic parameter +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. They are updated through +an executable graph which contains the node associated with a dynamic parameter object. Dynamic parameters may represent either a `sycl::buffer` or a pointer to a USM @@ -512,6 +519,45 @@ with them. This can provide meaningful names to help track what the parameters are when updating them. These labels are not used by the SYCL runtime in any way. +Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +template +dynamic_parameter(T param, std::string label = "") const; +---- +|Constructs a dynamic parameter used to reference `param` as a node argument +during update. + +See <> for more information +about updating node parameters. + +Parameters: + +* `param` - The parameter used in the node which should be registered for +updating. It can be either a `sycl::buffer` or a pointer to a USM allocation. + +* `label` - An optional string label which can be used to help identify the +parameter which has been registered as dynamic. + +Exceptions: + +* Throws with error code `invalid` if `param` is not a `sycl::buffer` or a + pointer to a USM allocation. + +| +[source,c++] +---- +const std::string& get_label() const; +---- +|Returns the label associated with the dynamic parameter. + +|=== + ==== Node Member Functions Table {counter: tableNumber}. Member functions of the `node` class. @@ -522,52 +568,68 @@ Table {counter: tableNumber}. Member functions of the `node` class. | [source,c++] ---- -node_type getType() const; +node_type get_type() const; ---- |Returns a value representing the type of command this node represents. | [source,c++] ---- -std::vector getDependencies() const; +std::vector get_dependencies() const; ---- |Returns a list of the nodes which this node depends on. | [source,c++] ---- -std::vector getDependentNodes() const; +std::vector get_dependent_nodes() const; ---- |Returns a list of the nodes which depend on this node. | [source,c++] ---- -template -dynamic_parameter registerDynamicParameter(T param, std::string label = "") const; +void register_dynamic_parameter(T dynamicParam) const; ---- -|Registers `param` as a dynamic parameter which can be updated in the future. +|Associate a dynamic parameter with this node, so that an updated to the + dynamic parameter See <> for more information about updating node parameters. Parameters: -* `param` - The parameter used in the node which should be registered for -updating. It can be either a `sycl::buffer` or a pointer to a USM allocation. +* `dynamicParam` - Dynamic parameter used in the node which is registered for + updating. -* `label` - An optional string label which can be used to help identify the -parameter which has been registered as dynamic. +Exceptions: -Returns: A `dynamic_parameter` object which can be used to reference this -parameter in future when updating it. +* Throws with error code `invalid` if the memory wrapped by `dynamicParam` can not + be tied to a parameter of the node. + +* Throws with error code `invalid` if the type of the node is not a kernel + execution. + +| +[source,c++] +---- +static node get_node_from_event(event nodeEvent); +---- +| Finds the node associated with an event created from a submission to a queue + in the recording state. + +Parameters: + +* `nodeEvent` - Event returned from a submission to a queue in the recording + state. + +Returns: Graph node associated to `nodeEvent`. Exceptions: -* Throws with error code `invalid` if `param` has not been used in the node. +* Throws with error code `invalid` if `nodeEvent` is not associated with a + graph node. -* Throws with error code `invalid` if `param` is not a `sycl::buffer` or a - pointer to a USM allocation. |=== ==== Depends-On Property @@ -649,17 +711,50 @@ graph LR ==== Executable Graph Update -Memory parameters to individual nodes in a graph in the `executable` state -can be updated between graph executions using dynamic parameters. Updates -to these parameters will take effect from the next submission of a graph -and will not affect any previous submissions or in-flight executions of the -same graph. +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 node take effect from the next +submission of a graph and will not affect any previous submissions or in-flight +executions of the same graph. + +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. -Parameters are associated with a `dynamic_parameter` object which is later -used to reference it when calling -`command_graph::updateDynamicParameter()`. This -prevents the user from having to keep track of the old parameter itself -until the update is performed. +===== Whole Graph Update + +===== 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 it is associated with a memory object. +By registering the `dynamic_parameter` with a node using +`node::register_dynamic_parameter()` the memory object underlying dynamic parameter +is matched to a node argument, and the `dynamic_parameter` can then be +used to update that node argument in future. This registration is necessary +because SYCL kernel arguments don't have indices and ordering of lambda captures +isn't defined, so to we need a way to tie the underlying node argument to +a `dynamic_parameter` object. + +The update itself is then performed on the `dynamic_parameter` by calling +`command_graph::update_dynamic_parameter()`. If +a user tries to register a dynamic parameter with a node, after that dynamic +parameter has been updated with a new value, then it is the new value that +must match one of the nodes arguments for the registration to be successful, +rather than the value originally used to construct the dynamic parameter. + +The other node configuration that can be updated is the ND-Range, this can +be set through `command_graph::update_nd_range` but +does not require any prior registration. Since the structure of the graph became fixed when finalizing, updating parameters on a node in a graph in the `executable` state will not change @@ -905,16 +1000,27 @@ Exceptions: * Throws synchronously with error code `invalid` if the path is invalid or the file extension is not supported or if the write operation failed. +|=== + +Table {counter: tableNumber}. Member functions of the `command_graph` class for +graph update +[cols="2a,a"] +|=== +|Member function|Description + | [source,c++] ---- template -void updateParameter(node target, dynamic_parameter dynamicParam, T newValue); +void update_dynamic_parameter(const std::vector& nodes, + dynamic_parameter dynamicParam, + T newValue); ---- -| Updates a parameter in a target node with a new value. This new value will +| Updates a parameter in target nodes with a new value. This new value will not affect any prior submissions of this graph and will take affect only -for future submissions. +for future submissions. See <> +for more information about updating node parameters. Preconditions: @@ -923,7 +1029,7 @@ Preconditions: Parameters: -* `target` - The node in this graph which will have its parameter value updated. +* `nodes` - The node in this graph which will have their parameter value updated. * `dynamicParam` - The object which represents the parameter to be updated. @@ -931,8 +1037,10 @@ Parameters: Exceptions: -* Throws with error code `invalid` if `target` is not part of the graph. -* Throws with error code `invalid` if `dynamicParam` was not registered with `target`. +* Throws with error code `invalid` if any node in `nodes` is not part of the + graph. +* Throws with error code `invalid` if `dynamicParam` is not registered with any + node in `nodes`. * Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. Parameters are considered incompatible if: @@ -940,6 +1048,37 @@ Exceptions: ** For buffers `newValue` must have a total memory size that is greater than or equal to the parameter represented by `dynamicParam`. + +| +[source,c++] +---- +template +void update_nd_range(const std::vector& nodes, + nd_range executionRange); +---- + +| Updates the ND-Range for the nodes with a new value. This new value will +not affect any prior submissions of this graph and will take affect only +for future submissions. See <> +for more information about updating node parameters. + +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 have their ND-Range values updated. + +* `executionRange` - The new value for the ND-Range. + +Exceptions: + +* Throws with error code `invalid` if any node in `nodes` is not part of the + graph. +* Throws with error code `nd_range` if `executionRange` is an invalid ND-Range + for any node in `nodes`. |=== Table {counter: tableNumber}. Member functions of the `command_graph` class for queue recording. @@ -1046,14 +1185,14 @@ Exceptions: | [source,c++] ---- -std::vector getNodes() const; +std::vector get_nodes() const; ---- |Returns a list of all the nodes present in the graph. | [source,c++] ---- -std::vector getRootNodes() const; +std::vector get_root_nodes() const; ---- |Returns a list of all nodes in the graph which have no dependencies. From fd65441c201bbe704208a513e6f373103336dd2d Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 29 Nov 2023 08:50:57 +0000 Subject: [PATCH 08/22] Add finalization property to allow a graph to be updated. --- .../sycl_ext_oneapi_graph.asciidoc | 31 ++++++++++++++++--- 1 file changed, 26 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 0097dc36283db..a61fad07aa932 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -338,6 +338,11 @@ class assume_buffer_outlives_graph { public: assume_buffer_outlives_graph() = default; }; + +class updatable_graph { + public: + updatable_graph() = default; +}; } // namespace graph namespace node { @@ -730,6 +735,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 ===== Individual Node Update @@ -969,8 +986,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. @@ -1037,6 +1055,9 @@ Parameters: Exceptions: +* Throws synchronously with error code `invalid` if + `property::graph::updatable_graph` 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. * Throws with error code `invalid` if `dynamicParam` is not registered with any @@ -1378,9 +1399,9 @@ handler::ext_oneapi_graph(command_graph& graph) ---- |Invokes the execution of a graph. Only one instance of `graph` will -execute at any time. If `graph` is submitted multiple times, dependencies -are automatically added by the runtime to prevent concurrent executions of -an identical graph. +execute at any time. If `graph` is submitted multiple times and +`property::graph::updatable_graph` is not set, dependencies are automatically +added by the runtime to prevent concurrent executions of an identical graph. Parameters: From e0f882df99469d84f139fa3b8d0f88b7b76c058f Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 29 Nov 2023 14:26:07 +0000 Subject: [PATCH 09/22] Address PR comments --- .../sycl_ext_oneapi_graph.asciidoc | 48 +++++++++++-------- 1 file changed, 27 insertions(+), 21 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index a61fad07aa932..aa8a05ee2bc32 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -732,8 +732,8 @@ The aspects of a kernel execution node that can be configured during update are: * 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. +An API for updating the complete graph object, which is most 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 @@ -758,9 +758,9 @@ By registering the `dynamic_parameter` with a node using `node::register_dynamic_parameter()` the memory object underlying dynamic parameter is matched to a node argument, and the `dynamic_parameter` can then be used to update that node argument in future. This registration is necessary -because SYCL kernel arguments don't have indices and ordering of lambda captures -isn't defined, so to we need a way to tie the underlying node argument to -a `dynamic_parameter` object. +because SYCL kernel arguments don't have indices and the argument ordering of +lambda captures isn't defined, so we need a way to tie the underlying node +argument to a `dynamic_parameter` object. The update itself is then performed on the `dynamic_parameter` by calling `command_graph::update_dynamic_parameter()`. If @@ -783,7 +783,10 @@ of buffer parameters do not change the behaviour 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 behaviour. +data dependency, and prevents unexpected behaviour. 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 `update_dynamic_parameter()` call +with this list of nodes will maintain the graphs data dependencies. ==== Graph Properties [[graph-properties]] @@ -1018,10 +1021,26 @@ Exceptions: * Throws synchronously with error code `invalid` if the path is invalid or the file extension is not supported or if the write operation failed. +| +[source,c++] +---- +std::vector get_nodes() const; +---- +|Returns a list of all the nodes present in the graph in the order that they +were added. + +| +[source,c++] +---- +std::vector get_root_nodes() const; +---- +|Returns a list of all nodes in the graph which have no dependencies in the +order they were added to the graph. + |=== Table {counter: tableNumber}. Member functions of the `command_graph` class for -graph update +graph update. [cols="2a,a"] |=== |Member function|Description @@ -1138,6 +1157,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++] ---- @@ -1203,20 +1223,6 @@ Exceptions: * Throws synchronously with error code `invalid` if `recordingQueue` is recording to a different graph. -| -[source,c++] ----- -std::vector get_nodes() const; ----- -|Returns a list of all the nodes present in the graph. - -| -[source,c++] ----- -std::vector get_root_nodes() const; ----- -|Returns a list of all nodes in the graph which have no dependencies. - | [source, c++] ---- From 52448a502ff4f0d59dbaecb48e458d460eca482d Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 29 Nov 2023 15:45:43 +0000 Subject: [PATCH 10/22] Remove nodes list from `update_dynamic_parameter()` --- .../sycl_ext_oneapi_graph.asciidoc | 29 ++++++++----------- 1 file changed, 12 insertions(+), 17 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index aa8a05ee2bc32..edc00468bb046 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -435,8 +435,7 @@ class command_graph { public: command_graph() = delete; - void update_dynamic_parameter(const std::vector& nodes, - dynamic_parameter dynamicParam, + void update_dynamic_parameter(dynamic_parameter dynamicParam, T newValue); template void update_nd_range(const std::vector& nodes, @@ -596,8 +595,11 @@ std::vector get_dependent_nodes() const; ---- void register_dynamic_parameter(T dynamicParam) const; ---- -|Associate a dynamic parameter with this node, so that an updated to the - dynamic parameter +|Associate a dynamic parameter with a kernel parameter of this node, so that an +update to the dynamic parameter modifies the kernel parameter. All matching +kernel arguments are bound, therefore if a kernel has identical parameters +which reference the same memory object as `dynamicParam`, all those parameters +are associated to `dynamicParam` and will be affected by an update to it. See <> for more information about updating node parameters. @@ -786,7 +788,7 @@ this buffer parameter updated to the new value. This maintains the correct data dependency, and prevents unexpected behaviour. 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 `update_dynamic_parameter()` call -with this list of nodes will maintain the graphs data dependencies. +will maintain the graphs data dependencies. ==== Graph Properties [[graph-properties]] @@ -1049,15 +1051,14 @@ graph update. [source,c++] ---- template -void update_dynamic_parameter(const std::vector& nodes, - dynamic_parameter dynamicParam, +void update_dynamic_parameter(dynamic_parameter dynamicParam, T newValue); ---- -| Updates a parameter in target nodes with a new value. This new value will -not affect any prior submissions of this graph and will take affect only -for future submissions. See <> -for more information about updating node parameters. +| Updates parameters in nodes associated with `dynamicParam` to a new value. +This new value will not affect any prior submissions of this graph and will +take affect only for future submissions. See <> for more information about updating node parameters. Preconditions: @@ -1066,8 +1067,6 @@ Preconditions: Parameters: -* `nodes` - The node in this graph which will have their parameter value updated. - * `dynamicParam` - The object which represents the parameter to be updated. * `newValue` - The new value for the parameter being updated. @@ -1077,10 +1076,6 @@ Exceptions: * Throws synchronously with error code `invalid` if `property::graph::updatable_graph` 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. -* Throws with error code `invalid` if `dynamicParam` is not registered with any - node in `nodes`. * Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. Parameters are considered incompatible if: From adc0a4d4999433b12a6623962f3c6ceff0d221b0 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 30 Nov 2023 18:32:02 +0000 Subject: [PATCH 11/22] Enable update by argument index. --- .../sycl_ext_oneapi_graph.asciidoc | 186 ++++++++++++------ 1 file changed, 122 insertions(+), 64 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index edc00468bb046..e1194396517ee 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -314,7 +314,7 @@ enum class graph_support_level { }; enum class node_type { - empty = 0, + empty, subgraph, kernel, memcpy, @@ -373,7 +373,9 @@ struct graphs_support; class dynamic_parameter { public: template - explicit dynamic_parameter(T param, std::string label = ""); + dynamic_parameter(T param, std::string label = ""); + + dynamic_parameter(std::string label = ""); const std::string& get_label() const; }; @@ -382,11 +384,20 @@ class node { public: node() = delete; + /* -- parameter update interface -- */ + + void register_dynamic_parameter(dynamic_parameter dynamicParam) const; + + void register_dynamic_parameter(int argIndex, + dynamic_parameter dynamicParam) const; + + /* -- query interface -- */ + node_type get_type() const; + std::vector get_dependencies() const; - std::vector get_dependent_nodes() const; - void register_dynamic_parameter(dynamic_parameter dynamicParam) const; + std::vector get_dependent_nodes() const; static node get_node_from_event(event nodeEvent); }; @@ -435,9 +446,12 @@ class command_graph { public: command_graph() = delete; - void update_dynamic_parameter(dynamic_parameter dynamicParam, - T newValue); - template + /* -- update interface -- */ + + template + void update_dynamic_parameter(T newValue, dynamic_parameter dynamicParam); + + template void update_nd_range(const std::vector& nodes, nd_range executionRange); }; @@ -508,15 +522,16 @@ The `node` class provides the {crs}[common reference semantics]. 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. They are updated through -an executable graph which contains the node associated with a dynamic parameter +an executable graph which contains the nodes associated with a dynamic parameter object. -Dynamic parameters may represent either a `sycl::buffer` or a pointer to a USM -allocation. When dynamic parameters are updated, the new value must be -compatible with the existing value. For USM allocations this must be a pointer -to a valid USM allocation in the same context, and for buffers the new value -must have a total memory size that is greater than or equal to the original -value. +When dynamic parameters are created they are either empty or represent an +underlying memory object. This underlying memory object can be a `sycl::buffer`, +a pointer to a USM allocation, otherwise it is treated as a scalar. When dynamic +parameters are registered with nodes or updated, the new value must be +compatible with any existing value. Empty nodes are registered with kernel +parameters using the argument index, after which they are initialized to +reference the underlying memory object at that index. Dynamic parameters can optionally have a user-specified string label associated with them. This can provide meaningful names to help track what the parameters @@ -543,15 +558,26 @@ about updating node parameters. Parameters: * `param` - The parameter used in the node which should be registered for -updating. It can be either a `sycl::buffer` or a pointer to a USM allocation. +updating. It can be either a `sycl::buffer`, a pointer to a USM allocation, +or else is treated as a scalar. * `label` - An optional string label which can be used to help identify the parameter which has been registered as dynamic. -Exceptions: +| +[source,c++] +---- +dynamic_parameter(std::string label = "") const; +---- +|Constructs an empty dynamic parameter. The registration of an empty dynamic +parameter must use the kernel argument index, as there is no underlying memory +object to match kernel parameters against. -* Throws with error code `invalid` if `param` is not a `sycl::buffer` or a - pointer to a USM allocation. +See <> for more information +about updating node parameters. + +* `label` - An optional string label which can be used to help identify the +parameter which has been registered as dynamic. | [source,c++] @@ -572,58 +598,88 @@ Table {counter: tableNumber}. Member functions of the `node` class. | [source,c++] ---- -node_type get_type() const; +void register_dynamic_parameter(dynamic_parameter dynamicParam) const; ---- -|Returns a value representing the type of command this node represents. +|Associate a dynamic parameter with kernel parameter of this node. +All kernel arguments matching the memory object referenced by `dynamicParam` +are bound to `dynamicParam`. -| -[source,c++] ----- -std::vector get_dependencies() const; ----- -|Returns a list of the nodes which this node depends on. +See <> for more information +about updating node parameters. -| -[source,c++] ----- -std::vector get_dependent_nodes() const; ----- -|Returns a list of the nodes which depend on this node. +Parameters: + +* `dynamicParam` - Dynamic parameter which is to be associated with the kernel + arguments matching `param`. + +Exceptions: + +* Throws with error code `invalid` if `dynamicParam` is empty or does not match + any arguments in the node. + +* Throws with error code `invalid` if the type of the node is not a kernel + execution. | [source,c++] ---- -void register_dynamic_parameter(T dynamicParam) const; +void register_dynamic_parameter(int argIndex, + dynamic_parameter dynamicParam) const; ---- -|Associate a dynamic parameter with a kernel parameter of this node, so that an -update to the dynamic parameter modifies the kernel parameter. All matching -kernel arguments are bound, therefore if a kernel has identical parameters -which reference the same memory object as `dynamicParam`, all those parameters -are associated to `dynamicParam` and will be affected by an update to it. +|Associate a dynamic parameter with a kernel parameter of this node via the +argument index. Although argument ordering is not defined for lambda captures, +the ordering is reliable for arguments set using `handler::set_arg()` or +`handler::set_args()`. See <> for more information about updating node parameters. Parameters: -* `dynamicParam` - Dynamic parameter used in the node which is registered for - updating. +* `argIndex` - Kernel argument index starting from zero to associate + `dynamicParam` with. +* `dynamicParam` - Dynamic parameter which is associated with the kernel + argument at index `argIndex`. Exceptions: -* Throws with error code `invalid` if the memory wrapped by `dynamicParam` can not - be tied to a parameter of the node. +* Throws with error code `invalid` if `argIndex` is not a valid argument index + for the kernel. + +* Throws with error code `invalid` if `dynamicParam` wraps a memory object that + cannot be tied to the parameter of the node at `argIndex`. * Throws with error code `invalid` if the type of the node is not a kernel execution. +| +[source,c++] +---- +node_type get_type() const; +---- +|Returns a value representing the type of command this node represents. + +| +[source,c++] +---- +std::vector get_dependencies() const; +---- +|Returns a list of the nodes which this node depends on. + +| +[source,c++] +---- +std::vector get_dependent_nodes() const; +---- +|Returns a list of the nodes which depend on this node. + | [source,c++] ---- static node get_node_from_event(event nodeEvent); ---- -| Finds the node associated with an event created from a submission to a queue - in the recording state. +|Finds the node associated with an event created from a submission to a queue + in the recording state. Parameters: @@ -755,24 +811,23 @@ otherwise exist if the same executable graph was executed concurrently. 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 it is associated with a memory object. -By registering the `dynamic_parameter` with a node using -`node::register_dynamic_parameter()` the memory object underlying dynamic parameter -is matched to a node argument, and the `dynamic_parameter` can then be -used to update that node argument in future. This registration is necessary -because SYCL kernel arguments don't have indices and the argument ordering of -lambda captures isn't defined, so we need a way to tie the underlying node -argument to a `dynamic_parameter` object. - -The update itself is then performed on the `dynamic_parameter` by calling -`command_graph::update_dynamic_parameter()`. If -a user tries to register a dynamic parameter with a node, after that dynamic -parameter has been updated with a new value, then it is the new value that -must match one of the nodes arguments for the registration to be successful, -rather than the value originally used to construct the dynamic parameter. +`dynamic_parameter` instance is created it is either empty or associated with a +memory object. By registering the `dynamic_parameter` with a node using +`node::register_dynamic_parameter()`, the dynamic parameter is matched to a node +argument, and the `dynamic_parameter` can then be used to update that node +argument in future. + +The update itself is then performed using a `dynamic_parameter` instance by +calling `command_graph::update_dynamic_parameter()` to +update all the parameters of nodes in an executable graph to which the +`dynamic_parameter` is registered. If a user tries to register a dynamic +parameter with a node, after that dynamic parameter has been updated with a new +value, then it is the new value that must match one of the nodes arguments for +the registration to be successful, rather than the value originally used to +construct the dynamic parameter. The other node configuration that can be updated is the ND-Range, this can -be set through `command_graph::update_nd_range` but +be set through `command_graph::update_nd_range()` but does not require any prior registration. Since the structure of the graph became fixed when finalizing, updating @@ -1051,8 +1106,7 @@ graph update. [source,c++] ---- template -void update_dynamic_parameter(dynamic_parameter dynamicParam, - T newValue); +void update_dynamic_parameter(T newValue, dynamic_parameter dynamicParam); ---- | Updates parameters in nodes associated with `dynamicParam` to a new value. @@ -1067,10 +1121,11 @@ Preconditions: Parameters: -* `dynamicParam` - The object which represents the parameter to be updated. - * `newValue` - The new value for the parameter being updated. +* `dynamicParam` - The object which represents the node parameters to be + updated. + Exceptions: * Throws synchronously with error code `invalid` if @@ -1084,6 +1139,9 @@ Exceptions: ** For buffers `newValue` must have a total memory size that is greater than or equal to the parameter represented by `dynamicParam`. +** For scalars `newValue` must have a memory size that is equal to the +parameter represented by `dynamicParam`. + | [source,c++] ---- From 4419ea93a9d82b01fcf0004483434e3cadc23187 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 30 Nov 2023 22:30:31 +0000 Subject: [PATCH 12/22] Template dynamic parameter type --- .../sycl_ext_oneapi_graph.asciidoc | 273 +++++++++++++++--- 1 file changed, 226 insertions(+), 47 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e1194396517ee..434b21761dbdb 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -370,26 +370,47 @@ struct graphs_support; } // namespace device } // namespace info +enum class dynamic_parameter_type { + usm, + buffer, + scalar +}; + +template class dynamic_parameter { public: - template - dynamic_parameter(T param, std::string label = ""); - dynamic_parameter(std::string label = ""); const std::string& get_label() const; }; +template <> +class dynamic_parameter { +public: + dynamic_parameter(size_t size, std::string label = ""); + + size_t get_size() const; + const std::string& get_label() const; +}; + class node { public: node() = delete; /* -- parameter update interface -- */ - void register_dynamic_parameter(dynamic_parameter dynamicParam) const; + void register_dynamic_parameter(void* param, + dynamic_parameter dynamicParam) const; + template + void register_dynamic_parameter(T param, + dynamic_parameter dynamicParam) const; + template + void register_dynamic_parameter(T param, + dynamic_parameter dynamicParam) const; + template void register_dynamic_parameter(int argIndex, - dynamic_parameter dynamicParam) const; + dynamic_parameter dynamicParam) const; /* -- query interface -- */ @@ -448,8 +469,17 @@ public: /* -- update interface -- */ + void + update_dynamic_parameter(void* newValue, + dynamic_parameter dynamicParam); + + template + void update_dynamic_parameter(T newValue, + dynamic_parameter dynamicParam); + template - void update_dynamic_parameter(T newValue, dynamic_parameter dynamicParam); + void update_dynamic_parameter(T newValue, + dynamic_parameter dynamicParam); template void update_nd_range(const std::vector& nodes, @@ -525,13 +555,14 @@ by the user after the node has been added to a graph. They are updated through an executable graph which contains the nodes associated with a dynamic parameter object. -When dynamic parameters are created they are either empty or represent an -underlying memory object. This underlying memory object can be a `sycl::buffer`, -a pointer to a USM allocation, otherwise it is treated as a scalar. When dynamic -parameters are registered with nodes or updated, the new value must be -compatible with any existing value. Empty nodes are registered with kernel -parameters using the argument index, after which they are initialized to -reference the underlying memory object at that index. +The type of the underlying object a dynamic parameter represents is set at compile +time using a template parameter. This underlying memory object can be a +`sycl::buffer`, a pointer to a USM allocation, or a scalar passed by value. + +Dynamic parameters are registered with nodes, with each registration associating +a node argument to the dynamic parameter instance. An update of the dynamic +parameter instance to a new value, will then update all the parameter which have +been registered to the node. Dynamic parameters can optionally have a user-specified string label associated with them. This can provide meaningful names to help track what the parameters @@ -546,20 +577,20 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. | [source,c++] ---- -template -dynamic_parameter(T param, std::string label = "") const; +dynamic_parameter(std::string label = ""); ---- -|Constructs a dynamic parameter used to reference `param` as a node argument -during update. +|Constructs a dynamic parameter object that can be registered with node +arguments. See <> for more information about updating node parameters. -Parameters: +Preconditions: + +* This constructor is only available when the templated `dynamic_parameter_type` + is `dynamic_parameter_type::usm` or `dynamic_parameter_type::buffer`. -* `param` - The parameter used in the node which should be registered for -updating. It can be either a `sycl::buffer`, a pointer to a USM allocation, -or else is treated as a scalar. +Parameters: * `label` - An optional string label which can be used to help identify the parameter which has been registered as dynamic. @@ -567,15 +598,23 @@ parameter which has been registered as dynamic. | [source,c++] ---- -dynamic_parameter(std::string label = "") const; +dynamic_parameter(size_t size, std::string label = ""); ---- -|Constructs an empty dynamic parameter. The registration of an empty dynamic -parameter must use the kernel argument index, as there is no underlying memory -object to match kernel parameters against. +|Constructs a scalar dynamic parameter object that can be registered with node +arguments. See <> for more information about updating node parameters. +Preconditions: + +* This constructor is only available when the templated `dynamic_parameter_type` + is `dynamic_parameter_type::scalar`. + +Parameters: + +* `size` - Size in bytes of the scalar object. + * `label` - An optional string label which can be used to help identify the parameter which has been registered as dynamic. @@ -586,6 +625,18 @@ const std::string& get_label() const; ---- |Returns the label associated with the dynamic parameter. +| +[source,c++] +---- +size_t get_size() const; +---- +|Returns the size associated with the scalar dynamic parameter. + +Preconditions: + +* This constructor is only available when the templated `dynamic_parameter_type` + is `dynamic_parameter_type::scalar`. + |=== ==== Node Member Functions @@ -598,24 +649,56 @@ Table {counter: tableNumber}. Member functions of the `node` class. | [source,c++] ---- -void register_dynamic_parameter(dynamic_parameter dynamicParam) const; +void register_dynamic_parameter( + void* param, + dynamic_parameter dynamicParam) const; +---- +|Associate a USM dynamic parameter with USM kernel parameter of this node. +All USM kernel arguments matching `param` are bound to `dynamicParam`. + +See <> for more information +about updating node parameters. + +Parameters: + +* `param` - USM pointer to used as a node argument. +* `dynamicParam` - Dynamic parameter which is to be associated with the kernel + arguments matching `param`. + +Exceptions: + +* Throws with error code `invalid` if `param` does not match any arguments in + the node. + +* Throws with error code `invalid` if the type of the node is not a kernel + execution. + +| +[source,c++] +---- +template +void register_dynamic_parameter( + T param, + dynamic_parameter dynamicParam) const; ---- -|Associate a dynamic parameter with kernel parameter of this node. -All kernel arguments matching the memory object referenced by `dynamicParam` -are bound to `dynamicParam`. +|Associate a buffer dynamic parameter with buffer kernel parameter of this node. +All buffer kernel arguments matching `param` are bound to `dynamicParam`. See <> for more information about updating node parameters. Parameters: +* `param` - Buffer accessor used as a node argument. * `dynamicParam` - Dynamic parameter which is to be associated with the kernel arguments matching `param`. Exceptions: -* Throws with error code `invalid` if `dynamicParam` is empty or does not match - any arguments in the node. +* Throws with error code `invalid` if `param` does not match any arguments in + the node. + +* Throws with error code `invalid` if `param` is not a buffer accessor. * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -623,8 +706,42 @@ Exceptions: | [source,c++] ---- -void register_dynamic_parameter(int argIndex, - dynamic_parameter dynamicParam) const; +template +void register_dynamic_parameter( + T param, + dynamic_parameter dynamicParam) const; +---- +|Associate a scalar dynamic parameter with a scalar kernel parameter of this node. +All scalar kernel arguments matching `param` are bound to `dynamicParam`. + +See <> for more information +about updating node parameters. + +Parameters: + +* `param` - Scalar used as a node argument. +* `dynamicParam` - Dynamic parameter which is to be associated with the kernel + arguments matching `param`. + +Exceptions: + +* Throws with error code `invalid` if `param` does not match any arguments in + the node. + +* Throws with error code `invalid` if `param` does not match the specified on + creation of `dynamicParam`. + +* Throws with error code `invalid` if the type of the node is not a kernel + execution. + + +| +[source,c++] +---- +template +void register_dynamic_parameter( + int argIndex, + dynamic_parameter dynamicParam) const; ---- |Associate a dynamic parameter with a kernel parameter of this node via the argument index. Although argument ordering is not defined for lambda captures, @@ -646,8 +763,8 @@ Exceptions: * Throws with error code `invalid` if `argIndex` is not a valid argument index for the kernel. -* Throws with error code `invalid` if `dynamicParam` wraps a memory object that - cannot be tied to the parameter of the node at `argIndex`. +* Throws with error code `invalid` if the type of `dynamicParam` does not + correspond to the type of the node at `argIndex`. * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -811,8 +928,8 @@ otherwise exist if the same executable graph was executed concurrently. 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 it is either empty or associated with a -memory object. By registering the `dynamic_parameter` with a node using +`dynamic_parameter` instance is created it is empty with no associated +kernel arguments. By registering the `dynamic_parameter` with a node using `node::register_dynamic_parameter()`, the dynamic parameter is matched to a node argument, and the `dynamic_parameter` can then be used to update that node argument in future. @@ -1102,14 +1219,50 @@ graph update. |=== |Member function|Description +| +[source,c++] +---- +void update_dynamic_parameter( + void* newValue, + dynamic_parameter dynamicParam); +---- + +| Updates USM parameters in nodes associated with `dynamicParam` to a new value. +This new value will not affect any prior submissions of this graph and will +take affect only for future submissions. See <> for more information about updating node parameters. + +Preconditions: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `newValue` - The new value for the parameter being updated. + +* `dynamicParam` - The object which represents the USM node parameters to be + updated. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable_graph` was not set when the executable graph + was created. + +* Throws synchronously with error code `invalid` if `newValue` is not a pointer + to a valid USM allocation. + | [source,c++] ---- template -void update_dynamic_parameter(T newValue, dynamic_parameter dynamicParam); +void update_dynamic_parameter( + T newValue, + dynamic_parameter dynamicParam); ---- -| Updates parameters in nodes associated with `dynamicParam` to a new value. +| Updates buffer parameters in nodes associated with `dynamicParam` to a new value. This new value will not affect any prior submissions of this graph and will take affect only for future submissions. See <> for more information about updating node parameters. @@ -1123,7 +1276,7 @@ Parameters: * `newValue` - The new value for the parameter being updated. -* `dynamicParam` - The object which represents the node parameters to be +* `dynamicParam` - The object which represents the buffer node parameters to be updated. Exceptions: @@ -1131,16 +1284,42 @@ Exceptions: * Throws synchronously with error code `invalid` if `property::graph::updatable_graph` was not set when the executable graph was created. -* Throws with error code `invalid` if `newValue` is not compatible with `dynamicParam`. - Parameters are considered incompatible if: -** For pointers to USM allocations `newValue` must be a pointer type. +* Throws with error code `invalid` if `newValue` is not a valid buffer. + +| +[source,c++] +---- +template +void update_dynamic_parameter( + T newValue, + dynamic_parameter dynamicParam); +---- + +| Updates scalar parameters in nodes associated with `dynamicParam` to a new value. +This new value will not affect any prior submissions of this graph and will +take affect only for future submissions. See <> for more information about updating node parameters. + +Preconditions: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `newValue` - The new value for the parameter being updated. + +* `dynamicParam` - The object which represents the scalar node parameters to be + updated. -** For buffers `newValue` must have a total memory size that is greater -than or equal to the parameter represented by `dynamicParam`. +Exceptions: -** For scalars `newValue` must have a memory size that is equal to the -parameter represented by `dynamicParam`. +* Throws synchronously with error code `invalid` if + `property::graph::updatable_graph` was not set when the executable graph + was created. +* Throws with error code `invalid` if `newValue` does not have the same size + as specified on creation of `dynamicParam`. | [source,c++] From 83f8852af155d05d5304c9a956560941ff0ba89a Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Fri, 1 Dec 2023 10:09:31 +0000 Subject: [PATCH 13/22] Clarifications Clarify that it's buffer accessors, not a `sycl::buffer` that are used for registration and update of buffer-type dynamic parameters. Other clarifications based on PR feedback. --- .../sycl_ext_oneapi_graph.asciidoc | 25 ++++++++++--------- 1 file changed, 13 insertions(+), 12 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 434b21761dbdb..738f72fa2a4cb 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -397,7 +397,7 @@ class node { public: node() = delete; - /* -- parameter update interface -- */ + /* -- parameter registration interface -- */ void register_dynamic_parameter(void* param, dynamic_parameter dynamicParam) const; @@ -555,14 +555,15 @@ by the user after the node has been added to a graph. They are updated through an executable graph which contains the nodes associated with a dynamic parameter object. -The type of the underlying object a dynamic parameter represents is set at compile -time using a template parameter. This underlying memory object can be a -`sycl::buffer`, a pointer to a USM allocation, or a scalar passed by value. +The type of the underlying object a dynamic parameter represents is set at +compile time using a template parameter. This underlying memory object can be a +`sycl::buffer` accessor, a pointer to a USM allocation, or a scalar passed by +value. Dynamic parameters are registered with nodes, with each registration associating a node argument to the dynamic parameter instance. An update of the dynamic -parameter instance to a new value, will then update all the parameter which have -been registered to the node. +parameter instance to a new value, will then update all the node parameter which +have been registered to the node. Dynamic parameters can optionally have a user-specified string label associated with them. This can provide meaningful names to help track what the parameters @@ -681,8 +682,9 @@ void register_dynamic_parameter( T param, dynamic_parameter dynamicParam) const; ---- -|Associate a buffer dynamic parameter with buffer kernel parameter of this node. -All buffer kernel arguments matching `param` are bound to `dynamicParam`. +|Associate a buffer dynamic parameter with buffer accessor kernel parameter of +this node. All buffer accessor kernel arguments matching `param` are bound to +`dynamicParam`. See <> for more information about updating node parameters. @@ -734,7 +736,6 @@ Exceptions: * Throws with error code `invalid` if the type of the node is not a kernel execution. - | [source,c++] ---- @@ -907,8 +908,8 @@ The aspects of a kernel execution node that can be configured during update are: * 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, which is most useful when the -graph was created from a library, and an individual node update API. +An API for updating the whole graph object, which is most useful when the +graph was recorded, 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 @@ -1285,7 +1286,7 @@ Exceptions: `property::graph::updatable_graph` was not set when the executable graph was created. -* Throws with error code `invalid` if `newValue` is not a valid buffer. +* Throws with error code `invalid` if `newValue` is not a valid buffer accessor. | [source,c++] From 89e8a243290ad961b363a1e9e4c41cf6a2c28463 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Fri, 1 Dec 2023 11:18:24 +0000 Subject: [PATCH 14/22] Change semantics of dyn param retaining values --- .../sycl_ext_oneapi_graph.asciidoc | 248 ++++++++++++++---- 1 file changed, 201 insertions(+), 47 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 738f72fa2a4cb..b30c5b86073ed 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -339,9 +339,9 @@ class assume_buffer_outlives_graph { assume_buffer_outlives_graph() = default; }; -class updatable_graph { +class updatable { public: - updatable_graph() = default; + updatable() = default; }; } // namespace graph @@ -379,7 +379,8 @@ enum class dynamic_parameter_type { template class dynamic_parameter { public: - dynamic_parameter(std::string label = ""); + dynamic_parameter(command_graph graph, + std::string label = ""); const std::string& get_label() const; }; @@ -387,7 +388,8 @@ public: template <> class dynamic_parameter { public: - dynamic_parameter(size_t size, std::string label = ""); + dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); size_t get_size() const; const std::string& get_label() const; @@ -401,9 +403,11 @@ public: void register_dynamic_parameter(void* param, dynamic_parameter dynamicParam) const; + template void register_dynamic_parameter(T param, dynamic_parameter dynamicParam) const; + template void register_dynamic_parameter(T param, dynamic_parameter dynamicParam) const; @@ -561,15 +565,37 @@ compile time using a template parameter. This underlying memory object can be a value. Dynamic parameters are registered with nodes, with each registration associating -a node argument to the dynamic parameter instance. An update of the dynamic -parameter instance to a new value, will then update all the node parameter which -have been registered to the node. +one or more node arguments to the dynamic parameter instance. When registration +is done by matching the node argument value, it must be the original argument value. +It is valid for a node argument to be registered with more than one dynamic +argument instance. + +Attempting to register the dynamic parameter again to the same node argument by +matching against an updated value will not work, as the dynamic parameter does +not retain the `param` value used for registration, only which kernel argument it +corresponds to. For example, + +[source,c++] +---- +// Update a node with a single USM parameter 'PtrA', to 'PtrB'. +dynamic_parameter DynParam +Node.register_dynamic_parameter(PtrA, DynParam); +ExecGraph.update_dynamic_param(PtrB, DynParam); + +// Exception thrown here as no match against original value. +Node.register_dynamic_parameter(PtrB, DynParam); + +// No-op as dynamic parameter already registered with 'PtrA'. +Node.register_dynamic_parameter(PtrA, DynParam); +---- Dynamic parameters can optionally have a user-specified string label associated with them. This can provide meaningful names to help track what the parameters are when updating them. These labels are not used by the SYCL runtime in any way. +The `dynamic_parameter` class provides the {crs}[common reference semantics]. + Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. [cols="2a,a"] |=== @@ -578,7 +604,8 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. | [source,c++] ---- -dynamic_parameter(std::string label = ""); +dynamic_parameter(command_graph graph, + std::string label = ""); ---- |Constructs a dynamic parameter object that can be registered with node arguments. @@ -593,13 +620,15 @@ Preconditions: Parameters: +* `graph` - Graph containing the nodes which will have dynamic parameters. * `label` - An optional string label which can be used to help identify the parameter which has been registered as dynamic. | [source,c++] ---- -dynamic_parameter(size_t size, std::string label = ""); +dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); ---- |Constructs a scalar dynamic parameter object that can be registered with node arguments. @@ -614,6 +643,7 @@ Preconditions: Parameters: +* `graph` - Graph containing the nodes which will have dynamic parameters. * `size` - Size in bytes of the scalar object. * `label` - An optional string label which can be used to help identify the @@ -655,7 +685,9 @@ void register_dynamic_parameter( dynamic_parameter dynamicParam) const; ---- |Associate a USM dynamic parameter with USM kernel parameter of this node. -All USM kernel arguments matching `param` are bound to `dynamicParam`. +All USM kernel arguments matching `param` are bound to `dynamicParam`. If +`param` has already registered with `dynamicParam` for this node, this +operation is a no-op. See <> for more information about updating node parameters. @@ -670,9 +702,10 @@ Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in the node. - * Throws with error code `invalid` if the type of the node is not a kernel execution. +* Throws with error code `invalid` if the node is not a member of the graph set + on creation of `dynamicParam`. | [source,c++] @@ -684,7 +717,8 @@ void register_dynamic_parameter( ---- |Associate a buffer dynamic parameter with buffer accessor kernel parameter of this node. All buffer accessor kernel arguments matching `param` are bound to -`dynamicParam`. +`dynamicParam`. If `param` has already registered with `dynamicParam` for this +node, this operation is a no-op. See <> for more information about updating node parameters. @@ -699,11 +733,11 @@ Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in the node. - * Throws with error code `invalid` if `param` is not a buffer accessor. - * Throws with error code `invalid` if the type of the node is not a kernel execution. +* Throws with error code `invalid` if the node is not a member of the graph set + on creation of `dynamicParam`. | [source,c++] @@ -713,8 +747,10 @@ void register_dynamic_parameter( T param, dynamic_parameter dynamicParam) const; ---- -|Associate a scalar dynamic parameter with a scalar kernel parameter of this node. -All scalar kernel arguments matching `param` are bound to `dynamicParam`. +|Associate a scalar dynamic parameter with a scalar kernel parameter of this +node. All scalar kernel arguments matching `param` are bound to `dynamicParam`. +If `param` has already registered with `dynamicParam` for this node, this +operation is a no-op. See <> for more information about updating node parameters. @@ -729,12 +765,12 @@ Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in the node. - * Throws with error code `invalid` if `param` does not match the specified on creation of `dynamicParam`. - * Throws with error code `invalid` if the type of the node is not a kernel execution. +* Throws with error code `invalid` if the node is not a member of the graph set + on creation of `dynamicParam`. | [source,c++] @@ -747,7 +783,8 @@ void register_dynamic_parameter( |Associate a dynamic parameter with a kernel parameter of this node via the argument index. Although argument ordering is not defined for lambda captures, the ordering is reliable for arguments set using `handler::set_arg()` or -`handler::set_args()`. +`handler::set_args()`. If `argIndex` has already registered with `dynamicParam` +for this node, this operation is a no-op. See <> for more information about updating node parameters. @@ -763,12 +800,12 @@ Exceptions: * Throws with error code `invalid` if `argIndex` is not a valid argument index for the kernel. - * Throws with error code `invalid` if the type of `dynamicParam` does not correspond to the type of the node at `argIndex`. - * Throws with error code `invalid` if the type of the node is not a kernel execution. +* Throws with error code `invalid` if the node is not a member of the graph set + on creation of `dynamicParam`. | [source,c++] @@ -782,14 +819,14 @@ node_type get_type() const; ---- std::vector get_dependencies() const; ---- -|Returns a list of the nodes which this node depends on. +|Returns a list of the predecessor nodes which this node depends on. | [source,c++] ---- std::vector get_dependent_nodes() const; ---- -|Returns a list of the nodes which depend on this node. +|Returns a list of the successor nodes which depend on this node. | [source,c++] @@ -804,7 +841,8 @@ Parameters: * `nodeEvent` - Event returned from a submission to a queue in the recording state. -Returns: Graph node associated to `nodeEvent`. +Returns: Graph node that was created when the command that returned +`nodeEvent` was submitted. Exceptions: @@ -911,17 +949,17 @@ Two methods are provided by the API to the user for performing this update. An API for updating the whole graph object, which is most useful when the graph was recorded, 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 +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. -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. +The `property::graph::updatable` 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 @@ -1165,8 +1203,8 @@ Preconditions: Parameters: * `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 <>. + 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. @@ -1230,8 +1268,11 @@ void update_dynamic_parameter( | Updates USM parameters in nodes associated with `dynamicParam` to a new value. This new value will not affect any prior submissions of this graph and will -take affect only for future submissions. See <> for more information about updating node parameters. +take affect only for future submissions. It is not an error if `newValue` is +set to the current parameter value, i.e. no update occurred. + +See <> for more information +about updating node parameters. Preconditions: @@ -1248,11 +1289,15 @@ Parameters: Exceptions: * Throws synchronously with error code `invalid` if - `property::graph::updatable_graph` was not set when the executable graph - was created. - + `property::graph::updatable` was not set when the executable graph was + created. * Throws synchronously with error code `invalid` if `newValue` is not a pointer to a valid USM allocation. +* Throws synchronously with error code `invalid` if the graph set on creation of + `dynamicParam` is not the same graph as used to create the executable graph + object. +* Throws synchronously with error code `invalid` if `dynamicParam` is empty, + i.e. not associated with any node parameters. | [source,c++] @@ -1265,8 +1310,11 @@ void update_dynamic_parameter( | Updates buffer parameters in nodes associated with `dynamicParam` to a new value. This new value will not affect any prior submissions of this graph and will -take affect only for future submissions. See <> for more information about updating node parameters. +take affect only for future submissions. It is not an error if `newValue` is +set to the current parameter value, i.e. no update occurred. + +See <> for more information +about updating node parameters. Preconditions: @@ -1283,10 +1331,15 @@ Parameters: Exceptions: * Throws synchronously with error code `invalid` if - `property::graph::updatable_graph` was not set when the executable graph - was created. - + `property::graph::updatable` was not set when the executable graph was + created. * Throws with error code `invalid` if `newValue` is not a valid buffer accessor. +* Throws synchronously with error code `invalid` if the graph set on creation of + `dynamicParam` is not the same graph as used to create the executable graph + object. +* Throws synchronously with error code `invalid` if `dynamicParam` is empty, + i.e. not associated with any node parameters. + | [source,c++] @@ -1299,8 +1352,11 @@ void update_dynamic_parameter( | Updates scalar parameters in nodes associated with `dynamicParam` to a new value. This new value will not affect any prior submissions of this graph and will -take affect only for future submissions. See <> for more information about updating node parameters. +take affect only for future submissions. It is not an error if `newValue` is +set to the current parameter value, i.e. no update occurred. + +See <> for more information +about updating node parameters. Preconditions: @@ -1317,10 +1373,15 @@ Parameters: Exceptions: * Throws synchronously with error code `invalid` if - `property::graph::updatable_graph` was not set when the executable graph - was created. + `property::graph::updatable` was not set when the executable graph was + created. * Throws with error code `invalid` if `newValue` does not have the same size as specified on creation of `dynamicParam`. +* Throws synchronously with error code `invalid` if the graph set on creation of + `dynamicParam` is not the same graph as used to create the executable graph + object. +* Throws synchronously with error code `invalid` if `dynamicParam` is empty, + i.e. not associated with any node parameters. | [source,c++] @@ -2095,6 +2156,99 @@ submitted in its entirety for execution via ... ---- +=== Dynamic Parameter Update + +Example showing a graph with two kernel nodes, the first created using a kernel +bundle with `handler::set_args()`, 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); + +// Initialize USM allocations + +// 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; +dynamic_parameter dynParamOutput; +dynamic_parameter dynParamScalar; + +// 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, myScalar); + cgh.parallel_for(range {n}, builtinKernel); +}); + +// Register nodeA dynamic parameters +nodeA.register_dynamic_parameter(0, dynParamInput); // Argument index 0 is ptrX +nodeA.register_dynamic_parameter(1, dynParamScalar); // Argument index 1 is myScalar + +// Second node uses ptrX as an inputs, and ptrY as input/output/ +node nodeB = myGraph.add[&](handler& cgh) { + cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) { + const size_t i = it[0]; + ptrY[i] += ptrX[i]; + }); + }, + {sycl_ext::property::node::depends_on(nodeA)}); + +// Register nodeB dynamic parameters +nodeB.register_dynamic_parameter(ptrX, dynParamInput); +nodeB.register_dynamic_parameter(ptrY, dynParamOutput); + +// 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 & B to ptrZ +execGraph.update_dynamic_parameter(ptrZ, dynParamInput); + +// Change myScalar argument to node A to newScalar +int newScalar = 12; +execGraph.update_dynamic_parameter(newScalar, dynParamScalar); + +// Change ptrY argument to node B to ptrQ +execGraph.update_dynamic_parameter(ptrQ, dynParamOutput); + +// Execute graph again, without needing to wait on completion of first submission. +myQueue.ext_oneapi_graph(exec); +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 From 9c6265302bbbec2dcfe3c22bdd9800613ebde081 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Fri, 8 Dec 2023 10:09:08 +0000 Subject: [PATCH 15/22] Make registration and update methods on the dynamic_parameter --- .../sycl_ext_oneapi_graph.asciidoc | 625 ++++++++++-------- 1 file changed, 350 insertions(+), 275 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index b30c5b86073ed..3d97472b09a50 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -372,51 +372,77 @@ struct graphs_support; enum class dynamic_parameter_type { usm, - buffer, + accessor, scalar }; template -class dynamic_parameter { +class dynamic_parameter; + +template <> +class dynamic_parameter { public: dynamic_parameter(command_graph graph, std::string label = ""); + void register(void* param, node node); + void register(void* param); + + void register(int argIndex, node node); + + void update(void* newValue, + command_graph execGraph); + const std::string& get_label() const; }; template <> -class dynamic_parameter { +class dynamic_parameter { public: dynamic_parameter(command_graph graph, - size_t size, std::string label = ""); + std::string label = ""); + + template + void register(T param, node node); + + template + void register(T param); + + void register(int argIndex, node node); + + template + void update(T newValue, + command_graph execGraph); - size_t get_size() const; const std::string& get_label() const; }; -class node { -public: - node() = delete; - /* -- parameter registration interface -- */ +template <> +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); - void register_dynamic_parameter(void* param, - dynamic_parameter dynamicParam) const; + template + void register(T param, node node); template - void register_dynamic_parameter(T param, - dynamic_parameter dynamicParam) const; + void register(T param); + + void register(int argIndex, node node); template - void register_dynamic_parameter(T param, - dynamic_parameter dynamicParam) const; + void update(T newValue, + command_graph execGraph); - template - void register_dynamic_parameter(int argIndex, - dynamic_parameter dynamicParam) const; + size_t get_size() const; + const std::string& get_label() const; +}; - /* -- query interface -- */ +class node { +public: + node() = delete; node_type get_type() const; @@ -471,19 +497,6 @@ class command_graph { public: command_graph() = delete; - /* -- update interface -- */ - - void - update_dynamic_parameter(void* newValue, - dynamic_parameter dynamicParam); - - template - void update_dynamic_parameter(T newValue, - dynamic_parameter dynamicParam); - - template - void update_dynamic_parameter(T newValue, - dynamic_parameter dynamicParam); template void update_nd_range(const std::vector& nodes, @@ -560,33 +573,32 @@ an executable graph which contains the nodes associated with a dynamic parameter object. The type of the underlying object a dynamic parameter represents is set at -compile time using a template parameter. This underlying memory object can be a -`sycl::buffer` accessor, a pointer to a USM allocation, or a scalar passed by -value. +compile time using a template parameter. This underlying memory object can be an +accessor, a pointer to a USM allocation, or a scalar passed by value. -Dynamic parameters are registered with nodes, with each registration associating -one or more node arguments to the dynamic parameter instance. When registration -is done by matching the node argument value, it must be the original argument value. -It is valid for a node argument to be registered with more than one dynamic -argument instance. +Dynamic parameters are registered with nodes in a modifiable graph, with each +registration associating one or more node arguments to the dynamic parameter +instance. It is valid for a node argument to be registered with more than one +dynamic argument instance. -Attempting to register the dynamic parameter again to the same node argument by -matching against an updated value will not work, as the dynamic parameter does -not retain the `param` value used for registration, only which kernel argument it -corresponds to. For example, +When registration is done by matching the node argument value, it must +be the original argument value. Attempting to register the dynamic parameter +again to the same node argument by matching against an updated value will not +work. This is because registration operates on a modifiable graph, while update +operates on an executable graph. For example, [source,c++] ---- // Update a node with a single USM parameter 'PtrA', to 'PtrB'. -dynamic_parameter DynParam -Node.register_dynamic_parameter(PtrA, DynParam); -ExecGraph.update_dynamic_param(PtrB, DynParam); +dynamic_parameter DynParam(Graph); +DynParam.register(PtrA, Node); +DynParam.update(PtrB, ExecGraph); // Exception thrown here as no match against original value. -Node.register_dynamic_parameter(PtrB, DynParam); +DynParam.register(PtrB, Node); // No-op as dynamic parameter already registered with 'PtrA'. -Node.register_dynamic_parameter(PtrA, DynParam); +DynParam.register(PtrA, Node); ---- Dynamic parameters can optionally have a user-specified string label associated @@ -594,6 +606,9 @@ with them. This can provide meaningful names to help track what the parameters are when updating them. These labels are not used by the SYCL runtime in any way. +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. @@ -610,13 +625,10 @@ dynamic_parameter(command_graph graph, |Constructs a dynamic parameter object that can be registered with node arguments. -See <> for more information -about updating node parameters. - Preconditions: * This constructor is only available when the templated `dynamic_parameter_type` - is `dynamic_parameter_type::usm` or `dynamic_parameter_type::buffer`. + is `dynamic_parameter_type::usm` or `dynamic_parameter_type::accessor`. Parameters: @@ -633,9 +645,6 @@ dynamic_parameter(command_graph graph, |Constructs a scalar dynamic parameter object that can be registered with node arguments. -See <> for more information -about updating node parameters. - Preconditions: * This constructor is only available when the templated `dynamic_parameter_type` @@ -645,7 +654,6 @@ Parameters: * `graph` - Graph containing the nodes which will have dynamic parameters. * `size` - Size in bytes of the scalar object. - * `label` - An optional string label which can be used to help identify the parameter which has been registered as dynamic. @@ -670,9 +678,8 @@ Preconditions: |=== -==== Node Member Functions - -Table {counter: tableNumber}. Member functions of the `node` class. +Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class +for registration. [cols="2a,a"] |=== |Member Function|Description @@ -680,132 +687,325 @@ Table {counter: tableNumber}. Member functions of the `node` class. | [source,c++] ---- -void register_dynamic_parameter( - void* param, - dynamic_parameter dynamicParam) const; +void register(void* param, node node); ---- -|Associate a USM dynamic parameter with USM kernel parameter of this node. -All USM kernel arguments matching `param` are bound to `dynamicParam`. If -`param` has already registered with `dynamicParam` for this node, this +|Associate the dynamic parameter with USM kernel parameter for `node`. All USM +kernel arguments matching `param` are bound to the dynamic parameter. If `param` +has already been registered with the dynamic-parameter for `node`, this operation is a no-op. -See <> for more information -about updating node parameters. +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::usm`. Parameters: * `param` - USM pointer to used as a node argument. -* `dynamicParam` - Dynamic parameter which is to be associated with the kernel - arguments matching `param`. +* `node` - Kernel command node to look for a argument matching `param`. Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in - the node. -* Throws with error code `invalid` if the type of the node is not a kernel + `node`. +* Throws with error code `invalid` if the type of `node` is not a kernel execution. -* Throws with error code `invalid` if the node is not a member of the graph set - on creation of `dynamicParam`. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic-parameter was created with. + +| +[source,c++] +---- +void register(void* param); +---- +|Associate the dynamic parameter with all the USM kernel parameters which match +`param` in the graph used on creation of the dynamic parameter. For every kernel +node in the graph, the dynamic parameter binds to any kernel arguments matching +`param`. It is not an error if the dynamic parameter fails to bind to a single +node, only if no nodes in the graph have an argument matching `param`. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::usm`. + +Parameters: + +* `param` - USM pointer to used as a node argument. + +Exceptions: + +* Throws with error code `invalid` if `param` no node arguments are matched + in the graph. | [source,c++] ---- template -void register_dynamic_parameter( - T param, - dynamic_parameter dynamicParam) const; +void register(T param, node node); ---- -|Associate a buffer dynamic parameter with buffer accessor kernel parameter of -this node. All buffer accessor kernel arguments matching `param` are bound to -`dynamicParam`. If `param` has already registered with `dynamicParam` for this -node, this operation is a no-op. +|Associate an accessor dynamic parameter with an accessor kernel parameter of +`node`. All accessor kernel arguments matching `param` are bound to the dynamic +parameter. If `param` has already been registered the dynamic parameter for +`node`, this operation is a no-op. -See <> for more information -about updating node parameters. +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::accessor`. Parameters: -* `param` - Buffer accessor used as a node argument. -* `dynamicParam` - Dynamic parameter which is to be associated with the kernel - arguments matching `param`. +* `param` - Accessor used as a node argument. +* `node` - Kernel command node to look for a argument matching `param`. Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in - the node. + `node`. * Throws with error code `invalid` if `param` is not a buffer accessor. -* Throws with error code `invalid` if the type of the node is not a kernel +* Throws with error code `invalid` if the type of `node` is not a kernel execution. -* Throws with error code `invalid` if the node is not a member of the graph set - on creation of `dynamicParam`. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic-parameter was created with. | [source,c++] ---- template -void register_dynamic_parameter( - T param, - dynamic_parameter dynamicParam) const; +void register(T param); ---- -|Associate a scalar dynamic parameter with a scalar kernel parameter of this -node. All scalar kernel arguments matching `param` are bound to `dynamicParam`. -If `param` has already registered with `dynamicParam` for this node, this -operation is a no-op. +|Associate the dynamic parameter with all the accessor kernel parameters which +match `param` in the graph used on creation of the dynamic parameter. For +every kernel node in the graph, the dynamic parameter binds to any kernel +arguments matching `param`. It is not an error if the dynamic parameter fails +to bind to a single node, only if no nodes in the graph have an argument +matching `param`. -See <> for more information -about updating node parameters. +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::accessor`. + +Parameters: + +* `param` - Scalar used as a node argument. + +Exceptions: + +* Throws with error code `invalid` if `param` no node arguments are matched + in the graph. + +| +[source,c++] +---- +template +void register(T param, node node); +---- +|Associate a scalar dynamic parameter with a scalar kernel parameter of `node`. +All scalar kernel arguments matching `param` are bound to the dynamic parameter. +If `param` has already been registered with the dynamic parameter for this node, +then this operation is a no-op. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::scalar`. Parameters: * `param` - Scalar used as a node argument. -* `dynamicParam` - Dynamic parameter which is to be associated with the kernel - arguments matching `param`. +* `node` - Kernel command node to look for a argument matching `param`. Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in - the node. -* Throws with error code `invalid` if `param` does not match the specified on - creation of `dynamicParam`. -* Throws with error code `invalid` if the type of the node is not a kernel + `node`. +* Throws with error code `invalid` if the size of `param` does not match the + size specified on creation of the dynamic parameter. +* Throws with error code `invalid` if the type of `node` is not a kernel execution. -* Throws with error code `invalid` if the node is not a member of the graph set - on creation of `dynamicParam`. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic-parameter was created with. | [source,c++] ---- -template -void register_dynamic_parameter( - int argIndex, - dynamic_parameter dynamicParam) const; +template +void register(T param); ---- -|Associate a dynamic parameter with a kernel parameter of this node via the +|Associate the dynamic parameter with all the scalar kernel parameters which +match `param` in the graph used on creation of the dynamic parameter. For +every kernel node in the graph, the dynamic parameter binds to any kernel +arguments matching `param`. It is not an error if the dynamic parameter fails +to bind to a single node, only if no nodes in the graph have an argument +matching `param`. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::scalar`. + +Parameters: + +* `param` - Scalar used as a node argument. + +Exceptions: + +* Throws with error code `invalid` if `param` no node arguments are matched + in the graph. + +* Throws with error code `invalid` if the size of `param` does not match the + size specified on creation of the dynamic parameter. + +| +[source,c++] +---- +template +void register_dynamic_parameter(int argIndex, node node); +---- +|Associate the dynamic parameter with a kernel parameter of `node` via the argument index. Although argument ordering is not defined for lambda captures, the ordering is reliable for arguments set using `handler::set_arg()` or -`handler::set_args()`. If `argIndex` has already registered with `dynamicParam` -for this node, this operation is a no-op. - -See <> for more information -about updating node parameters. +`handler::set_args()`. If `argIndex` is already registered with the dynamic +parameter for `node`, this operation is a no-op. Parameters: -* `argIndex` - Kernel argument index starting from zero to associate - `dynamicParam` with. -* `dynamicParam` - Dynamic parameter which is associated with the kernel - argument at index `argIndex`. +* `argIndex` - Kernel argument index starting from zero to associate with + the dynamic parameter. +* `node` - Kernel command node to associate with argument at index `argIndex`. Exceptions: * Throws with error code `invalid` if `argIndex` is not a valid argument index - for the kernel. -* Throws with error code `invalid` if the type of `dynamicParam` does not - correspond to the type of the node at `argIndex`. -* Throws with error code `invalid` if the type of the node is not a kernel + for the `node` kernel. +* Throws with error code `invalid` if the type of the dynamic parameter does not + correspond to the type of the `node` argument at `argIndex`. +* Throws with error code `invalid` if the type of `node` is not a kernel execution. -* Throws with error code `invalid` if the node is not a member of the graph set - on creation of `dynamicParam`. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic-parameter was created with. + +|=== + +Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class +for update. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +void update(void* newValue, + command_graph execGraph); +---- + +|Updates USM parameters in `execGraph` nodes associated with the dynamic +parameter to `newValue`. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value in nodes of +`execGraph`, i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::usm`. + +Parameters: + +* `newValue` - The new value for the parameters being updated. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws synchronously with error code `invalid` if `newValue` is not a pointer + to a valid USM allocation. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +| +[source,c++] +---- +template +void update(T newValue, + command_graph execGraph); +---- + +|Updates accessor parameters in `execGraph` nodes associated with the dynamic +parameter to `newValue`. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value in nodes of `execGraph`, +i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::accessor`. + +Parameters: + +* `newValue` - The new value for the parameter being updated. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws with error code `invalid` if `newValue` is not a valid accessor. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +| +[source,c++] +---- +template +void update(T newValue, + command_graph execGraph); +---- + +|Updates scalar parameters in `execGraph` nodes associated with the dynamic +parameter to `newValue`. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value of nodes in +`execGraph`, i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::scalar`. + +Parameters: + +* `newValue` - The new value for the parameter being updated. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws with error code `invalid` if `newValue` does not have the same size + as specified on creation of the dynamic parameter. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +|=== + +==== Node Member Functions + +Table {counter: tableNumber}. Member functions of the `node` class. +[cols="2a,a"] +|=== +|Member Function|Description | [source,c++] @@ -967,20 +1167,20 @@ exist if the same executable graph was executed concurrently. 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 it is empty with no associated -kernel arguments. By registering the `dynamic_parameter` with a node using -`node::register_dynamic_parameter()`, the dynamic parameter is matched to a node -argument, and the `dynamic_parameter` can then be used to update that node -argument in future. +`dynamic_parameter` instance is created with a modifiable state graph it is +empty, with no associated kernel arguments. By registering the +`dynamic_parameter` with a node using `dynamic_parameter::register()`, the +dynamic parameter is matched to a node argument, and the `dynamic_parameter` +can then be used to update that node argument in future. + +If a user doesn't have the individual `node` handles to register parameters +with, they can use the `dynamic_parameter::register()` API that doesn't take +a `node`. This entry-point iterates through the modifiable graph nodes, and +tries to register the parameter against every kernel node in the graph. The update itself is then performed using a `dynamic_parameter` instance by -calling `command_graph::update_dynamic_parameter()` to -update all the parameters of nodes in an executable graph to which the -`dynamic_parameter` is registered. If a user tries to register a dynamic -parameter with a node, after that dynamic parameter has been updated with a new -value, then it is the new value that must match one of the nodes arguments for -the registration to be successful, rather than the value originally used to -construct the dynamic parameter. +calling `dynamic_parameter::update()` to update all the parameters of nodes in +an executable graph to which the `dynamic_parameter` is registered. The other node configuration that can be updated is the ND-Range, this can be set through `command_graph::update_nd_range()` but @@ -998,7 +1198,7 @@ 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 behaviour. 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 `update_dynamic_parameter()` call +use the buffer as a parameter. Then a single `dynamic_parameter::update()` call will maintain the graphs data dependencies. ==== Graph Properties [[graph-properties]] @@ -1258,131 +1458,6 @@ graph update. |=== |Member function|Description -| -[source,c++] ----- -void update_dynamic_parameter( - void* newValue, - dynamic_parameter dynamicParam); ----- - -| Updates USM parameters in nodes associated with `dynamicParam` to a new value. -This new value will not affect any prior submissions of this graph and will -take affect only for future submissions. It is not an error if `newValue` is -set to the current parameter value, i.e. no update occurred. - -See <> for more information -about updating node parameters. - -Preconditions: - -* This member function is only available when the `command_graph` state is - `graph_state::executable`. - -Parameters: - -* `newValue` - The new value for the parameter being updated. - -* `dynamicParam` - The object which represents the USM node parameters to be - updated. - -Exceptions: - -* Throws synchronously with error code `invalid` if - `property::graph::updatable` was not set when the executable graph was - created. -* Throws synchronously with error code `invalid` if `newValue` is not a pointer - to a valid USM allocation. -* Throws synchronously with error code `invalid` if the graph set on creation of - `dynamicParam` is not the same graph as used to create the executable graph - object. -* Throws synchronously with error code `invalid` if `dynamicParam` is empty, - i.e. not associated with any node parameters. - -| -[source,c++] ----- -template -void update_dynamic_parameter( - T newValue, - dynamic_parameter dynamicParam); ----- - -| Updates buffer parameters in nodes associated with `dynamicParam` to a new value. -This new value will not affect any prior submissions of this graph and will -take affect only for future submissions. It is not an error if `newValue` is -set to the current parameter value, i.e. no update occurred. - -See <> for more information -about updating node parameters. - -Preconditions: - -* This member function is only available when the `command_graph` state is - `graph_state::executable`. - -Parameters: - -* `newValue` - The new value for the parameter being updated. - -* `dynamicParam` - The object which represents the buffer node parameters to 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 `newValue` is not a valid buffer accessor. -* Throws synchronously with error code `invalid` if the graph set on creation of - `dynamicParam` is not the same graph as used to create the executable graph - object. -* Throws synchronously with error code `invalid` if `dynamicParam` is empty, - i.e. not associated with any node parameters. - - -| -[source,c++] ----- -template -void update_dynamic_parameter( - T newValue, - dynamic_parameter dynamicParam); ----- - -| Updates scalar parameters in nodes associated with `dynamicParam` to a new value. -This new value will not affect any prior submissions of this graph and will -take affect only for future submissions. It is not an error if `newValue` is -set to the current parameter value, i.e. no update occurred. - -See <> for more information -about updating node parameters. - -Preconditions: - -* This member function is only available when the `command_graph` state is - `graph_state::executable`. - -Parameters: - -* `newValue` - The new value for the parameter being updated. - -* `dynamicParam` - The object which represents the scalar node parameters to 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 `newValue` does not have the same size - as specified on creation of `dynamicParam`. -* Throws synchronously with error code `invalid` if the graph set on creation of - `dynamicParam` is not the same graph as used to create the executable graph - object. -* Throws synchronously with error code `invalid` if `dynamicParam` is empty, - i.e. not associated with any node parameters. - | [source,c++] ---- @@ -2193,9 +2268,9 @@ kernel builtinKernel = myBundle.get_kernel(builtinKernelIds[0]); sycl_ext::command_graph myGraph(myContext, myDevice); // Create graph dynamic parameters -dynamic_parameter dynParamInput; -dynamic_parameter dynParamOutput; -dynamic_parameter dynParamScalar; +dynamic_parameter dynParamInput(myGraph); +dynamic_parameter dynParamOutput(myGraph); +dynamic_parameter dynParamScalar(myGraph); // First node uses ptrX as an input & output parameter, with operand // mySclar as another argument. @@ -2206,8 +2281,8 @@ node nodeA = myGraph.add[&](handler& cgh) { }); // Register nodeA dynamic parameters -nodeA.register_dynamic_parameter(0, dynParamInput); // Argument index 0 is ptrX -nodeA.register_dynamic_parameter(1, dynParamScalar); // Argument index 1 is myScalar +dynParamInput.register(0, nodeA); // Argument index 0 is ptrX +dynParamInput.register(1, nodeA); // Argument index 1 is myScalar // Second node uses ptrX as an inputs, and ptrY as input/output/ node nodeB = myGraph.add[&](handler& cgh) { @@ -2219,8 +2294,8 @@ node nodeB = myGraph.add[&](handler& cgh) { {sycl_ext::property::node::depends_on(nodeA)}); // Register nodeB dynamic parameters -nodeB.register_dynamic_parameter(ptrX, dynParamInput); -nodeB.register_dynamic_parameter(ptrY, dynParamOutput); +dynParamInput.register(ptrX, nodeB); +dynParamInput.register(ptrY, nodeB); // Create an executable graph with the updatable property. auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); @@ -2229,14 +2304,14 @@ auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); myQueue.ext_oneapi_graph(execGraph); // Change ptrX argument to node A & B to ptrZ -execGraph.update_dynamic_parameter(ptrZ, dynParamInput); +dynParamInput.update(ptrZ, execGraph); // Change myScalar argument to node A to newScalar int newScalar = 12; -execGraph.update_dynamic_parameter(newScalar, dynParamScalar); +dynParamScalar.update(newScalar, execGraph); // Change ptrY argument to node B to ptrQ -execGraph.update_dynamic_parameter(ptrQ, dynParamOutput); +dynParamOutput.update(ptrQ, execGraph); // Execute graph again, without needing to wait on completion of first submission. myQueue.ext_oneapi_graph(exec); From 5d4c8a44c9a6a42c0b02358a52eb99acf327a258 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 11 Dec 2023 09:51:53 +0000 Subject: [PATCH 16/22] Minor cleanup --- .../sycl_ext_oneapi_graph.asciidoc | 40 +++++++++---------- 1 file changed, 20 insertions(+), 20 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 3d97472b09a50..105933d10006d 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -244,7 +244,7 @@ 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. |=== @@ -323,6 +323,7 @@ enum class node_type { prefetch, memadvise, ext_oneapi_barrier, + host_task, }; namespace property { @@ -497,7 +498,6 @@ class command_graph { public: command_graph() = delete; - template void update_nd_range(const std::vector& nodes, nd_range executionRange); @@ -568,8 +568,8 @@ The `node` class provides the {crs}[common reference semantics]. ==== 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. They are updated through -an executable graph which contains the nodes associated with a dynamic parameter +by the user after the node has been added to a graph. They are updated on an +executable graph which contains the nodes associated with the dynamic parameter object. The type of the underlying object a dynamic parameter represents is set at @@ -691,7 +691,7 @@ void register(void* param, node node); ---- |Associate the dynamic parameter with USM kernel parameter for `node`. All USM kernel arguments matching `param` are bound to the dynamic parameter. If `param` -has already been registered with the dynamic-parameter for `node`, this +has already been registered with the dynamic parameter for `node`, this operation is a no-op. Preconditions: @@ -701,8 +701,8 @@ Preconditions: Parameters: -* `param` - USM pointer to used as a node argument. -* `node` - Kernel command node to look for a argument matching `param`. +* `param` - USM pointer used as a node argument. +* `node` - Kernel command to look for an argument matching `param`. Exceptions: @@ -711,7 +711,7 @@ Exceptions: * Throws with error code `invalid` if the type of `node` is not a kernel execution. * Throws with error code `invalid` if `node` is not a member of the graph that - this dynamic-parameter was created with. + this dynamic parameter was created with. | [source,c++] @@ -731,7 +731,7 @@ Preconditions: Parameters: -* `param` - USM pointer to used as a node argument. +* `param` - USM pointer used as a node argument. Exceptions: @@ -757,17 +757,17 @@ Preconditions: Parameters: * `param` - Accessor used as a node argument. -* `node` - Kernel command node to look for a argument matching `param`. +* `node` - Kernel command to look for an argument matching `param`. Exceptions: * Throws with error code `invalid` if `param` does not match any arguments in `node`. -* Throws with error code `invalid` if `param` is not a buffer accessor. +* Throws with error code `invalid` if `param` is not an accessor. * Throws with error code `invalid` if the type of `node` is not a kernel execution. * Throws with error code `invalid` if `node` is not a member of the graph that - this dynamic-parameter was created with. + this dynamic parameter was created with. | [source,c++] @@ -789,7 +789,7 @@ Preconditions: Parameters: -* `param` - Scalar used as a node argument. +* `param` - Accessor used as a node argument. Exceptions: @@ -815,7 +815,7 @@ Preconditions: Parameters: * `param` - Scalar used as a node argument. -* `node` - Kernel command node to look for a argument matching `param`. +* `node` - Kernel command to look for an argument matching `param`. Exceptions: @@ -826,7 +826,7 @@ Exceptions: * Throws with error code `invalid` if the type of `node` is not a kernel execution. * Throws with error code `invalid` if `node` is not a member of the graph that - this dynamic-parameter was created with. + this dynamic parameter was created with. | [source,c++] @@ -874,7 +874,7 @@ Parameters: * `argIndex` - Kernel argument index starting from zero to associate with the dynamic parameter. -* `node` - Kernel command node to associate with argument at index `argIndex`. +* `node` - Kernel command to associate with argument at index `argIndex`. Exceptions: @@ -885,7 +885,7 @@ Exceptions: * Throws with error code `invalid` if the type of `node` is not a kernel execution. * Throws with error code `invalid` if `node` is not a member of the graph that - this dynamic-parameter was created with. + this dynamic parameter was created with. |=== @@ -915,7 +915,7 @@ Preconditions: Parameters: -* `newValue` - The new value for the parameters being updated. +* `newValue` - Value to the parameters being updated to. * `execGraph` - The executable graph to update the node parameters for. Exceptions: @@ -950,7 +950,7 @@ Preconditions: Parameters: -* `newValue` - The new value for the parameter being updated. +* `newValue` - Value to the parameters being updated to. * `execGraph` - The executable graph to update the node parameters for. Exceptions: @@ -984,7 +984,7 @@ Preconditions: Parameters: -* `newValue` - The new value for the parameter being updated. +* `newValue` - Value to the parameters being updated to. * `execGraph` - The executable graph to update the node parameters for. Exceptions: From 35f81f0084b76f27fe331bcb75bf0180d8570fa4 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 11 Dec 2023 13:54:07 +0000 Subject: [PATCH 17/22] updatable must be set for nd-range update --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 17 ++++++++++++++--- 1 file changed, 14 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 105933d10006d..c176288c72f8e 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -1484,13 +1484,16 @@ Parameters: 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. * Throws with error code `nd_range` if `executionRange` is an invalid ND-Range for any node in `nodes`. |=== -Table {counter: tableNumber}. Member functions of the `command_graph` class for queue recording. +Table {counter: tableNumber}. Member functions of the `command_graph` class for +queue recording. [cols="2a,a"] |=== |Member function|Description @@ -2528,6 +2531,14 @@ runtime. == Issues +=== Update More Command Types + +Support updating arguments to types of nodes other that kernel execution +commands. + +**RESOLVED** Should be added be for at least memory copy nodes, however +fully scope of support needs designed and implemented. + === Multi Device Graph Allow an executable graph to contain nodes targeting different devices. @@ -2592,8 +2603,8 @@ if used in application code. `sycl::ext::intel::property::queue::no_immediate_command_list` should be set on construction to any queues an executable graph is submitted to. -. Synchronization between multiple executions of the same command-buffer - must be handled in the host for level-zero backend, which may involve +. Synchronization between multiple executions of the same command-buffer + must be handled in the host for level-zero backend, which may involve extra latency for subsequent submissions. == Revision History From 47ce7bf57ea45bc14236b916ffe7920ccfd5e058 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Tue, 12 Dec 2023 13:21:59 +0000 Subject: [PATCH 18/22] Add raw dynamic parameter type --- .../sycl_ext_oneapi_graph.asciidoc | 139 ++++++++++++++++-- 1 file changed, 127 insertions(+), 12 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index c176288c72f8e..a99bb3b286c09 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -374,7 +374,8 @@ struct graphs_support; enum class dynamic_parameter_type { usm, accessor, - scalar + scalar, + raw }; template @@ -387,6 +388,7 @@ public: std::string label = ""); void register(void* param, node node); + void register(void* param); void register(int argIndex, node node); @@ -441,6 +443,25 @@ public: const std::string& get_label() const; }; +template <> +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); + + void register(void* param, node node); + + void register(void* param); + + void register(int argIndex, node node); + + void update(void* newValue, + command_graph execGraph); + + size_t get_size() const; + const std::string& get_label() const; +}; + class node { public: node() = delete; @@ -574,7 +595,10 @@ object. The type of the underlying object a dynamic parameter represents is set at compile time using a template parameter. This underlying memory object can be an -accessor, a pointer to a USM allocation, or a scalar passed by value. +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 @@ -642,18 +666,18 @@ parameter which has been registered as dynamic. dynamic_parameter(command_graph graph, size_t size, std::string label = ""); ---- -|Constructs a scalar dynamic parameter object that can be registered with node -arguments. +|Constructs a dynamic parameter object representing defined number of bytes +that can be registered with node arguments. Preconditions: * This constructor is only available when the templated `dynamic_parameter_type` - is `dynamic_parameter_type::scalar`. + is `dynamic_parameter_type::scalar` or `dynamic_parameter_type::raw`. Parameters: * `graph` - Graph containing the nodes which will have dynamic parameters. -* `size` - Size in bytes of the scalar object. +* `size` - Size in bytes of the parameter. * `label` - An optional string label which can be used to help identify the parameter which has been registered as dynamic. @@ -674,7 +698,7 @@ size_t get_size() const; Preconditions: * This constructor is only available when the templated `dynamic_parameter_type` - is `dynamic_parameter_type::scalar`. + is `dynamic_parameter_type::scalar` or `dynamic_parameter_type::raw`. |=== @@ -735,7 +759,7 @@ Parameters: Exceptions: -* Throws with error code `invalid` if `param` no node arguments are matched +* Throws with error code `invalid` if no node arguments are matched to `param` in the graph. | @@ -793,7 +817,7 @@ Parameters: Exceptions: -* Throws with error code `invalid` if `param` no node arguments are matched +* Throws with error code `invalid` if no node arguments are matched to `param` in the graph. | @@ -852,7 +876,7 @@ Parameters: Exceptions: -* Throws with error code `invalid` if `param` no node arguments are matched +* Throws with error code `invalid` if no node arguments are matched to `param` in the graph. * Throws with error code `invalid` if the size of `param` does not match the @@ -861,8 +885,65 @@ Exceptions: | [source,c++] ---- -template -void register_dynamic_parameter(int argIndex, node node); +void register(void* param, node node); +---- +|Associate the dynamic parameter with kernel parameter for `node`. All +kernel arguments matching the bytes pointed to by `param` are bound to the +dynamic parameter. If `param` has already been registered with the dynamic +parameter for `node`, this operation is a no-op. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::raw`. + +Parameters: + +* `param` - Byte array of size set on the creation of the dynamic parameter. +* `node` - Kernel command to look for an argument matching `param`. + +Exceptions: + +* Throws with error code `invalid` if `param` is a null pointer. +* Throws with error code `invalid` if `param` does not match any arguments in + `node`. +* Throws with error code `invalid` if the type of `node` is not a kernel + execution. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic parameter was created with. + +| +[source,c++] +---- +void register(void* param); +---- +|Associate the dynamic parameter with all the kernel parameters which match +the bytes pointed to by `param` in the graph used on creation of the dynamic +parameter. For every kernel node in the graph, the dynamic parameter binds to +any kernel arguments matching the bytes of `param`. It is not an error if the +dynamic parameter fails to bind to a single node, only if no nodes in the graph +have an argument matching `param`. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::raw`. + +Parameters: + +* `param` - Byte array of size set on the creation of the dynamic parameter. + +Exceptions: + +* Throws with error code `invalid` if `param` is a null pointer. +* Throws with error code `invalid` if no node arguments are matched to `param` + in the graph. + +| +[source,c++] +---- +template +void register(int argIndex, node node); ---- |Associate the dynamic parameter with a kernel parameter of `node` via the argument index. Although argument ordering is not defined for lambda captures, @@ -998,6 +1079,40 @@ Exceptions: * Throws synchronously with error code `invalid` if the dynamic parameter is not associated with any nodes in `execGraph`. +| +[source,c++] +---- +void update(void* newValue, + command_graph execGraph); +---- + +|Updates raw byte parameters in `execGraph` nodes associated with the dynamic +parameter to a new value represented by array `newValue` of the size set when +the dynamic parameter was created. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value in nodes of +`execGraph`, i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::raw`. + +Parameters: + +* `newValue` - Byte array containing the value to update the parameters to. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if `newValue` is a null pointer. +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + |=== ==== Node Member Functions From c324d0ac4447be39d81d23ad85e17c404d7e5190 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 13 Dec 2023 10:20:54 +0000 Subject: [PATCH 19/22] Expand on multiple registration --- .../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 a99bb3b286c09..be9d096772a0d 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -603,7 +603,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. It is valid for a node argument to be registered with more than one -dynamic argument instance. +dynamic parameter instance, the semantics of this are that the last update to +any of the dynamic parameters on an executable graph is the new node argument +on the next submissions of that executable graph. When registration is done by matching the node argument value, it must be the original argument value. Attempting to register the dynamic parameter From 360fa1b021e862c3a681ad8ca94bd152cf86994c Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Fri, 15 Dec 2023 09:59:06 +0000 Subject: [PATCH 20/22] Remove scalar/raw registration by matching --- .../sycl_ext_oneapi_graph.asciidoc | 144 ++---------------- 1 file changed, 12 insertions(+), 132 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index be9d096772a0d..932988e675aef 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -427,12 +427,6 @@ public: dynamic_parameter(command_graph graph, size_t size, std::string label = ""); - template - void register(T param, node node); - - template - void register(T param); - void register(int argIndex, node node); template @@ -449,10 +443,6 @@ public: dynamic_parameter(command_graph graph, size_t size, std::string label = ""); - void register(void* param, node node); - - void register(void* param); - void register(int argIndex, node node); void update(void* newValue, @@ -607,6 +597,14 @@ dynamic parameter instance, the semantics of this are that the last update to any of the dynamic parameters on an executable graph is the new node argument on the next submissions of that executable graph. +Registration is done either using the argument index (known when argument +originally set using `set_arg()`/`set_args()`), or by matching the argument +value. Matching for USM typed dynamic parameters is done by direct comparison +of the pointer values. Matching for accessors is done by verifying the +underlying managed data is the same. Raw and scalar typed dynamic parameters +can't be registered by matching however, as the chances of unintentional +collisions in the comparison are too high. + When registration is done by matching the node argument value, it must be the original argument value. Attempting to register the dynamic parameter again to the same node argument by matching against an updated value will not @@ -819,125 +817,6 @@ Parameters: Exceptions: -* Throws with error code `invalid` if no node arguments are matched to `param` - in the graph. - -| -[source,c++] ----- -template -void register(T param, node node); ----- -|Associate a scalar dynamic parameter with a scalar kernel parameter of `node`. -All scalar kernel arguments matching `param` are bound to the dynamic parameter. -If `param` has already been registered with the dynamic parameter for this node, -then this operation is a no-op. - -Preconditions: - -* This method is only available when the templated `dynamic_parameter_type` - of the `dynamic_parameter` instance is `dynamic_parameter_type::scalar`. - -Parameters: - -* `param` - Scalar used as a node argument. -* `node` - Kernel command to look for an argument matching `param`. - -Exceptions: - -* Throws with error code `invalid` if `param` does not match any arguments in - `node`. -* Throws with error code `invalid` if the size of `param` does not match the - size specified on creation of the dynamic parameter. -* Throws with error code `invalid` if the type of `node` is not a kernel - execution. -* Throws with error code `invalid` if `node` is not a member of the graph that - this dynamic parameter was created with. - -| -[source,c++] ----- -template -void register(T param); ----- -|Associate the dynamic parameter with all the scalar kernel parameters which -match `param` in the graph used on creation of the dynamic parameter. For -every kernel node in the graph, the dynamic parameter binds to any kernel -arguments matching `param`. It is not an error if the dynamic parameter fails -to bind to a single node, only if no nodes in the graph have an argument -matching `param`. - -Preconditions: - -* This method is only available when the templated `dynamic_parameter_type` - of the `dynamic_parameter` instance is `dynamic_parameter_type::scalar`. - -Parameters: - -* `param` - Scalar used as a node argument. - -Exceptions: - -* Throws with error code `invalid` if no node arguments are matched to `param` - in the graph. - -* Throws with error code `invalid` if the size of `param` does not match the - size specified on creation of the dynamic parameter. - -| -[source,c++] ----- -void register(void* param, node node); ----- -|Associate the dynamic parameter with kernel parameter for `node`. All -kernel arguments matching the bytes pointed to by `param` are bound to the -dynamic parameter. If `param` has already been registered with the dynamic -parameter for `node`, this operation is a no-op. - -Preconditions: - -* This method is only available when the templated `dynamic_parameter_type` - of the `dynamic_parameter` instance is `dynamic_parameter_type::raw`. - -Parameters: - -* `param` - Byte array of size set on the creation of the dynamic parameter. -* `node` - Kernel command to look for an argument matching `param`. - -Exceptions: - -* Throws with error code `invalid` if `param` is a null pointer. -* Throws with error code `invalid` if `param` does not match any arguments in - `node`. -* Throws with error code `invalid` if the type of `node` is not a kernel - execution. -* Throws with error code `invalid` if `node` is not a member of the graph that - this dynamic parameter was created with. - -| -[source,c++] ----- -void register(void* param); ----- -|Associate the dynamic parameter with all the kernel parameters which match -the bytes pointed to by `param` in the graph used on creation of the dynamic -parameter. For every kernel node in the graph, the dynamic parameter binds to -any kernel arguments matching the bytes of `param`. It is not an error if the -dynamic parameter fails to bind to a single node, only if no nodes in the graph -have an argument matching `param`. - -Preconditions: - -* This method is only available when the templated `dynamic_parameter_type` - of the `dynamic_parameter` instance is `dynamic_parameter_type::raw`. - -Parameters: - -* `param` - Byte array of size set on the creation of the dynamic parameter. - -Exceptions: - -* Throws with error code `invalid` if `param` is a null pointer. * Throws with error code `invalid` if no node arguments are matched to `param` in the graph. @@ -1291,9 +1170,10 @@ dynamic parameter is matched to a node argument, and the `dynamic_parameter` can then be used to update that node argument in future. If a user doesn't have the individual `node` handles to register parameters -with, they can use the `dynamic_parameter::register()` API that doesn't take -a `node`. This entry-point iterates through the modifiable graph nodes, and -tries to register the parameter against every kernel node in the graph. +with, and the dynamic parameter type isn't raw or scalar, they can use the +`dynamic_parameter::register()` API that doesn't take a `node` argument. +This entry-point iterates through the modifiable graph nodes, and tries to +register the parameter against every kernel node in the graph. The update itself is then performed using a `dynamic_parameter` instance by calling `dynamic_parameter::update()` to update all the parameters of nodes in From 358412c101ac6edfa5d71e644a2feed8ea73b0fe Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 18 Dec 2023 08:47:14 +0000 Subject: [PATCH 21/22] Rename query functions to get predecessors/successors --- .../sycl_ext_oneapi_graph.asciidoc | 18 ++++++++---------- 1 file changed, 8 insertions(+), 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 932988e675aef..5802a438ab332 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -458,9 +458,9 @@ public: node_type get_type() const; - std::vector get_dependencies() const; + std::vector get_predecessors() const; - std::vector get_dependent_nodes() const; + std::vector get_successors() const; static node get_node_from_event(event nodeEvent); }; @@ -1013,14 +1013,14 @@ node_type get_type() const; | [source,c++] ---- -std::vector get_dependencies() const; +std::vector get_predecessors() const; ---- |Returns a list of the predecessor nodes which this node depends on. | [source,c++] ---- -std::vector get_dependent_nodes() const; +std::vector get_successors() const; ---- |Returns a list of the successor nodes which depend on this node. @@ -1188,12 +1188,12 @@ parameters on a node in a graph in the `executable` state 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 behaviour of a graph when executed. +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 behaviour. To achieve this, one +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. @@ -2255,8 +2255,6 @@ int *ptrY = malloc_device(n, myQueue); int *ptrZ = malloc_shared(n, myQueue); int *ptrQ = malloc_device(n, myQueue); -// Initialize USM allocations - // Kernel loaded from kernel bundle const std::vector builtinKernelIds = myDevice.get_info(); @@ -2275,7 +2273,7 @@ dynamic_parameter dynParamScalar(myGra // First node uses ptrX as an input & output parameter, with operand // mySclar as another argument. int myScalar = 42; -node nodeA = myGraph.add[&](handler& cgh) { +node nodeA = myGraph.add([&](handler& cgh) { cgh.set_args(ptrX, myScalar); cgh.parallel_for(range {n}, builtinKernel); }); @@ -2285,7 +2283,7 @@ dynParamInput.register(0, nodeA); // Argument index 0 is ptrX dynParamInput.register(1, nodeA); // Argument index 1 is myScalar // Second node uses ptrX as an inputs, and ptrY as input/output/ -node nodeB = myGraph.add[&](handler& cgh) { +node nodeB = myGraph.add([&](handler& cgh) { cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) { const size_t i = it[0]; ptrY[i] += ptrX[i]; From 6406b2a39dd7fc95c4939a50d4428eea90050a3e Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 21 Dec 2023 12:35:56 +0000 Subject: [PATCH 22/22] Remove queries from this PR --- .../sycl_ext_oneapi_graph.asciidoc | 80 +------------------ 1 file changed, 1 insertion(+), 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 5802a438ab332..9464b8369be66 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -313,19 +313,6 @@ enum class graph_support_level { emulated }; -enum class node_type { - empty, - subgraph, - kernel, - memcpy, - memset, - memfill, - prefetch, - memadvise, - ext_oneapi_barrier, - host_task, -}; - namespace property { namespace graph { @@ -452,18 +439,7 @@ public: const std::string& get_label() const; }; -class node { -public: - node() = delete; - - node_type get_type() const; - - std::vector get_predecessors() const; - - std::vector get_successors() const; - - static node get_node_from_event(event nodeEvent); -}; +class node {}; // State of a graph enum class graph_state { @@ -499,9 +475,6 @@ public: void make_edge(node& src, node& dest); void print_graph(std::string path, bool verbose = false) const; - - std::vector get_nodes() const; - std::vector get_root_nodes() const; }; template<> @@ -996,57 +969,6 @@ Exceptions: |=== -==== Node Member Functions - -Table {counter: tableNumber}. Member functions of the `node` class. -[cols="2a,a"] -|=== -|Member Function|Description - -| -[source,c++] ----- -node_type get_type() const; ----- -|Returns a value representing the type of command this node represents. - -| -[source,c++] ----- -std::vector get_predecessors() const; ----- -|Returns a list of the predecessor nodes which this node depends on. - -| -[source,c++] ----- -std::vector get_successors() const; ----- -|Returns a list of the successor nodes which depend on this node. - -| -[source,c++] ----- -static node get_node_from_event(event nodeEvent); ----- -|Finds the node associated with an event created from a submission to a queue - in the recording state. - -Parameters: - -* `nodeEvent` - Event returned from a submission to a queue in the recording - state. - -Returns: Graph node that was created when the command that returned -`nodeEvent` was submitted. - -Exceptions: - -* Throws with error code `invalid` if `nodeEvent` is not associated with a - graph node. - -|=== - ==== Depends-On Property The API for explicitly adding nodes to a `command_graph` includes a