From 7858877e796a7f0e42ffa3eb89e676027f20d89a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 17 Jul 2024 12:51:29 +0100 Subject: [PATCH 1/6] [SYCL][Graph] Add specification for kernel binary updates Adds the kernel binary update feature to the sycl graph specification. This introduces a new dynamic_command_group class which can be used to update the command-group function of a kernel nodes in graphs. --- .../sycl_ext_oneapi_graph.asciidoc | 253 ++++++++++++++++-- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 45 ++++ 2 files changed, 275 insertions(+), 23 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 56f09c04d3055..9751dddcbe49c 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -556,6 +556,114 @@ Parameters: |=== +==== Dynamic Command Groups + +[source,c++] +---- +namespace ext::oneapi::experimental { +class dynamic_command_group { +public: + dynamic_command_group( + command_graph graph, + const std::vector>& cgfList); + + size_t get_active_cgf(); + void set_active_cgf(size_t cgfIndex); +}; +---- + +Dynamic command-groups can be added as nodes to a graph. They provide a mechanism that +allows updating the command-group function of a node after the graph is finalized. +There is always one command-group function in the dynamic command-group that is set +as active. When a dynamic command-group node is executed, the kernel of the active +command-group function will be run and all the other command-group functions in +`cgfList` will be ignored. + +See <> for more information +about updating command-groups. + +===== Limitations + +Dynamic command-groups can only be used to update kernels. Trying to update a command-group +function that contains other operations will result in an error. + +All the command-group functions in a dynamic command-group must have identical dependencies. +It is not allowed for a dynamic command-group to have command-group functions that would +result in a change to the graph topology when set to active. In practice, this means that +any calls to `handler.depends_on()` must be identical for all the command-group functions +in a dynamic command-group. + +Table {counter: tableNumber}. Member functions of the `dynamic_command_group` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +dynamic_command_group( +command_graph graph, +const std::vector>& cgfList); +---- + +|Constructs a dynamic command-group object that can be added as a node to a `command_graph`. + +Parameters: + +* `graph` - Graph to be associated with this `dynamic_command_group`. +* `cgfList` - The list of command-group functions that can be activated for this dynamic command-group. + The command-group function at index 0 will be active by default. + +Exceptions: + +* Throws synchronously with error code `invalid` if the graph wasn't created with + the `property::graph::assume_buffer_outlives_graph` property and the `dynamic_command_group` + is created with command-group functions that use buffers. See the + <> + property for more information. + +* Throws with error code `invalid` if the `dynamic_command_group` is created with + command-group functions that are not kernel executions. + +* Throws with error code `invalid` if the command-group functions in `cgfList` have + event dependencies that are incompatible with each other and would result in + different graph topologies when set to active. + +| +[source,c++] +---- +size_t get_active_cgf(); +---- +|Returns the index of the currently active command-group function in this +`dynamic_command_group`. + +| +[source,c++] +---- +void set_active_cgf(size_t cgfIndex); +---- +| Sets the command-group function with index `cgfIndex` as active. The index of the +command-group function in a `dynamic_command_group` is identical to its index in the +`cgfList` vector when it was passed to the `dynamic_command_group` constructor. + +This change will be reflected immediately in the modifiable graph which contains this +`dynamic_command_group`. The new value will not be reflected in any executable graphs +created from that modifiable graph until `command_graph::update()` is called, passing +the modified nodes, or a new executable graph is finalized from the modifiable graph. + +Setting `cgfIndex` to the index of the currently active command-group function is +a no-op. + +Parameters: + +* `cgfIndex` - The index of the command-group function that should be set as active. + +Exceptions: + +* Throw with error code `invalid` if `cgfIndex` is not a valid index. + +|=== + ==== Depends-On Property [source,c++] @@ -631,6 +739,8 @@ public: template node add(T cgf, const property_list& propList = {}); + node add(dynamic_command_group& dynamicCG, const property_list& propList = {}); + void make_edge(node& src, node& dest); void print_graph(std::string path, bool verbose = false) const; @@ -711,21 +821,39 @@ Updates to a graph will be scheduled after any in-flight executions of the same graph and will not affect previous submissions of the same graph. The user is not required to wait on any previous submissions of a graph before updating it. -The only type of nodes that are currently able to be updated in a graph are -kernel execution nodes. - -The aspects of a kernel execution node that can be configured during update are: - -* Parameters to the kernel. -* Execution ND-Range of the kernel. - To update an executable graph, the `property::graph::updatable` property must have been set when the graph was created during finalization. Otherwise, an exception will be thrown if a user tries to update an executable graph. This guarantee allows the backend to provide a more optimized implementation, if possible. -===== Individual Node Update +===== Supported Features + +The only types of nodes that are currently able to be updated in a graph are +kernel execution nodes. + +There are two different API's that can be used to update a graph: + +* <> which allows updating +individual nodes of a command-graph. +* <> which allows updating the +entirety of the graph simultaneously by using another graph as a +reference. + +The aspects of a kernel execution node that can be changed during update are +different depending on the API used to perform the update: + +* For the <> API it's possible to update +the kernel function, the parameters to the kernel, and the ND-Range. +* For the <> API, only the parameters of the kernel +and the ND-Range can be updated. + +===== Individual Node Update [[individual-node-update]] + +Individual nodes of an executable graph can be updated directly. Depending on the attribute +of the node that requires updating, different API's should be used: + +====== Parameter Updates Parameters to individual nodes in a graph in the `executable` state can be updated between graph executions using dynamic parameters. A `dynamic_parameter` @@ -739,14 +867,6 @@ Parameter updates are performed using a `dynamic_parameter` instance by calling not registered, even if they use the same parameter value as a `dynamic_parameter`. -The other node configuration that can be updated is the execution range of the -kernel, this can be set through `node::update_nd_range()` or -`node::update_range()` but does not require any prior registration. - -The executable graph can then be updated by passing the updated nodes to -`command_graph::update(node& node)` or -`command_graph::update(const std::vector& nodes)`. - Since the structure of the graph became fixed when finalizing, updating parameters on a node will not change the already defined dependencies between nodes. This is important to note when updating buffer parameters to a node, @@ -762,6 +882,41 @@ 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. +====== Execution Range Updates + +Another configuration that can be updated is the execution range of the +kernel, this can be set through `node::update_nd_range()` or +`node::update_range()` but does not require any prior registration. + +An alternative way to update the execution range of a node is to do so while +updating command groups as described in the next section. + +====== Command Group Updates + +The command-groups of a kernel node can be updated using dynamic command-groups. +Dynamic command-groups allow replacing the command-group function of a kernel +node with a different one. This effectively allows updating the kernel function +and/or the kernel execution range. + +Command-group updates are performed by creating an instance of the +`dynamic_command_group` class. A dynamic command-group is created with a modifiable +state graph and a list of possible command-group functions. Command-group functions +within a dynamic command-group can then be set to active by using the member function +`dynamic_command_group::set_active_cgf()`. + +Dynamic command-groups are compatible with dynamic parameters. This means that +dynamic parameters can be used in command-group functions that are part of +dynamic command-groups. Updates to such dynamic parameters will be reflected +in the command-group functions once they are activated. + +====== Committing Updates + +Updating a node using the methods mentioned above will take effect immediately +for nodes in modifiable command-graphs. However, for graphs that are in the executable +state, in order to commit the update, the updated nodes must be passed to +`command_graph::update(node& node)` or +`command_graph::update(const std::vector& nodes)`. + ===== Whole Graph Update [[whole-graph-update]] A graph in the executable state can have all of its nodes updated using the @@ -1042,6 +1197,42 @@ Exceptions: | [source,c++] ---- +node add(dynamic_command_group& dynamicCG, const property_list& propList = {}); +---- + +| Adds the dynamic command-group `dynamicCG` as a node to the graph and sets the +current active command-group function in `dynamicCG` as the executable for future +executions of this graph node. + +The current active command-group function in `dynamicCG` will be executed asynchronously +when the graph is submitted to a queue. The requisites of this command-group +function will be used to identify any dependent nodes in the graph +to form edges with. The other command-group functions in `dynamicCG` will be captured +into the graph but will not be executed in a graph submission unless they are +set to active. + +Constraints: + +* This member function is only available when the `command_graph` state is + `graph_state::modifiable`. + +Parameters: + +* `dynamicCG` - Dynamic command-group object to be added as a node. + +* `propList` - Zero or more properties can be provided to the constructed node + via an instance of `property_list`. The `property::node::depends_on` property + can be passed here with a list of nodes to create dependency edges on. + +Returns: The dynamic command-group object node which has been added to the graph. + +Exceptions: + +* Throws synchronously with error code `invalid` if a queue is recording + commands to the graph. +| +[source,c++] +---- void make_edge(node& src, node& dest); ---- @@ -1157,8 +1348,9 @@ void update(node& node); ---- | Updates an executable graph node that corresponds to `node`. `node` must be a -kernel execution node. Kernel arguments and the ND-range of the node will be -updated inside the executable graph to reflect the current values in `node`. +kernel execution node. The command-group function of the node will be updated, +inside the executable graph, to reflect the current values in `node`. This +includes the kernel function, the kernel nd-range and the kernel parameters. Updating these values will not change the structure of the graph. @@ -1190,9 +1382,10 @@ void update(const std::vector& nodes); ---- | Updates all executable graph nodes that corresponds to the nodes contained in -`nodes`. All nodes must be kernel nodes. Kernel arguments and the ND-range of -each node will be updated inside the executable graph to reflect the current -values in each node in `nodes`. +`nodes`. All nodes must be kernel nodes. The command-group function of each node +will be updated, inside the executable graph, to reflect the current values in +`nodes`. This includes the kernel function, the kernel nd-range and the kernel +parameters". Updating these values will not change the structure of the graph. @@ -1749,6 +1942,10 @@ the call to `queue::submit()` or `command_graph::add()` along with the calls to handler functions and this will not be reflected on future executions of the graph. +Similarly, any command-group function inside a `dynamic_command_group` will be +evaluated once, in index order, when submitted to the graph using +`command_graph::add()`. + Any code like this should be moved to a separate host-task and added to the graph via the recording or explicit APIs in order to be compatible with this extension. @@ -1980,7 +2177,7 @@ can be used to add nodes to a graph when creating a graph from queue recording. == Examples and Usage Guide Detailed code examples and usage guidelines are provided in the -link:../../SYCLGraphUsageGuide.md[SYCL Graph Usage Guide]. +link:../../syclgraph/SYCLGraphUsageGuide.md[SYCL Graph Usage Guide]. == Future Direction [[future-direction]] @@ -2120,6 +2317,16 @@ to ensure this is desired and makes sense to users. **UNRESOLVED** Needs more discussion. +=== Updatable command-groups in the Record & Replay API: + +Currently the only way to update command-groups in a graph is to use the +Explicit API. There is a limitation in some backends that requires all +the command-groups used for updating to be specified before the graph +is finalized. This restriction makes it hard to implement the +Record & Replay API in a performant manner. + +**UNRESOLVED** Needs more discussion. + === Multi Device Graph Allow an executable graph to contain nodes targeting different devices. diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index f51d441174ef6..c7c824b5191ee 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -453,6 +453,51 @@ node kernelNode = myGraph.add([&](handler& cgh) { dynParamAccessor.update(bufferB.get_access()); ``` +### Dynamic Command Groups + +Example showing how a graph with a dynamic command group node can be updated. + +```cpp +queue Queue{}; +exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + +int *PtrA = malloc_device(1024, Queue); +int *PtrB = malloc_device(1024, Queue)​ + +auto CgfA = [&](handler &cgh) { + cgh.parallel_for(1024, [=](item<1> Item) { + PtrA[Item.get_id()] = 1;​ + }); +}; + +auto CgfB = [&](handler &cgh) { + cgh.parallel_for(512, [=](item<1> Item) { + PtrB[Item.get_id()] = 2; + }); +}; + +// Construct a dynamic command-group with CgfA as the active cgf (index 0). +auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CgfA, CgfB}); + +// Create a dynamic command-group graph node. +auto DynamicCGNode = Graph.add(DynamicCG); + +auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); + +// The graph will execute CgfA. +Queue.ext_oneapi_graph(ExecGraph).wait(); + +// Sets CgfB as active in the dynamic command-group (index 1). +DynamicCG.set_active_cgf(1); + +// Calls update to update the executable graph node with the changes to DynamicCG. +ExecGraph.update(DynamicCGNode); + +// The graph will execute CgfB. +Queue.ext_oneapi_graph(ExecGraph).wait(); +``` + + ### Whole Graph Update Example that shows recording and updating several nodes with different From 119ba28c0ffa934b40ae6bf46b3a90e2ea835b56 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Tue, 15 Oct 2024 17:21:35 +0100 Subject: [PATCH 2/6] Refine based on implementation findings --- .../sycl_ext_oneapi_graph.asciidoc | 69 +++++++++++-------- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 11 ++- 2 files changed, 47 insertions(+), 33 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 9751dddcbe49c..466cea13aae2f 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -425,7 +425,7 @@ Exceptions: template void update_nd_range(nd_range executionRange); ---- -| Updates the ND-Range for this node with a new value. This new value will not +| Updates the ND-range for this node with a new value. This new value will not affect any executable graphs this node is part of until it is passed to the executable graph's update function. See <> for more information @@ -433,15 +433,12 @@ about updating kernel nodes. Parameters: -* `executionRange` - The new value for the ND-Range. +* `executionRange` - The new value for the ND-range. Exceptions: * Throws with error code `invalid` if `Dimensions` does not match the dimensions - of the nd_range the kernel was originally created with. - -* Throws with error code `invalid` if the kernel node was originally created - with a `sycl::range`. + of the existing kernel execution range. * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -452,7 +449,7 @@ Exceptions: template void update_range(range executionRange); ---- -| Updates the execution Range for this node with a new value. This new value +| Updates the execution range for this node with a new value. This new value will not affect any executable graphs this node is part of until it is passed to the executable graph's update function. See <> for more information @@ -460,15 +457,12 @@ about updating kernel nodes. Parameters: -* `executionRange` - The new value for the Range. +* `executionRange` - The new value for the range. Exceptions: * Throws with error code `invalid` if `Dimensions` does not match the dimensions - of the range the kernel was originally created with. - -* Throws with error code `invalid` if the kernel node was originally created - with a `sycl::nd_range`. + of the existing kernel execution range. * Throws with error code `invalid` if the type of the node is not a kernel execution. @@ -523,7 +517,7 @@ Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. | [source,c++] ---- -dynamic_parameter(command_graph graph, +dynamic_parameter(command_graph graph, const ValueT &initialValue); ---- |Constructs a dynamic parameter object that can be registered with command graph @@ -564,10 +558,10 @@ namespace ext::oneapi::experimental { class dynamic_command_group { public: dynamic_command_group( - command_graph graph, + command_graph &graph, const std::vector>& cgfList); - size_t get_active_cgf(); + size_t get_active_cgf() const; void set_active_cgf(size_t cgfIndex); }; ---- @@ -579,6 +573,8 @@ as active. When a dynamic command-group node is executed, the kernel of the acti command-group function will be run and all the other command-group functions in `cgfList` will be ignored. +The `dynamic_command_group` class provides the {crs}[common reference semantics]. + See <> for more information about updating command-groups. @@ -591,7 +587,8 @@ All the command-group functions in a dynamic command-group must have identical d It is not allowed for a dynamic command-group to have command-group functions that would result in a change to the graph topology when set to active. In practice, this means that any calls to `handler.depends_on()` must be identical for all the command-group functions -in a dynamic command-group. +in a dynamic command-group. The dependencies created by buffer accessors must also create +identical node dependencies across all of the command-group functions. Table {counter: tableNumber}. Member functions of the `dynamic_command_group` class. [cols="2a,a"] @@ -602,7 +599,7 @@ Table {counter: tableNumber}. Member functions of the `dynamic_command_group` cl [source,c++] ---- dynamic_command_group( -command_graph graph, +command_graph &graph, const std::vector>& cgfList); ---- @@ -618,21 +615,19 @@ Exceptions: * Throws synchronously with error code `invalid` if the graph wasn't created with the `property::graph::assume_buffer_outlives_graph` property and the `dynamic_command_group` - is created with command-group functions that use buffers. See the + is created with any command-group function that uses buffers. See the <> property for more information. * Throws with error code `invalid` if the `dynamic_command_group` is created with command-group functions that are not kernel executions. -* Throws with error code `invalid` if the command-group functions in `cgfList` have - event dependencies that are incompatible with each other and would result in - different graph topologies when set to active. +* Throws with error code `invalid` if `cgfList` is empty. | [source,c++] ---- -size_t get_active_cgf(); +size_t get_active_cgf() const; ---- |Returns the index of the currently active command-group function in this `dynamic_command_group`. @@ -844,9 +839,9 @@ The aspects of a kernel execution node that can be changed during update are different depending on the API used to perform the update: * For the <> API it's possible to update -the kernel function, the parameters to the kernel, and the ND-Range. +the kernel function, the parameters to the kernel, and the ND-range. * For the <> API, only the parameters of the kernel -and the ND-Range can be updated. +and the ND-range can be updated. ===== Individual Node Update [[individual-node-update]] @@ -889,7 +884,11 @@ kernel, this can be set through `node::update_nd_range()` or `node::update_range()` but does not require any prior registration. An alternative way to update the execution range of a node is to do so while -updating command groups as described in the next section. +updating command groups as described in the next section. Using this mechanism +lifts the restriction from `node::update_nd_range()` / `node::update_range()` +of only being to update the execution range in the same dimension. As the +update being tied to a change in command-group means that the updated kernel +code may be defined as operating in a different dimension. ====== Command Group Updates @@ -909,6 +908,14 @@ dynamic parameters can be used in command-group functions that are part of dynamic command-groups. Updates to such dynamic parameters will be reflected in the command-group functions once they are activated. +Note that the execution range is tied to the command-group, therefore updating +the range of a node which uses a dynamic command-group shared by another node +will also update the execution range of the nodes sharing the dynamic +command-group. Activating a command-group with `set_active_cgf` to a +command-group that previously had its execution range updated with +`node::update_range()` or `node::update_nd_range()` will not reset the execution +range to the original, but instead use the most recently updated value. + ====== Committing Updates Updating a node using the methods mentioned above will take effect immediately @@ -1230,6 +1237,14 @@ Exceptions: * Throws synchronously with error code `invalid` if a queue is recording commands to the graph. + +* Throws synchronously with error code `invalid` if the graph does not match + the graph used on construction of `dynamicCG`. + +* Throws with error code `invalid` if the command-group functions in `cgfList` have + event or accessor dependencies that are incompatible with each other and + would result in different graph topologies when set to active. + | [source,c++] ---- @@ -2381,8 +2396,8 @@ if used in application code. . Using reductions in a graph node. . Using sycl streams in a graph node. -. 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 diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index c7c824b5191ee..50d36dd2b5b80 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -394,12 +394,12 @@ sycl_ext::command_graph myGraph(myContext, myDevice); int myScalar = 42; // Create graph dynamic parameters -dynamic_parameter dynParamInput(myGraph, ptrX); -dynamic_parameter dynParamScalar(myGraph, myScalar); +sycl_ext::dynamic_parameter dynParamInput(myGraph, ptrX); +sycl_ext::dynamic_parameter dynParamScalar(myGraph, myScalar); // The node uses ptrX as an input & output parameter, with operand // mySclar as another argument. -node kernelNode = myGraph.add([&](handler& cgh) { +sycl_ext::node kernelNode = myGraph.add([&](handler& cgh) { cgh.set_args(dynParamInput, ptrY, dynParamScalar); cgh.parallel_for(range {n}, builtinKernel); }); @@ -438,9 +438,9 @@ sycl::buffer bufferB{...}; // Create graph dynamic parameter using a placeholder accessor, since the // sycl::handler is not available here outside of the command-group scope. -dynamic_parameter dynParamAccessor(myGraph, bufferA.get_access()); +sycl_ext::dynamic_parameter dynParamAccessor(myGraph, bufferA.get_access()); -node kernelNode = myGraph.add([&](handler& cgh) { +sycl_ext::node kernelNode = myGraph.add([&](handler& cgh) { // Require the accessor contained in the dynamic paramter cgh.require(dynParamAccessor); // Set the arg on the kernel using the dynamic parameter directly @@ -497,7 +497,6 @@ ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); ``` - ### Whole Graph Update Example that shows recording and updating several nodes with different From 6655435f62f4bf60b6ec1f8b54cb119d921b7931 Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 29 Oct 2024 14:54:04 +0000 Subject: [PATCH 3/6] Add example using dynamic CGF and parameters together --- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 61 +++++++++++++++++++++++ 1 file changed, 61 insertions(+) diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index 50d36dd2b5b80..5a13f2b987262 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -497,6 +497,67 @@ ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); ``` +### Dynamic Command Groups With Dynamic Parameters + +Example showing how a graph with a dynamic command group that uses dynamic +parameters in a node can be updated. + +```cpp +size_t n = 1024; +queue Queue{}; +exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + +int *PtrA = malloc_device(n, Queue); +int *PtrB = malloc_device(n, Queue)​; + +// Kernels loaded from kernel bundle +const std::vector builtinKernelIds = + myDevice.get_info(); +kernel_bundle myBundle = + get_kernel_bundle(myContext, { myDevice }, builtinKernelIds); + +kernel builtinKernelA = myBundle.get_kernel(builtinKernelIds[0]); +kernel builtinKernelB = myBundle.get_kernel(builtinKernelIds[1]); + +// Create a dynamic parameter with an initial value of PtrA +exp_ext::dynamic_parameter DynamicPointerArg{Graph, PtrA}; + +// Create command groups for both kernels which use DynamicPointerArg +auto CgfA = [&](handler &cgh) { + cgh.set_arg(0, DynamicPointerArg); + cgh.parallel_for(range {n}, builtinKernelA); +}; + +auto CgfB = [&](handler &cgh) { + cgh.set_arg(0, DynamicPointerArg); + cgh.parallel_for(range {n / 2}, builtinKernelB); +}; + +// Construct a dynamic command-group with CgfA as the active cgf (index 0). +auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CgfA, CgfB}); + +// Create a dynamic command-group graph node. +auto DynamicCGNode = Graph.add(DynamicCG); + +auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); + +// The graph will execute CgfA with PtrA. +Queue.ext_oneapi_graph(ExecGraph).wait(); + +//Update DynamicPointerArg with a new value +DynamicPointerArg.update(PtrB); + +// Sets CgfB as active in the dynamic command-group (index 1). +DynamicCG.set_active_cgf(1); + +// Calls update to update the executable graph node with the changes to +// DynamicCG and DynamicPointerArg. +ExecGraph.update(DynamicCGNode); + +// The graph will execute CgfB with PtrB. +Queue.ext_oneapi_graph(ExecGraph).wait(); +``` + ### Whole Graph Update Example that shows recording and updating several nodes with different From c5b9fa2219fd1b8f01b9d3905834d9c4101dd798 Mon Sep 17 00:00:00 2001 From: Benjamin Tracy Date: Wed, 30 Oct 2024 14:54:44 +0000 Subject: [PATCH 4/6] Apply suggestions from code review Co-authored-by: Pablo Reble --- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index 5a13f2b987262..e25774953add4 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -462,7 +462,7 @@ queue Queue{}; exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; int *PtrA = malloc_device(1024, Queue); -int *PtrB = malloc_device(1024, Queue)​ +int *PtrB = malloc_device(1024, Queue)​; auto CgfA = [&](handler &cgh) { cgh.parallel_for(1024, [=](item<1> Item) { @@ -514,7 +514,7 @@ int *PtrB = malloc_device(n, Queue)​; const std::vector builtinKernelIds = myDevice.get_info(); kernel_bundle myBundle = - get_kernel_bundle(myContext, { myDevice }, builtinKernelIds); + get_kernel_bundle(myContext, { myDevice }, builtinKernelIds); kernel builtinKernelA = myBundle.get_kernel(builtinKernelIds[0]); kernel builtinKernelB = myBundle.get_kernel(builtinKernelIds[1]); From 6fa91735ab99a455e418fefc48fe8f0550134238 Mon Sep 17 00:00:00 2001 From: Benjamin Tracy Date: Wed, 30 Oct 2024 14:55:33 +0000 Subject: [PATCH 5/6] Update sycl/doc/syclgraph/SYCLGraphUsageGuide.md Co-authored-by: Pablo Reble --- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index e25774953add4..d402b107a1502 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -505,7 +505,9 @@ parameters in a node can be updated. ```cpp size_t n = 1024; queue Queue{}; -exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; +auto myContext = Queue.get_context(); +auto myDevice = Queue.get_device(); +exp_ext::command_graph Graph{myContext, myDevice}; int *PtrA = malloc_device(n, Queue); int *PtrB = malloc_device(n, Queue)​; From 3edf87064abdb35ea49dc464af39602e8aceced4 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 2 Dec 2024 09:07:19 +0000 Subject: [PATCH 6/6] Update get/setters to "index" rather than "cgf" Also refine spec wording based on PR feedback --- .../sycl_ext_oneapi_graph.asciidoc | 44 +++++++----- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 56 ++++++++------- .../sycl/ext/oneapi/experimental/graph.hpp | 4 +- sycl/source/detail/graph_impl.cpp | 4 +- .../Graph/Update/dyn_cgf_accessor.cpp | 2 +- .../Graph/Update/dyn_cgf_accessor_deps.cpp | 2 +- .../Graph/Update/dyn_cgf_accessor_deps2.cpp | 2 +- .../Graph/Update/dyn_cgf_accessor_spv.cpp | 2 +- .../Update/dyn_cgf_different_arg_nums.cpp | 6 +- .../Graph/Update/dyn_cgf_event_deps.cpp | 2 +- .../Graph/Update/dyn_cgf_get_active_index.cpp | 70 +++++++++++++++++++ .../test-e2e/Graph/Update/dyn_cgf_ndrange.cpp | 2 +- .../Graph/Update/dyn_cgf_ndrange_3D.cpp | 4 +- .../Graph/Update/dyn_cgf_overwrite_range.cpp | 4 +- .../Graph/Update/dyn_cgf_parameters.cpp | 2 +- .../Graph/Update/dyn_cgf_shared_nodes.cpp | 2 +- .../Update/dyn_cgf_update_before_finalize.cpp | 2 +- sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp | 2 +- .../Update/dyn_cgf_with_all_dyn_params.cpp | 4 +- ...dyn_cgf_with_different_type_dyn_params.cpp | 6 +- .../Update/dyn_cgf_with_some_dyn_params.cpp | 4 +- .../Graph/Update/whole_update_dynamic_cgf.cpp | 4 +- sycl/test/abi/sycl_symbols_linux.dump | 12 ++-- sycl/test/abi/sycl_symbols_windows.dump | 4 +- 24 files changed, 165 insertions(+), 81 deletions(-) create mode 100644 sycl/test-e2e/Graph/Update/dyn_cgf_get_active_index.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e579384363784..7c410bdc473f9 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -562,17 +562,20 @@ public: command_graph &graph, const std::vector>& cgfList); - size_t get_active_cgf() const; - void set_active_cgf(size_t cgfIndex); + size_t get_active_index() const; + void set_active_index(size_t cgfIndex); }; ---- -Dynamic command-groups can be added as nodes to a graph. They provide a mechanism that -allows updating the command-group function of a node after the graph is finalized. -There is always one command-group function in the dynamic command-group that is set -as active. When a dynamic command-group node is executed, the kernel of the active -command-group function will be run and all the other command-group functions in -`cgfList` will be ignored. +Dynamic command-groups can be added as nodes to a graph. They provide a +mechanism that allows updating the command-group function of a node after the +graph is finalized. There is always one command-group function in the dynamic +command-group that is set as active, this is the kernel which will execute for +the node when the graph is finalized into an executable state `command_graph`, +and all the other command-group functions in `cgfList` will be ignored. The +executable `command_graph` node can then be updated to a different kernel in +`cgfList`, by selecting a new active index on the dynamic command-group object +and calling the `update(node& node)` method on the executable `command_graph`. The `dynamic_command_group` class provides the {crs}[common reference semantics]. @@ -581,8 +584,9 @@ about updating command-groups. ===== Limitations -Dynamic command-groups can only be used to update kernels. Trying to update a command-group -function that contains other operations will result in an error. +Dynamic command-groups can only contain kernel operations. Trying to construct +a dynamic command-group with functions that contain other operations will +result in an error. All the command-group functions in a dynamic command-group must have identical dependencies. It is not allowed for a dynamic command-group to have command-group functions that would @@ -628,7 +632,7 @@ Exceptions: | [source,c++] ---- -size_t get_active_cgf() const; +size_t get_active_index() const; ---- |Returns the index of the currently active command-group function in this `dynamic_command_group`. @@ -636,7 +640,7 @@ size_t get_active_cgf() const; | [source,c++] ---- -void set_active_cgf(size_t cgfIndex); +void set_active_index(size_t cgfIndex); ---- | Sets the command-group function with index `cgfIndex` as active. The index of the command-group function in a `dynamic_command_group` is identical to its index in the @@ -902,7 +906,7 @@ Command-group updates are performed by creating an instance of the `dynamic_command_group` class. A dynamic command-group is created with a modifiable state graph and a list of possible command-group functions. Command-group functions within a dynamic command-group can then be set to active by using the member function -`dynamic_command_group::set_active_cgf()`. +`dynamic_command_group::set_active_index()`. Dynamic command-groups are compatible with dynamic parameters. This means that dynamic parameters can be used in command-group functions that are part of @@ -910,12 +914,14 @@ dynamic command-groups. Updates to such dynamic parameters will be reflected in the command-group functions once they are activated. Note that the execution range is tied to the command-group, therefore updating -the range of a node which uses a dynamic command-group shared by another node -will also update the execution range of the nodes sharing the dynamic -command-group. Activating a command-group with `set_active_cgf` to a -command-group that previously had its execution range updated with -`node::update_range()` or `node::update_nd_range()` will not reset the execution -range to the original, but instead use the most recently updated value. +the range of a node which uses a dynamic command-group will update the +execution range of the currently active command-group. If the dynamic +command-group is shared by another node, it will also update the execution +range of the other nodes sharing that dynamic command-group. Activating a +command-group with `set_active_index` to a command-group that previously had +its execution range updated with `node::update_range()` or +`node::update_nd_range()` will not reset the execution range to the original +value, but instead use the most recently updated value. ====== Committing Updates diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index d402b107a1502..a2ca77258bdd8 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -458,15 +458,19 @@ dynParamAccessor.update(bufferB.get_access()); Example showing how a graph with a dynamic command group node can be updated. ```cpp +... +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + queue Queue{}; -exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; +sycl_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; int *PtrA = malloc_device(1024, Queue); -int *PtrB = malloc_device(1024, Queue)​; +int *PtrB = malloc_device(1024, Queue); auto CgfA = [&](handler &cgh) { cgh.parallel_for(1024, [=](item<1> Item) { - PtrA[Item.get_id()] = 1;​ + PtrA[Item.get_id()] = 1; }); }; @@ -477,18 +481,18 @@ auto CgfB = [&](handler &cgh) { }; // Construct a dynamic command-group with CgfA as the active cgf (index 0). -auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CgfA, CgfB}); +auto DynamicCG = sycl_ext::dynamic_command_group(Graph, {CgfA, CgfB}); // Create a dynamic command-group graph node. auto DynamicCGNode = Graph.add(DynamicCG); -auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); +auto ExecGraph = Graph.finalize(sycl_ext::property::graph::updatable{}); // The graph will execute CgfA. Queue.ext_oneapi_graph(ExecGraph).wait(); // Sets CgfB as active in the dynamic command-group (index 1). -DynamicCG.set_active_cgf(1); +DynamicCG.set_active_index(1); // Calls update to update the executable graph node with the changes to DynamicCG. ExecGraph.update(DynamicCGNode); @@ -503,45 +507,49 @@ Example showing how a graph with a dynamic command group that uses dynamic parameters in a node can be updated. ```cpp -size_t n = 1024; +... +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + +size_t N = 1024; queue Queue{}; -auto myContext = Queue.get_context(); -auto myDevice = Queue.get_device(); -exp_ext::command_graph Graph{myContext, myDevice}; +auto MyContext = Queue.get_context(); +auto MyDevice = Queue.get_device(); +sycl_ext::command_graph Graph{MyContext, MyDevice}; -int *PtrA = malloc_device(n, Queue); -int *PtrB = malloc_device(n, Queue)​; +int *PtrA = malloc_device(N, Queue); +int *PtrB = malloc_device(N, Queue); // Kernels loaded from kernel bundle -const std::vector builtinKernelIds = - myDevice.get_info(); -kernel_bundle myBundle = - get_kernel_bundle(myContext, { myDevice }, builtinKernelIds); +const std::vector BuiltinKernelIds = + MyDevice.get_info(); +kernel_bundle MyBundle = + get_kernel_bundle(MyContext, { MyDevice }, BuiltinKernelIds); -kernel builtinKernelA = myBundle.get_kernel(builtinKernelIds[0]); -kernel builtinKernelB = myBundle.get_kernel(builtinKernelIds[1]); +kernel BuiltinKernelA = MyBundle.get_kernel(BuiltinKernelIds[0]); +kernel BuiltinKernelB = MyBundle.get_kernel(BuiltinKernelIds[1]); // Create a dynamic parameter with an initial value of PtrA -exp_ext::dynamic_parameter DynamicPointerArg{Graph, PtrA}; +sycl_ext::dynamic_parameter DynamicPointerArg{Graph, PtrA}; // Create command groups for both kernels which use DynamicPointerArg auto CgfA = [&](handler &cgh) { cgh.set_arg(0, DynamicPointerArg); - cgh.parallel_for(range {n}, builtinKernelA); + cgh.parallel_for(range {N}, BuiltinKernelA); }; auto CgfB = [&](handler &cgh) { cgh.set_arg(0, DynamicPointerArg); - cgh.parallel_for(range {n / 2}, builtinKernelB); + cgh.parallel_for(range {N / 2}, BuiltinKernelB); }; // Construct a dynamic command-group with CgfA as the active cgf (index 0). -auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CgfA, CgfB}); +auto DynamicCG = sycl_ext::dynamic_command_group(Graph, {CgfA, CgfB}); // Create a dynamic command-group graph node. auto DynamicCGNode = Graph.add(DynamicCG); -auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); +auto ExecGraph = Graph.finalize(sycl_ext::property::graph::updatable{}); // The graph will execute CgfA with PtrA. Queue.ext_oneapi_graph(ExecGraph).wait(); @@ -550,7 +558,7 @@ Queue.ext_oneapi_graph(ExecGraph).wait(); DynamicPointerArg.update(PtrB); // Sets CgfB as active in the dynamic command-group (index 1). -DynamicCG.set_active_cgf(1); +DynamicCG.set_active_index(1); // Calls update to update the executable graph node with the changes to // DynamicCG and DynamicPointerArg. diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index e205d939b3b30..f12e22396f448 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -223,8 +223,8 @@ class __SYCL_EXPORT dynamic_command_group { const command_graph &Graph, const std::vector> &CGFList); - size_t get_active_cgf() const; - void set_active_cgf(size_t Index); + size_t get_active_index() const; + void set_active_index(size_t Index); private: template diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 10624c7c5fd61..8ebb56f713444 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -2075,10 +2075,10 @@ dynamic_command_group::dynamic_command_group( impl->finalizeCGFList(CGFList); } -size_t dynamic_command_group::get_active_cgf() const { +size_t dynamic_command_group::get_active_index() const { return impl->getActiveIndex(); } -void dynamic_command_group::set_active_cgf(size_t Index) { +void dynamic_command_group::set_active_index(size_t Index) { return impl->setActiveIndex(Index); } } // namespace experimental diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp index 803f296d9f71a..318f1290ecc14 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp @@ -45,7 +45,7 @@ int main() { assert(HostData[i] == PatternA); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp index a8018190ab741..b7f9813b82e64 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp @@ -60,7 +60,7 @@ int main() { assert(HostData[i] == Ref); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp index 4e9ada8a3c246..32a861cb347fc 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp @@ -69,7 +69,7 @@ int main() { assert(HostData[i] == (InitA + InitB + PatternA)); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp index 08b5fa293cf80..aaaced3e7ce21 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp @@ -73,7 +73,7 @@ int main(int, char **argv) { assert(check_value(i, 0, HostDataB[i], "HostDataB")); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp index 7288fba3a73d1..5aa691b9c36ae 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp @@ -90,7 +90,7 @@ int main() { // CHECK-SAME: .argIndex = 0 // CHECK-SAME: .stype = UR_STRUCTURE_TYPE_EXP_COMMAND_BUFFER_UPDATE_VALUE_ARG_DESC // CHECK-SAME: .argIndex = 1 - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); Queue.copy(Ptr, HostData.data(), Size).wait(); @@ -107,7 +107,7 @@ int main() { // CHECK-SAME: .numNewValueArgs = 0 // CHECK-SAME: .stype = UR_STRUCTURE_TYPE_EXP_COMMAND_BUFFER_UPDATE_POINTER_ARG_DESC // CHECK-SAME: .argIndex = 0 - DynamicCG.set_active_cgf(2); + DynamicCG.set_active_index(2); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); Queue.copy(Ptr, HostData.data(), Size).wait(); @@ -130,7 +130,7 @@ int main() { // CHECK-SAME: .argIndex = 2 // CHECK-SAME: .stype = UR_STRUCTURE_TYPE_EXP_COMMAND_BUFFER_UPDATE_VALUE_ARG_DESC // CHECK-SAME: .argIndex = 3 - DynamicCG.set_active_cgf(3); + DynamicCG.set_active_index(3); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); Queue.copy(Ptr, HostData.data(), Size).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp index 9556f97de69f1..0379243737911 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp @@ -52,7 +52,7 @@ int main() { assert(HostData[i] == PatternA * PatternB); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_get_active_index.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_get_active_index.cpp new file mode 100644 index 0000000000000..882121539165f --- /dev/null +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_get_active_index.cpp @@ -0,0 +1,70 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// Tests the `get_active_index()` query + +#include "../graph_common.hpp" + +int main() { + queue Queue{}; + exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + + int *Ptr = malloc_device(Size, Queue); + std::vector HostData(Size); + + int PatternA = 42; + auto CGFA = [&](handler &CGH) { + CGH.parallel_for(Size, + [=](item<1> Item) { Ptr[Item.get_id()] = PatternA; }); + }; + + int PatternB = 0xA; + auto CGFB = [&](handler &CGH) { + CGH.parallel_for(Size, + [=](item<1> Item) { Ptr[Item.get_id()] = PatternB; }); + }; + + auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB}); + size_t ActiveIndex = DynamicCG.get_active_index(); + assert(0 == ActiveIndex); // Active index is zero by default + + // Set active index to 1 before adding node to graph + DynamicCG.set_active_index(1); + ActiveIndex = DynamicCG.get_active_index(); + assert(1 == ActiveIndex); + + auto DynamicCGNode = Graph.add(DynamicCG); + + // Set active index to 0 before finalizing the graph + DynamicCG.set_active_index(0); + ActiveIndex = DynamicCG.get_active_index(); + assert(0 == ActiveIndex); + auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); + + Queue.ext_oneapi_graph(ExecGraph).wait(); + Queue.copy(Ptr, HostData.data(), Size).wait(); + for (size_t i = 0; i < Size; i++) { + assert(HostData[i] == PatternA); + } + + // Set active index to 1 before updating the graph + DynamicCG.set_active_index(1); + ActiveIndex = DynamicCG.get_active_index(); + assert(1 == ActiveIndex); + + ExecGraph.update(DynamicCGNode); + + Queue.ext_oneapi_graph(ExecGraph).wait(); + Queue.copy(Ptr, HostData.data(), Size).wait(); + for (size_t i = 0; i < Size; i++) { + assert(HostData[i] == PatternB); + } + + sycl::free(Ptr, Queue); + + return 0; +} diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp index cbe1c2c3e117a..141e8708fd67d 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp @@ -51,7 +51,7 @@ int main() { } } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp index 3fd32ef575cf4..87ff509114a73 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp @@ -54,7 +54,7 @@ int main() { assert(HostData[i] == PatternA); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); @@ -63,7 +63,7 @@ int main() { assert(HostData[i] == PatternB); } - DynamicCG.set_active_cgf(2); + DynamicCG.set_active_index(2); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_overwrite_range.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_overwrite_range.cpp index 04697077bec36..28bb1dc8a958d 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_overwrite_range.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_overwrite_range.cpp @@ -36,10 +36,10 @@ int main() { sycl::range<1> UpdateRange(NewRange); DynamicCGNode.update_range(UpdateRange); - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); // Check that the UpdateRange from active CGF 0 is preserved - DynamicCG.set_active_cgf(0); + DynamicCG.set_active_index(0); auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp index 7f00d0f8750ce..364ed21709672 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp @@ -49,7 +49,7 @@ int main() { assert(HostDataB[i] == 0); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp index eab640b45b258..245ece9ea08b6 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp @@ -55,7 +55,7 @@ int main() { assert(HostData[i] == Ref); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(Node1); ExecGraph.update(Node3); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_update_before_finalize.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_update_before_finalize.cpp index 8c0c705960ef6..c3092cd97a610 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_update_before_finalize.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_update_before_finalize.cpp @@ -31,7 +31,7 @@ int main() { auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB}); auto DynamicCGNode = Graph.add(DynamicCG); - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); auto ExecGraph = Graph.finalize(); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp index 97c454b6db92a..21fb5a2135d99 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp @@ -38,7 +38,7 @@ int main() { assert(HostData[i] == PatternA); } - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); Queue.ext_oneapi_graph(ExecGraph).wait(); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp index 28a55ecfeceeb..06050d8750141 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp @@ -99,11 +99,11 @@ int main() { ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(false, true, false); - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(false, true, false); - DynamicCG.set_active_cgf(2); + DynamicCG.set_active_index(2); // Should be ignored as DynParam1 not used in active node DynParam1.update(PtrA); ExecGraph.update(DynamicCGNode); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp index 925839729cce8..65d27070a1b0c 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp @@ -111,17 +111,17 @@ int main() { ExecuteGraphAndVerifyResults(0, UpdatedScalarValue, 0); // CGFB using PtrB in its dynamic parameter and immutable ScalarValue - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(0, ScalarValue, false); // CGFC using immutable PtrC and UpdatedScalarValue in its dynamic parameter - DynamicCG.set_active_cgf(2); + DynamicCG.set_active_index(2); ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(0, 0, UpdatedScalarValue); // CGFD using immutable PtrA and immutable ScalarValue for arguments - DynamicCG.set_active_cgf(3); + DynamicCG.set_active_index(3); ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(ScalarValue, 0, 0); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp index 6ee6dafaaea60..c6c639a83cf91 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp @@ -86,12 +86,12 @@ int main() { ExecuteGraphAndVerifyResults(false, true, false); // CGFB with DynParam using PtrB - DynamicCG.set_active_cgf(1); + DynamicCG.set_active_index(1); ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(false, true, false); // CGFC unconditionally using PtrC - DynamicCG.set_active_cgf(2); + DynamicCG.set_active_index(2); ExecGraph.update(DynamicCGNode); ExecuteGraphAndVerifyResults(false, false, true); diff --git a/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp b/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp index bf40f1baf7661..34e334b86861b 100644 --- a/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp +++ b/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp @@ -34,7 +34,7 @@ int main() { auto DynamicCGB = exp_ext::dynamic_command_group(GraphB, {CGFA, CGFB}); auto DynamicCGNodeB = GraphB.add(DynamicCGB); - DynamicCGB.set_active_cgf(1); // Check if doesn't affect GraphA + DynamicCGB.set_active_index(1); // Check if doesn't affect GraphA auto ExecGraph = GraphA.finalize(exp_ext::property::graph::updatable{}); @@ -57,7 +57,7 @@ int main() { // Both ExecGraph and Graph B have CGFB as active, so // whole graph update should be valid as graphs match. - DynamicCGA.set_active_cgf(1); + DynamicCGA.set_active_index(1); ExecGraph.update(DynamicCGNodeA); ExecGraph.update(GraphB); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 3e7673d95b236..e3316b2305641 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3026,7 +3026,7 @@ _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmRKNS3_16image_desc _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmRKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmmmjRKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmmmjRKNS0_6deviceERKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental21dynamic_command_group14set_active_cgfEm +_ZN4sycl3_V13ext6oneapi12experimental21dynamic_command_group16set_active_indexEm _ZN4sycl3_V13ext6oneapi12experimental21dynamic_command_groupC1ERKNS3_13command_graphILNS3_11graph_stateE0EEERKSt6vectorISt8functionIFvRNS0_7handlerEEESaISF_EE _ZN4sycl3_V13ext6oneapi12experimental21dynamic_command_groupC2ERKNS3_13command_graphILNS3_11graph_stateE0EEERKSt6vectorISt8functionIFvRNS0_7handlerEEESaISF_EE _ZN4sycl3_V13ext6oneapi12experimental21get_composite_devicesEv @@ -3245,6 +3245,10 @@ _ZN4sycl3_V16detail13lgamma_r_implEfPi _ZN4sycl3_V16detail13make_platformEmNS0_7backendE _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE +_ZN4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv +_ZN4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv +_ZN4sycl3_V16detail14SubmissionInfoC1Ev +_ZN4sycl3_V16detail14SubmissionInfoC2Ev _ZN4sycl3_V16detail14addCounterInitERNS0_7handlerERSt10shared_ptrINS1_10queue_implEERS4_IiE _ZN4sycl3_V16detail14getBorderColorENS0_19image_channel_orderE _ZN4sycl3_V16detail14tls_code_loc_t5queryEv @@ -3254,10 +3258,6 @@ _ZN4sycl3_V16detail14tls_code_loc_tC2ERKNS1_13code_locationE _ZN4sycl3_V16detail14tls_code_loc_tC2Ev _ZN4sycl3_V16detail14tls_code_loc_tD1Ev _ZN4sycl3_V16detail14tls_code_loc_tD2Ev -_ZN4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv -_ZN4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv -_ZN4sycl3_V16detail14SubmissionInfoC1Ev -_ZN4sycl3_V16detail14SubmissionInfoC2Ev _ZN4sycl3_V16detail16AccessorBaseHost10getAccDataEv _ZN4sycl3_V16detail16AccessorBaseHost14getAccessRangeEv _ZN4sycl3_V16detail16AccessorBaseHost14getMemoryRangeEv @@ -3612,7 +3612,7 @@ _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem10get_deviceEv _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem11get_contextEv _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem3mapEmmNS3_19address_access_modeEm _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem4sizeEv -_ZNK4sycl3_V13ext6oneapi12experimental21dynamic_command_group14get_active_cgfEv +_ZNK4sycl3_V13ext6oneapi12experimental21dynamic_command_group16get_active_indexEv _ZNK4sycl3_V13ext6oneapi12experimental4node14get_successorsEv _ZNK4sycl3_V13ext6oneapi12experimental4node16get_predecessorsEv _ZNK4sycl3_V13ext6oneapi12experimental4node8get_typeEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index a84e393e5b006..61ba627e864f8 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -331,13 +331,13 @@ ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBX@Z -?get_active_cgf@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ +?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ ??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z ?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEAVdynamic_command_group@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV?$command_graph@$0A@@12345@AEBV?$vector@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V?$allocator@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@@2@@std@@@Z ??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z ??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z -?set_active_cgf@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z +?set_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z ??1dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??0event@_V1@sycl@@AEAA@V?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@@Z