Merge pull request #853 from ROCm-Developer-Tools/hip-module-hash-fix

Introduce hash key to HIP module implementation

[ROCm/clr commit: d3a5502f17]

projects/clr/hipamd/src/hip_hcc_internal.h

@@ -374,6 +374,7 @@ struct ihipModule_t {
     std::string fileName;
     hsa_executable_t executable = {};
     hsa_code_object_reader_t coReader = {};
+    std::string hash;
 
     ~ihipModule_t() {
         if (executable.handle) hsa_executable_destroy(executable);

projects/clr/hipamd/src/hip_module.cpp

@@ -55,6 +55,21 @@ using namespace ELFIO;
 using namespace hip_impl;
 using namespace std;
 
+// calculate MD5 checksum
+inline std::string checksum(size_t size, const char *source) {
+    // FNV-1a hashing, 64-bit version
+    const uint64_t FNV_prime = 0x100000001b3;
+    const uint64_t FNV_basis = 0xcbf29ce484222325;
+    uint64_t hash = FNV_basis;
+
+    const char *str = static_cast<const char *>(source);
+    for (auto i = 0; i < size; ++i) {
+        hash ^= *str++;
+        hash *= FNV_prime;
+    }
+    return std::to_string(hash);
+}
+
 inline uint64_t alignTo(uint64_t Value, uint64_t Align, uint64_t Skew = 0) {
     assert(Align != 0u && "Align can't be 0.");
     Skew %= Align;
@@ -332,21 +347,21 @@ pair<hipDeviceptr_t, size_t> read_global_description(ForwardIterator f, ForwardI
 
 hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
                                          const char* name) {
-    static unordered_map<hipModule_t, vector<Agent_global>> agent_globals;
+    // the key of the map would the hash of code object associated with the
+    // hipModule_t instance
+    static unordered_map<std::string, vector<Agent_global>> agent_globals;
+    auto key = hmod->hash;
 
-    // TODO: this is not particularly robust.
-    if (agent_globals.count(hmod) == 0) {
+    if (agent_globals.count(key) == 0) {
         static mutex mtx;
         lock_guard<mutex> lck{mtx};
 
-        if (agent_globals.count(hmod) == 0) {
-            agent_globals.emplace(hmod, read_agent_globals(this_agent(), hmod->executable));
+        if (agent_globals.count(key) == 0) {
+            agent_globals.emplace(key, read_agent_globals(this_agent(), hmod->executable));
         }
     }
 
-    // TODO: This is unsafe iff some other emplacement triggers rehashing.
-    //       It will have to be properly fleshed out in the future.
-    const auto it0 = agent_globals.find(hmod);
+    const auto it0 = agent_globals.find(key);
     if (it0 == agent_globals.cend()) {
         throw runtime_error{"agent_globals data structure corrupted."};
     }
@@ -568,8 +583,14 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
 
     auto tmp = code_object_blob_for_agent(image, this_agent());
 
-    (*module)->executable = hip_impl::load_executable(
-        tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent());
+    auto content = tmp.empty() ? read_elf_file_as_string(image) : tmp;
+
+    (*module)->executable = hip_impl::load_executable(content,
+                                                      (*module)->executable,
+                                                      this_agent());
+
+    // compute the hash of the code object
+    (*module)->hash = checksum(content.length(), content.data());
 
     return (*module)->executable.handle ? hipSuccess : hipErrorUnknown;
 }