From 2f64c7678914b483d4b1454b33371cbb6755b0d4 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Tue, 8 Jan 2019 17:18:06 +0000 Subject: [PATCH] Introduce hash key to HIP module implementation A hash calculated via FNV-1a algorithm is introduced in ihipModule_t, the internal of hipModule_t. The hash is used by HIP module APIs such as - read_agent_global_from_module to determine whether the agent-scope globals for a module have been iterated. This commit fixes one issue that applications which load / unload modules frequently would occasionally fail. After deep investigation of the issue it turns out the old implementation in read_agent_global_from_module uses hipModule_t as the key, which is not robust enough, as hipModule_t instances are allocated dynamically so there are cases that one memory address may be used by multiple hipModule_t instances. The real solution is to introduce a uniquely identifiable hash for the code object associated with the HIP module. And that's the rationale behind this commit. [ROCm/clr commit: 6e68d44220a53d07a26a89df11116f18ef502d7e] --- projects/clr/hipamd/src/hip_hcc_internal.h | 1 + projects/clr/hipamd/src/hip_module.cpp | 41 ++++++++++++++++------ 2 files changed, 32 insertions(+), 10 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index d71ad850e3..be794956f6 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/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); diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 657fb06b5e..780240c067 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/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(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 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> agent_globals; + // the key of the map would the hash of code object associated with the + // hipModule_t instance + static unordered_map> 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 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; }