diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 7d26e61f9c4d8..101807a4ef3fc 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -176,28 +176,31 @@ std::shared_ptr graph_impl::addNodesToExits( std::shared_ptr graph_impl::addSubgraphNodes( const std::shared_ptr &SubGraphExec) { std::map, std::shared_ptr> NodesMap; - std::list> NewNodeList; + std::list> NewNodesList; - std::list> NodeList = SubGraphExec->getSchedule(); + std::list> NodesList = SubGraphExec->getSchedule(); + // Duplication of nodes for (std::list>::const_iterator NodeIt = - NodeList.end(); - NodeIt != NodeList.begin();) { + NodesList.end(); + NodeIt != NodesList.begin();) { --NodeIt; auto Node = *NodeIt; std::shared_ptr NodeCopy; duplicateNode(Node, NodeCopy); - NewNodeList.push_back(NodeCopy); + NewNodesList.push_back(NodeCopy); NodesMap.insert({Node, NodeCopy}); for (auto &NextNode : Node->MSuccessors) { if (NodesMap.find(NextNode) != NodesMap.end()) { auto Successor = NodesMap[NextNode]; NodeCopy->registerSuccessor(Successor, NodeCopy); + } else { + assert("Node duplication failed. A duplicated node is missing."); } } } - return addNodesToExits(NewNodeList); + return addNodesToExits(NewNodesList); } void graph_impl::addRoot(const std::shared_ptr &Root) { @@ -523,6 +526,11 @@ void exec_graph_impl::createCommandBuffers(sycl::device Device) { // TODO extract kernel bundle logic from enqueueImpKernel for (auto Node : MSchedule) { + // Empty nodes are not processed as other nodes, but only their + // dependencies are propagated in findRealDeps + if (Node->isEmpty()) + continue; + sycl::detail::CG::CGTYPE type = Node->MCGType; // If the node is a kernel with no special requirements we can enqueue it // directly. @@ -662,8 +670,9 @@ exec_graph_impl::enqueue(const std::shared_ptr &Queue, "Error during emulated graph command group submission."); } ScheduledEvents.push_back(NewEvent); - } else { - + } else if (!NodeImpl->isEmpty()) { + // Empty nodes are node processed as other nodes, but only their + // dependencies are propagated in findRealDeps sycl::detail::EventImplPtr EventImpl = sycl::detail::Scheduler::getInstance().addCG(NodeImpl->getCGCopy(), Queue); diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index 7b961e78d85af..c30875078b440 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -142,10 +142,8 @@ class node_impl { if (std::find(Schedule.begin(), Schedule.end(), Next) == Schedule.end()) Next->sortTopological(Next, Schedule); } - // We don't need to schedule empty nodes as they are only used when - // calculating dependencies - if (!NodeImpl->isEmpty()) - Schedule.push_front(NodeImpl); + + Schedule.push_front(NodeImpl); } /// Checks if this node has a given requirement. diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index f5ad662f20c53..6ea0f6aa2172f 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -838,13 +838,160 @@ TEST_F(CommandGraphTest, SubGraph) { auto MainGraphExecImpl = sycl::detail::getSyclObjImpl(MainGraphExec); auto Schedule = MainGraphExecImpl->getSchedule(); auto ScheduleIt = Schedule.begin(); - ASSERT_EQ(Schedule.size(), 4ul); + // The schedule list must contain 5 nodes: 4 regulars + 1 empty. + // Indeed an empty node is added as an exit point of the added subgraph to + // facilitate the handling of dependencies + ASSERT_EQ(Schedule.size(), 5ul); ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node1MainGraph)); ScheduleIt++; ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node1Graph))); ScheduleIt++; ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node2Graph))); ScheduleIt++; + ASSERT_TRUE((*ScheduleIt)->isEmpty()); + ScheduleIt++; + ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node3MainGraph)); + ASSERT_EQ(Queue.get_context(), MainGraphExecImpl->getContext()); +} + +TEST_F(CommandGraphTest, SubGraphWithEmptyNode) { + // Add sub-graph with two nodes + auto Node1Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Empty1Graph = + Graph.add([&](sycl::handler &cgh) { /*empty node */ }, + {experimental::property::node::depends_on(Node1Graph)}); + auto Node2Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(Empty1Graph)}); + + auto GraphExec = Graph.finalize(); + + // Add node to main graph followed by sub-graph and another node + experimental::command_graph MainGraph(Queue.get_context(), Dev); + auto Node1MainGraph = MainGraph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Node2MainGraph = + MainGraph.add([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }, + {experimental::property::node::depends_on(Node1MainGraph)}); + auto Node3MainGraph = MainGraph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(Node2MainGraph)}); + + // Assert order of the added sub-graph + ASSERT_NE(sycl::detail::getSyclObjImpl(Node2MainGraph), nullptr); + ASSERT_TRUE(sycl::detail::getSyclObjImpl(Node2MainGraph)->isEmpty()); + // Check the structure of the main graph. + // 1 root connected to 1 successor (the single root of the subgraph) + ASSERT_EQ(sycl::detail::getSyclObjImpl(MainGraph)->MRoots.size(), 1lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.size(), + 1lu); + // Subgraph nodes are duplicated when inserted to parent graph. + // we thus check the node content only. + ASSERT_TRUE( + *(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.front()) == + *(sycl::detail::getSyclObjImpl(Node1Graph))); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.size(), + 1lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node2MainGraph)->MSuccessors.size(), + 1lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MPredecessors.size(), + 0lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node2MainGraph)->MPredecessors.size(), + 1lu); + + // Finalize main graph and check schedule + auto MainGraphExec = MainGraph.finalize(); + auto MainGraphExecImpl = sycl::detail::getSyclObjImpl(MainGraphExec); + auto Schedule = MainGraphExecImpl->getSchedule(); + auto ScheduleIt = Schedule.begin(); + // The schedule list must contain 6 nodes: 5 regulars + 1 empty. + // Indeed an empty node is added as an exit point of the added subgraph to + // facilitate the handling of dependencies + ASSERT_EQ(Schedule.size(), 6ul); + ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node1MainGraph)); + ScheduleIt++; + ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node1Graph))); + ScheduleIt++; + ASSERT_TRUE((*ScheduleIt)->isEmpty()); // empty node inside the subgraph + ScheduleIt++; + ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node2Graph))); + ScheduleIt++; + ASSERT_TRUE( + (*ScheduleIt)->isEmpty()); // empty node added by the impl to handle + // depenendcies w.r.t. the added subgraph + ScheduleIt++; + ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node3MainGraph)); + ASSERT_EQ(Queue.get_context(), MainGraphExecImpl->getContext()); +} + +TEST_F(CommandGraphTest, SubGraphWithEmptyNodeLast) { + // Add sub-graph with two nodes + auto Node1Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Node2Graph = Graph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(Node1Graph)}); + auto EmptyGraph = + Graph.add([&](sycl::handler &cgh) { /*empty node */ }, + {experimental::property::node::depends_on(Node2Graph)}); + + auto GraphExec = Graph.finalize(); + + // Add node to main graph followed by sub-graph and another node + experimental::command_graph MainGraph(Queue.get_context(), Dev); + auto Node1MainGraph = MainGraph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto Node2MainGraph = + MainGraph.add([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }, + {experimental::property::node::depends_on(Node1MainGraph)}); + auto Node3MainGraph = MainGraph.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(Node2MainGraph)}); + + // Assert order of the added sub-graph + ASSERT_NE(sycl::detail::getSyclObjImpl(Node2MainGraph), nullptr); + ASSERT_TRUE(sycl::detail::getSyclObjImpl(Node2MainGraph)->isEmpty()); + // Check the structure of the main graph. + // 1 root connected to 1 successor (the single root of the subgraph) + ASSERT_EQ(sycl::detail::getSyclObjImpl(MainGraph)->MRoots.size(), 1lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.size(), + 1lu); + // Subgraph nodes are duplicated when inserted to parent graph. + // we thus check the node content only. + ASSERT_TRUE( + *(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.front()) == + *(sycl::detail::getSyclObjImpl(Node1Graph))); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.size(), + 1lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node2MainGraph)->MSuccessors.size(), + 1lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MPredecessors.size(), + 0lu); + ASSERT_EQ(sycl::detail::getSyclObjImpl(Node2MainGraph)->MPredecessors.size(), + 1lu); + + // Finalize main graph and check schedule + auto MainGraphExec = MainGraph.finalize(); + auto MainGraphExecImpl = sycl::detail::getSyclObjImpl(MainGraphExec); + auto Schedule = MainGraphExecImpl->getSchedule(); + auto ScheduleIt = Schedule.begin(); + // The schedule list must contain 6 nodes: 5 regulars + 1 empty. + // Indeed an empty node is added as an exit point of the added subgraph to + // facilitate the handling of dependencies + ASSERT_EQ(Schedule.size(), 6ul); + ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node1MainGraph)); + ScheduleIt++; + ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node1Graph))); + ScheduleIt++; + ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node2Graph))); + ScheduleIt++; + ASSERT_TRUE((*ScheduleIt)->isEmpty()); // empty node inside the subgraph + ScheduleIt++; + ASSERT_TRUE( + (*ScheduleIt)->isEmpty()); // empty node added by the impl to handle + // depenendcies w.r.t. the added subgraph + ScheduleIt++; ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node3MainGraph)); ASSERT_EQ(Queue.get_context(), MainGraphExecImpl->getContext()); } @@ -881,7 +1028,10 @@ TEST_F(CommandGraphTest, RecordSubGraph) { auto MainGraphExecImpl = sycl::detail::getSyclObjImpl(MainGraphExec); auto Schedule = MainGraphExecImpl->getSchedule(); auto ScheduleIt = Schedule.begin(); - ASSERT_EQ(Schedule.size(), 4ul); + // The schedule list must contain 5 nodes: 4 regulars + 1 empty. + // Indeed an empty node is added as an exit point of the added subgraph to + // facilitate the handling of dependencies + ASSERT_EQ(Schedule.size(), 5ul); // The first and fourth nodes should have events associated with MainGraph but // not graph. The second and third nodes were added as a sub-graph and @@ -904,6 +1054,9 @@ TEST_F(CommandGraphTest, RecordSubGraph) { ASSERT_ANY_THROW( sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt)); + ScheduleIt++; + ASSERT_TRUE((*ScheduleIt)->isEmpty()); + ScheduleIt++; ASSERT_ANY_THROW( sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt)); @@ -1023,9 +1176,12 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) { auto GraphExecImpl = sycl::detail::getSyclObjImpl(GraphExec); auto Schedule = GraphExecImpl->getSchedule(); auto ScheduleIt = Schedule.begin(); - ASSERT_EQ(Schedule.size(), 2ul); + // the schedule list contains all types of nodes (even empty nodes) + ASSERT_EQ(Schedule.size(), 3ul); ASSERT_EQ(*ScheduleIt, PtrNode1); ScheduleIt++; + ASSERT_TRUE((*ScheduleIt)->isEmpty()); + ScheduleIt++; ASSERT_EQ(*ScheduleIt, PtrNode3); ASSERT_EQ(InOrderQueue.get_context(), GraphExecImpl->getContext()); } @@ -1080,7 +1236,10 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) { auto GraphExecImpl = sycl::detail::getSyclObjImpl(GraphExec); auto Schedule = GraphExecImpl->getSchedule(); auto ScheduleIt = Schedule.begin(); - ASSERT_EQ(Schedule.size(), 2ul); + // the schedule list contains all types of nodes (even empty nodes) + ASSERT_EQ(Schedule.size(), 3ul); + ASSERT_TRUE((*ScheduleIt)->isEmpty()); + ScheduleIt++; ASSERT_EQ(*ScheduleIt, PtrNode2); ScheduleIt++; ASSERT_EQ(*ScheduleIt, PtrNode3); @@ -1137,10 +1296,13 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) { auto GraphExecImpl = sycl::detail::getSyclObjImpl(GraphExec); auto Schedule = GraphExecImpl->getSchedule(); auto ScheduleIt = Schedule.begin(); - ASSERT_EQ(Schedule.size(), 2ul); + // the schedule list contains all types of nodes (even empty nodes) + ASSERT_EQ(Schedule.size(), 3ul); ASSERT_EQ(*ScheduleIt, PtrNode1); ScheduleIt++; ASSERT_EQ(*ScheduleIt, PtrNode2); + ScheduleIt++; + ASSERT_TRUE((*ScheduleIt)->isEmpty()); ASSERT_EQ(InOrderQueue.get_context(), GraphExecImpl->getContext()); }