diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index c4c704b1b3..0fa2a56d3c 100644 --- a/hipamd/src/hip_device.cpp +++ b/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()); } // ================================================================================================ diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index 8a49bfd8e9..636b72d40a 100644 --- a/hipamd/src/hip_event.cpp +++ b/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); } diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 6ab73234e7..66b567cb45 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -27,6 +27,8 @@ #include "hip_mempool_impl.hpp" namespace hip { +extern std::unordered_map> GraphExecStatus_; +extern amd::Monitor GraphExecStatusLock_; std::vector 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(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); } diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index a571a1d62f..49e8f795d2 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -49,6 +49,9 @@ const char* GetGraphNodeTypeString(uint32_t op) { }; namespace hip { +std::unordered_map> GraphExecStatus_; +amd::Monitor GraphExecStatusLock_{"Guards graph execution state"}; + int GraphNode::nextID = 0; int Graph::nextID = 0; std::unordered_set 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(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 diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 8fe5de22cc..d69d986b3d 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -544,7 +544,7 @@ struct Graph { } }; struct GraphKernelNode; -struct GraphExec { +struct GraphExec : public amd::ReferenceCountedObject { std::vector> parallelLists_; // Topological order of the graph doesn't include nodes embedded as part of the child graph std::vector topoOrder_; @@ -573,7 +573,8 @@ struct GraphExec { GraphExec(std::vector& topoOrder, std::vector>& lists, std::unordered_map>& nodeWaitLists, struct Graph*& clonedGraph, std::unordered_map& 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; }; diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 1c3f230442..0ae7c393e4 100644 --- a/hipamd/src/hip_internal.hpp +++ b/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; diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 2e80e8e906..f73a9275cd 100644 --- a/hipamd/src/hip_stream.cpp +++ b/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;