Added Graph free function kernels test without update

Author: Konrad Kusiak

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

@@ -0,0 +1,13 @@
+// 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 %}
+//
+// XFAIL: cuda
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16004
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/free_function_kernels.cpp"

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

@@ -0,0 +1,44 @@
+// Tests compatibility with free function kernels extension
+
+#include "../graph_common.hpp"
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::single_task_kernel))
+void ff_0(int *ptr) {
+  for (size_t i{0}; i < Size; ++i) {
+    ptr[i] = i;
+  }
+}
+
+int main() {
+  queue Queue{};
+  context ctxt{Queue.get_context()};
+
+  exp_ext::command_graph Graph{ctxt, Queue.get_device()};
+
+  int *PtrA = malloc_device<int>(Size, Queue);
+
+  std::vector<int> HostDataA(Size);
+
+  Queue.memset(PtrA, 0, Size * sizeof(int)).wait();
+
+#ifndef __SYCL_DEVICE_ONLY__
+  kernel_bundle Bundle = get_kernel_bundle<bundle_state::executable>(ctxt);
+  kernel_id Kernel_id = exp_ext::get_kernel_id<ff_0>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, PtrA);
+    cgh.single_task(Kernel);
+  });
+
+  auto ExecGraph = Graph.finalize();
+
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  Queue.copy(PtrA, HostDataA.data(), Size).wait();
+  for (size_t i = 0; i < Size; i++) {
+    assert(HostDataA[i] == i);
+  }
+  sycl::free(PtrA, Queue);
+#endif
+  return 0;
+}

sycl/test-e2e/Graph/RecordReplay/free_function_kernels.cpp

@@ -0,0 +1,13 @@
+// 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 %}
+//
+// XFAIL: cuda
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16004
+
+#define GRAPH_E2E_RECORD_REPLAY
+
+#include "../Inputs/free_function_kernels.cpp"