[SYCL] Add graph record and replay entry points

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -9,11 +9,17 @@
 #pragma once
 
 #include <sycl/detail/defines_elementary.hpp>
+#include <sycl/queue.hpp>
+
+#include "graph_defines.hpp"
 
 #include <list>
 #include <set>
 
 namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+class queue;
+
 namespace ext {
 namespace oneapi {
 namespace experimental {
@@ -142,8 +148,6 @@ struct node {
   void set_root() { MGraph->add_root(MNode); }
 };
 
-enum class graph_state { modifiable, executable };
-
 template <graph_state State = graph_state::modifiable> class command_graph {
 public:
   // Adding empty node with [0..n] predecessors:
@@ -219,5 +223,6 @@ void command_graph<graph_state::executable>::exec_and_wait(sycl::queue q) {
 } // namespace experimental
 } // namespace oneapi
 } // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl
 

sycl/include/sycl/ext/oneapi/experimental/graph_defines.hpp

@@ -0,0 +1,28 @@
+//==--------- graph_defines.hpp --- SYCL graph extension
+//---------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+enum class graph_state {
+  modifiable,
+  executable,
+};
+
+/// Forward declaration for command_graph
+template <graph_state State> class command_graph;
+}
+} // namespace oneapi
+} // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/queue.hpp

@@ -23,6 +23,8 @@
 #include <sycl/property_list.hpp>
 #include <sycl/stl.hpp>
 
+#include <sycl/ext/oneapi/experimental/graph_defines.hpp>
+
 // Explicitly request format macros
 #ifndef __STDC_FORMAT_MACROS
 #define __STDC_FORMAT_MACROS 1
@@ -72,6 +74,16 @@ static event submitAssertCapture(queue &, event &, queue *,
 #endif
 } // namespace detail
 
+namespace ext {
+namespace oneapi {
+namespace experimental {
+// State of a queue with regards to graph recording,
+// returned by info::queue::state
+enum class queue_state { executing, recording };
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+
 /// Encapsulates a single SYCL queue which schedules kernels on a SYCL device.
 ///
 /// A SYCL queue can be used to submit command groups to be executed by the SYCL
@@ -1068,6 +1080,26 @@ class __SYCL_EXPORT queue {
   /// \return the backend associated with this queue.
   backend get_backend() const noexcept;
 
+public:
+  /// Places the queue into command_graph recording mode.
+  ///
+  /// \return true if the queue was not already in recording mode.
+  bool
+  begin_recording(ext::oneapi::experimental::command_graph<
+                  ext::oneapi::experimental::graph_state::modifiable> &graph);
+
+  /// Ends recording mode on the queue and returns to the normal state.
+  ///
+  /// \return true if the queue was already in recording mode.
+  bool end_recording();
+
+  /// Submits an executable command_graph for execution on this queue
+  ///
+  /// \return an event representing the execution of the command_graph
+  event submit(ext::oneapi::experimental::command_graph<
+               ext::oneapi::experimental::graph_state::executable>
+                   graph);
+
 private:
   pi_native_handle getNative() const;
 

sycl/source/queue.cpp

@@ -15,6 +15,8 @@
 #include <sycl/queue.hpp>
 #include <sycl/stl.hpp>
 
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+
 #include <algorithm>
 
 namespace sycl {
@@ -212,5 +214,24 @@ bool queue::device_has(aspect Aspect) const {
   // avoid creating sycl object from impl
   return impl->getDeviceImplPtr()->has(Aspect);
 }
+
+bool queue::begin_recording(
+    ext::oneapi::experimental::command_graph<
+        ext::oneapi::experimental::graph_state::modifiable> &graph) {
+  // Empty Implementation
+  return true;
+}
+
+bool queue::end_recording() {
+  // Empty Implementation
+  return true;
+}
+
+event queue::submit(ext::oneapi::experimental::command_graph<
+             ext::oneapi::experimental::graph_state::executable>
+                 graph) {
+  graph.exec_and_wait(*this);
+  return {};
+}
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/test/graph/graph-record-simple.cpp

@@ -0,0 +1,53 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+
+int main() {
+  const size_t n = 10;
+  const float expectedValue = 1.f;
+
+  sycl::property_list properties{
+      sycl::property::queue::in_order(),
+      sycl::ext::oneapi::property::queue::lazy_execution{}};
+
+  sycl::queue q{sycl::default_selector_v, properties};
+
+  sycl::ext::oneapi::experimental::command_graph g;
+
+  float *arr = sycl::malloc_shared<float>(n, q);
+
+  q.begin_recording(g);
+
+  q.submit([&](sycl::handler &h) {
+    h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx) {
+      size_t i = idx;
+      arr[i] = expectedValue;
+    });
+  });
+
+  q.end_recording();
+
+  auto exec_graph = g.finalize(q.get_context());
+
+  q.submit(exec_graph);
+
+  int errors = 0;
+  // Verify results
+  for (size_t i = 0; i < n; i++) {
+    if (arr[i] != expectedValue) {
+      std::cout << "Unexpected result at index: " << i
+                << ", expected: " << expectedValue << " actual: " << arr[i]
+                << "\n";
+      errors++;
+    }
+  }
+
+  std::cout << "done.\n";
+
+  sycl::free(arr, q.get_context());
+
+  return errors;
+}