From f278ac9ea0cdbf7cd5f78f74cb91113645336c3c Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Tue, 17 Sep 2024 15:13:23 +0000 Subject: [PATCH] SWDEV-485904 - Fix virtual,physical mem obj leaks Change-Id: Ie0456b5dcfec206ae54a6aabfc2a15a620cac693 [ROCm/clr commit: 870842201d42944b5aff4ef30a1d4ddfbd40f3f8] --- projects/clr/hipamd/src/hip_memory.cpp | 20 ++++++++++---------- projects/clr/hipamd/src/hip_vm.cpp | 13 ++++++------- projects/clr/hipamd/src/hip_vm.hpp | 5 ++++- projects/clr/rocclr/device/device.cpp | 6 +----- projects/clr/rocclr/platform/memory.cpp | 4 +++- 5 files changed, 24 insertions(+), 24 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 84ddee4c13..3e9f96b2ef 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -426,16 +426,16 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, return hipErrorUnknown; } - // Size validation - if (sizeBytes > (srcMemory->getSize() - sOffset)) { - return hipErrorInvalidValue; - } - // If the mem object is a VMM sub buffer (subbuffer has parent set), // then use parent's size for validation. if (srcMemory->parent() && (srcMemory->getMemFlags() & CL_MEM_VA_RANGE_AMD)) { srcMemory = srcMemory->parent(); } + + // Size validation + if (sizeBytes > (srcMemory->getSize() - sOffset)) { + return hipErrorInvalidValue; + } } if (dstMemory != nullptr) { @@ -444,16 +444,16 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, return hipErrorUnknown; } - // Size validation - if (sizeBytes > (dstMemory->getSize() - dOffset)) { - return hipErrorInvalidValue; - } - // If the mem object is a VMM sub buffer (subbuffer has parent set), // then use parent's size for validation. if (dstMemory->parent() && (dstMemory->getMemFlags() & CL_MEM_VA_RANGE_AMD)) { dstMemory = dstMemory->parent(); } + + // Size validation + if (sizeBytes > (dstMemory->getSize() - dOffset)) { + return hipErrorInvalidValue; + } } //If src and dst ptr are null then kind must be either h2h or def. diff --git a/projects/clr/hipamd/src/hip_vm.cpp b/projects/clr/hipamd/src/hip_vm.cpp index aa1bbb1cf3..bf29686a6a 100644 --- a/projects/clr/hipamd/src/hip_vm.cpp +++ b/projects/clr/hipamd/src/hip_vm.cpp @@ -39,10 +39,13 @@ hipError_t hipMemAddressFree(void* devPtr, size_t size) { if (devPtr == nullptr || size == 0) { HIP_RETURN(hipErrorInvalidValue); } - + amd::Memory* memObj = amd::MemObjMap::FindVirtualMemObj(devPtr); + if (memObj == nullptr) { + LogPrintfError("Cannot find the Virtual MemObj entry for this addr 0x%x", devPtr); + } // Single call frees address range for all devices. g_devices[0]->devices()[0]->virtualFree(devPtr); - + memObj->release(); HIP_RETURN(hipSuccess); } @@ -126,10 +129,6 @@ hipError_t hipMemCreate(hipMemGenericAllocationHandle_t* handle, size_t size, phys_mem_obj->getUserData().data = new hip::GenericAllocation(*phys_mem_obj, size, *prop); *handle = reinterpret_cast(phys_mem_obj->getUserData().data); - // Remove because the entry of 0x1 is not needed in MemObjMap. - // We save the copy of Phy mem obj in virtual mem obj during mapping. - amd::MemObjMap::RemoveMemObj(ptr); - HIP_RETURN(hipSuccess); } @@ -261,7 +260,6 @@ hipError_t hipMemMap(void* ptr, size_t size, size_t offset, hipMemGenericAllocat ga->retain(); auto& queue = *g_devices[ga->GetProperties().location.id]->NullStream(); - // Map the physical address to virtual address amd::Command* cmd = new amd::VirtualMapCommand(queue, amd::Command::EventWaitList{}, ptr, size, &ga->asAmdMemory()); @@ -366,6 +364,7 @@ hipError_t hipMemUnmap(void* ptr, size_t size) { cmd->enqueue(); cmd->awaitCompletion(); cmd->release(); + vaddr_sub_obj->release(); // restore the original pa of the generic allocation hip::GenericAllocation* ga diff --git a/projects/clr/hipamd/src/hip_vm.hpp b/projects/clr/hipamd/src/hip_vm.hpp index 678c480175..e37b2b4b81 100644 --- a/projects/clr/hipamd/src/hip_vm.hpp +++ b/projects/clr/hipamd/src/hip_vm.hpp @@ -38,7 +38,10 @@ class GenericAllocation : public amd::RuntimeObject { public: GenericAllocation(amd::Memory& phys_mem_ref, size_t size, const hipMemAllocationProp& prop) : phys_mem_ref_(phys_mem_ref), size_(size), properties_(prop) {} - ~GenericAllocation() {} + ~GenericAllocation() { + amd::Context* amdContext = g_devices[properties_.location.id]->asContext(); + amd::SvmBuffer::free(*amdContext, phys_mem_ref_.getSvmPtr()); + } const hipMemAllocationProp& GetProperties() const { return properties_; } hipMemGenericAllocationHandle_t asMemGenericAllocationHandle() { diff --git a/projects/clr/rocclr/device/device.cpp b/projects/clr/rocclr/device/device.cpp index f8479c3817..7661f3e9d3 100644 --- a/projects/clr/rocclr/device/device.cpp +++ b/projects/clr/rocclr/device/device.cpp @@ -479,11 +479,7 @@ bool Device::DestroyVirtualBuffer(amd::Memory* vaddr_mem_obj) { return false; } - if (vaddr_mem_obj->parent() == nullptr) { - // If parent is nullptr, then vaddr_mem_obj is the parent. - amd::MemObjMap::RemoveVirtualMemObj(vaddr_mem_obj->getSvmPtr()); - return true; - } else { + if (vaddr_mem_obj->parent() != nullptr) { // If parent is not nullptr, this is the sub-buffer object. amd::Memory* vaddr_base_obj = amd::MemObjMap::FindVirtualMemObj(vaddr_mem_obj->getSvmPtr()); if (vaddr_base_obj == nullptr) { diff --git a/projects/clr/rocclr/platform/memory.cpp b/projects/clr/rocclr/platform/memory.cpp index 97cbac1fe3..a3a1e2103f 100644 --- a/projects/clr/rocclr/platform/memory.cpp +++ b/projects/clr/rocclr/platform/memory.cpp @@ -456,7 +456,9 @@ Memory::~Memory() { } hostMemRef_.deallocateMemory(context_()); if (getMemFlags() & CL_MEM_VA_RANGE_AMD) { - amd::MemObjMap::RemoveVirtualMemObj(getSvmPtr()); + if (parent_ == nullptr) { + amd::MemObjMap::RemoveVirtualMemObj(getSvmPtr()); + } // If runtime executes graph mempool with VM, then VA can be mapped in space // for graph validation logic during execution. And the reason it's not unmaped // in graph itself because the app can have a graph without a free node