[SYCL][Graph] Duplicate sub-graph nodes

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<node_impl> CurrentNode,
+                             std::vector<std::shared_ptr<node_impl>> &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_impl> Node,
+                   std::shared_ptr<node_impl> &NodeCopy) {
+  if (Node->MCGType == sycl::detail::CG::None) {
+    NodeCopy = std::make_shared<node_impl>();
+    NodeCopy->MCGType = sycl::detail::CG::None;
+  } else {
+    NodeCopy = std::make_shared<node_impl>(Node->MCGType, Node->getCGCopy());
+  }
+}
+
 } // anonymous namespace
 
 void exec_graph_impl::schedule() {
@@ -125,7 +150,7 @@ graph_impl::~graph_impl() {
   }
 }
 
-std::shared_ptr<node_impl> graph_impl::addSubgraphNodes(
+std::shared_ptr<node_impl> graph_impl::addNodesToExits(
     const std::list<std::shared_ptr<node_impl>> &NodeList) {
   // Find all input and output nodes from the node list
   std::vector<std::shared_ptr<node_impl>> Inputs;
@@ -148,6 +173,33 @@ std::shared_ptr<node_impl> graph_impl::addSubgraphNodes(
   return this->add(Outputs);
 }
 
+std::shared_ptr<node_impl> graph_impl::addSubgraphNodes(
+    const std::shared_ptr<exec_graph_impl> &SubGraphExec) {
+  std::map<std::shared_ptr<node_impl>, std::shared_ptr<node_impl>> NodesMap;
+  std::list<std::shared_ptr<node_impl>> NewNodeList;
+
+  std::list<std::shared_ptr<node_impl>> NodeList = SubGraphExec->getSchedule();
+
+  for (std::list<std::shared_ptr<node_impl>>::const_iterator NodeIt =
+           NodeList.end();
+       NodeIt != NodeList.begin();) {
+    --NodeIt;
+    auto Node = *NodeIt;
+    std::shared_ptr<node_impl> 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<node_impl> &Root) {
   MRoots.insert(Root);
 }

sycl/source/detail/graph_impl.hpp

@@ -92,6 +92,46 @@ class node_impl {
             std::unique_ptr<sycl::detail::CG> &&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<sycl::detail::CGExecKernel *>(MCommandGroup.get());
+      sycl::detail::CGExecKernel *ExecKernelB =
+          static_cast<sycl::detail::CGExecKernel *>(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<sycl::detail::CGCopyUSM *>(MCommandGroup.get());
+      sycl::detail::CGCopyUSM *CopyB =
+          static_cast<sycl::detail::CGCopyUSM *>(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<sycl::detail::CGCopy *>(MCommandGroup.get());
+      sycl::detail::CGCopy *CopyB =
+          static_cast<sycl::detail::CGCopy *>(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_impl> 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<sycl::detail::CGExecKernel *>(MCommandGroup.get());
-      sycl::detail::CGExecKernel *ExecKernelB =
-          static_cast<sycl::detail::CGExecKernel *>(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<node_impl>
-  addSubgraphNodes(const std::list<std::shared_ptr<node_impl>> &NodeList);
+  addSubgraphNodes(const std::shared_ptr<exec_graph_impl> &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<node_impl> &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<node_impl> &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<node_impl>
+  addNodesToExits(const std::list<std::shared_ptr<node_impl>> &NodeList);
 };
 
 /// Class representing the implementation of command_graph<executable>.
@@ -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<graph_impl> &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) {

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.

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"

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"

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

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<float>(i); });
+    CGH.parallel_for(N, [=](id<1> it) { X[it] = static_cast<float>(it); });
   });
 
   auto B2 = add_node(

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"

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"