From 71f5314dbea6ca0763fe9b89a2f3f0d9cad6ff9b Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Mon, 14 Aug 2023 11:15:28 +0100 Subject: [PATCH] [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 | 81 ++++++++--- 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 - .../subgraph_interleaved_submit.cpp | 91 ------------- .../subgraph_two_parent_graphs.cpp | 128 ------------------ sycl/unittests/Extensions/CommandGraph.cpp | 24 ++-- 12 files changed, 132 insertions(+), 266 deletions(-) delete mode 100644 sycl/test-e2e/Graph/RecordReplay/subgraph_interleaved_submit.cpp delete mode 100644 sycl/test-e2e/Graph/RecordReplay/subgraph_two_parent_graphs.cpp diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 4582ded1ab2ce..7d26e61f9c4d8 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -108,6 +108,31 @@ bool visitNodeDepthFirst( NodeStack.pop_back(); return false; } + +/// 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() { @@ -125,7 +150,7 @@ graph_impl::~graph_impl() { } } -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; @@ -148,6 +173,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 114022fa6cd9e..7b961e78d85af 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -92,6 +92,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. @@ -293,26 +333,16 @@ class node_impl { /// Tests is the caller is similar to Node /// @return True if the two nodes are similars bool isSimilar(std::shared_ptr Node) { - if (MCGType != Node->MCGType) - return false; - if (MSuccessors.size() != Node->MSuccessors.size()) return false; if (MPredecessors.size() != Node->MPredecessors.size()) return false; - if ((MCGType == sycl::detail::CG::CGTYPE::Kernel) && - (MCGType == sycl::detail::CG::CGTYPE::Kernel)) { - sycl::detail::CGExecKernel *ExecKernelA = - static_cast(MCommandGroup.get()); - sycl::detail::CGExecKernel *ExecKernelB = - static_cast(Node->MCommandGroup.get()); + if (*this == *Node.get()) + return true; - if (ExecKernelA->MKernelName.compare(ExecKernelB->MKernelName) != 0) - return false; - } - return true; + return false; } /// Recursive traversal of successor nodes checking for @@ -485,11 +515,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. @@ -666,10 +697,6 @@ class graph_impl { /// @return True if a cycle is detected, false if not. bool checkForCycles(); - /// Insert node into list of root nodes. - /// @param Root Node to add to list of root nodes. - void addRoot(const std::shared_ptr &Root); - /// Context associated with this graph. sycl::context MContext; /// Device associated with this graph. All graph nodes will execute on this @@ -703,6 +730,16 @@ class graph_impl { /// Controls whether we allow buffers to be used in the graph. Set by the /// presence of the assume_buffer_outlives_graph property. bool MAllowBuffers = false; + + /// 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. @@ -754,6 +791,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; } + /// Prints the contents of the graph to a text file in DOT format /// @param GraphName is a string appended to the output file name void printGraphAsDot(const std::string GraphName) { diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 1181947c70cd6..a31b9ded69dfd 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -1367,7 +1367,9 @@ void handler::ext_oneapi_graph( } // 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/test-e2e/Graph/RecordReplay/subgraph_interleaved_submit.cpp b/sycl/test-e2e/Graph/RecordReplay/subgraph_interleaved_submit.cpp deleted file mode 100644 index cb0c10aba3d40..0000000000000 --- a/sycl/test-e2e/Graph/RecordReplay/subgraph_interleaved_submit.cpp +++ /dev/null @@ -1,91 +0,0 @@ -// REQUIRES: level_zero, gpu -// RUN: %{build} -o %t.out -// RUN: %{run} %t.out -// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %} -// -// CHECK-NOT: LEAK - -// XFAIL:* -// Submit a graph as a subgraph more than once doesn't yet work. - -// Tests creating a parent graph with the same sub-graph interleaved with -// other nodes. -// The second run is to check that there are no leaks reported with the embedded -// ZE_DEBUG=4 testing capability. - -#include "../graph_common.hpp" - -int main() { - queue Queue{gpu_selector_v}; - - exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; - exp_ext::command_graph SubGraph{Queue.get_context(), Queue.get_device()}; - - const size_t N = 10; - float *X = malloc_device(N, Queue); - - SubGraph.begin_recording(Queue); - - auto S1 = Queue.submit([&](handler &CGH) { - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] *= 2.0f; - }); - }); - - Queue.submit([&](handler &CGH) { - CGH.depends_on(S1); - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] += 0.5f; - }); - }); - - SubGraph.end_recording(Queue); - - auto ExecSubGraph = SubGraph.finalize(); - - Graph.begin_recording(Queue); - - auto P1 = Queue.submit([&](handler &CGH) { - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] = 1.0f; - }); - }); - - auto P2 = Queue.submit([&](handler &CGH) { - CGH.depends_on(P1); - CGH.ext_oneapi_graph(ExecSubGraph); - }); - - auto P3 = Queue.submit([&](handler &CGH) { - CGH.depends_on(P2); - CGH.parallel_for(range<1>{N}, [=](id<1> it) { - const size_t i = it[0]; - X[i] *= -1.0f; - }); - }); - - Queue.submit([&](handler &CGH) { - CGH.depends_on(P3); - CGH.ext_oneapi_graph(ExecSubGraph); - }); - - Graph.end_recording(); - - auto ExecGraph = Graph.finalize(); - - auto E = Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }); - - std::vector Output(N); - Queue.memcpy(Output.data(), X, N * sizeof(float), E).wait(); - - for (size_t i = 0; i < N; i++) { - assert(Output[i] == -6.25f); - } - - sycl::free(X, Queue); - - return 0; -} diff --git a/sycl/test-e2e/Graph/RecordReplay/subgraph_two_parent_graphs.cpp b/sycl/test-e2e/Graph/RecordReplay/subgraph_two_parent_graphs.cpp deleted file mode 100644 index 5e2abe0dbf0f9..0000000000000 --- a/sycl/test-e2e/Graph/RecordReplay/subgraph_two_parent_graphs.cpp +++ /dev/null @@ -1,128 +0,0 @@ -// REQUIRES: level_zero, gpu -// RUN: %{build} -o %t.out -// RUN: %{run} %t.out -// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %} -// -// CHECK-NOT: LEAK - -// XFAIL: * -// Subgraph doesn't work properly in second parent graph - -// Tests adding an executable graph object as a sub-graph of two different -// parent graphs. -// The second run is to check that there are no leaks reported with the embedded -// ZE_DEBUG=4 testing capability. - -#include "../graph_common.hpp" - -int main() { - queue Queue{gpu_selector_v}; - - exp_ext::command_graph GraphA{Queue.get_context(), Queue.get_device()}; - exp_ext::command_graph GraphB{Queue.get_context(), Queue.get_device()}; - exp_ext::command_graph SubGraph{Queue.get_context(), Queue.get_device()}; - - const size_t N = 10; - float *X = malloc_device(N, Queue); - - SubGraph.begin_recording(Queue); - - auto S1 = Queue.submit([&](handler &CGH) { - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] *= 2.0f; - }); - }); - - Queue.submit([&](handler &CGH) { - CGH.depends_on(S1); - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] += 0.5f; - }); - }); - - SubGraph.end_recording(Queue); - - auto ExecSubGraph = SubGraph.finalize(); - - GraphA.begin_recording(Queue); - - auto A1 = Queue.submit([&](handler &CGH) { - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] = 1.0f; - }); - }); - - auto A2 = Queue.submit([&](handler &CGH) { - CGH.depends_on(A1); - CGH.ext_oneapi_graph(ExecSubGraph); - }); - - Queue.submit([&](handler &CGH) { - CGH.depends_on(A2); - CGH.parallel_for(range<1>{N}, [=](id<1> it) { - const size_t i = it[0]; - X[i] *= -1.0f; - }); - }); - - GraphA.end_recording(); - - auto ExecGraphA = GraphA.finalize(); - - GraphB.begin_recording(Queue); - - auto B1 = Queue.submit([&](handler &CGH) { - CGH.parallel_for(N, [=](id<1> it) { - const size_t i = it[0]; - X[i] = static_cast(i); - }); - }); - - auto B2 = Queue.submit([&](handler &CGH) { - CGH.depends_on(B1); - CGH.ext_oneapi_graph(ExecSubGraph); - }); - - Queue.submit([&](handler &CGH) { - CGH.depends_on(B2); - CGH.parallel_for(range<1>{N}, [=](id<1> it) { - const size_t i = it[0]; - X[i] *= X[i]; - }); - }); - - GraphB.end_recording(); - auto ExecGraphB = GraphB.finalize(); - - auto EventA1 = - Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraphA); }); - std::vector OutputA(N); - auto EventA2 = Queue.memcpy(OutputA.data(), X, N * sizeof(float), EventA1); - - auto EventB1 = Queue.submit([&](handler &CGH) { - CGH.depends_on(EventA2); - CGH.ext_oneapi_graph(ExecGraphB); - }); - std::vector OutputB(N); - Queue.memcpy(OutputB.data(), X, N * sizeof(float), EventB1); - Queue.wait(); - - auto refB = [](size_t i) { - float result = static_cast(i); - result = result * 2.0f + 0.5f; - result *= result; - return result; - }; - - for (size_t i = 0; i < N; i++) { - assert(OutputA[i] == -2.5f); - assert(OutputB[i] == refB(i)); - } - - sycl::free(X, Queue); - - return 0; -} diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 6c8e7ce2dd3e5..f5ad662f20c53 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -548,7 +548,6 @@ bool checkExecGraphSchedule( } return true; } - } // anonymous namespace class CommandGraphTest : public ::testing::Test { @@ -822,8 +821,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(), @@ -839,9 +841,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()); @@ -882,8 +884,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( @@ -893,14 +895,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(