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

sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.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_get_active_index.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();

sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.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();

sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.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();

sycl/test-e2e/Graph/Update/dyn_cgf_overwrite_range.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();

sycl/test-e2e/Graph/Update/dyn_cgf_parameters.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);
 

sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.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();

sycl/test-e2e/Graph/Update/dyn_cgf_update_before_finalize.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();