[SYCL][Graph] Command Graph PoC

sycl/include/sycl/detail/pi.h

@@ -52,9 +52,10 @@
 // 10.13 Added new PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS queue property.
 // 10.14 Add PI_EXT_INTEL_DEVICE_INFO_FREE_MEMORY as an extension for
 // piDeviceGetInfo.
+// 10.15 Add new PI_EXT_ONEAPI_QUEUE_LAZY_EXECUTION queue property
 
 #define _PI_H_VERSION_MAJOR 10
-#define _PI_H_VERSION_MINOR 14
+#define _PI_H_VERSION_MINOR 15
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -569,6 +570,14 @@ constexpr pi_queue_properties PI_QUEUE_PROFILING_ENABLE = (1 << 1);
 constexpr pi_queue_properties PI_QUEUE_ON_DEVICE = (1 << 2);
 constexpr pi_queue_properties PI_QUEUE_ON_DEVICE_DEFAULT = (1 << 3);
 constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS = (1 << 4);
+// In a lazy queue, enqueued commands are not submitted for execution
+// immediately, instead they are submitted for execution once the queue is
+// flushed.
+//
+// This is to enable prototyping of the SYCL_EXT_ONEAPI_GRAPH extension,
+// before a native command-list interface in PI can be designed and
+// implemented.
+constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_LAZY_EXECUTION = (1 << 5);
 
 using pi_result = _pi_result;
 using pi_platform_info = _pi_platform_info;

sycl/include/sycl/detail/property_helper.hpp

@@ -35,8 +35,9 @@ enum DataLessPropKind {
   UseDefaultStream = 8,
   DiscardEvents = 9,
   DeviceReadOnly = 10,
+  LazyExecution = 11,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 10,
+  LastKnownDataLessPropKind = 11,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

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

@@ -0,0 +1,144 @@
+//==--------- graph.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
+
+#include <memory>
+#include <vector>
+
+#include <sycl/detail/common.hpp>
+#include <sycl/detail/defines_elementary.hpp>
+#include <sycl/property_list.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+
+class handler;
+class queue;
+namespace ext {
+namespace oneapi {
+namespace experimental {
+
+namespace detail {
+struct node_impl;
+struct graph_impl;
+
+using node_ptr = std::shared_ptr<node_impl>;
+using graph_ptr = std::shared_ptr<graph_impl>;
+} // namespace detail
+
+enum class graph_state {
+  modifiable,
+  executable,
+};
+
+class __SYCL_EXPORT node {
+private:
+  node(detail::node_ptr Impl) : impl(Impl) {}
+
+  template <class Obj>
+  friend decltype(Obj::impl)
+  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
+  template <class T>
+  friend T sycl::detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
+
+  detail::node_ptr impl;
+  detail::graph_ptr MGraph;
+};
+
+template <graph_state State = graph_state::modifiable>
+class __SYCL_EXPORT command_graph {
+public:
+  command_graph(const property_list &propList = {});
+
+  // Adding empty node with [0..n] predecessors:
+  node add(const std::vector<node> &dep = {});
+
+  // Adding device node:
+  template <typename T> node add(T cgf, const std::vector<node> &dep = {}) {
+    return add_impl(cgf, dep);
+  }
+
+  // Adding dependency between two nodes.
+  void make_edge(node sender, node receiver);
+
+  command_graph<graph_state::executable>
+  finalize(const sycl::context &syclContext,
+           const property_list &propList = {}) const;
+
+  /// Change the state of a queue to be recording and associate this graph with
+  /// it.
+  /// @param recordingQueue The queue to change state on and associate this
+  /// graph with.
+  /// @return True if the queue had its state changed from executing to
+  /// recording.
+  bool begin_recording(queue recordingQueue);
+
+  /// Change the state of multiple queues to be recording and associate this
+  /// graph with each of them.
+  /// @param recordingQueues The queues to change state on and associate this
+  /// graph with.
+  /// @return True if any queue had its state changed from executing to
+  /// recording.
+  bool begin_recording(const std::vector<queue> &recordingQueues);
+
+  /// Set all queues currently recording to this graph to the executing state.
+  /// @return True if any queue had its state changed from recording to
+  /// executing.
+  bool end_recording();
+
+  /// Set a queues currently recording to this graph to the executing state.
+  /// @param recordingQueue The queue to change state on.
+  /// @return True if the queue had its state changed from recording to
+  /// executing.
+  bool end_recording(queue recordingQueue);
+
+  /// Set multiple queues currently recording to this graph to the executing
+  /// state.
+  /// @param recordingQueue The queues to change state on.
+  /// @return True if any queue had its state changed from recording to
+  /// executing.
+  bool end_recording(const std::vector<queue> &recordingQueues);
+
+private:
+  command_graph(detail::graph_ptr Impl) : impl(Impl) {}
+
+  // Template-less implementation of add()
+  node add_impl(std::function<void(handler &)> cgf,
+                const std::vector<node> &dep);
+
+  template <class Obj>
+  friend decltype(Obj::impl)
+  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
+  template <class T>
+  friend T sycl::detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
+
+  detail::graph_ptr impl;
+};
+
+template <> class __SYCL_EXPORT command_graph<graph_state::executable> {
+public:
+  command_graph() = delete;
+
+  command_graph(detail::graph_ptr g, const sycl::context &ctx)
+      : MTag(rand()), MCtx(ctx), impl(g) {}
+
+private:
+  template <class Obj>
+  friend decltype(Obj::impl)
+  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
+
+  int MTag;
+  const sycl::context &MCtx;
+  detail::graph_ptr impl;
+};
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/feature_test.hpp.in

@@ -47,6 +47,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES 1
 #define SYCL_EXT_ONEAPI_GROUP_ALGORITHMS 1
 #define SYCL_EXT_ONEAPI_GROUP_SORT 1
+#define SYCL_EXT_ONEAPI_LAZY_QUEUE 1
 #define SYCL_EXT_ONEAPI_MAX_WORK_GROUP_QUERY 1
 #define SYCL_EXT_ONEAPI_ND_RANGE_REDUCTIONS 1
 #define SYCL_EXT_ONEAPI_DEFAULT_CONTEXT 1

sycl/include/sycl/handler.hpp

@@ -30,6 +30,8 @@
 #include <sycl/sampler.hpp>
 #include <sycl/stl.hpp>
 
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+
 #include <functional>
 #include <limits>
 #include <memory>
@@ -2516,6 +2518,13 @@ class __SYCL_EXPORT handler {
   /// \param Advice is a device-defined advice for the specified allocation.
   void mem_advise(const void *Ptr, size_t Length, int Advice);
 
+  /// Executes a command_graph.
+  ///
+  /// \param Graph Executable command_graph to run
+  void ext_oneapi_graph(ext::oneapi::experimental::command_graph<
+                        ext::oneapi::experimental::graph_state::executable>
+                            Graph);
+
 private:
   std::shared_ptr<detail::handler_impl> MImpl;
   std::shared_ptr<detail::queue_impl> MQueue;

sycl/include/sycl/properties/queue_properties.hpp

@@ -28,6 +28,8 @@ namespace property {
 namespace queue {
 class discard_events
     : public ::sycl::detail::DataLessProperty<::sycl::detail::DiscardEvents> {};
+class lazy_execution
+    : public ::sycl::detail::DataLessProperty<::sycl::detail::LazyExecution> {};
 } // namespace queue
 } // namespace property
 
@@ -65,6 +67,9 @@ template <>
 struct is_property<ext::oneapi::property::queue::discard_events>
     : std::true_type {};
 template <>
+struct is_property<ext::oneapi::property::queue::lazy_execution>
+    : std::true_type {};
+template <>
 struct is_property<property::queue::cuda::use_default_stream> : std::true_type {
 };
 template <>
@@ -80,6 +85,9 @@ template <>
 struct is_property_of<ext::oneapi::property::queue::discard_events, queue>
     : std::true_type {};
 template <>
+struct is_property_of<ext::oneapi::property::queue::lazy_execution, queue>
+    : std::true_type {};
+template <>
 struct is_property_of<property::queue::cuda::use_default_stream, queue>
     : std::true_type {};
 template <>

sycl/include/sycl/queue.hpp

@@ -23,6 +23,7 @@
 #include <sycl/property_list.hpp>
 #include <sycl/stl.hpp>
 
+
 // Explicitly request format macros
 #ifndef __STDC_FORMAT_MACROS
 #define __STDC_FORMAT_MACROS 1
@@ -72,6 +73,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
@@ -1058,6 +1069,53 @@ class __SYCL_EXPORT queue {
 // Clean KERNELFUNC macros.
 #undef _KERNELFUNCPARAM
 
+  /// Shortcut for executing a graph of commands.
+  ///
+  /// \param Graph the graph of commands to execute
+  /// \return an event representing graph execution operation.
+  event ext_oneapi_graph(ext::oneapi::experimental::command_graph<
+                         ext::oneapi::experimental::graph_state::executable>
+                             Graph) {
+    const detail::code_location CodeLoc = {};
+    return submit([&](handler &CGH) { CGH.ext_oneapi_graph(Graph); }, CodeLoc);
+  }
+
+  /// Shortcut for executing a graph of commands.
+  ///
+  /// \param Graph the graph of commands to execute
+  /// \param DepEvent is an event that specifies the graph execution
+  /// dependencies.
+  /// \return an event representing graph execution operation.
+  event ext_oneapi_graph(ext::oneapi::experimental::command_graph<
+                             ext::oneapi::experimental::graph_state::executable>
+                             Graph,
+                         event DepEvent _CODELOCPARAM(&CodeLoc)) {
+    return submit(
+        [&](handler &CGH) {
+          CGH.depends_on(DepEvent);
+          CGH.ext_oneapi_graph(Graph);
+        } _CODELOCFW(CodeLoc));
+  }
+
+  /// Shortcut for executing a graph of commands.
+  ///
+  /// \param Graph the graph of commands to execute
+  /// \param DepEvents is a vector of events that specifies the graph
+  /// execution dependencies.
+  /// \return an event representing graph execution operation.
+  event ext_oneapi_graph(ext::oneapi::experimental::command_graph<
+                             ext::oneapi::experimental::graph_state::executable>
+                             Graph,
+                         const std::vector<event> &DepEvents) {
+    const detail::code_location CodeLoc = {};
+    return submit(
+        [&](handler &CGH) {
+          CGH.depends_on(DepEvents);
+          CGH.ext_oneapi_graph(Graph);
+        },
+        CodeLoc);
+  }
+
   /// Returns whether the queue is in order or OoO
   ///
   /// Equivalent to has_property<property::queue::in_order>()