Re-sync with upstream and re-factor platform global management for texture references.

[ROCm/clr commit: 6a0efb7ed2]

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

@@ -23,6 +23,7 @@ THE SOFTWARE.
 #pragma once
 
 #include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
 
 #include <cstddef>
 #include <istream>
@@ -68,12 +69,15 @@ 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();
 
     hsa_executable_t load_executable(
         hsa_executable_t executable, hsa_agent_t agent, std::istream& file);

projects/clr/hipamd/src/hip_module.cpp

@@ -725,9 +725,9 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
         if(name == NULL || hmod == NULL){
             ret = hipErrorNotInitialized;
         } else{
-            const auto it = hmod->coGlobals.find(name);
-		    if (it == hmod->coGlobals.end()) return ihipLogStatus(hipErrorInvalidValue);
-		    *texRef = reinterpret_cast<textureReference*>(it->second);
+            const auto it = hip_impl::globals().find(name);
+		    if (it == hip_impl::globals().end()) return ihipLogStatus(hipErrorInvalidValue);
+		    *texRef = reinterpret_cast<textureReference*>(it->second.get());
             ret = hipSuccess;
         }
     }

projects/clr/hipamd/src/program_state.cpp

@@ -158,14 +158,7 @@ namespace
             symbol_section_accessor{reader, code_object_dynsym});
 
         for (auto&& x : undefined_symbols) {
-            using RAII_global =
-                unique_ptr<void, decltype(hsa_amd_memory_unlock)*>;
-
-            static unordered_map<string, RAII_global> globals;
-            static once_flag f;
-            call_once(f, [=]() { globals.reserve(symbol_addresses().size()); });
-
-            if (globals.find(x) != globals.cend()) return;
+            if (globals().find(x) != globals().cend()) return;
 
             const auto it1 = symbol_addresses().find(x);
 
@@ -176,7 +169,7 @@ namespace
             static mutex mtx;
             lock_guard<mutex> lck{mtx};
 
-            if (globals.find(x) != globals.cend()) return;
+            if (globals().find(x) != globals().cend()) return;
 
             void* p = nullptr;
             hsa_amd_memory_lock(
@@ -186,12 +179,10 @@ namespace
                 0,
                 &p);
 
-            if (!p) { cerr << it1->first << endl; assert(false); }
-
             hsa_executable_agent_global_variable_define(
                 executable, agent, x.c_str(), p);
 
-            globals.emplace(x, RAII_global{p, hsa_amd_memory_unlock});
+            globals().emplace(x, RAII_global{p, hsa_amd_memory_unlock});
         }
     }
 
@@ -534,6 +525,15 @@ namespace hip_impl
         return r;
     }
 
+    unordered_map<string, RAII_global>& globals()
+    {
+        static unordered_map<string, RAII_global> r;
+        static once_flag f;
+        call_once(f, []() { r.reserve(symbol_addresses().size()); });
+
+        return r;
+    }
+
     hsa_executable_t load_executable(
         hsa_executable_t executable, hsa_agent_t agent, istream& file)
     {