AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-539710 - Defer allocation of managed variable (#652)

Просмотр исходного кода 
Co-authored-by: Rahul Manocha <rmanocha@amd.com>
Этот коммит содержится в:
Manocha, Rahul
2025-07-31 08:30:23 -07:00
коммит произвёл GitHub
родитель c88f345229
Коммит 3f6f9d6081
11 изменённых файлов: 177 добавлений и 87 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
hipamd/src/hip_code_object.cpp
+18 -13
Просмотреть файл
@@ -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<void**>(managedVar->getManagedVarPtr())));
if (managedVar->getAllocFlag()) { // check if it is a managed or host alloc
err = ihipFree(*(static_cast<void**>(managedVar->getManagedVarPtr())));
} else {
void** pointer = static_cast<void**>(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<address>(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<address>(dvar->device_ptr()), var->getManagedVarPtr(),
dvar->size(), hipMemcpyHostToDevice, *stream);
}
}
managedVarsDevicePtrInitalized_[deviceId] = true;
hipamd/src/hip_global.cpp
+29 -5
Просмотреть файл
@@ -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<void**>(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
hipamd/src/hip_global.hpp
+14 -6
Просмотреть файл
@@ -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<DeviceVar*> 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
hipamd/src/hip_hmm.cpp
Обычный файл → Исполняемый файл
+12 -6
Просмотреть файл
@@ -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;
}
hipamd/src/hip_platform.cpp
+39 -16
Просмотреть файл
@@ -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<hip::FatBinaryInfo**>(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());
}
rocclr/device/rocm/rocdevice.cpp
+33 -32
Просмотреть файл
@@ -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<void*>(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<void*>(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) {
rocclr/device/rocm/rocmemory.cpp
+26 -2
Просмотреть файл
@@ -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)
rocclr/platform/context.cpp
+2 -2
Просмотреть файл
@@ -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_);
rocclr/platform/context.hpp
+1 -1
Просмотреть файл
@@ -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
rocclr/platform/memory.cpp
+2 -2
Просмотреть файл
@@ -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;
rocclr/platform/memory.hpp
+1 -2
Просмотреть файл
@@ -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);