Update get/setters to "index" rather than "cgf"

Also refine spec wording based on PR feedback

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -562,17 +562,20 @@ public:
       command_graph<graph_state::modifiable> &graph,
       const std::vector<std::function<void(handler &)>>& 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,15 +632,15 @@ 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`.
 
 |
 [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,20 +906,22 @@ 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
 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
 

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<int>(1024, Queue);
-int *PtrB = malloc_device<int>(1024, Queue)​;
+int *PtrB = malloc_device<int>(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<int>(n, Queue);
-int *PtrB = malloc_device<int>(n, Queue)​;
+int *PtrA = malloc_device<int>(N, Queue);
+int *PtrB = malloc_device<int>(N, Queue);
 
 // Kernels loaded from kernel bundle
-const std::vector<kernel_id> builtinKernelIds =
-      myDevice.get_info<info::device::built_in_kernel_ids>();
-kernel_bundle<bundle_state::executable> myBundle =
-      get_kernel_bundle<sycl::bundle_state::executable>(myContext, { myDevice }, builtinKernelIds);
+const std::vector<kernel_id> BuiltinKernelIds =
+      MyDevice.get_info<info::device::built_in_kernel_ids>();
+kernel_bundle<bundle_state::executable> MyBundle =
+      get_kernel_bundle<sycl::bundle_state::executable>(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.

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -223,8 +223,8 @@ class __SYCL_EXPORT dynamic_command_group {
       const command_graph<graph_state::modifiable> &Graph,
       const std::vector<std::function<void(handler &)>> &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 <class Obj>

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

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();
 

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();

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();

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();

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();

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();

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<int>(Size, Queue);
+  std::vector<int> 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;
+}

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();