[SYCL][GRAPH] Add graph E2E tests which use free function extension (#16159)

Dynamic parameters in SYCL Graphs should be used for now with
free_function_kernels extension since then the order of kernel arguments
is well defined.

The tests are a copy of few selected ones from `Graph/Update` and
modified such that they use the extension. They live in
`Graph/Update/FreeFunctionKernels`.

Since the free function kernels are not yet implemented on CUDA, we
should keep for now both copies of the tests (with and without free
functions).

Author: Konrad Kusiak

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -58,6 +58,7 @@ Jack Kirk, Codeplay +
 Ronan Keryell, AMD +
 Andrey Alekseenko, KTH Royal Institute of Technology +
 Fábio Mestre, Codeplay +
+Konrad Kusiak, Codeplay +
 
 == Dependencies
 
@@ -1979,10 +1980,16 @@ can be used adding nodes to a graph when creating a graph from queue recording.
 New methods are also defined that enable submitting an executable graph, 
 e.g. directly to a queue without returning an event.
 
+==== sycl_ext_oneapi_free_function_kernels
+
+`sycl_ext_oneapi_free_function_kernels`, defined in 
+link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[sycl_ext_oneapi_free_function_kernels]
+can be used with SYCL Graphs. 
+
 == Examples and Usage Guide
 
 Detailed code examples and usage guidelines are provided in the
-link:../../SYCLGraphUsageGuide.md[SYCL Graph Usage Guide].
+link:../../syclgraph/SYCLGraphUsageGuide.md[SYCL Graph Usage Guide].
 
 == Future Direction [[future-direction]]
 

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,45 @@
+// 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);
+  }
+#endif
+  sycl::free(PtrA, Queue);
+
+  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"

sycl/test-e2e/Graph/Update/FreeFunctionKernels/free_function_kernels.hpp

@@ -0,0 +1,57 @@
+#pragma once
+
+#include "../../graph_common.hpp"
+#include "sycl/ext/oneapi/kernel_properties/properties.hpp"
+#include "sycl/kernel_bundle.hpp"
+#include <sycl/ext/oneapi/free_function_queries.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;
+  }
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::single_task_kernel))
+void ff_1(int *Ptr) {
+  for (size_t i{0}; i < Size; ++i) {
+    Ptr[i] += i;
+  }
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::single_task_kernel))
+void ff_2(int *Ptr, size_t Size, size_t NumKernelLoops) {
+  for (size_t j{0}; j < NumKernelLoops; j++) {
+    for (size_t i{0}; i < Size; i++) {
+      Ptr[i] += i;
+    }
+  }
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::nd_range_kernel<3>))
+void ff_3(int *Ptr) {
+  size_t GlobalID =
+      ext::oneapi::this_work_item::get_nd_item<3>().get_global_linear_id();
+  Ptr[GlobalID] = GlobalID;
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::nd_range_kernel<3>))
+void ff_4(int *Ptr) {
+  size_t GlobalID =
+      ext::oneapi::this_work_item::get_nd_item<3>().get_global_linear_id();
+  Ptr[GlobalID] *= 2;
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::nd_range_kernel<1>))
+void ff_5(int *PtrA, int *PtrB, int *PtrC) {
+  size_t GlobalID =
+      ext::oneapi::this_work_item::get_nd_item<1>().get_global_id();
+  PtrC[GlobalID] += PtrA[GlobalID] * PtrB[GlobalID];
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::single_task_kernel))
+void ff_6(int *Ptr, int ScalarValue) {
+  for (size_t i{0}; i < Size; ++i) {
+    Ptr[i] = ScalarValue;
+  }
+}

sycl/test-e2e/Graph/Update/FreeFunctionKernels/update_before_finalize.cpp

@@ -0,0 +1,60 @@
+// 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
+
+// Tests updating a graph node before finalization
+
+#include "../../graph_common.hpp"
+#include "free_function_kernels.hpp"
+
+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);
+  int *PtrB = malloc_device<int>(Size, Queue);
+
+  std::vector<int> HostDataA(Size);
+  std::vector<int> HostDataB(Size);
+
+  Queue.memset(PtrA, 0, Size * sizeof(int)).wait();
+  Queue.memset(PtrB, 0, Size * sizeof(int)).wait();
+
+  exp_ext::dynamic_parameter InputParam(Graph, PtrA);
+
+#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, InputParam);
+    cgh.single_task(Kernel);
+  });
+  // Swap PtrB to be the input
+  InputParam.update(PtrB);
+
+  auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // Only PtrB should be filled with values
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  Queue.copy(PtrA, HostDataA.data(), Size).wait();
+  Queue.copy(PtrB, HostDataB.data(), Size).wait();
+  for (size_t i = 0; i < Size; i++) {
+    assert(HostDataA[i] == 0);
+    assert(HostDataB[i] == i);
+  }
+#endif
+  sycl::free(PtrA, Queue);
+  sycl::free(PtrB, Queue);
+
+  return 0;
+}

sycl/test-e2e/Graph/Update/FreeFunctionKernels/update_with_indices_multiple_exec_graphs.cpp

@@ -0,0 +1,77 @@
+// 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
+
+// Tests creating multiple executable graphs from the same modifiable graph and
+// only updating one of them.
+
+#include "../../graph_common.hpp"
+#include "free_function_kernels.hpp"
+
+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);
+  int *PtrB = malloc_device<int>(Size, Queue);
+
+  std::vector<int> HostDataA(Size);
+  std::vector<int> HostDataB(Size);
+
+  Queue.memset(PtrA, 0, Size * sizeof(int)).wait();
+  Queue.memset(PtrB, 0, Size * sizeof(int)).wait();
+
+  exp_ext::dynamic_parameter InputParam(Graph, PtrA);
+
+#ifndef __SYCL_DEVICE_ONLY__
+  kernel_bundle Bundle = get_kernel_bundle<bundle_state::executable>(Ctxt);
+  kernel_id Kernel_id = exp_ext::get_kernel_id<ff_1>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, InputParam);
+    cgh.single_task(Kernel);
+  });
+
+  auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
+  auto ExecGraph2 = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // PtrA values should be modified twice
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+  Queue.ext_oneapi_graph(ExecGraph2).wait();
+
+  Queue.copy(PtrA, HostDataA.data(), Size).wait();
+  Queue.copy(PtrB, HostDataB.data(), Size).wait();
+  for (size_t i = 0; i < Size; i++) {
+    assert(HostDataA[i] == i * 2);
+    assert(HostDataB[i] == 0);
+  }
+
+  // Swap PtrB to be the input
+  InputParam.update(PtrB);
+  // Only update ExecGraph, which should now modify PtrB while ExecGraph2
+  // modifies PtrA still
+  ExecGraph.update(KernelNode);
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+  Queue.ext_oneapi_graph(ExecGraph2).wait();
+
+  Queue.copy(PtrA, HostDataA.data(), Size).wait();
+  Queue.copy(PtrB, HostDataB.data(), Size).wait();
+  for (size_t i = 0; i < Size; i++) {
+    // A should have been modified 3 times by now, B only once
+    assert(HostDataA[i] == i * 3);
+    assert(HostDataB[i] == i);
+  }
+#endif
+  sycl::free(PtrA, Queue);
+  sycl::free(PtrB, Queue);
+
+  return 0;
+}

sycl/test-e2e/Graph/Update/FreeFunctionKernels/update_with_indices_ordering.cpp

@@ -0,0 +1,80 @@
+// 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
+
+// Tests that updating a graph is ordered with respect to previous executions of
+// the graph which may be in flight.
+
+#include "../../graph_common.hpp"
+#include "free_function_kernels.hpp"
+
+int main() {
+  queue Queue{};
+  context Ctxt{Queue.get_context()};
+
+  // Use a large N to try and make the kernel slow
+  const size_t N = 1 << 16;
+  // Loop inside kernel to make even slower (too large N runs out of memory)
+  const size_t NumKernelLoops = 4;
+  const size_t NumSubmitLoops = 8;
+
+  exp_ext::command_graph Graph{Ctxt, Queue.get_device()};
+
+  int *PtrA = malloc_device<int>(N, Queue);
+  int *PtrB = malloc_device<int>(N, Queue);
+
+  std::vector<int> HostDataA(N);
+  std::vector<int> HostDataB(N);
+
+  Queue.memset(PtrA, 0, N * sizeof(int)).wait();
+  Queue.memset(PtrB, 0, N * sizeof(int)).wait();
+
+  exp_ext::dynamic_parameter InputParam(Graph, PtrA);
+
+#ifndef __SYCL_DEVICE_ONLY__
+  kernel_bundle Bundle = get_kernel_bundle<bundle_state::executable>(Ctxt);
+  kernel_id Kernel_id = exp_ext::get_kernel_id<ff_2>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, InputParam);
+    cgh.set_arg(1, N);
+    cgh.set_arg(2, NumKernelLoops);
+    cgh.single_task(Kernel);
+  });
+
+  auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // Submit a bunch of graphs without waiting
+  for (size_t i = 0; i < NumSubmitLoops; i++) {
+    Queue.ext_oneapi_graph(ExecGraph);
+  }
+
+  // Swap PtrB to be the input
+  InputParam.update(PtrB);
+
+  ExecGraph.update(KernelNode);
+
+  // Submit another set of graphs then wait on all submissions
+  for (size_t i = 0; i < NumSubmitLoops; i++) {
+    Queue.ext_oneapi_graph(ExecGraph);
+  }
+  Queue.wait_and_throw();
+
+  Queue.copy(PtrA, HostDataA.data(), N).wait();
+  Queue.copy(PtrB, HostDataB.data(), N).wait();
+  for (size_t i = 0; i < N; i++) {
+    assert(HostDataA[i] == i * NumKernelLoops * NumSubmitLoops);
+    assert(HostDataB[i] == i * NumKernelLoops * NumSubmitLoops);
+  }
+#endif
+  sycl::free(PtrA, Queue);
+  sycl::free(PtrB, Queue);
+
+  return 0;
+}