[SYCL][Graph] Implement dynamic command-groups

Implement Dynamic Command-Group feature specified in
PR [[SYCL][Graph] Add specification for kernel binary updates](intel#14896)

This feature enables updating `ur_kernel_handle_t` objects in graph nodes
between executions as well as parameters and execution range of nodes.

This functionality is currently supported on CUDA & HIP which are used
for testing in the new E2E tests. Level Zero support will follow
shortly, resulting in the removal of the `XFAIL` labels from the E2E
tests.

The code for adding nodes to a graph has been refactored to split out
verification of edges, and marking memory objects used in a node, as
separate helper functions. This allows path for adding a command-group
node to do this functions over each CG in the list before creating the
node itself.

The `dynamic_parameter_impl` code has also been refactored so the code
is shared for updating a dynamic parameter used in both a regular kernel
node and a dynamic command-group node.

See the addition to the design doc for further details on the
implementation.

sycl/doc/design/CommandGraph.md

@@ -282,6 +282,23 @@ requirements for these new accessors to correctly trigger allocations before
 updating. This is similar to how individual graph commands are enqueued when
 accessors are used in a graph node.
 
+### Dynamic Command-Group
+
+To implement the `dynamic_command_group` class for updating the command-groups (CG)
+associated with nodes, the CG member of the node implementation class changes
+from a `std::unique_ptr` to a `std::shared_ptr` so that multiple nodes and the
+`dynamic_command_group_impl` object can share the same CG object. This avoids
+the overhead of having to allocate and free copies of the CG when a new active
+CG is selected.
+
+The `dynamic_command_group_impl` class contains weak pointers to the nodes which
+have been created with it, so that when a new active CG is selected it can
+propagate the change to those nodes. The `node_impl` class also contains a
+reference to the dynamic command-group that created it, so that when the graph
+is finalized each node can use the list of kernels in its dynamic command-group
+as part of the `urCommandBufferAppendKernelLaunchExp` call to pass the possible
+alternative kernels.
+
 ## Optimizations
 ### Interactions with Profiling
 

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

@@ -93,6 +93,7 @@ class node_impl;
 class graph_impl;
 class exec_graph_impl;
 class dynamic_parameter_impl;
+class dynamic_command_group_impl;
 } // namespace detail
 
 enum class node_type {
@@ -213,6 +214,23 @@ class depends_on_all_leaves : public ::sycl::detail::DataLessProperty<
 } // namespace node
 } // namespace property
 
+class __SYCL_EXPORT dynamic_command_group {
+public:
+  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);
+
+private:
+  template <class Obj>
+  friend const decltype(Obj::impl) &
+  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
+
+  std::shared_ptr<detail::dynamic_command_group_impl> impl;
+};
+
 namespace detail {
 // Templateless modifiable command-graph base class.
 class __SYCL_EXPORT modifiable_command_graph {
@@ -269,6 +287,28 @@ class __SYCL_EXPORT modifiable_command_graph {
     return Node;
   }
 
+  /// Add a Dynamic command-group node to the graph.
+  /// @param DynamicCG Dynamic command-group function to create node with.
+  /// @param PropList Property list used to pass [0..n] predecessor nodes.
+  /// @return Constructed node which has been added to the graph.
+  node add(dynamic_command_group &DynamicCG,
+           const property_list &PropList = {}) {
+    if (PropList.has_property<property::node::depends_on>()) {
+      auto Deps = PropList.get_property<property::node::depends_on>();
+      node Node = addImpl(DynamicCG, Deps.get_dependencies());
+      if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+        addGraphLeafDependencies(Node);
+      }
+      return Node;
+    }
+
+    node Node = addImpl(DynamicCG, {});
+    if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+      addGraphLeafDependencies(Node);
+    }
+    return Node;
+  }
+
   /// Add a dependency between two nodes.
   /// @param Src Node which will be a dependency of \p Dest.
   /// @param Dest Node which will be dependent on \p Src.
@@ -328,6 +368,12 @@ class __SYCL_EXPORT modifiable_command_graph {
   modifiable_command_graph(const std::shared_ptr<detail::graph_impl> &Impl)
       : impl(Impl) {}
 
+  /// Template-less implementation of add() for dynamic command-group nodes.
+  /// @param DynCGF Dynamic Command-group function object to add.
+  /// @param Dep List of predecessor nodes.
+  /// @return Node added to the graph.
+  node addImpl(dynamic_command_group &DynCGF, const std::vector<node> &Dep);
+
   /// Template-less implementation of add() for CGF nodes.
   /// @param CGF Command-group function to add.
   /// @param Dep List of predecessor nodes.

sycl/include/sycl/handler.hpp

@@ -3376,6 +3376,7 @@ class __SYCL_EXPORT handler {
                                  size_t Size, bool Block = false);
   friend class ext::oneapi::experimental::detail::graph_impl;
   friend class ext::oneapi::experimental::detail::dynamic_parameter_impl;
+  friend class ext::oneapi::experimental::detail::dynamic_command_group_impl;
 
   bool DisableRangeRounding();