Fixed host allocated globals address lookup for host usage

Fixed texture driver APIs failure


[ROCm/hip commit: 24ab820a11]

projects/hip/include/hip/hcc_detail/program_state.hpp

@@ -69,18 +69,16 @@ namespace hip_impl
         }
     };
 
-    using RAII_global = std::unique_ptr<void, decltype(hsa_amd_memory_unlock)*>;
-
     const std::unordered_map<
         hsa_agent_t, std::vector<hsa_executable_t>>& executables();
     const std::unordered_map<
         std::uintptr_t,
         std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>& functions();
     const std::unordered_map<std::uintptr_t, std::string>& function_names();
-    std::unordered_map<std::string, RAII_global>& globals();
+    std::unordered_map<std::string, void*>& globals();
 
     hsa_executable_t load_executable(
         const std::string& file,
         hsa_executable_t executable,
         hsa_agent_t agent);
-} // Namespace hip_impl.
+} // Namespace hip_impl.

projects/hip/src/hip_memory.cpp

@@ -453,6 +453,7 @@ hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAlloc
         array[0]->width = pAllocateArray->width;
         array[0]->height = pAllocateArray->height;
         array[0]->isDrv = true;
+        array[0]->textureType = hipTextureType2D;
         void ** ptr = &array[0]->data;
 	    if (ctx) {
             const unsigned am_flags = 0;

projects/hip/src/hip_module.cpp

@@ -568,7 +568,6 @@ hipError_t hipModuleGetTexRef(
     const auto it = globals().find(name);
     if (it == globals().end()) return ihipLogStatus(hipErrorInvalidValue);
 
-    *texRef = static_cast<textureReference*>(it->second.get());
-
+    *texRef = reinterpret_cast<textureReference*>(it->second);
     return ihipLogStatus(hipSuccess);
 }

projects/hip/src/program_state.cpp

@@ -169,7 +169,7 @@ namespace
             lock_guard<mutex> lck{mtx};
 
             if (globals().find(x) != globals().cend()) return;
-
+            globals().emplace(x, (void*)(it1->second.first));
             void* p = nullptr;
             hsa_amd_memory_lock(
                 reinterpret_cast<void*>(it1->second.first),
@@ -181,7 +181,6 @@ namespace
             hsa_executable_agent_global_variable_define(
                 executable, agent, x.c_str(), p);
 
-            globals().emplace(x, RAII_global{p, hsa_amd_memory_unlock});
         }
     }
 
@@ -462,9 +461,9 @@ namespace hip_impl
         return r;
     }
 
-    unordered_map<string, RAII_global>& globals()
+    unordered_map<string, void*>& globals()
     {
-        static unordered_map<string, RAII_global> r;
+        static unordered_map<string, void*> r;
         static once_flag f;
         call_once(f, []() { r.reserve(symbol_addresses().size()); });
 
@@ -491,4 +490,4 @@ namespace hip_impl
 
         return executable;
     }
-} // Namespace hip_impl.
+} // Namespace hip_impl.