SWDEV-469331 - Fix issue of graph sync.

If graph has multiple branches, End command is enqueued on launch stream which
makes sure all the internal parallel streams are finsihed.

When node is removed from the graph, indegree and outdegree are not getting update correctly for parent, child nodes and
resulting in endNode not having deps on parallel commands. Resulting in graph sync issues.

Change-Id: I33cc2f21220e1c017d88099b29b542e05b683f73


[ROCm/clr commit: 9ad7e79e50]

projects/clr/hipamd/src/hip_graph.cpp

@@ -65,7 +65,7 @@ inline hipError_t ihipGraphAddNode(hip::GraphNode* graphNode, hip::Graph* graph,
       return hipErrorInvalidValue;
     }
     DuplicateDep.insert(pDependencies[i]);
-    pDependencies[i]->AddEdge(graphNode);
+    pDependencies[i]->AddEdgeDep(graphNode);
   }
   if (capture == false) {
     {
@@ -1017,18 +1017,16 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) {
     return hipErrorStreamCaptureInvalidated;
   }
 
-  // check if all parallel streams have joined
-  // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do
-  // not result in hipErrorStreamCaptureUnjoined
-  // add temporary node to check if all parallel streams have joined
+  // Add temporary node to check if all parallel streams have joined
   hip::GraphNode* pGraphNode;
   pGraphNode = reinterpret_cast<hip::GraphNode*>(new hip::GraphEmptyNode());
   hipError_t status =
       ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(),
                        s->GetLastCapturedNodes().size());
-
   if (s->GetCaptureGraph()->GetLeafNodeCount() > 1) {
     std::vector<hip::GraphNode*> leafNodes = s->GetCaptureGraph()->GetLeafNodes();
+    // Manually added nodes doesnt result in hipErrorStreamCaptureUnjoined. Remove them from leaf
+    // nodes.
     std::unordered_set<hip::GraphNode*> nodes = s->GetCaptureGraph()->GetManualNodesDuringCapture();
     for (auto node : nodes) {
       const auto& fnode = std::find(leafNodes.begin(), leafNodes.end(), node);
@@ -1036,6 +1034,8 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) {
         leafNodes.erase(fnode);
       }
     }
+    // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do
+    // not result in hipErrorStreamCaptureUnjoined
     const std::vector<hip::GraphNode*>& removedDepNodes = s->GetRemovedDependencies();
     bool foundInRemovedDep = false;
     for (auto leafNode : leafNodes) {
@@ -1055,6 +1055,7 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) {
     // remove temporary node
     s->GetCaptureGraph()->RemoveNode(pGraphNode);
   }
+
   *pGraph = s->GetCaptureGraph();
   // end capture on all streams/events part of graph capture
   return s->EndCapture();
@@ -1681,7 +1682,7 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
         HIP_RETURN(hipErrorInvalidValue);
       }
     }
-    fromNode[i]->AddEdge(toNode[i]);
+    fromNode[i]->AddEdgeDep(toNode[i]);
   }
   HIP_RETURN(hipSuccess);
 }
@@ -1930,7 +1931,7 @@ hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* fr
   hip::Graph* g = reinterpret_cast<hip::Graph*>(graph);
   for (size_t i = 0; i < numDependencies; i++) {
     if (toNode[i]->GetParentGraph() != g || fromNode[i]->GetParentGraph() != g ||
-        fromNode[i]->RemoveUpdateEdge(toNode[i]) == false) {
+        fromNode[i]->RemoveEdgeDep(toNode[i]) == false) {
       HIP_RETURN(hipErrorInvalidValue);
     }
   }
@@ -2053,16 +2054,8 @@ hipError_t hipGraphDestroyNode(hipGraphNode_t node) {
       n->GetType() == hipGraphNodeTypeMemFree) {
     HIP_RETURN(hipErrorNotSupported);
   }
-  // First remove all the edges both incoming and outgoing from node.
-  for (auto& edge : n->GetEdges()) {
-    n->RemoveUpdateEdge(edge);
-  }
-  const std::vector<hip::GraphNode*>& dependencies = n->GetDependencies();
-  for (auto& parent : dependencies) {
-    parent->RemoveEdge(n);
-    parent->SetOutDegree(parent->GetOutDegree() - 1);
-  }
-  // Remove the node from graph.
+  // Remove the node from graph should takecare of updating edges of parent and deps of child nodes
+  // as part of graph node destructor
   n->GetParentGraph()->RemoveNode(n);
   HIP_RETURN(hipSuccess);
 }

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -273,32 +273,31 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
     }
   }
   /// Add graph node dependency
-  void AddDependency(const Node& node) { dependencies_.push_back(node); }
+  void AddDependency(const Node& node) {
+    dependencies_.push_back(node);
+    inDegree_++;
+  }
   /// Remove graph node dependency
   void RemoveDependency(const Node& node) {
     dependencies_.erase(std::remove(dependencies_.begin(), dependencies_.end(), node),
                         dependencies_.end());
+    inDegree_--;
   }
   void RemoveEdge(const Node& childNode) {
     edges_.erase(std::remove(edges_.begin(), edges_.end(), childNode), edges_.end());
+    outDegree_--;
   }
-  /// Return graph node children
-  const std::vector<Node>& GetEdges() const { return edges_; }
-  /// Updates graph node children
-  void SetEdges(std::vector<Node>& edges) {
-    for (auto entry : edges) {
-      edges_.push_back(entry);
-    }
-  }
-  /// Add edge, update parent node outdegree, child node indegree and dependency
   void AddEdge(const Node& childNode) {
     edges_.push_back(childNode);
     outDegree_++;
-    childNode->SetInDegree(childNode->GetInDegree() + 1);
+  }
+  /// Add edge, update parent node outdegree, child node indegree and dependency
+  void AddEdgeDep(const Node& childNode) {
+    AddEdge(childNode);
     childNode->AddDependency(this);
   }
   /// Remove edge, update parent node outdegree, child node indegree and dependency
-  bool RemoveUpdateEdge(const Node& childNode) {
+  bool RemoveEdgeDep(const Node& childNode) {
     // std::remove changes the end() hence saving it before hand for validation
     auto currEdgeEnd = edges_.end();
     auto it = std::remove(edges_.begin(), edges_.end(), childNode);
@@ -308,10 +307,17 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
     }
     edges_.erase(it, edges_.end());
     outDegree_--;
-    childNode->SetInDegree(childNode->GetInDegree() - 1);
     childNode->RemoveDependency(this);
     return true;
   }
+  /// Return graph node children
+  const std::vector<Node>& GetEdges() const { return edges_; }
+  /// Updates graph node children
+  void SetEdges(std::vector<Node>& edges) {
+    for (auto entry : edges) {
+      edges_.push_back(entry);
+    }
+  }
   /// Get Runlist of the nodes embedded as part of the graphnode(e.g. ChildGraph)
   virtual void GetRunList(std::vector<std::vector<Node>>& parallelList,
                           std::unordered_map<Node, std::vector<Node>>& dependencies) {}