[SYCL][Graph] E2E test updating a SPIR-V Kernel (#13890)

Add an Graph E2E test which checks for updating an accessor argument to
a kernel created from SPIR-V.

`update_with_indices_accessor.spv` was created by using
`-fsycl-dump-device-code` on existing E2E test
`update_with_indices_accessor.cpp`, and can be dissembled using
`spirv-dis` if anyone is interested in looking at it further.

Author: EwanC

sycl/test-e2e/Graph/Inputs/Kernels/update_with_indices_accessor.spv

@@ -3,32 +3,8 @@
 // are identical to the ones used in KernelCompiler/Kernels/kernels.spv
 
 #include "../graph_common.hpp"
-#include <fstream>
-
-sycl::kernel_bundle<sycl::bundle_state::executable>
-loadKernelsFromFile(sycl::queue &Q, std::string FileName) {
-
-  // Read the SPIR-V module from disk.
-  std::ifstream SpvStream(FileName, std::ios::binary);
-  SpvStream.seekg(0, std::ios::end);
-  size_t sz = SpvStream.tellg();
-  SpvStream.seekg(0);
-  std::vector<std::byte> Spv(sz);
-  SpvStream.read(reinterpret_cast<char *>(Spv.data()), sz);
-
-  // Create a kernel bundle from the binary SPIR-V.
-  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> KernelBundleSrc =
-      exp_ext::create_kernel_bundle_from_source(
-          Q.get_context(), exp_ext::source_language::spirv, Spv);
-
-  // Build the SPIR-V module for our device.
-  sycl::kernel_bundle<sycl::bundle_state::executable> KernelBundleExe =
-      exp_ext::build(KernelBundleSrc);
-  return KernelBundleExe;
-}
 
 int main(int, char **argv) {
-
   const sycl::device Dev{sycl::default_selector_v};
   const sycl::context Ctx{Dev};
 

sycl/test-e2e/Graph/Inputs/kernel_bundle_spirv.cpp

@@ -16,7 +16,9 @@ option `config.required_features` to the graph aspect required to run the tests:
 
 Most of the tests are written in a similar manner to other `e2e` tests. The
 exception to this are tests in the `Inputs` directory which are meant to be used
-as inputs to other tests.
+as inputs to other tests. The `Kernels` subdirectory of `Inputs` contains SPIR-V
+kernels for testing. These can be generated from SYCL kernels by using the
+`-fsycl-dump-device-code=<dir>` option to the DPC++ compiler.
 
 Often, the same feature, needs to be tested for both the `Explicit`
 and `Record and Replay` APIs. To avoid code duplication, such tests are added to
@@ -29,7 +31,7 @@ common code.
 
 The other directories are used to group similar tests together. Tests that
 require a specific `aspect` are also grouped together in order to use the
-`lit.local.cfg` file. Directories might themselves contain subdirectories named 
+`lit.local.cfg` file. Directories might themselves contain subdirectories named
 `Explicit` and `RecordReplay` if they make use of the framework described above.
 
 In addition, in order to help identify specific tests, the matching files

sycl/test-e2e/Graph/README.md

@@ -0,0 +1,74 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv
+// 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 %S/../Inputs/Kernels/update_with_indices_accessor.spv 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 %S/../Inputs/Kernels/update_with_indices_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+// REQUIRES: level_zero
+
+// Tests updating an accessor argument to a graph node created from SPIR-V
+// using index-based explicit update
+
+#include "../graph_common.hpp"
+
+int main(int, char **argv) {
+  queue Queue{};
+  sycl::kernel_bundle KernelBundle = loadKernelsFromFile(Queue, argv[1]);
+  const auto getKernel =
+      [](sycl::kernel_bundle<sycl::bundle_state::executable> &bundle,
+         const std::string &name) {
+        return bundle.ext_oneapi_get_kernel(name);
+      };
+
+  kernel kernel = getKernel(
+      KernelBundle, "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_");
+
+  const size_t N = 1024;
+
+  exp_ext::command_graph Graph{
+      Queue.get_context(),
+      Queue.get_device(),
+      {exp_ext::property::graph::assume_buffer_outlives_graph{}}};
+  std::vector<int> HostDataA(N, 0);
+  std::vector<int> HostDataB(N, 0);
+
+  buffer BufA{HostDataA};
+  buffer BufB{HostDataB};
+  BufA.set_write_back(false);
+  BufB.set_write_back(false);
+  // Initial accessor for use in kernel and dynamic parameter
+  auto Acc = BufA.get_access();
+  exp_ext::dynamic_parameter InputParam(Graph, Acc);
+
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.require(InputParam);
+    cgh.set_arg(0, InputParam);
+    cgh.single_task(kernel);
+  });
+
+  auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // BufA should be filled with values
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  Queue.copy(BufA.get_access(), HostDataA.data()).wait();
+  Queue.copy(BufB.get_access(), HostDataB.data()).wait();
+  for (size_t i = 0; i < N; i++) {
+    assert(HostDataA[i] == i);
+    assert(HostDataB[i] == 0);
+  }
+
+  // Swap BufB to be the input
+  InputParam.update(BufB.get_access());
+  ExecGraph.update(KernelNode);
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  Queue.copy(BufA.get_access(), HostDataA.data()).wait();
+  Queue.copy(BufB.get_access(), HostDataB.data()).wait();
+  for (size_t i = 0; i < N; i++) {
+    assert(HostDataA[i] == i);
+    assert(HostDataB[i] == i);
+  }
+  return 0;
+}

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

@@ -4,7 +4,8 @@
 #include <sycl/ext/oneapi/experimental/graph.hpp>
 
 #include <condition_variable> // std::conditional_variable
-#include <mutex>              // std::mutex, std::unique_lock
+#include <fstream>
+#include <mutex> // std::mutex, std::unique_lock
 #include <numeric>
 
 // Test constants.
@@ -458,3 +459,24 @@ bool inline check_value(const size_t index, const T &Ref, const T &Got,
 
   return true;
 }
+
+kernel_bundle<bundle_state::executable>
+loadKernelsFromFile(queue &Q, std::string FileName) {
+  // Read the SPIR-V module from disk.
+  std::ifstream SpvStream(FileName, std::ios::binary);
+  SpvStream.seekg(0, std::ios::end);
+  size_t sz = SpvStream.tellg();
+  SpvStream.seekg(0);
+  std::vector<std::byte> Spv(sz);
+  SpvStream.read(reinterpret_cast<char *>(Spv.data()), sz);
+
+  // Create a kernel bundle from the binary SPIR-V.
+  kernel_bundle<bundle_state::ext_oneapi_source> KernelBundleSrc =
+      exp_ext::create_kernel_bundle_from_source(
+          Q.get_context(), exp_ext::source_language::spirv, Spv);
+
+  // Build the SPIR-V module for our device.
+  kernel_bundle<bundle_state::executable> KernelBundleExe =
+      exp_ext::build(KernelBundleSrc);
+  return KernelBundleExe;
+}