diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 1b254f3029..3dc9ee7614 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -25,6 +25,21 @@ #include "hip_event.hpp" thread_local std::vector g_captureStreams; +static std::unordered_set 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; diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index ae42d798a1..cd5487c2d9 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -46,6 +46,9 @@ const char* GetGraphNodeTypeString(uint32_t op) { }; int hipGraphNode::nextID = 0; +std::unordered_set 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& levelOrder) { } } -ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const { +const ihipGraph* ihipGraph::getOriginalGraph() const { + return pOriginalGraph_; +} + +void ihipGraph::setOriginalGraph(const ihipGraph* pOriginalGraph) { + pOriginalGraph_ = pOriginalGraph; +} + +ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const{ ihipGraph* newGraph = new ihipGraph(); for (auto entry : vertices_) { hipGraphNode* node = entry->clone(); @@ -640,10 +651,11 @@ ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const { } clonedNodes[node]->SetDependencies(clonedDependencies); } + newGraph->setOriginalGraph(this); return newGraph; } -ihipGraph* ihipGraph::clone() const { +ihipGraph* ihipGraph::clone() const{ std::unordered_map clonedNodes; return clone(clonedNodes); } diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 39319cfe91..e26be80d8d 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -41,6 +41,7 @@ hipError_t FillCommands(std::vector>& parallelLists, amd::Command*& endCommand, amd::HostQueue* queue); void UpdateQueue(std::vector>& 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 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 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& GetNodes() const { return vertices_; } /// returns all the edges in the graph std::vector> 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& visited, std::vector& singleList, std::vector>& parallelLists, std::unordered_map>& dependencies);