From e77419eb2a3ec5f676d1cbbb1519dccca8d9f323 Mon Sep 17 00:00:00 2001 From: foreman Date: Mon, 4 Nov 2019 10:13:20 -0500 Subject: [PATCH] P4 to Git Change 2024251 by kjayapra@0_HIPWS_LNX1_ROCM on 2019/11/04 10:07:42 SWDEV-206759 - Adding support for duplicate global vars Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#45 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#81 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#46 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#48 edit ... //depot/stg/opencl/drivers/opencl/make/hip.git/tests/build/Makefile.hip_tests#30 edit [ROCm/hip commit: 33038437b3eda6998d04ac7a4cf29ead945e3ac8] --- projects/hip/api/hip/hip_internal.hpp | 11 +-- projects/hip/api/hip/hip_memory.cpp | 16 ++--- projects/hip/api/hip/hip_module.cpp | 16 ++++- projects/hip/api/hip/hip_platform.cpp | 96 +++++++++++++++++++-------- 4 files changed, 96 insertions(+), 43 deletions(-) diff --git a/projects/hip/api/hip/hip_internal.hpp b/projects/hip/api/hip/hip_internal.hpp index ba4a300739..1ce259f708 100644 --- a/projects/hip/api/hip/hip_internal.hpp +++ b/projects/hip/api/hip/hip_internal.hpp @@ -155,7 +155,7 @@ public: }; private: std::unordered_map functions_; - std::unordered_map vars_; + std::unordered_multimap vars_; static PlatformState* platform_; @@ -168,16 +168,19 @@ public: std::vector< std::pair >* unregisterVar(hipModule_t hmod); + + PlatformState::DeviceVar* findVar(std::string hostVar, int deviceId, hipModule_t hmod); void registerVar(const void* hostvar, const DeviceVar& var); void registerFunction(const void* hostFunction, const DeviceFunction& func); hipFunction_t getFunc(const void* hostFunction, int deviceId); bool getFuncAttr(const void* hostFunction, hipFuncAttributes* func_attr); - bool getGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, - size_t* size_ptr); + bool getGlobalVar(const void* hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr); bool getTexRef(const char* hostVar, textureReference** texRef); - bool getShadowVarInfo(std::string var_name, void** var_addr, size_t* var_size); + bool getShadowVarInfo(std::string var_name, hipModule_t hmod, + void** var_addr, size_t* var_size); void setupArgument(const void *arg, size_t size, size_t offset); void configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index 6e07ddd30b..8d35a5aef7 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/projects/hip/api/hip/hip_memory.cpp @@ -587,8 +587,8 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou hipDeviceptr_t device_ptr = nullptr; /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, - &sym_size)) { + if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -611,8 +611,8 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, hipDeviceptr_t device_ptr = nullptr; /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, - &sym_size)) { + if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -635,8 +635,8 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_ hipDeviceptr_t device_ptr = nullptr; /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, - &sym_size)) { + if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -659,8 +659,8 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co hipDeviceptr_t device_ptr = nullptr; /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &device_ptr, - &sym_size)) { + if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } diff --git a/projects/hip/api/hip/hip_module.cpp b/projects/hip/api/hip/hip_module.cpp index 77a2359614..280e804570 100644 --- a/projects/hip/api/hip/hip_module.cpp +++ b/projects/hip/api/hip/hip_module.cpp @@ -183,13 +183,25 @@ hipError_t ihipModuleLoadData(hipModule_t *module, const void *image) program->setVarInfoCallBack(&getSvarInfo); if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentContext()->devices()[0], image, ElfSize(image))) { +>>>> ORIGINAL //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#44 + return hipErrorUnknown; +==== THEIRS //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#45 return hipErrorInvalidKernelFile; +==== YOURS //0_HIPWS_LNX1_ROCM/main/drivers/opencl/api/hip/hip_module.cpp + return hipErrorUnknown; +<<<< } *module = reinterpret_cast(as_cl(program)); if (!ihipModuleRegisterGlobal(program, module)) { +>>>> ORIGINAL //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#44 + return hipErrorUnknown; +==== THEIRS //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#45 return hipErrorSharedObjectSymbolNotFound; +==== YOURS //0_HIPWS_LNX1_ROCM/main/drivers/opencl/api/hip/hip_module.cpp + return hipErrorUnknown; +<<<< } if (!ihipModuleRegisterUndefined(program, module)) { @@ -230,8 +242,8 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), dptr, - bytes)) { + if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), hmod, + dptr, bytes)) { HIP_RETURN(hipErrorNotFound); } diff --git a/projects/hip/api/hip/hip_platform.cpp b/projects/hip/api/hip/hip_platform.cpp index bd6cd55317..0902ce1a08 100644 --- a/projects/hip/api/hip/hip_platform.cpp +++ b/projects/hip/api/hip/hip_platform.cpp @@ -186,6 +186,42 @@ std::vector< std::pair >* PlatformState::unregisterVar(hipMod return rmodules; } +PlatformState::DeviceVar* PlatformState::findVar(std::string hostVar, int deviceId, hipModule_t hmod) { + DeviceVar* dvar = nullptr; + + if (hmod != nullptr) { + // If module is provided, then get the var only from that module + auto var_range = vars_.equal_range(hostVar); + for (auto it = var_range.first; it != var_range.second; ++it) { + if ((*it->second.modules)[deviceId].first == hmod) { + dvar = &(it->second); + break; + } + } + } else { + // If var count is < 2, return the var + if (vars_.count(hostVar) < 2) { + auto it = vars_.find(hostVar); + dvar = ((it == vars_.end()) ? nullptr : &(it->second)); + } else { + // If var count is > 2, return the original var, + // if original var count != 1, return vars_.end()/Invalid + size_t orig_global_count = 0; + auto var_range = vars_.equal_range(hostVar); + for (auto it = var_range.first; it != var_range.second; ++it) { + // when dyn_undef is set, it is a shadow var + if (it->second.dyn_undef == false) { + ++orig_global_count; + dvar = &(it->second); + } + } + dvar = ((orig_global_count == 1) ? dvar : nullptr); + } + } + + return dvar; +} + void PlatformState::registerVar(const void* hostvar, const DeviceVar& rvar) { amd::ScopedLock lock(lock_); @@ -216,12 +252,12 @@ bool ihipGetFuncAttributes(const char* func_name, amd::Program* program, hipFunc return true; } -bool PlatformState::getShadowVarInfo(std::string var_name, void** var_addr, size_t* var_size) { - const auto it = vars_.find(var_name); - if (it != vars_.cend()) { - DeviceVar& dvar = it->second; - *var_addr = dvar.shadowVptr; - *var_size = dvar.size; +bool PlatformState::getShadowVarInfo(std::string var_name, hipModule_t hmod, + void** var_addr, size_t* var_size) { + DeviceVar* dvar = findVar(var_name, ihipGetDevice(), hmod); + if (dvar != nullptr) { + *var_addr = dvar->shadowVptr; + *var_size = dvar->size; return true; } else { return false; @@ -230,7 +266,8 @@ bool PlatformState::getShadowVarInfo(std::string var_name, void** var_addr, size bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, size_t* var_size) { - return PlatformState::instance().getShadowVarInfo(var_name, var_addr, var_size); + return PlatformState::instance().getShadowVarInfo(var_name, reinterpret_cast(program), + var_addr, var_size); } hipFunction_t PlatformState::getFunc(const void* hostFunction, int deviceId) { @@ -294,52 +331,50 @@ bool PlatformState::getFuncAttr(const void* hostFunction, bool PlatformState::getTexRef(const char* hostVar, textureReference** texRef) { amd::ScopedLock lock(lock_); - const auto it = vars_.find(std::string(reinterpret_cast(hostVar))); - if (it == vars_.cend()) { + DeviceVar* dvar = findVar(std::string(hostVar), ihipGetDevice(), nullptr); + if (dvar == nullptr) { return false; } - DeviceVar& dvar = it->second; - if (!dvar.dyn_undef) { + if (!dvar->dyn_undef) { return false; } - *texRef = reinterpret_cast(dvar.shadowVptr); + *texRef = reinterpret_cast(dvar->shadowVptr); return true; } -bool PlatformState::getGlobalVar(const void* hostVar, int deviceId, +bool PlatformState::getGlobalVar(const void* hostVar, int deviceId, hipModule_t hmod, hipDeviceptr_t* dev_ptr, size_t* size_ptr) { amd::ScopedLock lock(lock_); - const auto it = vars_.find(std::string(reinterpret_cast(hostVar))); - if (it != vars_.cend()) { - DeviceVar& dvar = it->second; - if (dvar.rvars[deviceId].getdeviceptr() == nullptr) { + DeviceVar* dvar = findVar(std::string(reinterpret_cast(hostVar)), deviceId, hmod); + if (dvar != nullptr) { + if (dvar->rvars[deviceId].getdeviceptr() == nullptr) { size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; amd::Memory* amd_mem_obj = nullptr; - if (!(*dvar.modules)[deviceId].second) { - amd::Program* program = as_amd(reinterpret_cast((*dvar.modules)[deviceId].first)); + if (!(*dvar->modules)[deviceId].second) { + amd::Program* program = as_amd(reinterpret_cast((*dvar->modules)[deviceId].first)); program->setVarInfoCallBack(&getSvarInfo); if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr)) { return false; } - (*dvar.modules)[deviceId].second = true; + (*dvar->modules)[deviceId].second = true; } - if((hipSuccess == ihipCreateGlobalVarObj(dvar.hostVar.c_str(), (*dvar.modules)[deviceId].first, + if((hipSuccess == ihipCreateGlobalVarObj(dvar->hostVar.c_str(), (*dvar->modules)[deviceId].first, &amd_mem_obj, &device_ptr, &sym_size)) && (device_ptr != nullptr)) { - dvar.rvars[deviceId].size_ = sym_size; - dvar.rvars[deviceId].devicePtr_ = device_ptr; - dvar.rvars[deviceId].amd_mem_obj_ = amd_mem_obj; + dvar->rvars[deviceId].size_ = sym_size; + dvar->rvars[deviceId].devicePtr_ = device_ptr; + dvar->rvars[deviceId].amd_mem_obj_ = amd_mem_obj; amd::MemObjMap::AddMemObj(device_ptr, amd_mem_obj); } else { LogError("[HIP] __hipRegisterVar cannot find kernel for device \n"); } } - *size_ptr = dvar.rvars[deviceId].getvarsize(); - *dev_ptr = dvar.rvars[deviceId].getdeviceptr(); + *size_ptr = dvar->rvars[deviceId].getvarsize(); + *dev_ptr = dvar->rvars[deviceId].getdeviceptr(); return true; } else { return false; @@ -523,7 +558,8 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { size_t size = 0; - if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), devPtr, &size)) { + if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + devPtr, &size)) { HIP_RETURN(hipErrorInvalidSymbol); } HIP_RETURN(hipSuccess); @@ -531,13 +567,15 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbolName) { hipDeviceptr_t devPtr = nullptr; - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), &devPtr, sizePtr)) { + if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + &devPtr, sizePtr)) { HIP_RETURN(hipErrorInvalidSymbol); } HIP_RETURN(hipSuccess); } -hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memory** amd_mem_obj, hipDeviceptr_t* dptr, size_t* bytes) +hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memory** amd_mem_obj, + hipDeviceptr_t* dptr, size_t* bytes) { HIP_INIT();