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

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
 
@@ -479,6 +484,20 @@ 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

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
 };

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.
@@ -239,6 +263,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);

sycl/source/detail/graph_impl.cpp

@@ -726,6 +726,21 @@ 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) {
+      if (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);

sycl/test/abi/sycl_symbols_linux.dump

@@ -3731,6 +3731,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_graph15dependsOnLeavesENS3_4nodeE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC2ERKNS3_16image_descriptorERKNS0_5queueE

sycl/unittests/Extensions/CommandGraph.cpp

@@ -1515,6 +1515,84 @@ 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, FusionExtensionExceptionCheck) {
   queue Q{ext::codeplay::experimental::property::queue::enable_fusion{}};