[SYCL][Graph] Add specification for kernel binary updates (#14896)

Adds the kernel binary update feature to the sycl graph specification.
This introduces a new dynamic_command_group class which can be used to
update the command-group function of a kernel nodes in graphs.

Implemented in:
* #16154
* #15700

---------

Co-authored-by: Ewan Crawford <[email protected]>
Co-authored-by: Ben Tracy <[email protected]>
Co-authored-by: Pablo Reble <[email protected]>

Author: 4 people

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -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<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;
+  });
+};
+
+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<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);
+
+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

sycl/doc/syclgraph/SYCLGraphUsageGuide.md

@@ -226,8 +226,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/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -2078,10 +2078,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/source/detail/graph_impl.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.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_deps.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_deps2.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_accessor_spv.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_different_arg_nums.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();