diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 16267a7937..5657fd6371 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -98,11 +98,11 @@ bool Device::FreeMemory(amd::Memory* memory, Stream* stream) { } // ================================================================================================ -void Device::ReleaseFreedMemory(Stream* stream) { +void Device::ReleaseFreedMemory() { amd::ScopedLock lock(lock_); // Search for memory in the entire list of pools for (auto it : mem_pools_) { - it->ReleaseFreedMemory(stream); + it->ReleaseFreedMemory(); } } diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index 00d6b83bdd..d233e42438 100644 --- a/hipamd/src/hip_event.cpp +++ b/hipamd/src/hip_event.cpp @@ -71,20 +71,22 @@ hipError_t Event::synchronize() { return hipSuccess; } + auto hip_device = g_devices[deviceId()]; // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status static constexpr bool kWaitCompletion = true; - if (!g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, kWaitCompletion)) { + if (!hip_device->devices()[0]->IsHwEventReady(*event_, kWaitCompletion)) { if (event_->HwEvent() != nullptr) { amd::Command* command = nullptr; hipError_t status = recordCommand(command, event_->command().queue(), flags); command->enqueue(); - g_devices[deviceId()]->devices()[0]->IsHwEventReady(command->event(), kWaitCompletion); + hip_device->devices()[0]->IsHwEventReady(command->event(), kWaitCompletion); command->release(); } else { event_->awaitCompletion(); } } - + // Release freed memory for all memory pools on the device + hip_device->ReleaseFreedMemory(); return hipSuccess; } diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 7afe92d6e0..88b12d8451 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -529,7 +529,7 @@ public: bool FreeMemory(amd::Memory* memory, Stream* stream); /// Release freed memory from all pools on the current device - void ReleaseFreedMemory(Stream* stream); + void ReleaseFreedMemory(); /// Removes a destroyed stream from the safe list of memory pools void RemoveStreamFromPools(Stream* stream); diff --git a/hipamd/src/hip_mempool_impl.cpp b/hipamd/src/hip_mempool_impl.cpp index de620e1ee6..ca480da551 100644 --- a/hipamd/src/hip_mempool_impl.cpp +++ b/hipamd/src/hip_mempool_impl.cpp @@ -115,7 +115,7 @@ bool Heap::ReleaseAllMemory(size_t min_bytes_to_hold, bool safe_release) { } // ================================================================================================ -bool Heap::ReleaseAllMemory(hip::Stream* stream) { +bool Heap::ReleaseAllMemory() { for (auto it = allocations_.begin(); it != allocations_.end();) { // Make sure the heap holds the minimum number of bytes if (total_size_ <= release_threshold_) { @@ -270,10 +270,10 @@ void MemoryPool::ReleaseAllMemory() { } // ================================================================================================ -void MemoryPool::ReleaseFreedMemory(hip::Stream* stream) { +void MemoryPool::ReleaseFreedMemory() { amd::ScopedLock lock(lock_pool_ops_); - free_heap_.ReleaseAllMemory(stream); + free_heap_.ReleaseAllMemory(); } // ================================================================================================ diff --git a/hipamd/src/hip_mempool_impl.hpp b/hipamd/src/hip_mempool_impl.hpp index 2541e7a490..c097e19dcf 100644 --- a/hipamd/src/hip_mempool_impl.hpp +++ b/hipamd/src/hip_mempool_impl.hpp @@ -110,10 +110,10 @@ public: bool RemoveMemory(amd::Memory* memory, MemoryTimestamp* ts = nullptr); /// Releases all memory, until the threshold value is met - bool ReleaseAllMemory(size_t min_bytes_to_hold = std::numeric_limits::max(), bool safe_release = false); + bool ReleaseAllMemory(size_t min_bytes_to_hold, bool safe_release = false); /// Releases all memory, safe to the provided stream, until the threshold value is met - bool ReleaseAllMemory(hip::Stream* stream); + bool ReleaseAllMemory(); /// Remove the provided stream from the safe list void RemoveStream(hip::Stream* stream); @@ -218,7 +218,7 @@ class MemoryPool : public amd::ReferenceCountedObject { /// Releases all allocations from free_heap_. It can be called on Stream or Device synchronization /// @note The caller must make sure it's safe to release memory - void ReleaseFreedMemory(hip::Stream* stream = nullptr); + void ReleaseFreedMemory(); /// Removes a stream from tracking void RemoveStream(hip::Stream* stream); diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 6dc5b16aa6..fad39e1692 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -139,6 +139,8 @@ void Stream::SyncAllStreams(int deviceId, bool cpu_wait) { it->finish(cpu_wait); it->release(); } + // Release freed memory for all memory pools on the device + g_devices[deviceId]->ReleaseFreedMemory(); } // ================================================================================================ @@ -450,8 +452,14 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) { } bool wait = (stream == nullptr) ? true : false; constexpr bool kDontWaitForCpu = false; + + auto hip_stream = hip::getStream(stream, wait); + // Wait for the current host queue - hip::getStream(stream, wait)->finish(kDontWaitForCpu); + hip_stream->finish(kDontWaitForCpu); + + // Release freed memory for all memory pools on the device + hip_stream->GetDevice()->ReleaseFreedMemory(); return hipSuccess; }