From 3f6f9d608122a21e780db184a179b55079df9bd9 Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Thu, 31 Jul 2025 08:30:23 -0700 Subject: [PATCH] SWDEV-539710 - Defer allocation of managed variable (#652) Co-authored-by: Rahul Manocha --- hipamd/src/hip_code_object.cpp | 31 ++++++++------- hipamd/src/hip_global.cpp | 34 ++++++++++++++--- hipamd/src/hip_global.hpp | 20 +++++++--- hipamd/src/hip_hmm.cpp | 18 ++++++--- hipamd/src/hip_platform.cpp | 55 +++++++++++++++++++-------- rocclr/device/rocm/rocdevice.cpp | 65 ++++++++++++++++---------------- rocclr/device/rocm/rocmemory.cpp | 28 +++++++++++++- rocclr/platform/context.cpp | 4 +- rocclr/platform/context.hpp | 2 +- rocclr/platform/memory.cpp | 4 +- rocclr/platform/memory.hpp | 3 +- 11 files changed, 177 insertions(+), 87 deletions(-) mode change 100644 => 100755 hipamd/src/hip_hmm.cpp diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index be68b44441..5596feba7b 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -33,7 +33,7 @@ using namespace llvm::ELF; namespace hip { hipError_t ihipFree(void* ptr); // forward declaration of methods required for managed variables -hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0); +hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align = 0, bool use_host_ptr = 0); namespace { // In uncompressed mode constexpr char kOffloadBundleUncompressedMagicStr[] = "__CLANG_OFFLOAD_BUNDLE__"; @@ -718,7 +718,7 @@ hipError_t DynCO::initDynManagedVars(const std::string& managedVar) { return status; } // Allocate managed memory for these symbols - status = ihipMallocManaged(&pointer, dvar->size()); + status = ihipMallocManaged(&pointer, dvar->size(), 0, 0); guarantee(status == hipSuccess, "Status %d, failed to allocate managed memory", status); // update as manager variable and set managed memory pointer and size @@ -875,7 +875,7 @@ hipError_t StatCO::removeFatBinary(FatBinaryInfo** module) { auto managedVarsIter = managedVars_.find(module); if (managedVarsIter != managedVars_.end()) { for (auto& managedVar : managedVarsIter->second) { - hipError_t err; + hipError_t err = hipSuccess; for (auto dev : g_devices) { DeviceVar* dvar = nullptr; IHIP_RETURN_ONFAIL(managedVar->getDeviceVarPtr(&dvar, dev->deviceId())); @@ -885,7 +885,12 @@ hipError_t StatCO::removeFatBinary(FatBinaryInfo** module) { assert(err == hipSuccess); } } - err = ihipFree(*(static_cast(managedVar->getManagedVarPtr()))); + if (managedVar->getAllocFlag()) { // check if it is a managed or host alloc + err = ihipFree(*(static_cast(managedVar->getManagedVarPtr()))); + } else { + void** pointer = static_cast(managedVar->getManagedVarPtr()); + amd::Os::releaseMemory(*pointer, managedVar->getSize()); + } assert(err == hipSuccess); delete managedVar; } @@ -1038,21 +1043,21 @@ hipError_t StatCO::initStatManagedVarDevicePtr(int deviceId) { // Lazy load FatBinaryInfo **module = var->moduleInfo(); if (*(module) == nullptr) { - hipError_t err = digestFatBinary(module_to_hostModule_[module], *module); + err = digestFatBinary(module_to_hostModule_[module], *module); assert(err == hipSuccess); } - - DeviceVar* dvar = nullptr; - IHIP_RETURN_ONFAIL(var->getStatDeviceVar(&dvar, deviceId)); - hip::Stream* stream = g_devices.at(deviceId)->NullStream(); - if (stream != nullptr) { - err = ihipMemcpy(reinterpret_cast
(dvar->device_ptr()), var->getManagedVarPtr(), - dvar->size(), hipMemcpyHostToDevice, *stream); - } else { + if (stream == nullptr) { ClPrint(amd::LOG_ERROR, amd::LOG_API, "Host Queue is NULL"); return hipErrorInvalidResourceHandle; } + // Allocate managed var for deferred loading + IHIP_RETURN_ONFAIL(var->allocateManagedVarPtr()); + // Copy from managed var host to device ptr + DeviceVar* dvar = nullptr; + IHIP_RETURN_ONFAIL(var->getStatDeviceVar(&dvar, deviceId)); + err = ihipMemcpy(reinterpret_cast
(dvar->device_ptr()), var->getManagedVarPtr(), + dvar->size(), hipMemcpyHostToDevice, *stream); } } managedVarsDevicePtrInitalized_[deviceId] = true; diff --git a/hipamd/src/hip_global.cpp b/hipamd/src/hip_global.cpp index daa0a0eb41..756833b38a 100644 --- a/hipamd/src/hip_global.cpp +++ b/hipamd/src/hip_global.cpp @@ -61,6 +61,9 @@ void __hipGetPCH(const char** pch, unsigned int *size) { #endif namespace hip { +// forward declaration of methods required for managed variables +hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align = 0, bool use_host_ptr = 0); + //Device Vars DeviceVar::DeviceVar(std::string name, hipModule_t hmod, @@ -218,10 +221,17 @@ Var::Var(const std::string& name, DeviceVarKind dVarKind, size_t size, int type, dVar_.resize(g_devices.size()); } -Var::Var(const std::string& name, DeviceVarKind dVarKind, void *pointer, size_t size, - unsigned align, FatBinaryInfo** modules) : name_(name), dVarKind_(dVarKind), - size_(size), modules_(modules), managedVarPtr_(pointer), align_(align), - type_(0), norm_(0) { +Var::Var(const std::string& name, DeviceVarKind dVarKind, void* pointer, size_t size, + unsigned align, FatBinaryInfo** modules) + : name_(name), + dVarKind_(dVarKind), + size_(size), + modules_(modules), + managedVarPtr_(pointer), + allocFlag_(false), + align_(align), + type_(0), + norm_(0) { dVar_.resize(g_devices.size()); } @@ -270,5 +280,19 @@ hipError_t Var::getStatDeviceVar(DeviceVar** dvar, int deviceId) { *dvar = dVar_[deviceId]; return hipSuccess; } - +// this method is added for allocation of managed var +hipError_t Var::allocateManagedVarPtr() { + void** pointer = static_cast(managedVarPtr_); + // check if it is deffered allocation + if (!allocFlag_) { + // Allocate managed memory for this var + const bool use_host_ptr = true; + IHIP_RETURN_ONFAIL(ihipMallocManaged(pointer, size_, align_, use_host_ptr)); + allocFlag_ = true; + } + if (dVar_.empty()) { + resize_dVar(g_devices.size()); + } + return hipSuccess; +} }; //namespace: hip diff --git a/hipamd/src/hip_global.hpp b/hipamd/src/hip_global.hpp index dc8fb94132..8ae8f52608 100644 --- a/hipamd/src/hip_global.hpp +++ b/hipamd/src/hip_global.hpp @@ -111,7 +111,7 @@ public: Var(const std::string& name, DeviceVarKind dVarKind, size_t size, int type, int norm, FatBinaryInfo** modules = nullptr); - Var(const std::string& name, DeviceVarKind dVarKind, void *pointer, size_t size, unsigned align, + Var(const std::string& name, DeviceVarKind dVarKind, void* pointer, size_t size, unsigned align, FatBinaryInfo** modules = nullptr); ~Var(); @@ -124,20 +124,28 @@ public: hipError_t getDeviceVarPtr(DeviceVar** dvar, int deviceId); + hipError_t allocateManagedVarPtr(); + void resize_dVar(size_t size) { dVar_.resize(size); } + //bool isEmpty_dVar() const { return dVar_.empty(); } + FatBinaryInfo** moduleInfo() { return modules_; }; DeviceVarKind getVarKind() const { return dVarKind_; } size_t getSize() const { return size_; } + size_t getAlignment() const { return align_; } std::string getName() const { return name_; } - void* getManagedVarPtr() { return managedVarPtr_; }; + void* getManagedVarPtr() const { return managedVarPtr_; } void setManagedVarInfo(void* pointer, size_t size) { managedVarPtr_ = pointer; size_ = size; dVarKind_ = DVK_Managed; } -private: + bool getAllocFlag() const { return allocFlag_; } + void setAllocFlag(bool val) { allocFlag_ = val; } + + private: std::vector dVar_; // DeviceVarObj per Device std::string name_; // Variable name (not unique identifier) DeviceVarKind dVarKind_; // Variable kind @@ -145,9 +153,9 @@ private: int type_; // Type(Textures/Surfaces only) int norm_; // Type(Textures/Surfaces only) FatBinaryInfo** modules_; // static module where it is referenced - - void *managedVarPtr_; // Managed memory pointer with size_ & align_ - unsigned int align_; // Managed memory alignment + void* managedVarPtr_; // Managed memory pointer with size_ & align_ + size_t align_; // Managed memory alignment + bool allocFlag_; // 0 : host alloc, 1: managed alloc }; }; //namespace: hip diff --git a/hipamd/src/hip_hmm.cpp b/hipamd/src/hip_hmm.cpp old mode 100644 new mode 100755 index 15b9105bff..be7152edcf --- a/hipamd/src/hip_hmm.cpp +++ b/hipamd/src/hip_hmm.cpp @@ -28,7 +28,7 @@ namespace hip { // Forward declaraiton of a function -hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0); +hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align = 0, bool use_host_ptr = 0); // Make sure HIP defines match ROCclr to avoid double conversion static_assert(hipCpuDeviceId == amd::CpuDeviceId, "CPU device ID mismatch with ROCclr!"); @@ -76,7 +76,7 @@ hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) { HIP_RETURN(hipErrorStreamCaptureUnsupported); } - HIP_RETURN(ihipMallocManaged(dev_ptr, size), *dev_ptr); + HIP_RETURN(ihipMallocManaged(dev_ptr, size, 0, 0), *dev_ptr); } // ================================================================================================ @@ -284,7 +284,7 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr, } // ================================================================================================ -hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) { +hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align, bool use_host_ptr) { if (ptr == nullptr) { return hipErrorInvalidValue; } else if (size == 0) { @@ -299,9 +299,15 @@ hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) { // Allocate SVM fine grain buffer with the forced host pointer, avoiding explicit memory // allocation in the device driver - *ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR, - size, (align == 0) ? dev.info().memBaseAddrAlign_ : align); - + if (use_host_ptr) { + // If the host pointer is already allocated, map it to svm fine grain buffer + *ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR, size, + (align == 0) ? dev.info().memBaseAddrAlign_ : align, nullptr, + *ptr); + } else { + *ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR, size, + (align == 0) ? dev.info().memBaseAddrAlign_ : align); + } if (*ptr == nullptr) { return hipErrorMemoryAllocation; } diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp index 7d91f209cd..8f7cb99e3b 100644 --- a/hipamd/src/hip_platform.cpp +++ b/hipamd/src/hip_platform.cpp @@ -35,7 +35,7 @@ constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF" PlatformState* PlatformState::platform_; // Initiaized as nullptr by default // forward declaration of methods required for __hipRegisrterManagedVar -hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0); +hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align = 0, bool use_host_ptr = 0); struct __CudaFatBinaryWrapper { unsigned int magic; @@ -152,24 +152,42 @@ void __hipRegisterManagedVar( void* init_value, // Initial value to be copied into \p pointer const char* name, // Name of the variable in code object size_t size, unsigned align) { - HIP_INIT_VOID(); - hipError_t status = ihipMallocManaged(pointer, size, align); - if (status == hipSuccess) { - hip::Stream* stream = hip::getNullStream(); - if (stream != nullptr) { - status = ihipMemcpy(*pointer, init_value, size, hipMemcpyHostToDevice, *stream); - guarantee((status == hipSuccess), "Error during memcpy to managed memory, error:%d!", - status); - } else { - ClPrint(amd::LOG_ERROR, amd::LOG_API, "Host Queue is NULL"); - } - } else { - guarantee(false, "Error during allocation of managed memory!, error: %d", status); - } + + static int enable_deferred_loading{[]() { + #ifdef _WIN32 // Don't defer loading for windows + return 0; + #else + char* var = getenv("HIP_ENABLE_DEFERRED_LOADING"); + return var ? atoi(var) : 1; + #endif + }()}; + hipError_t hip_error = hipSuccess; hip::Var* var_ptr = new hip::Var(std::string(name), hip::Var::DeviceVarKind::DVK_Managed, pointer, size, align, reinterpret_cast(hipModule)); - status = PlatformState::instance().registerStatManagedVar(var_ptr); + hipError_t status = PlatformState::instance().registerStatManagedVar(var_ptr); guarantee((status == hipSuccess), "Cannot register Static Managed Var, error: %d", status); + + if (enable_deferred_loading) { + // Allocate temporary var on host and initialize + *pointer = amd::Os::reserveMemory(0, size, align, amd::Os::MEM_PROT_RW); + ::memcpy(*pointer, init_value, size); + } else { + HIP_INIT_VOID(); + hipError_t status = ihipMallocManaged(pointer, size, align, 0); + var_ptr->setAllocFlag(true); // set flag true for managed alloc + if (status == hipSuccess) { + hip::Stream* stream = hip::getNullStream(); + if (stream != nullptr) { + status = ihipMemcpy(*pointer, init_value, size, hipMemcpyHostToDevice, *stream); + guarantee((status == hipSuccess), "Error during memcpy to managed memory, error:%d!", + status); + } else { + ClPrint(amd::LOG_ERROR, amd::LOG_API, "Host Queue is NULL"); + } + } else { + guarantee(false, "Error during allocation of managed memory!, error: %d", status); + } + } } void __hipRegisterTexture( @@ -742,6 +760,11 @@ void PlatformState::init() { for (auto& it : statCO_.vars_) { it.second->resize_dVar(g_devices.size()); } + for (auto& it : statCO_.managedVars_) { + for (auto& var: it.second) { + var->resize_dVar(g_devices.size()); + } + } for (auto& it : statCO_.functions_) { it.second->resize_dFunc(g_devices.size()); } diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 312904e8c0..44f3822d33 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -2288,45 +2288,46 @@ void Device::updateFreeMemory(size_t size, bool free) { void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_mem_flags flags, void* svmPtr) const { amd::Memory* mem = nullptr; - - if (nullptr == svmPtr) { - // create a hidden buffer, which will allocated on the device later - mem = new (context) amd::Buffer(context, flags, size, - reinterpret_cast(amd::Memory::MemoryType::kSvmMemoryPtr)); - if (mem == nullptr) { - LogError("failed to create a svm mem object!"); - return nullptr; - } - - if (!mem->create(nullptr)) { - LogError("failed to create a svm hidden buffer!"); - mem->release(); - return nullptr; - } - // if the device supports SVM FGS, return the committed CPU address directly. - Memory* gpuMem = getRocMemory(mem); - if (gpuMem == nullptr) { - LogError("failed to create GPU memory from svm hidden buffer!"); - return nullptr; - } - - // add the information to context so that we can use it later. - if (mem->getSvmPtr() != nullptr) { - amd::MemObjMap::AddMemObj(mem->getSvmPtr(), mem); - } - svmPtr = mem->getSvmPtr(); - } else { + void* svmPtrUsed = reinterpret_cast(amd::Memory::MemoryType::kSvmMemoryPtr); + + if (nullptr != svmPtr) { // Find the existing amd::mem object mem = amd::MemObjMap::FindMemObj(svmPtr); - if (nullptr == mem) { + if (mem != nullptr) { + return mem->getSvmPtr(); + } + if (flags & CL_MEM_USE_HOST_PTR ) { + svmPtrUsed = svmPtr; + } else { DevLogPrintfError("Cannot find svm_ptr: 0x%x \n", svmPtr); return nullptr; } - - svmPtr = mem->getSvmPtr(); + } + + // create a hidden buffer, which will allocated on the device later + mem = new (context) amd::Buffer(context, flags, size, svmPtrUsed); + if (mem == nullptr) { + LogError("failed to create a svm mem object!"); + return nullptr; } - return svmPtr; + if (!mem->create(nullptr)) { + LogError("failed to create a svm hidden buffer!"); + mem->release(); + return nullptr; + } + // if the device supports SVM FGS, return the committed CPU address directly. + Memory* gpuMem = getRocMemory(mem); + if (gpuMem == nullptr) { + LogError("failed to create GPU memory from svm hidden buffer!"); + return nullptr; + } + + // add the information to context so that we can use it later. + if (mem->getSvmPtr() != nullptr) { + amd::MemObjMap::AddMemObj(mem->getSvmPtr(), mem); + } + return mem->getSvmPtr(); } void* Device::virtualAlloc(void* req_addr, size_t size, size_t alignment) { diff --git a/rocclr/device/rocm/rocmemory.cpp b/rocclr/device/rocm/rocmemory.cpp index 1b7c2fb3e3..e8d1ff3d50 100644 --- a/rocclr/device/rocm/rocmemory.cpp +++ b/rocclr/device/rocm/rocmemory.cpp @@ -656,7 +656,7 @@ void Buffer::destroy() { if (kind_ != MEMORY_KIND_PTRGIVEN) { if (isFineGrain) { - if (memFlags & CL_MEM_ALLOC_HOST_PTR) { + if (memFlags & (CL_MEM_ALLOC_HOST_PTR)) { if (dev().info().hmmSupported_) { // AMD HMM path. Release reserved system memory dev().releaseMemory(deviceMemory_, size()); @@ -675,6 +675,17 @@ void Buffer::destroy() { } else { dev().memFree(deviceMemory_, size()); } + } else { + if (memFlags & CL_MEM_USE_HOST_PTR) { + // unlock svm host pointer from memory pool + if (!dev().info().hmmSupported_) { + hsa_amd_memory_unlock(owner()->getSvmPtr()); + } + // destroy system memory + if (!(amd::Os::releaseMemory(deviceMemory_, size()))) { + ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "[ROCClr] munmap failed \n"); + } + } } if ((deviceMemory_ != nullptr) && @@ -812,7 +823,7 @@ bool Buffer::create(bool alloc_local) { if (deviceMemory_ == NULL) { return false; } - // Currently HMM requires cirtain initial calls to mark sysmem allocation as + // Currently HMM requires certain initial calls to mark sysmem allocation as // GPU accessible or prefetch memory into GPU if (!dev().SvmAllocInit(deviceMemory_, size())) { ClPrint(amd::LOG_ERROR, amd::LOG_MEM, "SVM init in ROCr failed!"); @@ -864,6 +875,19 @@ bool Buffer::create(bool alloc_local) { } else { kind_ = MEMORY_KIND_PTRGIVEN; } + if (memFlags & CL_MEM_USE_HOST_PTR) { + if (dev().info().hmmSupported_) { + // Currently HMM requires certain initial calls to mark sysmem allocation as + // GPU accessible or prefetch memory into GPU + if (!dev().SvmAllocInit(deviceMemory_, size())) { + ClPrint(amd::LOG_ERROR, amd::LOG_MEM, "SVM init in ROCr failed!"); + return false; + } + } else { + deviceMemory_ = dev().hostLock(owner()->getSvmPtr(), size(), + getHostMemorySegment(memFlags)); + } + } } if ((deviceMemory_ != nullptr) && (dev().settings().apuSystem_ || !isFineGrain) diff --git a/rocclr/platform/context.cpp b/rocclr/platform/context.cpp index 6d0bc27288..3378cf624b 100644 --- a/rocclr/platform/context.cpp +++ b/rocclr/platform/context.cpp @@ -324,13 +324,13 @@ void Context::hostFree(void* ptr) const { } void* Context::svmAlloc(size_t size, size_t alignment, cl_svm_mem_flags flags, - const amd::Device* curDev) { + const amd::Device* curDev, void* svmPtr) { unsigned int numSVMDev = svmAllocDevice_.size(); if (numSVMDev < 1) { return nullptr; } - void* svmPtrAlloced = nullptr; + void* svmPtrAlloced = svmPtr; amd::ScopedLock lock(&ctxLock_); diff --git a/rocclr/platform/context.hpp b/rocclr/platform/context.hpp index 904d92e0aa..af64598ced 100644 --- a/rocclr/platform/context.hpp +++ b/rocclr/platform/context.hpp @@ -153,7 +153,7 @@ class Context : public RuntimeObject { * @param curDev The current device */ void* svmAlloc(size_t size, size_t alignment, cl_svm_mem_flags flags = CL_MEM_READ_WRITE, - const amd::Device* curDev = nullptr); + const amd::Device* curDev = nullptr, void* svmPtr = nullptr); /** * Release SVM buffer diff --git a/rocclr/platform/memory.cpp b/rocclr/platform/memory.cpp index 24e6bcf567..95aca41e24 100644 --- a/rocclr/platform/memory.cpp +++ b/rocclr/platform/memory.cpp @@ -1553,8 +1553,8 @@ bool SvmBuffer::Contains(uintptr_t ptr) { // The allocation flags are ignored for now. void* SvmBuffer::malloc(Context& context, cl_svm_mem_flags flags, size_t size, size_t alignment, - const amd::Device* curDev) { - void* ret = context.svmAlloc(size, alignment, flags, curDev); + const amd::Device* curDev, void* hostptr) { + void* ret = context.svmAlloc(size, alignment, flags, curDev, hostptr); if (ret == nullptr) { LogError("Unable to allocate aligned memory"); return nullptr; diff --git a/rocclr/platform/memory.hpp b/rocclr/platform/memory.hpp index fe08a5985b..a084edfeaa 100644 --- a/rocclr/platform/memory.hpp +++ b/rocclr/platform/memory.hpp @@ -676,8 +676,7 @@ class SvmBuffer : AllStatic { public: //! Allocate a shared buffer that is accessible by all devices in the context static void* malloc(Context& context, cl_svm_mem_flags flags, size_t size, size_t alignment, - const amd::Device* curDev = nullptr); - + const amd::Device* curDev = nullptr, void* hostptr = nullptr); //! Release shared buffer static void free(const Context& context, void* ptr);