[SYCL][Graph] Adding experimental implementation of SYCL Command Graph extension

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/add_nodes_after_finalize.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/basic_buffer.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/basic_usm.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/basic_usm_host.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/basic_usm_mixed.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/basic_usm_shared.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/basic_usm_system.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_2d.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_host2target.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_host2target_2d.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_host2target_offset.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_offsets.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_target2host.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_target2host_2d.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_copy_target2host_offset.cpp"

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

@@ -0,0 +1,11 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/buffer_ordering.cpp"

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

@@ -0,0 +1,49 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+// Tests that an event returned from adding a graph node using the queue
+// recording API can be passed to `handler::depends_on` inside a node
+// added using the explicit API. This should create a graph edge.
+
+#include "../graph_common.hpp"
+
+int main() {
+
+  queue Queue{gpu_selector_v};
+
+  exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
+
+  const size_t N = 10;
+  float *Arr = malloc_device<float>(N, Queue);
+
+  Graph.begin_recording(Queue);
+  // `Event` corresponds to a graph node
+  event Event = Queue.submit([&](handler &CGH) {
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) { Arr[idx] = 42.0f; });
+  });
+  Graph.end_recording(Queue);
+
+  Graph.add([&](handler &CGH) {
+    CGH.depends_on(Event); // creates edge to recorded graph node
+    CGH.parallel_for(range<1>{N}, [=](id<1> idx) { Arr[idx] *= 2.0f; });
+  });
+
+  auto ExecGraph = Graph.finalize();
+
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }).wait();
+
+  constexpr float ref = 42.0f * 2.0f;
+  std::vector<float> Output(N);
+  Queue.memcpy(Output.data(), Arr, N * sizeof(float)).wait();
+  for (int i = 0; i < N; i++)
+    assert(Output[i] == ref);
+
+  sycl::free(Arr, Queue);
+
+  return 0;
+}

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

@@ -0,0 +1,14 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+// Expected fail as reduction support is not complete.
+// XFAIL: *
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/dotp_buffer_reduction.cpp"

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

@@ -0,0 +1,82 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+// Expected fail as reduction support is not complete.
+// XFAIL: *
+
+// Tests constructing a graph using the explicit API to perform a dotp
+// operation which uses a sycl reduction with USM memory.
+
+#include "../graph_common.hpp"
+
+int main() {
+  queue Queue{gpu_selector_v};
+
+  exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
+
+  float *Dotp = malloc_device<float>(1, Queue);
+
+  const size_t N = 10;
+  float *X = malloc_device<float>(N, Queue);
+  float *Y = malloc_device<float>(N, Queue);
+  float *Z = malloc_device<float>(N, Queue);
+
+  auto NodeI = Graph.add([&](handler &CGH) {
+    CGH.parallel_for(N, [=](id<1> it) {
+      const size_t i = it[0];
+      X[i] = 1.0f;
+      Y[i] = 2.0f;
+      Z[i] = 3.0f;
+    });
+  });
+
+  auto NodeA = Graph.add(
+      [&](handler &CGH) {
+        CGH.parallel_for(range<1>{N}, [=](id<1> it) {
+          const size_t i = it[0];
+          X[i] = Alpha * X[i] + Beta * Y[i];
+        });
+      },
+      {exp_ext::property::node::depends_on(NodeI)});
+
+  auto NodeB = Graph.add(
+      [&](handler &CGH) {
+        CGH.parallel_for(range<1>{N}, [=](id<1> it) {
+          const size_t i = it[0];
+          Z[i] = Gamma * Z[i] + Beta * Y[i];
+        });
+      },
+      {exp_ext::property::node::depends_on(NodeI)});
+
+  auto NodeC = Graph.add(
+      [&](handler &CGH) {
+        CGH.parallel_for(range<1>{N}, reduction(Dotp, 0.0f, std::plus()),
+                         [=](id<1> it, auto &Sum) {
+                           const size_t i = it[0];
+                           Sum += X[i] * Z[i];
+                         });
+      },
+      {exp_ext::property::node::depends_on(NodeA, NodeB)});
+
+  auto ExecGraph = Graph.finalize();
+
+  // Using shortcut for executing a graph of commands
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  float Output;
+  Queue.memcpy(&Output, Dotp, sizeof(float)).wait();
+
+  assert(Output == dotp_reference_result(N));
+
+  sycl::free(Dotp, Queue);
+  sycl::free(X, Queue);
+  sycl::free(Y, Queue);
+  sycl::free(Z, Queue);
+
+  return 0;
+}

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

@@ -0,0 +1,14 @@
+// REQUIRES: level_zero, gpu
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using ZE_DEBUG
+// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %}
+//
+// CHECK-NOT: LEAK
+
+// Expected fail as executable graph update isn't implemented yet
+// XFAIL: *
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/double_buffer.cpp"