Added graph tests which use free function extension

Author: Konrad Kusiak

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

@@ -0,0 +1,56 @@
+#include "sycl/ext/oneapi/kernel_properties/properties.hpp"
+#include "sycl/kernel_bundle.hpp"
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+
+namespace exp_ext = sycl::ext::oneapi::experimental;
+using namespace sycl;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::single_task_kernel))
+void ff_0(int *ptr, size_t size) {
+  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, size_t size) {
+  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 id = ext::oneapi::this_work_item::get_nd_item<1>().get_global_id();
+  ptrC[id] += ptrA[id] * ptrB[id];
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::single_task_kernel))
+void ff_6(int *ptr, int scalarValue, size_t size) {
+  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 %}
+//
+// The name mangling for free function kernels currently does not work with PTX.
+// UNSUPPORTED: cuda
+
+// 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()};
+
+  const size_t N = 1024;
+
+  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_0>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, InputParam);
+    cgh.set_arg(1, N);
+    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(), N).wait();
+  Queue.copy(PtrB, HostDataB.data(), N).wait();
+  for (size_t i = 0; i < N; i++) {
+    assert(HostDataA[i] == 0);
+    assert(HostDataB[i] == i);
+  }
+#endif
+  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 %}
+//
+// The name mangling for free function kernels currently does not work with PTX.
+// UNSUPPORTED: cuda
+
+// 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()};
+
+  const size_t N = 1024;
+
+  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_1>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, InputParam);
+    cgh.set_arg(1, N);
+    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(), N).wait();
+  Queue.copy(PtrB, HostDataB.data(), N).wait();
+  for (size_t i = 0; i < N; 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(), N).wait();
+  Queue.copy(PtrB, HostDataB.data(), N).wait();
+  for (size_t i = 0; i < N; i++) {
+    // A should have been modified 3 times by now, B only once
+    assert(HostDataA[i] == i * 3);
+    assert(HostDataB[i] == i);
+  }
+#endif
+  return 0;
+}

sycl/test-e2e/Graph/Update/FreeFunctionKernels/update_with_indices_ordering.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 %}
+//
+// The name mangling for free function kernels currently does not work with PTX.
+// UNSUPPORTED: cuda
+
+// 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
+  return 0;
+}

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

@@ -0,0 +1,70 @@
+// 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 %}
+//
+// The name mangling for free function kernels currently does not work with PTX.
+// UNSUPPORTED: cuda
+
+// Tests updating a graph node using index-based explicit update
+
+#include "../../graph_common.hpp"
+#include "free_function_kernels.hpp"
+
+int main() {
+  queue Queue{};
+  context ctxt{Queue.get_context()};
+
+  const size_t N = 1024;
+
+  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_0>();
+  kernel Kernel = Bundle.get_kernel(Kernel_id);
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, InputParam);
+    cgh.set_arg(0, N);
+    cgh.single_task(Kernel);
+  });
+
+  auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // PtrA should be filled with values
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  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);
+    assert(HostDataB[i] == 0);
+  }
+
+  // Swap PtrB to be the input
+  InputParam.update(PtrB);
+  ExecGraph.update(KernelNode);
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  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);
+    assert(HostDataB[i] == i);
+  }
+#endif
+  return 0;
+}