diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index aabab38b62cb1..f10f43e0430a8 100644 --- 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 index 81cf668f523fc..3009af8ee2890 100644 --- 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 index e15ac3eadb67f..15645d9884499 100644 --- 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()) { auto Deps = PropList.get_property(); - return addImpl(Deps.get_dependencies()); + node Node = addImpl(Deps.get_dependencies()); + if (PropList.has_property()) { + addGraphLeafDependencies(Node); + } + return Node; } - return addImpl({}); + node Node = addImpl({}); + if (PropList.has_property()) { + addGraphLeafDependencies(Node); + } + return Node; } /// Add a command-group node to the graph. @@ -171,9 +187,17 @@ class __SYCL_EXPORT modifiable_command_graph { template node add(T CGF, const property_list &PropList = {}) { if (PropList.has_property()) { auto Deps = PropList.get_property(); - return addImpl(CGF, Deps.get_dependencies()); + node Node = addImpl(CGF, Deps.get_dependencies()); + if (PropList.has_property()) { + addGraphLeafDependencies(Node); + } + return Node; + } + node Node = addImpl(CGF, {}); + if (PropList.has_property()) { + 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 &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 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 index 29ab2d20b8f6a..168d1bc83f253 100644 --- 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 CGF, return sycl::detail::createSyclObjFromImpl(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 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 SenderImpl = sycl::detail::getSyclObjImpl(Src); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 56cd6ebf90be1..5d1ca2b883caa 100644 --- 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 index ecc39d00ff59c..7e28aa4f96bbe 100644 --- 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 index 58ddec5c14f4e..b9a02706e9e4e 100644 --- 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>([]() {}); }); + auto Node2Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Node3Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + + 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>([]() {}); }); + auto Node2Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Node3Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Node4Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {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>([]() {}); }); + auto Node2Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + 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>([]() {}); }, + {experimental::property::node::depends_on(Node1Graph)}); + auto Node4Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {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>([]() {}); }); + auto Node2Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + 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>([]() {}); }); + 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{}};