[SYCL][Graph] Test WGU kernel mismatch (#14379)

We cannot currently update the kernel binary of a node in Whole Graph
Update. Rather than silently accepting inconsistent kernel functions,
which indicates the graphs aren't topologically identical, throw an
error when the kernel types of two nodes are mismatched.

This change requires removing the unittest for barrier nodes in Whole
Graph Update as the mock infrastructure does not
setup the internal `CG` class to the depth required to test working
functionality. This
functionality is already covered by
`test-e2e/Graph/Update/whole_update_barrier_node.cpp`

Author: Ewan Crawford

sycl/source/detail/graph_impl.cpp

@@ -1184,23 +1184,45 @@ void exec_graph_impl::update(std::shared_ptr<graph_impl> GraphImpl) {
     throw sycl::exception(sycl::make_error_code(errc::invalid),
                           "Cannot update using a graph with a different "
                           "topology. Mismatch found in the number of nodes.");
-  } else {
-    for (uint32_t i = 0; i < MNodeStorage.size(); ++i) {
-      if (MNodeStorage[i]->MSuccessors.size() !=
-              GraphImpl->MNodeStorage[i]->MSuccessors.size() ||
-          MNodeStorage[i]->MPredecessors.size() !=
-              GraphImpl->MNodeStorage[i]->MPredecessors.size()) {
-        throw sycl::exception(
-            sycl::make_error_code(errc::invalid),
-            "Cannot update using a graph with a different topology. Mismatch "
-            "found in the number of edges.");
-      }
+  }
+
+  for (uint32_t i = 0; i < MNodeStorage.size(); ++i) {
+    if (MNodeStorage[i]->MSuccessors.size() !=
+            GraphImpl->MNodeStorage[i]->MSuccessors.size() ||
+        MNodeStorage[i]->MPredecessors.size() !=
+            GraphImpl->MNodeStorage[i]->MPredecessors.size()) {
+      throw sycl::exception(
+          sycl::make_error_code(errc::invalid),
+          "Cannot update using a graph with a different topology. Mismatch "
+          "found in the number of edges.");
+    }
+    if (MNodeStorage[i]->MCGType != GraphImpl->MNodeStorage[i]->MCGType) {
+      throw sycl::exception(
+          sycl::make_error_code(errc::invalid),
+          "Cannot update using a graph with mismatched node types. Each pair "
+          "of nodes being updated must have the same type");
+    }
 
-      if (MNodeStorage[i]->MCGType != GraphImpl->MNodeStorage[i]->MCGType) {
-        throw sycl::exception(
-            sycl::make_error_code(errc::invalid),
-            "Cannot update using a graph with mismatched node types. Each pair "
-            "of nodes being updated must have the same type");
+    if (MNodeStorage[i]->MCGType == sycl::detail::CG::Kernel) {
+      sycl::detail::CGExecKernel *TargetCGExec =
+          static_cast<sycl::detail::CGExecKernel *>(
+              MNodeStorage[i]->MCommandGroup.get());
+      const std::string &TargetKernelName = TargetCGExec->getKernelName();
+
+      sycl::detail::CGExecKernel *SourceCGExec =
+          static_cast<sycl::detail::CGExecKernel *>(
+              GraphImpl->MNodeStorage[i]->MCommandGroup.get());
+      const std::string &SourceKernelName = SourceCGExec->getKernelName();
+
+      if (TargetKernelName.compare(SourceKernelName) != 0) {
+        std::stringstream ErrorStream(
+            "Cannot update using a graph with mismatched kernel "
+            "types. Source node type ");
+        ErrorStream << SourceKernelName;
+        ErrorStream << ", target node type ";
+        ErrorStream << TargetKernelName;
+        throw sycl::exception(sycl::make_error_code(errc::invalid),
+                              ErrorStream.str());
       }
     }
   }

sycl/test-e2e/Graph/Update/whole_update_dynamic_param.cpp

@@ -36,13 +36,15 @@ int main() {
   exp_ext::command_graph GraphA{Queue.get_context(), Queue.get_device()};
 
   exp_ext::dynamic_parameter InputParam(GraphA, InputDataDevice1);
+  auto KernelLambda = [=]() {
+    for (size_t i = 0; i < Size; i++) {
+      OutputDataDevice1[i] = InputDataDevice1[i];
+    }
+  };
+
   GraphA.add([&](handler &CGH) {
     CGH.set_arg(1, InputParam);
-    CGH.single_task([=]() {
-      for (size_t i = 0; i < Size; i++) {
-        OutputDataDevice1[i] = InputDataDevice1[i];
-      }
-    });
+    CGH.single_task(KernelLambda);
   });
 
   auto GraphExecA = GraphA.finalize();
@@ -59,13 +61,7 @@ int main() {
   InputParam.update(InputDataDevice2);
   exp_ext::command_graph GraphB{Queue.get_context(), Queue.get_device()};
 
-  GraphB.add([&](handler &CGH) {
-    CGH.single_task([=]() {
-      for (size_t i = 0; i < Size; i++) {
-        OutputDataDevice1[i] = InputDataDevice1[i];
-      }
-    });
-  });
+  GraphB.add([&](handler &CGH) { CGH.single_task(KernelLambda); });
 
   auto GraphExecB = GraphB.finalize(exp_ext::property::graph::updatable{});
   GraphExecB.update(GraphA);

sycl/test-e2e/Graph/Update/whole_update_kernel_type_mismatch.cpp

@@ -0,0 +1,103 @@
+// 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 %}
+
+// Test that an error is thrown when the types of kernels do not match in Whole
+// Graph Update
+
+#include "../graph_common.hpp"
+
+void testFunctors(queue Queue, int *Data) {
+  exp_ext::command_graph Graph{Queue};
+  exp_ext::command_graph UpdateGraph{Queue};
+  struct KernelFunctorA {
+    KernelFunctorA(int *Data) : Data(Data) {}
+
+    void operator()() const { Data[0] = 42; }
+
+    int *Data;
+  };
+
+  struct KernelFunctorB {
+    KernelFunctorB(int *Data) : Data(Data) {}
+    void operator()() const { Data[0] = 42; }
+
+    int *Data;
+  };
+
+  Graph.add([&](handler &CGH) { CGH.single_task(KernelFunctorA{Data}); });
+
+  UpdateGraph.add([&](handler &CGH) { CGH.single_task(KernelFunctorB{Data}); });
+
+  auto GraphExec = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // Check it's an error if kernel types don't match
+  std::error_code ErrorCode = make_error_code(sycl::errc::success);
+  try {
+    GraphExec.update(UpdateGraph);
+  } catch (const sycl::exception &e) {
+    ErrorCode = e.code();
+  }
+  assert(ErrorCode == sycl::errc::invalid);
+}
+
+void testUnNamedLambdas(queue Queue, int *Data) {
+  exp_ext::command_graph Graph{Queue};
+  exp_ext::command_graph UpdateGraph{Queue};
+
+  Graph.add([&](handler &CGH) { CGH.single_task([=]() { Data[0] = 42; }); });
+
+  UpdateGraph.add(
+      [&](handler &CGH) { CGH.single_task([=]() { Data[0] = 42; }); });
+
+  auto GraphExec = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // Check it's an error if kernel types don't match
+  std::error_code ErrorCode = make_error_code(sycl::errc::success);
+  try {
+    GraphExec.update(UpdateGraph);
+  } catch (const sycl::exception &e) {
+    ErrorCode = e.code();
+  }
+  assert(ErrorCode == sycl::errc::invalid);
+}
+void testNamedLambdas(queue Queue, int *Data) {
+  exp_ext::command_graph Graph{Queue};
+  exp_ext::command_graph UpdateGraph{Queue};
+
+  auto LambdaA = [=]() { Data[0] = 42; };
+
+  Graph.add([&](handler &CGH) { CGH.single_task<class TestLambdaA>(LambdaA); });
+
+  auto LambdaB = [=]() { Data[0] = 42; };
+
+  UpdateGraph.add(
+      [&](handler &CGH) { CGH.single_task<class TestLambdaB>(LambdaB); });
+
+  auto GraphExec = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // Check it's an error if kernel types don't match
+  std::error_code ErrorCode = make_error_code(sycl::errc::success);
+  try {
+    GraphExec.update(UpdateGraph);
+  } catch (const sycl::exception &e) {
+    ErrorCode = e.code();
+  }
+  assert(ErrorCode == sycl::errc::invalid);
+}
+
+int main() {
+  queue Queue{};
+  int *Data = malloc_device<int>(1, Queue);
+
+  testNamedLambdas(Queue, Data);
+  testUnNamedLambdas(Queue, Data);
+  testFunctors(Queue, Data);
+
+  sycl::free(Data, Queue);
+
+  return 0;
+}

sycl/unittests/Extensions/CommandGraph/Update.cpp

@@ -399,21 +399,3 @@ TEST_F(WholeGraphUpdateTest, EmptyNode) {
   auto GraphExec = Graph.finalize(experimental::property::graph::updatable{});
   GraphExec.update(UpdateGraph);
 }
-
-TEST_F(WholeGraphUpdateTest, BarrierNode) {
-  // Test that updating a graph that has a barrier node is not an error
-  Graph.begin_recording(Queue);
-  auto NodeKernel = Queue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
-  Queue.ext_oneapi_submit_barrier({NodeKernel});
-  Graph.end_recording(Queue);
-
-  UpdateGraph.begin_recording(Queue);
-  auto UpdateNodeKernel = Queue.submit(
-      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
-  Queue.ext_oneapi_submit_barrier({UpdateNodeKernel});
-  UpdateGraph.end_recording(Queue);
-
-  auto GraphExec = Graph.finalize(experimental::property::graph::updatable{});
-  GraphExec.update(UpdateGraph);
-}