[SYCL][Graph] Add test for using spirv kernels in graphs (#13532)

Adds new tests to check that the sycl_ext_oneapi_kernel_compiler_spirv
extension is compatible with sycl graphs.

Author: fabiomestre

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -1799,6 +1799,12 @@ code `invalid` if a user tries to add them to a graph.
 Removing this restriction is something we may look at for future revisions of
 `sycl_ext_oneapi_graph`.
 
+==== sycl_ext_oneapi_kernel_compiler_spirv
+
+The kernels loaded using
+link:../experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc[sycl_ext_oneapi_kernel_compiler_spirv]
+behave as normal when used in graph nodes.
+
 == Examples
 
 [NOTE]

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

@@ -0,0 +1,12 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out %S/../Inputs/Kernels/kernels.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/kernels.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// Extra run to check for immediate-command-list in Level Zero
+// RUN: %if level_zero && linux %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/kernels.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+// REQUIRES: level_zero
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/kernel_bundle_spirv.cpp"

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

@@ -0,0 +1,98 @@
+// Tests creating a node using a SPIR-V kernel imported with
+// sycl_ext_oneapi_kernel_compiler_spirv. The SPIR-V kernels used in this test
+// 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};
+
+  queue Queue{Ctx, Dev};
+
+  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);
+      };
+
+  sycl::kernel kernel = getKernel(KernelBundle, "my_kernel");
+  assert(kernel.get_backend() == backend::ext_oneapi_level_zero);
+
+  constexpr int N = 4;
+  std::array<int, N> input_array{0, 1, 2, 3};
+  std::array<int, N> output_array{};
+  std::array<int, N> output_array2{};
+
+  sycl::buffer input_buffer(input_array.data(), sycl::range<1>(N));
+  sycl::buffer output_buffer(output_array.data(), sycl::range<1>(N));
+  sycl::buffer output_buffer2(output_array2.data(), sycl::range<1>(N));
+
+  input_buffer.set_write_back(false);
+  output_buffer.set_write_back(false);
+  output_buffer2.set_write_back(false);
+
+  {
+    exp_ext::command_graph Graph{
+        Queue.get_context(),
+        Queue.get_device(),
+        {exp_ext::property::graph::assume_buffer_outlives_graph{}}};
+
+    add_node(Graph, Queue, ([&](sycl::handler &CGH) {
+               CGH.set_arg(
+                   0, input_buffer.get_access<sycl::access::mode::read>(CGH));
+               CGH.set_arg(
+                   1, output_buffer.get_access<sycl::access::mode::write>(CGH));
+               CGH.parallel_for(sycl::range<1>{N}, kernel);
+             }));
+
+    add_node(Graph, Queue, ([&](sycl::handler &CGH) {
+               CGH.set_arg(
+                   0, input_buffer.get_access<sycl::access::mode::read>(CGH));
+               CGH.set_arg(
+                   1,
+                   output_buffer2.get_access<sycl::access::mode::write>(CGH));
+               CGH.parallel_for(sycl::range<1>{N}, kernel);
+             }));
+
+    auto GraphExec = Graph.finalize();
+
+    Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); });
+    Queue.wait_and_throw();
+  }
+
+  host_accessor HostAccOutput(output_buffer);
+  host_accessor HostAccOutput2(output_buffer2);
+
+  for (int i = 0; i < N; i++) {
+    assert(HostAccOutput[i] == ((i * 2) + 100));
+    assert(HostAccOutput2[i] == ((i * 2) + 100));
+  }
+
+  return 0;
+}

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

@@ -0,0 +1,12 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out %S/../Inputs/Kernels/kernels.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/kernels.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// Extra run to check for immediate-command-list in Level Zero
+// RUN: %if level_zero && linux %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/kernels.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+// REQUIRES: level_zero
+
+#define GRAPH_E2E_RECORD_REPLAY
+
+#include "../Inputs/kernel_bundle_spirv.cpp"