[SYCL][Graph] Add a shortcut for adding leaves as dependencies (intel…

…#11855) Adds a node property that allows users to easily add all leaves of a graph as dependencies when creating a node with the explicit API. Updates the spec with this new feature. Adds unitests that check this behaviour.
reble · Nov 15, 2023 · 0b6b3b8 · 0b6b3b8
1 parent e15ebd0
commit 0b6b3b8
Show file tree

Hide file tree

Showing 7 changed files with 234 additions and 5 deletions.
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc
@@ -336,6 +336,11 @@ class depends_on {
     depends_on(NodeTN... nodes);
 };
 
+class depends_on_all_leaves {
+  public:
+    depends_on_all_leaves() = default;
+};
+
 } // namespace node
 } // namespace property
 
@@ -481,6 +486,21 @@ class depends_on {
 }
 ----
 
+==== Depends-On-All-Leaves Property 
+The API for explicitly adding nodes to a `command_graph` includes a
+`property_list` parameter. This extension defines the `depends_on_all_leaves` 
+property to be passed here. `depends_on_all_leaves` provides a shortcut for 
+adding all the current leaves of a graph as dependencies.
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental::property::node {
+class depends_on_all_leaves {
+  public:
+    depends_on_all_leaves();
+};
+}
+----
+
 === Graph
 
 This extension adds a new `command_graph` object which follows the

diff --git a/sycl/include/sycl/detail/property_helper.hpp b/sycl/include/sycl/detail/property_helper.hpp
@@ -46,8 +46,9 @@ enum DataLessPropKind {
   QueueSubmissionImmediate = 21,
   GraphAssumeDataOutlivesBuffer = 22,
   GraphAssumeBufferOutlivesGraph = 23,
+  GraphDependOnAllLeaves = 24,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 23,
+  LastKnownDataLessPropKind = 24,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp
@@ -137,6 +137,14 @@ class depends_on : public ::sycl::detail::PropertyWithData<
   const std::vector<::sycl::ext::oneapi::experimental::node> MDeps;
 };
 
+/// Property used to to add all previous graph leaves as dependencies when
+/// creating a new node with command_graph::add().
+class depends_on_all_leaves : public ::sycl::detail::DataLessProperty<
+                                  ::sycl::detail::GraphDependOnAllLeaves> {
+public:
+  depends_on_all_leaves() = default;
+};
+
 } // namespace node
 } // namespace property
 
@@ -159,9 +167,17 @@ class __SYCL_EXPORT modifiable_command_graph {
   node add(const property_list &PropList = {}) {
     if (PropList.has_property<property::node::depends_on>()) {
       auto Deps = PropList.get_property<property::node::depends_on>();
-      return addImpl(Deps.get_dependencies());
+      node Node = addImpl(Deps.get_dependencies());
+      if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+        addGraphLeafDependencies(Node);
+      }
+      return Node;
     }
-    return addImpl({});
+    node Node = addImpl({});
+    if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+      addGraphLeafDependencies(Node);
+    }
+    return Node;
   }
 
   /// Add a command-group node to the graph.
@@ -171,9 +187,17 @@ class __SYCL_EXPORT modifiable_command_graph {
   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>();
-      return addImpl(CGF, Deps.get_dependencies());
+      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 addImpl(CGF, {});
+    return Node;
   }
 
   /// Add a dependency between two nodes.
@@ -247,6 +271,11 @@ class __SYCL_EXPORT modifiable_command_graph {
   /// @return Node added to the graph.
   node addImpl(const std::vector<node> &Dep);
 
+  /// Adds all graph leaves as dependencies
+  /// @param Node Destination node to which the leaves of the graph will be
+  /// added as dependencies.
+  void addGraphLeafDependencies(node Node);
+
   template <class Obj>
   friend decltype(Obj::impl)
   sycl::detail::getSyclObjImpl(const Obj &SyclObject);

diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp
@@ -726,6 +726,19 @@ node modifiable_command_graph::addImpl(std::function<void(handler &)> CGF,
   return sycl::detail::createSyclObjFromImpl<node>(NodeImpl);
 }
 
+void modifiable_command_graph::addGraphLeafDependencies(node Node) {
+  // Find all exit nodes in the current graph and add them to the dependency
+  // vector
+  std::shared_ptr<detail::node_impl> DstImpl =
+      sycl::detail::getSyclObjImpl(Node);
+  graph_impl::WriteLock Lock(impl->MMutex);
+  for (auto &NodeImpl : impl->MNodeStorage) {
+    if ((NodeImpl->MSuccessors.size() == 0) && (NodeImpl != DstImpl)) {
+      impl->makeEdge(NodeImpl, DstImpl);
+    }
+  }
+}
+
 void modifiable_command_graph::make_edge(node &Src, node &Dest) {
   std::shared_ptr<detail::node_impl> SenderImpl =
       sycl::detail::getSyclObjImpl(Src);

diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3734,6 +3734,7 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplES
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph9make_edgeERNS3_4nodeES7_
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE
+_ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph24addGraphLeafDependenciesENS3_4nodeE
 _ZNK4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph11print_graphENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE

diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
@@ -873,6 +873,7 @@
 ?add@device_global_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@host_pipe_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA?AVnode@34567@AEBVproperty_list@67@@Z
+?addGraphLeafDependencies@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXVnode@34567@@Z
 ?addHostAccessorAndWait@detail@_V1@sycl@@YAXPEAVAccessorImplHost@123@@Z
 ?addHostSampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVSampledImageAccessorImplHost@123@@Z
 ?addHostUnsampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVUnsampledImageAccessorImplHost@123@@Z

diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp
@@ -1500,6 +1500,170 @@ TEST_F(CommandGraphTest, EnqueueMultipleBarrier) {
   }
 }
 
+TEST_F(CommandGraphTest, DependencyLeavesKeyword1) {
+  auto Node1Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node2Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node3Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+
+  auto EmptyNode =
+      Graph.add([&](sycl::handler &cgh) { /*empty node */ },
+                {experimental::property::node::depends_on_all_leaves()});
+
+  auto GraphImpl = sycl::detail::getSyclObjImpl(Graph);
+
+  // Check the graph structure
+  // (1) (2) (3)
+  //   \  |  /
+  //    \ | /
+  //     (E)
+  ASSERT_EQ(GraphImpl->MRoots.size(), 3lu);
+  auto EmptyImpl = sycl::detail::getSyclObjImpl(EmptyNode);
+  ASSERT_EQ(EmptyImpl->MPredecessors.size(), 3lu);
+  ASSERT_EQ(EmptyImpl->MSuccessors.size(), 0lu);
+
+  auto Node1Impl = sycl::detail::getSyclObjImpl(Node1Graph);
+  ASSERT_EQ(Node1Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node1Impl->MSuccessors[0].lock(), EmptyImpl);
+  auto Node2Impl = sycl::detail::getSyclObjImpl(Node2Graph);
+  ASSERT_EQ(Node2Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node2Impl->MSuccessors[0].lock(), EmptyImpl);
+  auto Node3Impl = sycl::detail::getSyclObjImpl(Node3Graph);
+  ASSERT_EQ(Node3Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node3Impl->MSuccessors[0].lock(), EmptyImpl);
+}
+
+TEST_F(CommandGraphTest, DependencyLeavesKeyword2) {
+  auto Node1Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node2Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node3Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node4Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); },
+      {experimental::property::node::depends_on(Node3Graph)});
+
+  auto EmptyNode =
+      Graph.add([&](sycl::handler &cgh) { /*empty node */ },
+                {experimental::property::node::depends_on_all_leaves()});
+
+  auto GraphImpl = sycl::detail::getSyclObjImpl(Graph);
+
+  // Check the graph structure
+  // (1) (2) (3)
+  //   \  |  /
+  //    \ | (4)
+  //     \| /
+  //     (E)
+  ASSERT_EQ(GraphImpl->MRoots.size(), 3lu);
+  auto EmptyImpl = sycl::detail::getSyclObjImpl(EmptyNode);
+  ASSERT_EQ(EmptyImpl->MPredecessors.size(), 3lu);
+  ASSERT_EQ(EmptyImpl->MSuccessors.size(), 0lu);
+
+  auto Node1Impl = sycl::detail::getSyclObjImpl(Node1Graph);
+  ASSERT_EQ(Node1Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node1Impl->MSuccessors[0].lock(), EmptyImpl);
+  auto Node2Impl = sycl::detail::getSyclObjImpl(Node2Graph);
+  ASSERT_EQ(Node2Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node2Impl->MSuccessors[0].lock(), EmptyImpl);
+  auto Node3Impl = sycl::detail::getSyclObjImpl(Node3Graph);
+  ASSERT_EQ(Node3Impl->MSuccessors.size(), 1lu);
+
+  auto Node4Impl = sycl::detail::getSyclObjImpl(Node4Graph);
+  ASSERT_EQ(Node4Impl->MPredecessors.size(), 1lu);
+  ASSERT_EQ(Node4Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node4Impl->MSuccessors[0].lock(), EmptyImpl);
+}
+
+TEST_F(CommandGraphTest, DependencyLeavesKeyword3) {
+  auto Node1Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node2Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto EmptyNode =
+      Graph.add([&](sycl::handler &cgh) { /*empty node */ },
+                {experimental::property::node::depends_on_all_leaves()});
+  auto Node3Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); },
+      {experimental::property::node::depends_on(Node1Graph)});
+  auto Node4Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); },
+      {experimental::property::node::depends_on(EmptyNode)});
+
+  auto GraphImpl = sycl::detail::getSyclObjImpl(Graph);
+
+  // Check the graph structure
+  // (1)(2)
+  //  |\ |
+  //  | (E)
+  // (3) |
+  //    (4)
+  ASSERT_EQ(GraphImpl->MRoots.size(), 2lu);
+  auto EmptyImpl = sycl::detail::getSyclObjImpl(EmptyNode);
+  ASSERT_EQ(EmptyImpl->MPredecessors.size(), 2lu);
+  ASSERT_EQ(EmptyImpl->MSuccessors.size(), 1lu);
+
+  auto Node1Impl = sycl::detail::getSyclObjImpl(Node1Graph);
+  auto Node2Impl = sycl::detail::getSyclObjImpl(Node2Graph);
+  ASSERT_EQ(Node1Impl->MSuccessors.size(), 2lu);
+  ASSERT_EQ(Node2Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node2Impl->MSuccessors[0].lock(), EmptyImpl);
+
+  auto Node3Impl = sycl::detail::getSyclObjImpl(Node3Graph);
+  ASSERT_EQ(Node3Impl->MPredecessors.size(), 1lu);
+  ASSERT_EQ(Node3Impl->MPredecessors[0].lock(), Node1Impl);
+
+  auto Node4Impl = sycl::detail::getSyclObjImpl(Node4Graph);
+  ASSERT_EQ(Node4Impl->MPredecessors.size(), 1lu);
+  ASSERT_EQ(Node4Impl->MPredecessors[0].lock(), EmptyImpl);
+}
+
+TEST_F(CommandGraphTest, DependencyLeavesKeyword4) {
+  auto Node1Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto Node2Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto EmptyNode =
+      Graph.add([&](sycl::handler &cgh) { /*empty node */ },
+                {experimental::property::node::depends_on_all_leaves()});
+  auto Node3Graph = Graph.add(
+      [&](sycl::handler &cgh) { cgh.single_task<TestKernel<>>([]() {}); });
+  auto EmptyNode2 =
+      Graph.add([&](sycl::handler &cgh) { /*empty node */ },
+                {experimental::property::node::depends_on_all_leaves()});
+
+  auto GraphImpl = sycl::detail::getSyclObjImpl(Graph);
+
+  // Check the graph structure
+  // (1)(2)
+  //   \/
+  //  (E1) (3)
+  //    \  /
+  //    (E2)
+  ASSERT_EQ(GraphImpl->MRoots.size(), 3lu);
+  auto EmptyImpl = sycl::detail::getSyclObjImpl(EmptyNode);
+  ASSERT_EQ(EmptyImpl->MPredecessors.size(), 2lu);
+  ASSERT_EQ(EmptyImpl->MSuccessors.size(), 1lu);
+
+  auto Node1Impl = sycl::detail::getSyclObjImpl(Node1Graph);
+  ASSERT_EQ(Node1Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node1Impl->MSuccessors[0].lock(), EmptyImpl);
+  auto Node2Impl = sycl::detail::getSyclObjImpl(Node2Graph);
+  ASSERT_EQ(Node2Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node2Impl->MSuccessors[0].lock(), EmptyImpl);
+
+  auto EmptyImpl2 = sycl::detail::getSyclObjImpl(EmptyNode2);
+  auto Node3Impl = sycl::detail::getSyclObjImpl(Node3Graph);
+  ASSERT_EQ(Node3Impl->MPredecessors.size(), 0lu);
+  ASSERT_EQ(Node3Impl->MSuccessors.size(), 1lu);
+  ASSERT_EQ(Node3Impl->MSuccessors[0].lock(), EmptyImpl2);
+
+  ASSERT_EQ(EmptyImpl2->MPredecessors.size(), 2lu);
+}
+
 TEST_F(CommandGraphTest, FusionExtensionExceptionCheck) {
   queue Q{ext::codeplay::experimental::property::queue::enable_fusion{}};