SWDEV-460770 - Handle Graph Exec release

Handle GraphExec instance is destroyed before async launch completes
GraphExec instance is destroyed after async launch completes
GraphExec instance is destroyed without a launch

Change-Id: I45a7c82295fea916c7559bd8f796df710513aea1

hipamd/src/hip_device.cpp

@@ -269,6 +269,8 @@ void Device::SyncAllStreams( bool cpu_wait) {
   }
   // Release freed memory for all memory pools on the device
   ReleaseFreedMemory();
+  // Release all graph exec objects destroyed by user.
+  ReleaseGraphExec(hip::getCurrentDevice()->deviceId());
 }
 
 // ================================================================================================

hipamd/src/hip_event.cpp

@@ -444,7 +444,8 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
   hipError_t status = e->synchronize();
   // Release freed memory for all memory pools on the device
   g_devices[e->deviceId()]->ReleaseFreedMemory();
-
+  // Release all graph exec objects destroyed by user.
+  ReleaseGraphExec(e->deviceId());
   HIP_RETURN(status);
 }
 

hipamd/src/hip_graph.cpp

@@ -27,6 +27,8 @@
 #include "hip_mempool_impl.hpp"
 
 namespace hip {
+extern std::unordered_map<GraphExec*, std::pair<hip::Stream*, bool>> GraphExecStatus_;
+extern amd::Monitor GraphExecStatusLock_;
 
 std::vector<hip::Stream*> g_captureStreams;
 amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"};
@@ -1396,8 +1398,16 @@ hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) {
   if (pGraphExec == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  amd::ScopedLock lock(GraphExecStatusLock_);
   hip::GraphExec* ge = reinterpret_cast<hip::GraphExec*>(pGraphExec);
-  delete ge;
+  // bool found = false;
+  if (GraphExecStatus_.find(ge) == GraphExecStatus_.end()) {
+    ge->release();
+  } else {
+    // graph execution is under progress. destroy graphExec during next sync point
+    auto pair = GraphExecStatus_[ge];
+    GraphExecStatus_[ge] = std::make_pair(pair.first, true);
+  }
   HIP_RETURN(hipSuccess);
 }
 

hipamd/src/hip_graph_internal.cpp

@@ -49,6 +49,9 @@ const char* GetGraphNodeTypeString(uint32_t op) {
 };
 
 namespace hip {
+std::unordered_map<GraphExec*, std::pair<hip::Stream*, bool>> GraphExecStatus_;
+amd::Monitor GraphExecStatusLock_{"Guards graph execution state"};
+
 int GraphNode::nextID = 0;
 int Graph::nextID = 0;
 std::unordered_set<GraphNode*> GraphNode::nodeSet_;
@@ -586,7 +589,6 @@ hipError_t GraphExec::Run(hipStream_t stream) {
         accumulate = new amd::AccumulateCommand(*hip_stream, {}, nullptr, lastCapturedPacket);
       }
     }
-
     for (int i = 0; i < topoOrder_.size() - 1; i++) {
       if (DEBUG_CLR_GRAPH_PACKET_CAPTURE && topoOrder_[i]->GetType() == hipGraphNodeTypeKernel &&
           !reinterpret_cast<hip::GraphKernelNode*>(topoOrder_[i])->HasHiddenHeap()) {
@@ -646,7 +648,41 @@ hipError_t GraphExec::Run(hipStream_t stream) {
       endCommand->release();
     }
   }
+  amd::ScopedLock lock(GraphExecStatusLock_);
+  GraphExecStatus_[this] = std::make_pair(hip_stream, false);
   ResetQueueIndex();
   return status;
 }
+void ReleaseGraphExec(int deviceId) {
+  // Release all graph exec objects destroyed by user.
+  amd::ScopedLock lock(GraphExecStatusLock_);
+  for (auto itr = GraphExecStatus_.begin(); itr != GraphExecStatus_.end();) {
+    auto pair = itr->second;
+    if (pair.first->DeviceId() == deviceId) {
+      if (pair.second == true) {
+        ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Release GraphExec");
+        (itr->first)->release();
+      }
+      GraphExecStatus_.erase(itr++);
+    } else {
+      itr++;
+    }
+  }
+}
+void ReleaseGraphExec(hip::Stream* stream) {
+  amd::ScopedLock lock(GraphExecStatusLock_);
+  for (auto itr = GraphExecStatus_.begin(); itr != GraphExecStatus_.end();) {
+    auto pair = itr->second;
+    if (pair.first == stream) {
+      if (pair.second == true) {
+        ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Release GraphExec");
+        (itr->first)->release();
+      }
+      GraphExecStatus_.erase(itr++);
+      break;
+    } else {
+      ++itr;
+    }
+  }
+}
 }  // namespace hip

hipamd/src/hip_graph_internal.hpp

@@ -544,7 +544,7 @@ struct Graph {
   }
 };
 struct GraphKernelNode;
-struct GraphExec {
+struct GraphExec : public amd::ReferenceCountedObject {
   std::vector<std::vector<Node>> parallelLists_;
   // Topological order of the graph doesn't include nodes embedded as part of the child graph
   std::vector<Node> topoOrder_;
@@ -573,7 +573,8 @@ struct GraphExec {
   GraphExec(std::vector<Node>& topoOrder, std::vector<std::vector<Node>>& lists,
             std::unordered_map<Node, std::vector<Node>>& nodeWaitLists, struct Graph*& clonedGraph,
             std::unordered_map<Node, Node>& clonedNodes, uint64_t flags = 0)
-      : parallelLists_(lists),
+      : ReferenceCountedObject(),
+        parallelLists_(lists),
         topoOrder_(topoOrder),
         nodeWaitLists_(nodeWaitLists),
         clonedGraph_(clonedGraph),
@@ -648,7 +649,6 @@ struct GraphExec {
   // Capture GPU Packets from graph commands
   hipError_t CaptureAQLPackets();
   hipError_t UpdateAQLPacket(hip::GraphKernelNode* node);
-
   using KernelArgImpl = device::Settings::KernelArgImpl;
 };
 

hipamd/src/hip_internal.hpp

@@ -73,6 +73,9 @@ struct Graph;
 struct GraphNode;
 struct GraphExec;
 struct UserObject;
+class Stream;
+extern void ReleaseGraphExec(int deviceId);
+extern void ReleaseGraphExec(hip::Stream* stream);
 typedef struct ihipIpcMemHandle_st {
   char ipc_handle[IHIP_IPC_MEM_HANDLE_SIZE];  ///< ipc memory handle on ROCr
   size_t psize;

hipamd/src/hip_stream.cpp

@@ -343,10 +343,14 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) {
   constexpr bool kDontWaitForCpu = false;
 
   auto hip_stream = hip::getStream(stream, wait);
-
   // Wait for the current host queue
   hip_stream->finish(kDontWaitForCpu);
-
+  if (stream == nullptr) {
+    // null stream will sync with other streams.
+    ReleaseGraphExec(hip_stream->DeviceId());
+  } else {
+    ReleaseGraphExec(hip_stream);
+  }
   // Release freed memory for all memory pools on the device
   hip_stream->GetDevice()->ReleaseFreedMemory();
   return hipSuccess;