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