graph.hpp

Go to the documentation of this file.

//==--------- 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 <sycl/accessor.hpp>               // for detail::AccessorBaseHost
#include <sycl/context.hpp>                // for context
#include <sycl/detail/export.hpp>          // for __SYCL_EXPORT
#include <sycl/detail/kernel_desc.hpp>     // for kernel_param_kind_t
#include <sycl/detail/property_helper.hpp> // for DataLessPropKind, PropWith...
#include <sycl/device.hpp>                 // for device
#include <sycl/nd_range.hpp>               // for range, nd_range
#include <sycl/properties/property_traits.hpp> // for is_property, is_property_of
#include <sycl/property_list.hpp>              // for property_list
 
#include <functional>  // for function
#include <memory>      // for shared_ptr
#include <type_traits> // for true_type
#include <vector>      // for vector
 
namespace sycl {
inline namespace _V1 {
 
class handler;
class queue;
class device;
namespace ext {
namespace oneapi {
namespace experimental {
 
enum class graph_state {
  modifiable, 
  executable, 
};
 
// Forward declare ext::oneapi::experimental classes
template <graph_state State> class command_graph;
class raw_kernel_arg;
 
namespace detail {
// List of sycl features and extensions which are not supported by graphs. Used
// for throwing errors when these features are used with graphs.
enum class UnsupportedGraphFeatures {
  sycl_reductions = 0,
  sycl_specialization_constants = 1,
  sycl_kernel_bundle = 2,
  sycl_ext_oneapi_kernel_properties = 3,
  sycl_ext_oneapi_enqueue_barrier = 4,
  sycl_ext_oneapi_memcpy2d = 5,
  sycl_ext_oneapi_device_global = 6,
  sycl_ext_oneapi_bindless_images = 7,
  sycl_ext_oneapi_experimental_cuda_cluster_launch = 8,
  sycl_ext_codeplay_enqueue_native_command = 9
};
 
inline const char *
UnsupportedFeatureToString(UnsupportedGraphFeatures Feature) {
  using UGF = UnsupportedGraphFeatures;
  switch (Feature) {
  case UGF::sycl_reductions:
    return "Reductions";
  case UGF::sycl_specialization_constants:
    return "Specialization Constants";
  case UGF::sycl_kernel_bundle:
    return "Kernel Bundles";
  case UGF::sycl_ext_oneapi_kernel_properties:
    return "sycl_ext_oneapi_kernel_properties";
  case UGF::sycl_ext_oneapi_enqueue_barrier:
    return "sycl_ext_oneapi_enqueue_barrier";
  case UGF::sycl_ext_oneapi_memcpy2d:
    return "sycl_ext_oneapi_memcpy2d";
  case UGF::sycl_ext_oneapi_device_global:
    return "sycl_ext_oneapi_device_global";
  case UGF::sycl_ext_oneapi_bindless_images:
    return "sycl_ext_oneapi_bindless_images";
  case UGF::sycl_ext_oneapi_experimental_cuda_cluster_launch:
    return "sycl_ext_oneapi_experimental_cuda_cluster_launch";
  case UGF::sycl_ext_codeplay_enqueue_native_command:
    return "sycl_ext_codeplay_enqueue_native_command";
  }
 
  assert(false && "Unhandled graphs feature");
  return {};
}
 
class node_impl;
class graph_impl;
class exec_graph_impl;
class dynamic_parameter_impl;
} // namespace detail
 
enum class node_type {
  empty = 0,
  subgraph = 1,
  kernel = 2,
  memcpy = 3,
  memset = 4,
  memfill = 5,
  prefetch = 6,
  memadvise = 7,
  ext_oneapi_barrier = 8,
  host_task = 9
};
 
class __SYCL_EXPORT node {
public:
  node() = delete;
 
  node_type get_type() const;
 
  std::vector<node> get_predecessors() const;
 
  std::vector<node> get_successors() const;
 
  static node get_node_from_event(event nodeEvent);
 
  template <int Dimensions>
  void update_nd_range(nd_range<Dimensions> executionRange);
 
  template <int Dimensions> void update_range(range<Dimensions> executionRange);
 
private:
  node(const std::shared_ptr<detail::node_impl> &Impl) : impl(Impl) {}
 
  template <class Obj>
  friend const decltype(Obj::impl) &
  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
  template <class T>
  friend T sycl::detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
 
  std::shared_ptr<detail::node_impl> impl;
};
 
namespace property {
namespace graph {
 
class no_cycle_check : public ::sycl::detail::DataLessProperty<
                           ::sycl::detail::GraphNoCycleCheck> {
public:
  no_cycle_check() = default;
};
 
class assume_buffer_outlives_graph
    : public ::sycl::detail::DataLessProperty<
          ::sycl::detail::GraphAssumeBufferOutlivesGraph> {
public:
  assume_buffer_outlives_graph() = default;
};
 
class updatable
    : public ::sycl::detail::DataLessProperty<::sycl::detail::GraphUpdatable> {
public:
  updatable() = default;
};
 
class enable_profiling : public ::sycl::detail::DataLessProperty<
                             ::sycl::detail::GraphEnableProfiling> {
public:
  enable_profiling() = default;
};
} // namespace graph
 
namespace node {
 
class depends_on : public ::sycl::detail::PropertyWithData<
                       ::sycl::detail::GraphNodeDependencies> {
public:
  template <typename... NodeTN> depends_on(NodeTN... nodes) : MDeps{nodes...} {}
 
  const std::vector<::sycl::ext::oneapi::experimental::node> &
  get_dependencies() const {
    return MDeps;
  }
 
private:
  const std::vector<::sycl::ext::oneapi::experimental::node> MDeps;
};
 
class depends_on_all_leaves : public ::sycl::detail::DataLessProperty<
                                  ::sycl::detail::GraphDependOnAllLeaves> {
public:
  depends_on_all_leaves() = default;
};
 
} // namespace node
} // namespace property
 
namespace detail {
// Templateless modifiable command-graph base class.
class __SYCL_EXPORT modifiable_command_graph {
public:
  modifiable_command_graph(const context &SyclContext, const device &SyclDevice,
                           const property_list &PropList = {});
 
  modifiable_command_graph(const queue &SyclQueue,
                           const property_list &PropList = {});
 
  node add(const property_list &PropList = {}) {
    if (PropList.has_property<property::node::depends_on>()) {
      auto Deps = PropList.get_property<property::node::depends_on>();
      node Node = addImpl(Deps.get_dependencies());
      if (PropList.has_property<property::node::depends_on_all_leaves>()) {
        addGraphLeafDependencies(Node);
      }
      return Node;
    }
    node Node = addImpl({});
    if (PropList.has_property<property::node::depends_on_all_leaves>()) {
      addGraphLeafDependencies(Node);
    }
    return Node;
  }
 
  template <typename T> node add(T CGF, const property_list &PropList = {}) {
    if (PropList.has_property<property::node::depends_on>()) {
      auto Deps = PropList.get_property<property::node::depends_on>();
      node Node = addImpl(CGF, Deps.get_dependencies());
      if (PropList.has_property<property::node::depends_on_all_leaves>()) {
        addGraphLeafDependencies(Node);
      }
      return Node;
    }
    node Node = addImpl(CGF, {});
    if (PropList.has_property<property::node::depends_on_all_leaves>()) {
      addGraphLeafDependencies(Node);
    }
    return Node;
  }
 
  void make_edge(node &Src, node &Dest);
 
  command_graph<graph_state::executable>
  finalize(const property_list &PropList = {}) const;
 
  void begin_recording(queue &RecordingQueue,
                       const property_list &PropList = {});
 
  void begin_recording(const std::vector<queue> &RecordingQueues,
                       const property_list &PropList = {});
 
  void end_recording();
 
  void end_recording(queue &RecordingQueue);
 
  void end_recording(const std::vector<queue> &RecordingQueues);
 
  void print_graph(const std::string path, bool verbose = false) const;
 
  std::vector<node> get_nodes() const;
 
  std::vector<node> get_root_nodes() const;
 
protected:
  modifiable_command_graph(const std::shared_ptr<detail::graph_impl> &Impl)
      : impl(Impl) {}
 
  node addImpl(std::function<void(handler &)> CGF,
               const std::vector<node> &Dep);
 
  node addImpl(const std::vector<node> &Dep);
 
  void addGraphLeafDependencies(node Node);
 
  template <class Obj>
  friend const decltype(Obj::impl) &
  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
  template <class T>
  friend T sycl::detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
 
  std::shared_ptr<detail::graph_impl> impl;
};
 
// Templateless executable command-graph base class.
class __SYCL_EXPORT executable_command_graph {
public:
  executable_command_graph() = delete;
 
  void update(const command_graph<graph_state::modifiable> &Graph);
 
  void update(const node &Node);
 
  void update(const std::vector<node> &Nodes);
 
protected:
  executable_command_graph(const std::shared_ptr<detail::graph_impl> &Graph,
                           const sycl::context &Ctx,
                           const property_list &PropList = {});
 
  template <class Obj>
  friend const decltype(Obj::impl) &
  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
 
  void finalizeImpl();
 
  std::shared_ptr<detail::exec_graph_impl> impl;
};
} // namespace detail
 
template <graph_state State = graph_state::modifiable>
class command_graph : public detail::modifiable_command_graph {
public:
  command_graph(const context &SyclContext, const device &SyclDevice,
                const property_list &PropList = {})
      : modifiable_command_graph(SyclContext, SyclDevice, PropList) {}
 
  command_graph(const queue &SyclQueue, const property_list &PropList = {})
      : modifiable_command_graph(SyclQueue, PropList) {}
 
private:
  command_graph(const std::shared_ptr<detail::graph_impl> &Impl)
      : modifiable_command_graph(Impl) {}
 
  template <class T>
  friend T sycl::detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
};
 
template <>
class command_graph<graph_state::executable>
    : public detail::executable_command_graph {
protected:
  friend command_graph<graph_state::executable>
  detail::modifiable_command_graph::finalize(const sycl::property_list &) const;
  using detail::executable_command_graph::executable_command_graph;
};
 
namespace detail {
class __SYCL_EXPORT dynamic_parameter_base {
public:
  dynamic_parameter_base(
      sycl::ext::oneapi::experimental::command_graph<graph_state::modifiable>
          Graph,
      size_t ParamSize, const void *Data);
 
protected:
  void updateValue(const void *NewValue, size_t Size);
 
  // Update a sycl_ext_oneapi_raw_kernel_arg parameter. Size parameter is
  // ignored as it represents sizeof(raw_kernel_arg), which doesn't represent
  // the number of underlying bytes.
  void updateValue(const raw_kernel_arg *NewRawValue, size_t Size);
 
  void updateAccessor(const sycl::detail::AccessorBaseHost *Acc);
  std::shared_ptr<dynamic_parameter_impl> impl;
 
  template <class Obj>
  friend const decltype(Obj::impl) &
  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
};
} // namespace detail
 
template <typename ValueT>
class dynamic_parameter : public detail::dynamic_parameter_base {
  static constexpr bool IsAccessor =
      std::is_base_of_v<sycl::detail::AccessorBaseHost, ValueT>;
  static constexpr sycl::detail::kernel_param_kind_t ParamType =
      IsAccessor ? sycl::detail::kernel_param_kind_t::kind_accessor
      : std::is_pointer_v<ValueT>
          ? sycl::detail::kernel_param_kind_t::kind_pointer
          : sycl::detail::kernel_param_kind_t::kind_std_layout;
 
public:
  dynamic_parameter(experimental::command_graph<graph_state::modifiable> Graph,
                    const ValueT &Param)
      : detail::dynamic_parameter_base(Graph, sizeof(ValueT), &Param) {}
 
  void update(const ValueT &NewValue) {
    if constexpr (IsAccessor) {
      detail::dynamic_parameter_base::updateAccessor(&NewValue);
    } else {
      detail::dynamic_parameter_base::updateValue(&NewValue, sizeof(ValueT));
    }
  }
};
 
template <typename ValueT>
dynamic_parameter(experimental::command_graph<graph_state::modifiable> Graph,
                  const ValueT &Param) -> dynamic_parameter<ValueT>;
template <graph_state State = graph_state::modifiable>
command_graph(const context &SyclContext, const device &SyclDevice,
              const property_list &PropList) -> command_graph<State>;
 
} // namespace experimental
} // namespace oneapi
} // namespace ext
 
template <>
struct is_property<ext::oneapi::experimental::property::graph::no_cycle_check>
    : std::true_type {};
 
template <>
struct is_property<ext::oneapi::experimental::property::node::depends_on>
    : std::true_type {};
 
template <>
struct is_property_of<
    ext::oneapi::experimental::property::graph::no_cycle_check,
    ext::oneapi::experimental::command_graph<
        ext::oneapi::experimental::graph_state::modifiable>> : std::true_type {
};
 
template <>
struct is_property_of<ext::oneapi::experimental::property::node::depends_on,
                      ext::oneapi::experimental::node> : std::true_type {};
 
} // namespace _V1
} // namespace sycl