SWDEV-241896 - Remove extern support in ROCclr/HIP since compiler added Texture reference support.

Change-Id: I1b0a7377b803b4e2b94ffef6ffd91e376fdb2b2f


[ROCm/clr commit: 19acb0bfe5]

projects/clr/hipamd/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

projects/clr/hipamd/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;

projects/clr/hipamd/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:

projects/clr/hipamd/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;

projects/clr/hipamd/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<hipModule_t>(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_;
 

projects/clr/hipamd/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);