From 064997de1be85506cea5c1a7c3787d9e360581e3 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Mon, 14 Aug 2023 11:15:28 +0100 Subject: [PATCH 1/6] [SYCL][Graph] Duplicate sub-graph nodes (#292) * [SYCL][Graph] Duplicate sub-graph nodes Duplicates the nodes of a sub-graph when added to its parents to enable multiple parents. As a result nodes of the initial sub-graph and the one of the main graph are no longer the exact same (shared_pointer) but two different nodes with similar content. A few unitests have been updated accordingly and an equal operator has been added to node_impl to ease comparison between nodes. --- sycl/source/detail/graph_impl.cpp | 54 ++++++++++++- sycl/source/detail/graph_impl.hpp | 80 +++++++++++++++++-- sycl/source/handler.cpp | 4 +- .../sub_graph_multiple_submission.cpp | 3 - .../Explicit/sub_graph_two_parent_graphs.cpp | 3 - .../Inputs/sub_graph_multiple_submission.cpp | 2 +- .../Inputs/sub_graph_two_parent_graphs.cpp | 2 +- .../sub_graph_multiple_submission.cpp | 3 - .../sub_graph_two_parent_graphs.cpp | 3 - sycl/unittests/Extensions/CommandGraph.cpp | 23 +++--- 10 files changed, 144 insertions(+), 33 deletions(-) diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 1a4f07285fe4b..79a9d2ba73a8c 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -71,6 +71,31 @@ bool checkForRequirement(sycl::detail::AccessorImplHost *Req, } return SuccessorAddedDep; } + +/// Recursively append CurrentNode to Outputs if a given node is an exit node +/// @param[in] CurrentNode Node to check as exit node. +/// @param[inout] Outputs list of exit nodes. +void appendExitNodesFromRoot(std::shared_ptr CurrentNode, + std::vector> &Outputs) { + if (CurrentNode->MSuccessors.size() > 0) { + for (auto Successor : CurrentNode->MSuccessors) { + appendExitNodesFromRoot(Successor, Outputs); + } + } else { + Outputs.push_back(CurrentNode); + } +} + +void duplicateNode(const std::shared_ptr Node, + std::shared_ptr &NodeCopy) { + if (Node->MCGType == sycl::detail::CG::None) { + NodeCopy = std::make_shared(); + NodeCopy->MCGType = sycl::detail::CG::None; + } else { + NodeCopy = std::make_shared(Node->MCGType, Node->getCGCopy()); + } +} + } // anonymous namespace void exec_graph_impl::schedule() { @@ -81,7 +106,7 @@ void exec_graph_impl::schedule() { } } -std::shared_ptr graph_impl::addSubgraphNodes( +std::shared_ptr graph_impl::addNodesToExits( const std::list> &NodeList) { // Find all input and output nodes from the node list std::vector> Inputs; @@ -104,6 +129,33 @@ std::shared_ptr graph_impl::addSubgraphNodes( return this->add(Outputs); } +std::shared_ptr graph_impl::addSubgraphNodes( + const std::shared_ptr &SubGraphExec) { + std::map, std::shared_ptr> NodesMap; + std::list> NewNodeList; + + std::list> NodeList = SubGraphExec->getSchedule(); + + for (std::list>::const_iterator NodeIt = + NodeList.end(); + NodeIt != NodeList.begin();) { + --NodeIt; + auto Node = *NodeIt; + std::shared_ptr NodeCopy; + duplicateNode(Node, NodeCopy); + NewNodeList.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); + } + } + } + + return addNodesToExits(NewNodeList); +} + void graph_impl::addRoot(const std::shared_ptr &Root) { MRoots.insert(Root); } diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index 5526aeaccef44..3c2629ab9ee8a 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -72,6 +72,46 @@ class node_impl { std::unique_ptr &&CommandGroup) : MCGType(CGType), MCommandGroup(std::move(CommandGroup)) {} + /// Tests if two nodes have the same content, + /// i.e. same command group + /// @param Node node to compare with + bool operator==(const node_impl &Node) { + if (MCGType != Node.MCGType) + return false; + + if (MCGType == sycl::detail::CG::CGTYPE::Kernel) { + sycl::detail::CGExecKernel *ExecKernelA = + static_cast(MCommandGroup.get()); + sycl::detail::CGExecKernel *ExecKernelB = + static_cast(Node.MCommandGroup.get()); + + if (ExecKernelA->MKernelName.compare(ExecKernelB->MKernelName) != 0) + return false; + } + if (MCGType == sycl::detail::CG::CGTYPE::CopyUSM) { + sycl::detail::CGCopyUSM *CopyA = + static_cast(MCommandGroup.get()); + sycl::detail::CGCopyUSM *CopyB = + static_cast(MCommandGroup.get()); + if ((CopyA->getSrc() != CopyB->getSrc()) || + (CopyA->getDst() != CopyB->getDst()) || + (CopyA->getLength() == CopyB->getLength())) + return false; + } + if ((MCGType == sycl::detail::CG::CGTYPE::CopyAccToAcc) || + (MCGType == sycl::detail::CG::CGTYPE::CopyAccToPtr) || + (MCGType == sycl::detail::CG::CGTYPE::CopyPtrToAcc)) { + sycl::detail::CGCopy *CopyA = + static_cast(MCommandGroup.get()); + sycl::detail::CGCopy *CopyB = + static_cast(MCommandGroup.get()); + if ((CopyA->getSrc() != CopyB->getSrc()) || + (CopyA->getDst() != CopyB->getDst())) + return false; + } + return true; + } + /// Recursively add nodes to execution stack. /// @param NodeImpl Node to schedule. /// @param Schedule Execution ordering to add node to. @@ -167,6 +207,21 @@ class node_impl { return nullptr; } + /// Tests is the caller is similar to Node + /// @return True if the two nodes are similar + bool isSimilar(std::shared_ptr Node) { + if (MSuccessors.size() != Node->MSuccessors.size()) + return false; + + if (MPredecessors.size() != Node->MPredecessors.size()) + return false; + + if (*this == *Node.get()) + return true; + + return false; + } + private: /// Creates a copy of the node's CG by casting to it's actual type, then using /// that to copy construct and create a new unique ptr from that copy. @@ -187,10 +242,6 @@ class graph_impl { : MContext(SyclContext), MDevice(SyclDevice), MRecordingQueues(), MEventsMap(), MInorderQueueMap() {} - /// Insert node into list of root nodes. - /// @param Root Node to add to list of root nodes. - void addRoot(const std::shared_ptr &Root); - /// Remove node from list of root nodes. /// @param Root Node to remove from list of root nodes. void removeRoot(const std::shared_ptr &Root); @@ -276,11 +327,12 @@ class graph_impl { "No event has been recorded for the specified graph node"); } - /// Adds sub-graph nodes from an executable graph to this graph. - /// @param NodeList List of nodes from sub-graph in schedule order. + /// Duplicates and Adds sub-graph nodes from an executable graph to this + /// graph. + /// @param SubGraphExec sub-graph to add to the parent. /// @return An empty node is used to schedule dependencies on this sub-graph. std::shared_ptr - addSubgraphNodes(const std::list> &NodeList); + addSubgraphNodes(const std::shared_ptr &SubGraphExec); /// Query for the context tied to this graph. /// @return Context associated with graph. @@ -333,6 +385,16 @@ class graph_impl { std::map, std::shared_ptr, std::owner_less>> MInorderQueueMap; + + /// Insert node into list of root nodes. + /// @param Root Node to add to list of root nodes. + void addRoot(const std::shared_ptr &Root); + + /// Adds nodes to the exit nodes of this graph. + /// @param NodeList List of nodes from sub-graph in schedule order. + /// @return An empty node is used to schedule dependencies on this sub-graph. + std::shared_ptr + addNodesToExits(const std::list> &NodeList); }; /// Class representing the implementation of command_graph. @@ -378,6 +440,10 @@ class exec_graph_impl { return MSchedule; } + /// Query the graph_impl. + /// @return pointer to the graph_impl MGraphImpl + const std::shared_ptr &getGraphImpl() const { return MGraphImpl; } + private: /// Create a command-group for the node and add it to command-buffer by going /// through the scheduler. diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 95db5c5eb66af..b9c13949acdb9 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -1313,7 +1313,9 @@ void handler::ext_oneapi_graph( if (ParentGraph) { // Store the node representing the subgraph in the handler so that we can // return it to the user later. - MSubgraphNode = ParentGraph->addSubgraphNodes(GraphImpl->getSchedule()); + // The nodes of the subgraph are duplicated when added to its parents. + // This avoids changing properties of the graph added as a subgraph. + MSubgraphNode = ParentGraph->addSubgraphNodes(GraphImpl); // If we are recording an in-order queue remember the subgraph node, so it // can be used as a dependency for any more nodes recorded from this queue. diff --git a/sycl/test-e2e/Graph/Explicit/sub_graph_multiple_submission.cpp b/sycl/test-e2e/Graph/Explicit/sub_graph_multiple_submission.cpp index 14e447c04104f..b8863b57c7290 100644 --- a/sycl/test-e2e/Graph/Explicit/sub_graph_multiple_submission.cpp +++ b/sycl/test-e2e/Graph/Explicit/sub_graph_multiple_submission.cpp @@ -6,9 +6,6 @@ // // CHECK-NOT: LEAK -// XFAIL:* -// Submit a graph as a subgraph more than once doesn't yet work. - #define GRAPH_E2E_EXPLICIT #include "../Inputs/sub_graph_multiple_submission.cpp" diff --git a/sycl/test-e2e/Graph/Explicit/sub_graph_two_parent_graphs.cpp b/sycl/test-e2e/Graph/Explicit/sub_graph_two_parent_graphs.cpp index cbf768203ec5a..4254a861fe344 100644 --- a/sycl/test-e2e/Graph/Explicit/sub_graph_two_parent_graphs.cpp +++ b/sycl/test-e2e/Graph/Explicit/sub_graph_two_parent_graphs.cpp @@ -6,9 +6,6 @@ // // CHECK-NOT: LEAK -// XFAIL: * -// Subgraph doesn't work properly in second parent graph - #define GRAPH_E2E_EXPLICIT #include "../Inputs/sub_graph_two_parent_graphs.cpp" diff --git a/sycl/test-e2e/Graph/Inputs/sub_graph_multiple_submission.cpp b/sycl/test-e2e/Graph/Inputs/sub_graph_multiple_submission.cpp index 81ad495da5268..a3a3f86b04895 100644 --- a/sycl/test-e2e/Graph/Inputs/sub_graph_multiple_submission.cpp +++ b/sycl/test-e2e/Graph/Inputs/sub_graph_multiple_submission.cpp @@ -62,7 +62,7 @@ int main() { Queue.memcpy(Output.data(), X, N * sizeof(float), E).wait(); for (size_t i = 0; i < N; i++) { - assert(Output[i] == -6.25f); + assert(Output[i] == -4.5f); } sycl::free(X, Queue); diff --git a/sycl/test-e2e/Graph/Inputs/sub_graph_two_parent_graphs.cpp b/sycl/test-e2e/Graph/Inputs/sub_graph_two_parent_graphs.cpp index 5c9732f4cc509..91ced1c25a9b0 100644 --- a/sycl/test-e2e/Graph/Inputs/sub_graph_two_parent_graphs.cpp +++ b/sycl/test-e2e/Graph/Inputs/sub_graph_two_parent_graphs.cpp @@ -50,7 +50,7 @@ int main() { auto ExecGraphA = GraphA.finalize(); auto B1 = add_node(GraphB, Queue, [&](handler &CGH) { - CGH.parallel_for(N, [=](id<1> it) { X[it] = static_cast(i); }); + CGH.parallel_for(N, [=](id<1> it) { X[it] = static_cast(it); }); }); auto B2 = add_node( diff --git a/sycl/test-e2e/Graph/RecordReplay/sub_graph_multiple_submission.cpp b/sycl/test-e2e/Graph/RecordReplay/sub_graph_multiple_submission.cpp index b8f4bfa3b3b83..aedd9e252e252 100644 --- a/sycl/test-e2e/Graph/RecordReplay/sub_graph_multiple_submission.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/sub_graph_multiple_submission.cpp @@ -6,9 +6,6 @@ // // CHECK-NOT: LEAK -// XFAIL:* -// Submit a graph as a subgraph more than once doesn't yet work. - #define GRAPH_E2E_RECORD_REPLAY #include "../Inputs/sub_graph_multiple_submission.cpp" diff --git a/sycl/test-e2e/Graph/RecordReplay/sub_graph_two_parent_graphs.cpp b/sycl/test-e2e/Graph/RecordReplay/sub_graph_two_parent_graphs.cpp index 2d2eb17cd7078..6eb49c39583da 100644 --- a/sycl/test-e2e/Graph/RecordReplay/sub_graph_two_parent_graphs.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/sub_graph_two_parent_graphs.cpp @@ -6,9 +6,6 @@ // // CHECK-NOT: LEAK -// XFAIL: * -// Subgraph doesn't work properly in second parent graph - #define GRAPH_E2E_RECORD_REPLAY #include "../Inputs/sub_graph_two_parent_graphs.cpp" diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 590d86b8e0019..f7a1fc3b6bbac 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -287,8 +287,11 @@ TEST_F(CommandGraphTest, SubGraph) { ASSERT_EQ(sycl::detail::getSyclObjImpl(MainGraph)->MRoots.size(), 1lu); ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.size(), 1lu); - ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MSuccessors.front(), - sycl::detail::getSyclObjImpl(Node1Graph)); + // 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(Node2MainGraph)->MSuccessors.size(), 1lu); ASSERT_EQ(sycl::detail::getSyclObjImpl(Node1MainGraph)->MPredecessors.size(), @@ -304,9 +307,9 @@ TEST_F(CommandGraphTest, SubGraph) { ASSERT_EQ(Schedule.size(), 4ul); ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node1MainGraph)); ScheduleIt++; - ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node1Graph)); + ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node1Graph))); ScheduleIt++; - ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node2Graph)); + ASSERT_TRUE(*(*ScheduleIt) == *(sycl::detail::getSyclObjImpl(Node2Graph))); ScheduleIt++; ASSERT_EQ(*ScheduleIt, sycl::detail::getSyclObjImpl(Node3MainGraph)); ASSERT_EQ(Queue.get_context(), MainGraphExecImpl->getContext()); @@ -347,8 +350,8 @@ TEST_F(CommandGraphTest, RecordSubGraph) { ASSERT_EQ(Schedule.size(), 4ul); // 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 should - // have events associated with Graph but not MainGraph. + // not graph. The second and third nodes were added as a sub-graph and + // duplicated. They should not have events associated with Graph or MainGraph. ASSERT_ANY_THROW( sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt)); ASSERT_EQ( @@ -358,14 +361,14 @@ TEST_F(CommandGraphTest, RecordSubGraph) { ScheduleIt++; ASSERT_ANY_THROW( sycl::detail::getSyclObjImpl(MainGraph)->getEventForNode(*ScheduleIt)); - ASSERT_EQ(sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt), - sycl::detail::getSyclObjImpl(Node1Graph)); + ASSERT_ANY_THROW( + sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt)); ScheduleIt++; ASSERT_ANY_THROW( sycl::detail::getSyclObjImpl(MainGraph)->getEventForNode(*ScheduleIt)); - ASSERT_EQ(sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt), - sycl::detail::getSyclObjImpl(Node2Graph)); + ASSERT_ANY_THROW( + sycl::detail::getSyclObjImpl(Graph)->getEventForNode(*ScheduleIt)); ScheduleIt++; ASSERT_ANY_THROW( From 824096f27c74c9933d523b97cf38b07fda0dadba Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Wed, 16 Aug 2023 17:01:00 +0100 Subject: [PATCH 2/6] [SYCL][Graph] Enable empty nodes in Subgraphs (#300) * [SYCL][Graph] Enable empty nodes in Subgraphs The implementation uses the list of scheduled nodes to add a subgraph to a main graph. However, since empty nodes are not scheduled, empty nodes were not listed in this list, resulting in inconsistent graphs when subgraph with empty node(s) were added to a main graph. This PR fixes this issue by forcing to list empty nodes when creating the list for inserting subgraph. It also adds unitests to check that subgraphs with empty nodes are correctly added to a main graph. * [SYCL][Graph] Enable empty nodes in Subgraphs Changes the definition of MSchedule as it now contains all types of nodes (including empty nodes). Empty nodes are now filtered out when creating the commandbuffer and/or enqueuing nodes to only keep their dependencies. This PR updates a few unitests to make them compliant to the new schedule list definition. * Update sycl/source/detail/graph_impl.cpp Co-authored-by: Ben Tracy --------- Co-authored-by: Ben Tracy --- sycl/source/detail/graph_impl.cpp | 25 ++- sycl/source/detail/graph_impl.hpp | 6 +- sycl/unittests/Extensions/CommandGraph.cpp | 172 ++++++++++++++++++++- 3 files changed, 186 insertions(+), 17 deletions(-) diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 79a9d2ba73a8c..fbb031c77e6d0 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -132,28 +132,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) { @@ -363,6 +366,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. @@ -490,8 +498,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 3c2629ab9ee8a..6f81ecddfdc6c 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -122,10 +122,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 f7a1fc3b6bbac..45920cff0e573 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -304,13 +304,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()); } @@ -347,7 +494,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 @@ -370,6 +520,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)); @@ -489,9 +642,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()); } @@ -546,7 +702,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); @@ -603,9 +762,12 @@ 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()); } From b3cf214cfdc8a60ec8e64d37d091bfbacb4bc94d Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 16 Aug 2023 17:21:02 +0100 Subject: [PATCH 3/6] Use TestKernel helper in unittests --- sycl/unittests/Extensions/CommandGraph.cpp | 47 +++++++++++----------- 1 file changed, 24 insertions(+), 23 deletions(-) diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 45920cff0e573..f8b40022060cd 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include @@ -57,7 +58,7 @@ TEST_F(CommandGraphTest, AddNode) { ASSERT_TRUE(GraphImpl->MRoots.empty()); auto Node1 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); ASSERT_NE(sycl::detail::getSyclObjImpl(Node1), nullptr); ASSERT_FALSE(sycl::detail::getSyclObjImpl(Node1)->isEmpty()); ASSERT_EQ(GraphImpl->MRoots.size(), 1lu); @@ -112,17 +113,17 @@ TEST_F(CommandGraphTest, Finalize) { sycl::buffer Buf(1); auto Node1 = Graph.add([&](sycl::handler &cgh) { sycl::accessor A(Buf, cgh, sycl::write_only, sycl::no_init); - cgh.single_task([=]() { A[0] = 1; }); + cgh.single_task>([]() {}); }); // Add independent node auto Node2 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); // Add a node that depends on Node1 due to the accessor auto Node3 = Graph.add([&](sycl::handler &cgh) { sycl::accessor A(Buf, cgh, sycl::write_only, sycl::no_init); - cgh.single_task([=]() { A[0] = 3; }); + cgh.single_task>([]() {}); }); // Guarantee order of independent nodes 1 and 2 @@ -148,7 +149,7 @@ TEST_F(CommandGraphTest, MakeEdge) { // Add two independent nodes auto Node1 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2 = Graph.add([&](sycl::handler &cgh) {}); ASSERT_EQ(GraphImpl->MRoots.size(), 2ul); ASSERT_TRUE(sycl::detail::getSyclObjImpl(Node1)->MSuccessors.empty()); @@ -242,7 +243,7 @@ TEST_F(CommandGraphTest, BeginEndRecording) { TEST_F(CommandGraphTest, GetCGCopy) { auto Node1 = Graph.add([&](sycl::handler &cgh) {}); auto Node2 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }, + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, {experimental::property::node::depends_on(Node1)}); // Get copy of CG of Node2 and check equality @@ -264,21 +265,21 @@ TEST_F(CommandGraphTest, GetCGCopy) { TEST_F(CommandGraphTest, SubGraph) { // Add sub-graph with two nodes auto Node1Graph = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2Graph = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }, + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, {experimental::property::node::depends_on(Node1Graph)}); 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([]() {}); }); + [&](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([]() {}); }, + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, {experimental::property::node::depends_on(Node2MainGraph)}); // Assert order of the added sub-graph @@ -466,10 +467,10 @@ TEST_F(CommandGraphTest, RecordSubGraph) { // Record sub-graph with two nodes Graph.begin_recording(Queue); auto Node1Graph = Queue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2Graph = Queue.submit([&](sycl::handler &cgh) { cgh.depends_on(Node1Graph); - cgh.single_task([]() {}); + cgh.single_task>([]() {}); }); Graph.end_recording(Queue); auto GraphExec = Graph.finalize(); @@ -478,14 +479,14 @@ TEST_F(CommandGraphTest, RecordSubGraph) { experimental::command_graph MainGraph(Queue.get_context(), Dev); MainGraph.begin_recording(Queue); auto Node1MainGraph = Queue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2MainGraph = Queue.submit([&](handler &cgh) { cgh.depends_on(Node1MainGraph); cgh.ext_oneapi_graph(GraphExec); }); auto Node3MainGraph = Queue.submit([&](sycl::handler &cgh) { cgh.depends_on(Node2MainGraph); - cgh.single_task([]() {}); + cgh.single_task>([]() {}); }); MainGraph.end_recording(Queue); @@ -541,7 +542,7 @@ TEST_F(CommandGraphTest, InOrderQueue) { // Record in-order queue with three nodes InOrderGraph.begin_recording(InOrderQueue); auto Node1Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode1 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -550,7 +551,7 @@ TEST_F(CommandGraphTest, InOrderQueue) { ASSERT_TRUE(PtrNode1->MPredecessors.empty()); auto Node2Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode2 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -563,7 +564,7 @@ TEST_F(CommandGraphTest, InOrderQueue) { ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1); auto Node3Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode3 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -601,7 +602,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) { // node InOrderGraph.begin_recording(InOrderQueue); auto Node1Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode1 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -622,7 +623,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) { ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1); auto Node3Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode3 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -669,7 +670,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) { ASSERT_TRUE(PtrNode1->MPredecessors.empty()); auto Node2Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode2 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -682,7 +683,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) { ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1); auto Node3Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode3 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -721,7 +722,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) { // Record in-order queue with two regular nodes then an empty node InOrderGraph.begin_recording(InOrderQueue); auto Node1Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode1 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -730,7 +731,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) { ASSERT_TRUE(PtrNode1->MPredecessors.empty()); auto Node2Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode2 = sycl::detail::getSyclObjImpl(InOrderGraph) From a722fdcfdf4eae42c8f25a4649937083d186d0d1 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Tue, 22 Aug 2023 09:29:16 +0100 Subject: [PATCH 4/6] [SYCL][Graph] Improves node_impl equal operator and new nodes list allocation policy (#303) --- sycl/source/detail/graph_impl.cpp | 7 ++++-- sycl/source/detail/graph_impl.hpp | 37 ++++++++++++++++++------------- 2 files changed, 26 insertions(+), 18 deletions(-) diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index fbb031c77e6d0..32ec162917370 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -132,19 +132,22 @@ 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> NewNodesList; std::list> NodesList = SubGraphExec->getSchedule(); + std::list> NewNodesList{NodesList.size()}; // Duplication of nodes + std::list>::iterator NewNodesIt = + NewNodesList.end(); for (std::list>::const_iterator NodeIt = NodesList.end(); NodeIt != NodesList.begin();) { --NodeIt; + --NewNodesIt; auto Node = *NodeIt; std::shared_ptr NodeCopy; duplicateNode(Node, NodeCopy); - NewNodesList.push_back(NodeCopy); + *NewNodesIt = NodeCopy; NodesMap.insert({Node, NodeCopy}); for (auto &NextNode : Node->MSuccessors) { if (NodesMap.find(NextNode) != NodesMap.end()) { diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index 6f81ecddfdc6c..ea880f7d1392c 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -74,42 +74,47 @@ class node_impl { /// Tests if two nodes have the same content, /// i.e. same command group + /// This function should only be used for internal purposes. + /// A true return from this operator is not a guarantee that the nodes are + /// equals according to the Common reference semantics. But this function is + /// an helper to verify that two nodes contain equivalent Command Groups. /// @param Node node to compare with + /// @return true if two nodes have equivament command groups. false otherwise. bool operator==(const node_impl &Node) { if (MCGType != Node.MCGType) return false; - if (MCGType == sycl::detail::CG::CGTYPE::Kernel) { + switch (MCGType) { + case sycl::detail::CG::CGTYPE::Kernel: { sycl::detail::CGExecKernel *ExecKernelA = static_cast(MCommandGroup.get()); sycl::detail::CGExecKernel *ExecKernelB = static_cast(Node.MCommandGroup.get()); - - if (ExecKernelA->MKernelName.compare(ExecKernelB->MKernelName) != 0) - return false; + return ExecKernelA->MKernelName.compare(ExecKernelB->MKernelName) == 0; } - if (MCGType == sycl::detail::CG::CGTYPE::CopyUSM) { + case sycl::detail::CG::CGTYPE::CopyUSM: { sycl::detail::CGCopyUSM *CopyA = static_cast(MCommandGroup.get()); sycl::detail::CGCopyUSM *CopyB = static_cast(MCommandGroup.get()); - if ((CopyA->getSrc() != CopyB->getSrc()) || - (CopyA->getDst() != CopyB->getDst()) || - (CopyA->getLength() == CopyB->getLength())) - return false; + return (CopyA->getSrc() == CopyB->getSrc()) && + (CopyA->getDst() == CopyB->getDst()) && + (CopyA->getLength() == CopyB->getLength()); } - if ((MCGType == sycl::detail::CG::CGTYPE::CopyAccToAcc) || - (MCGType == sycl::detail::CG::CGTYPE::CopyAccToPtr) || - (MCGType == sycl::detail::CG::CGTYPE::CopyPtrToAcc)) { + case sycl::detail::CG::CGTYPE::CopyAccToAcc: + case sycl::detail::CG::CGTYPE::CopyAccToPtr: + case sycl::detail::CG::CGTYPE::CopyPtrToAcc: { sycl::detail::CGCopy *CopyA = static_cast(MCommandGroup.get()); sycl::detail::CGCopy *CopyB = static_cast(MCommandGroup.get()); - if ((CopyA->getSrc() != CopyB->getSrc()) || - (CopyA->getDst() != CopyB->getDst())) - return false; + return (CopyA->getSrc() == CopyB->getSrc()) && + (CopyA->getDst() == CopyB->getDst()); + } + default: + assert(false && "Unexpected command group type!"); + return false; } - return true; } /// Recursively add nodes to execution stack. From 7c1c5df3b774309213d3f3739e694fcf43fd4c43 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 24 Aug 2023 17:19:26 +0100 Subject: [PATCH 5/6] Use `auto` for node list iterators --- sycl/source/detail/graph_impl.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index b77c941f42a83..fefb06a6280cf 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -137,10 +137,7 @@ std::shared_ptr graph_impl::addSubgraphNodes( std::list> NewNodesList{NodesList.size()}; // Duplication of nodes - std::list>::iterator NewNodesIt = - NewNodesList.end(); - for (std::list>::const_iterator NodeIt = - NodesList.end(); + for (auto NodeIt = NodesList.end(), NewNodesIt = NewNodesList.end(); NodeIt != NodesList.begin();) { --NodeIt; --NewNodesIt; From 0d65ddb89bdb071672503b80e522059efda4f439 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 24 Aug 2023 17:50:18 +0100 Subject: [PATCH 6/6] Remove unused function --- sycl/source/detail/graph_impl.cpp | 14 -------------- 1 file changed, 14 deletions(-) diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index fefb06a6280cf..aa51eb9d0c44c 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -72,20 +72,6 @@ bool checkForRequirement(sycl::detail::AccessorImplHost *Req, return SuccessorAddedDep; } -/// Recursively append CurrentNode to Outputs if a given node is an exit node -/// @param[in] CurrentNode Node to check as exit node. -/// @param[inout] Outputs list of exit nodes. -void appendExitNodesFromRoot(std::shared_ptr CurrentNode, - std::vector> &Outputs) { - if (CurrentNode->MSuccessors.size() > 0) { - for (auto Successor : CurrentNode->MSuccessors) { - appendExitNodesFromRoot(Successor, Outputs); - } - } else { - Outputs.push_back(CurrentNode); - } -} - void duplicateNode(const std::shared_ptr Node, std::shared_ptr &NodeCopy) { if (Node->MCGType == sycl::detail::CG::None) {