diff --git a/rocclr/hip_code_object.cpp b/rocclr/hip_code_object.cpp index c6a866c9c4..b0979c5246 100755 --- a/rocclr/hip_code_object.cpp +++ b/rocclr/hip_code_object.cpp @@ -202,19 +202,10 @@ hipError_t DynCO::populateDynGlobalVars() { return hipErrorSharedObjectSymbolNotFound; } - if (!dev_program->getUndefinedVarFromCodeObj(&undef_var_names)) { - DevLogPrintfError("Could not get undefined Variables for Module: 0x%x \n", module()); - return hipErrorSharedObjectSymbolNotFound; - } - for (auto& elem : var_names) { vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Variable, 0, 0, 0, nullptr))); } - for (auto& elem : undef_var_names) { - vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Texture, 0, 0, 0, nullptr))); - } - return hipSuccess; } @@ -377,20 +368,4 @@ hipError_t StatCO::getStatGlobalVar(const void* hostVar, int deviceId, hipDevice *size_ptr = dvar->size(); return hipSuccess; } - -hipError_t StatCO::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, - hipDeviceptr_t* dev_ptr, size_t* size_ptr) { - amd::ScopedLock lock(sclock_); - - for (auto& elem : vars_) { - if ((elem.second->name() == hostVar) - && (elem.second->module(deviceId) == hmod)) { - *dev_ptr = elem.second->device_ptr(deviceId); - *size_ptr = elem.second->device_size(deviceId); - return hipSuccess; - } - } - - return hipErrorNotFound; -} }; //namespace: hip diff --git a/rocclr/hip_code_object.hpp b/rocclr/hip_code_object.hpp index f5f179570b..0cc2a7051a 100755 --- a/rocclr/hip_code_object.hpp +++ b/rocclr/hip_code_object.hpp @@ -118,8 +118,6 @@ public: hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId); hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, size_t* size_ptr); - hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, - hipDeviceptr_t* dev_ptr, size_t* size_ptr); private: friend class ::PlatformState; diff --git a/rocclr/hip_global.hpp b/rocclr/hip_global.hpp index 3888daf30b..fd57ecfb50 100755 --- a/rocclr/hip_global.hpp +++ b/rocclr/hip_global.hpp @@ -95,11 +95,6 @@ public: hipError_t getStatDeviceVar(DeviceVar** dvar, int deviceId); void resize_dVar(size_t size) { dVar_.resize(size); } - //Accessor for device_ptrs. - std::string name() const { return name_; } - hipModule_t module(int deviceId) const { return nullptr; } - hipDeviceptr_t device_ptr(int deviceId) const { return dVar_[deviceId]->device_ptr(); } - size_t device_size(int deviceId) const { return dVar_[deviceId]->size(); } FatBinaryInfo** moduleInfo() { return modules_; }; private: diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index a950961ea7..7e0cc8b9a2 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -252,8 +252,6 @@ extern int ihipGetDevice(); extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags); extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size); -extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, - size_t* var_size); constexpr bool kOptionChangeable = true; constexpr bool kNewDevProg = false; diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 03bebeb7fb..6abea0df4e 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -80,27 +80,6 @@ extern "C" hip::FatBinaryInfo** __hipRegisterFatBinary(const void* data) return PlatformState::instance().addFatBinary(fbwrapper->binary); } -bool PlatformState::getShadowVarInfo(std::string var_name, hipModule_t hmod, - void** var_addr, size_t* var_size) { - - amd::ScopedLock lock(lock_); - if (hipSuccess == getDynGlobalVar(var_name.c_str(), ihipGetDevice(), hmod, var_addr, var_size)) { - return true; - } - - if (hipSuccess == getStatGlobalVarByName(var_name, ihipGetDevice(), hmod, var_addr, var_size)) { - return true; - } - - return false; -} - -bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, - size_t* var_size) { - return PlatformState::instance().getShadowVarInfo(var_name, reinterpret_cast(program), - var_addr, var_size); -} - extern "C" void __hipRegisterFunction( hip::FatBinaryInfo** modules, const void* hostFunction, @@ -879,11 +858,6 @@ hipError_t PlatformState::getStatGlobalVar(const void* hostVar, int deviceId, hi return statCO_.getStatGlobalVar(hostVar, deviceId, dev_ptr, size_ptr); } -hipError_t PlatformState::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, - hipDeviceptr_t* dev_ptr, size_t* size_ptr) { - return statCO_.getStatGlobalVarByName(hostVar, deviceId, hmod, dev_ptr, size_ptr); -} - void PlatformState::setupArgument(const void *arg, size_t size, size_t offset) { auto& arguments = execStack_.top().arguments_; diff --git a/rocclr/hip_platform.hpp b/rocclr/hip_platform.hpp index 2bcf620f6d..51fea0841e 100755 --- a/rocclr/hip_platform.hpp +++ b/rocclr/hip_platform.hpp @@ -77,11 +77,6 @@ public: hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId); hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, size_t* size_ptr); - hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, - hipDeviceptr_t* dev_ptr, size_t* size_ptr); - - bool getShadowVarInfo(std::string var_name, hipModule_t hmod, - void** var_addr, size_t* var_size); //Exec Functions void setupArgument(const void *arg, size_t size, size_t offset);