SWDEV-315443 - hipGraph and Node issues

Change-Id: I9fefd1a12d83bfa89d6704ee236c1409618ab02a

hipamd/src/hip_graph.cpp

@@ -25,6 +25,21 @@
 #include "hip_event.hpp"
 
 thread_local std::vector<hipStream_t> g_captureStreams;
+static std::unordered_set<ihipGraph*> graphSet;
+static amd::Monitor graphSetLock{"Guards global graph set"};
+
+ihipGraph::ihipGraph() {
+  amd::ScopedLock lock(graphSetLock);
+  graphSet.insert(this);
+}
+
+ihipGraph::~ihipGraph() {
+  for (auto node : vertices_) {
+    delete node;
+  }
+  amd::ScopedLock lock(graphSetLock);
+  graphSet.erase(this);
+}
 
 inline void ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph,
                              const hipGraphNode_t* pDependencies, size_t numDependencies) {
@@ -1374,7 +1389,7 @@ hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) {
 
 hipError_t hipGraphDestroyNode(hipGraphNode_t node) {
   HIP_INIT_API(hipGraphDestroyNode, node);
-  if (node == nullptr) {
+  if (node == nullptr || !hipGraphNode::isNodeValid(node)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   node->GetParentGraph()->RemoveNode(node);
@@ -1383,11 +1398,22 @@ hipError_t hipGraphDestroyNode(hipGraphNode_t node) {
   HIP_RETURN(hipSuccess);
 }
 
+bool isGraphValid(ihipGraph* pGraph) {
+  amd::ScopedLock lock(graphSetLock);
+  if (graphSet.find(pGraph) == graphSet.end()) {
+    return false;
+  }
+  return true;
+}
+
 hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph) {
   HIP_INIT_API(hipGraphClone, pGraphClone, originalGraph);
   if (originalGraph == nullptr || pGraphClone == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  if (!isGraphValid(originalGraph)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   *pGraphClone = originalGraph->clone();
   HIP_RETURN(hipSuccess);
 }
@@ -1398,6 +1424,10 @@ hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t origina
   if (pNode == nullptr || originalNode == nullptr || clonedGraph == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  if (clonedGraph->getOriginalGraph() == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
   for (auto node : clonedGraph->GetNodes()) {
     if (node->GetID() == originalNode->GetID()) {
       *pNode = node;

hipamd/src/hip_graph_internal.cpp

@@ -46,6 +46,9 @@ const char* GetGraphNodeTypeString(uint32_t op) {
 };
 
 int hipGraphNode::nextID = 0;
+std::unordered_set<hipGraphNode*> hipGraphNode::nodeSet_;
+amd::Monitor hipGraphNode::nodeSetLock_{"Guards global node set"};
+
 hipError_t hipGraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t count,
                                                 hipMemcpyKind kind) {
   hipError_t status = ihipMemcpy_validate(dst, src, count, kind);
@@ -614,7 +617,15 @@ void ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
   }
 }
 
-ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const {
+const ihipGraph* ihipGraph::getOriginalGraph() const {
+  return pOriginalGraph_;
+}
+
+void ihipGraph::setOriginalGraph(const ihipGraph* pOriginalGraph) {
+  pOriginalGraph_ = pOriginalGraph;
+}
+
+ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const{
   ihipGraph* newGraph = new ihipGraph();
   for (auto entry : vertices_) {
     hipGraphNode* node = entry->clone();
@@ -640,10 +651,11 @@ ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const {
     }
     clonedNodes[node]->SetDependencies(clonedDependencies);
   }
+  newGraph->setOriginalGraph(this);
   return newGraph;
 }
 
-ihipGraph* ihipGraph::clone() const {
+ihipGraph* ihipGraph::clone() const{
   std::unordered_map<Node, Node> clonedNodes;
   return clone(clonedNodes);
 }

hipamd/src/hip_graph_internal.hpp

@@ -41,6 +41,7 @@ hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
                         amd::Command*& endCommand, amd::HostQueue* queue);
 void UpdateQueue(std::vector<std::vector<Node>>& parallelLists, amd::HostQueue*& queue,
                  hipGraphExec* ptr);
+
 struct hipGraphNode {
  protected:
   amd::HostQueue* queue_;
@@ -57,6 +58,8 @@ struct hipGraphNode {
   size_t outDegree_;
   static int nextID;
   struct ihipGraph* parentGraph_;
+  static std::unordered_set<hipGraphNode*> nodeSet_;
+  static amd::Monitor nodeSetLock_;
 
  public:
   hipGraphNode(hipGraphNodeType type)
@@ -66,7 +69,10 @@ struct hipGraphNode {
         inDegree_(0),
         outDegree_(0),
         id_(nextID++),
-        parentGraph_(nullptr) {}
+        parentGraph_(nullptr) {
+    amd::ScopedLock lock(nodeSetLock_);
+    nodeSet_.insert(this);
+  }
   /// Copy Constructor
   hipGraphNode(const hipGraphNode& node) {
     level_ = node.level_;
@@ -76,6 +82,8 @@ struct hipGraphNode {
     visited_ = false;
     id_ = node.id_;
     parentGraph_ = nullptr;
+    amd::ScopedLock lock(nodeSetLock_);
+    nodeSet_.insert(this);
   }
 
   virtual ~hipGraphNode() {
@@ -85,6 +93,17 @@ struct hipGraphNode {
     for (auto node : dependencies_) {
       node->RemoveEdge(this);
     }
+    amd::ScopedLock lock(nodeSetLock_);
+    nodeSet_.erase(this);
+  }
+
+  // check node validity
+  static bool isNodeValid(hipGraphNode* pGraphNode) {
+    amd::ScopedLock lock(nodeSetLock_);
+    if (nodeSet_.find(pGraphNode) == nodeSet_.end()) {
+      return false;
+    }
+    return true;
   }
 
   amd::HostQueue* GetQueue() { return queue_; }
@@ -211,14 +230,14 @@ struct hipGraphNode {
 
 struct ihipGraph {
   std::vector<Node> vertices_;
+  const ihipGraph* pOriginalGraph_ = nullptr;
 
  public:
-  ihipGraph() {}
-  ~ihipGraph() {
-    for (auto node : vertices_) {
-      delete node;
-    }
-  }
+ ihipGraph();
+ ~ihipGraph();
+  // check graphs validity
+  bool isGraphValid(ihipGraph* pGraph);
+
   /// add node to the graph
   void AddNode(const Node& node);
   void RemoveNode(const Node& node);
@@ -234,6 +253,11 @@ struct ihipGraph {
   const std::vector<Node>& GetNodes() const { return vertices_; }
   /// returns all the edges in the graph
   std::vector<std::pair<Node, Node>> GetEdges() const;
+  // returns the original graph ptr if cloned
+  const ihipGraph* getOriginalGraph() const;
+  // saves the original graph ptr if cloned
+  void setOriginalGraph(const ihipGraph* pOriginalGraph);
+
   void GetRunListUtil(Node v, std::unordered_map<Node, bool>& visited,
                       std::vector<Node>& singleList, std::vector<std::vector<Node>>& parallelLists,
                       std::unordered_map<Node, std::vector<Node>>& dependencies);