diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index a5596b6e96158..7c410bdc473f9 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -426,7 +426,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 @@ -434,15 +434,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. @@ -453,7 +450,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 @@ -461,15 +458,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. @@ -524,7 +518,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 @@ -557,6 +551,119 @@ 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_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, 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]. + +See <> for more information +about updating command-groups. + +===== Limitations + +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 +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. 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"] +|=== +|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 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 `cgfList` is empty. + +| +[source,c++] +---- +size_t get_active_index() const; +---- +|Returns the index of the currently active command-group function in this +`dynamic_command_group`. + +| +[source,c++] +---- +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 +`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++] @@ -632,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; @@ -712,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` @@ -740,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, @@ -763,6 +882,55 @@ 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. 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 + +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_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 +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 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 + +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 @@ -1040,6 +1208,50 @@ Exceptions: * Throws with error code `invalid` if the type of the command-group is not a kernel execution and a `dynamic_parameter` was registered inside `cgf`. +| +[source,c++] +---- +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. + +* 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++] ---- @@ -1158,8 +1370,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. @@ -1191,9 +1404,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. @@ -1750,6 +1964,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. @@ -2129,6 +2347,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. @@ -2183,8 +2411,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 f51d441174ef6..a2ca77258bdd8 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 @@ -453,6 +453,121 @@ 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 +... +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + +queue Queue{}; +sycl_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 = sycl_ext::dynamic_command_group(Graph, {CgfA, CgfB}); + +// Create a dynamic command-group graph node. +auto DynamicCGNode = Graph.add(DynamicCG); + +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_index(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(); +``` + +### 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 +... +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(); +sycl_ext::command_graph Graph{MyContext, MyDevice}; + +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 +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); +}; + +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 = sycl_ext::dynamic_command_group(Graph, {CgfA, CgfB}); + +// Create a dynamic command-group graph node. +auto DynamicCGNode = Graph.add(DynamicCG); + +auto ExecGraph = Graph.finalize(sycl_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_index(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 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