From 23e9968752ddb2ceb78d9986182dcb503c13d9be Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 6 Mar 2019 18:27:28 +0000 Subject: [PATCH 1/7] Fix hash_for undefined reference in hipTestConstant test Issue: mismatch undefined symbols in different user env - Binary expects modified return value std::string& - Fails to match libhip_hcc.so: return value is std::string& but doesn't match modified C++ env Fix: Change return value to char*, create new key std::string in header from char* --- include/hip/hcc_detail/hip_runtime_api.h | 5 +++-- src/hip_module.cpp | 4 ++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 7b7ffb0638..72cfbf7d3b 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2518,7 +2518,7 @@ struct Agent_global { namespace hip_impl { hsa_executable_t executable_for(hipModule_t); -const std::string& hash_for(hipModule_t); +const char* hash_for(hipModule_t); template std::pair read_global_description( @@ -2543,7 +2543,8 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, // hipModule_t instance static std::unordered_map< std::string, std::vector> agent_globals; - auto key = hash_for(hmod); + const char* hash = hash_for(hmod); + std::string key(hash); if (agent_globals.count(key) == 0) { static std::mutex mtx; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 52d1fb38bd..178b0c9471 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -287,8 +287,8 @@ namespace hip_impl { return hmod->executable; } - const std::string& hash_for(hipModule_t hmod) { - return hmod->hash; + const char* hash_for(hipModule_t hmod) { + return hmod->hash.c_str(); } hsa_agent_t this_agent() { From 00d24d254dcc1097555f3d4e893ba4cb17e49969 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 6 Mar 2019 18:33:43 +0000 Subject: [PATCH 2/7] Fix Agent_global variables failing hipTestDeviceSymbol Issue: Header uses std::vector agent_globals which is created by hip_module.cpp - Move iterator fails to copy Agent_global from library source into header version - Due to different versions of std::string name in struct Agent_global Fix: Change Agent_global to use char* name instead of std::string name --- include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- src/hip_module.cpp | 7 +++++-- 2 files changed, 7 insertions(+), 4 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 72cfbf7d3b..f6ac507c6f 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2506,7 +2506,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); struct Agent_global { - std::string name; + char* name; hipDeviceptr_t address; uint32_t byte_cnt; }; @@ -2524,7 +2524,7 @@ template std::pair read_global_description( ForwardIterator f, ForwardIterator l, const char* name) { const auto it = std::find_if(f, l, [=](const Agent_global& x) { - return x.name == name; + return strcmp(x.name, name) == 0; }); return it == l ? diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 178b0c9471..7a76c4642f 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -310,7 +310,7 @@ namespace hip_impl { namespace { inline void track(const Agent_global& x, hsa_agent_t agent) { - tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name.c_str(), + tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name, x.address, x.byte_cnt); int deviceIndex =0; @@ -341,7 +341,10 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t); if (t == HSA_SYMBOL_KIND_VARIABLE) { - static_cast(out)->push_back(Agent_global{name(x), address(x), size(x)}); + Agent_global tmp = {nullptr, address(x), size(x)}; + tmp.name = (char *) malloc(sizeof(name(x).c_str())); + strcpy(tmp.name, name(x).c_str()); + static_cast(out)->push_back(tmp); track(static_cast(out)->back(),agent); } From bf1d48bf786e7a7d69b895bf51862eb0aa8ad21a Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 9 Mar 2019 13:50:26 -0500 Subject: [PATCH 3/7] Fix memory leak introduced by previous change to Agent_global. Make Agent_global manage the lifetime of the name string --- include/hip/hcc_detail/hip_runtime_api.h | 34 ++++++++++++++++++++++++ src/hip_module.cpp | 6 ++--- 2 files changed, 36 insertions(+), 4 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index f6ac507c6f..f1614c8ed2 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2506,6 +2506,40 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); struct Agent_global { + + Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {} + Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) + : name(nullptr), address(address), byte_cnt(byte_cnt) { + if (name) + this->name = strdup(name); + } + + Agent_global& operator=(Agent_global&& t) { + if (this == &t) return *this; + + if (name) free(name); + name = t.name; + address = t.address; + byte_cnt = t.byte_cnt; + + t.name = nullptr; + t.address = nullptr; + t.byte_cnt = 0; + + return *this; + } + + Agent_global(Agent_global&& t) + : name(nullptr), address(nullptr), byte_cnt(0) { + *this = std::move(t); + } + + // not needed, delete them to prevent bugs + Agent_global(const Agent_global&) = delete; + Agent_global& operator=(Agent_global& t) = delete; + + ~Agent_global() { if (name) free(name); } + char* name; hipDeviceptr_t address; uint32_t byte_cnt; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 7a76c4642f..aec9c58e7f 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -341,10 +341,8 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t); if (t == HSA_SYMBOL_KIND_VARIABLE) { - Agent_global tmp = {nullptr, address(x), size(x)}; - tmp.name = (char *) malloc(sizeof(name(x).c_str())); - strcpy(tmp.name, name(x).c_str()); - static_cast(out)->push_back(tmp); + Agent_global tmp(name(x).c_str(), address(x), size(x)); + static_cast(out)->push_back(std::move(tmp)); track(static_cast(out)->back(),agent); } From 824ee1aa72c8fcd5d13a66d88668d56ca8316cf1 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 9 Mar 2019 14:21:11 -0500 Subject: [PATCH 4/7] move triple_to_hsa_isa into the header --- include/hip/hcc_detail/code_object_bundle.hpp | 60 ++++++++++++++++++- src/code_object_bundle.cpp | 54 ----------------- 2 files changed, 59 insertions(+), 55 deletions(-) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 7b97503c16..67154026ef 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -33,7 +33,65 @@ THE SOFTWARE. #include namespace hip_impl { -hsa_isa_t triple_to_hsa_isa(const std::string& triple); + +inline +__attribute__((visibility("hidden"))) +std::string transmogrify_triple(const std::string& triple) +{ + static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"}; + static constexpr const char new_prefix[]{"hcc-amdgcn-amd-amdhsa--gfx"}; + + if (triple.find(old_prefix) == 0) { + return new_prefix + triple.substr(sizeof(old_prefix) - 1); + } + + return (triple.find(new_prefix) == 0) ? triple : ""; +} + +inline +__attribute__((visibility("hidden"))) +std::string isa_name(std::string triple) +{ + static constexpr const char offload_prefix[]{"hcc-"}; + + triple = transmogrify_triple(triple); + if (triple.empty()) return {}; + + triple.erase(0, sizeof(offload_prefix) - 1); + + static hsa_isa_t tmp{}; + static const bool is_old_rocr{ + hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS}; + + if (is_old_rocr) { + std::string tmp{triple.substr(triple.rfind('x') + 1)}; + triple.replace(0, std::string::npos, "AMD:AMDGPU"); + + for (auto&& x : tmp) { + triple.push_back(':'); + triple.push_back(x); + } + } + + return triple; +} + +inline +__attribute__((visibility("hidden"))) +hsa_isa_t triple_to_hsa_isa(const std::string& triple) { + const std::string isa{isa_name(std::move(triple))}; + + if (isa.empty()) return hsa_isa_t({}); + + hsa_isa_t r{}; + + if(HSA_STATUS_SUCCESS != hsa_isa_from_name(isa.c_str(), &r)) { + r.handle = 0; + } + + return r; +} + struct Bundled_code { union Header { diff --git a/src/code_object_bundle.cpp b/src/code_object_bundle.cpp index 91258f0c75..6285cb3828 100644 --- a/src/code_object_bundle.cpp +++ b/src/code_object_bundle.cpp @@ -10,60 +10,6 @@ using namespace std; -inline -std::string transmogrify_triple(const std::string& triple) -{ - static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"}; - static constexpr const char new_prefix[]{"hcc-amdgcn-amd-amdhsa--gfx"}; - - if (triple.find(old_prefix) == 0) { - return new_prefix + triple.substr(sizeof(old_prefix) - 1); - } - - return (triple.find(new_prefix) == 0) ? triple : ""; -} - -inline -std::string isa_name(std::string triple) -{ - static constexpr const char offload_prefix[]{"hcc-"}; - - triple = transmogrify_triple(triple); - if (triple.empty()) return {}; - - triple.erase(0, sizeof(offload_prefix) - 1); - - static hsa_isa_t tmp{}; - static const bool is_old_rocr{ - hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS}; - - if (is_old_rocr) { - std::string tmp{triple.substr(triple.rfind('x') + 1)}; - triple.replace(0, std::string::npos, "AMD:AMDGPU"); - - for (auto&& x : tmp) { - triple.push_back(':'); - triple.push_back(x); - } - } - - return triple; -} - -hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) { - const std::string isa{isa_name(std::move(triple))}; - - if (isa.empty()) return hsa_isa_t({}); - - hsa_isa_t r{}; - - if(HSA_STATUS_SUCCESS != hsa_isa_from_name(isa.c_str(), &r)) { - r.handle = 0; - } - - return r; -} - // DATA - STATICS constexpr const char hip_impl::Bundled_code_header::magic_string_[]; From d37f9e6b2de9daaf28beaca5fb1a585669ef6033 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 9 Mar 2019 14:23:34 -0500 Subject: [PATCH 5/7] remove old style triple name --- include/hip/hcc_detail/code_object_bundle.hpp | 14 -------------- 1 file changed, 14 deletions(-) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 67154026ef..b3f47ba542 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -59,20 +59,6 @@ std::string isa_name(std::string triple) triple.erase(0, sizeof(offload_prefix) - 1); - static hsa_isa_t tmp{}; - static const bool is_old_rocr{ - hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS}; - - if (is_old_rocr) { - std::string tmp{triple.substr(triple.rfind('x') + 1)}; - triple.replace(0, std::string::npos, "AMD:AMDGPU"); - - for (auto&& x : tmp) { - triple.push_back(':'); - triple.push_back(x); - } - } - return triple; } From cb9ea5cefc7c9fc56e2ba34c0e875697631d7fea Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sat, 9 Mar 2019 14:24:33 -0500 Subject: [PATCH 6/7] minor cleanup --- include/hip/hcc_detail/hip_runtime_api.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index f1614c8ed2..b011911a9d 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2577,8 +2577,7 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, // hipModule_t instance static std::unordered_map< std::string, std::vector> agent_globals; - const char* hash = hash_for(hmod); - std::string key(hash); + std::string key(hash_for(hmod)); if (agent_globals.count(key) == 0) { static std::mutex mtx; From 10d3084e202deb897a933be84e95fa815eb4d0c7 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 13 Mar 2019 11:58:32 -0400 Subject: [PATCH 7/7] remove visibility hidden attribute --- include/hip/hcc_detail/code_object_bundle.hpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index b3f47ba542..89b341afa7 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -35,7 +35,6 @@ THE SOFTWARE. namespace hip_impl { inline -__attribute__((visibility("hidden"))) std::string transmogrify_triple(const std::string& triple) { static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"}; @@ -49,7 +48,6 @@ std::string transmogrify_triple(const std::string& triple) } inline -__attribute__((visibility("hidden"))) std::string isa_name(std::string triple) { static constexpr const char offload_prefix[]{"hcc-"}; @@ -63,7 +61,6 @@ std::string isa_name(std::string triple) } inline -__attribute__((visibility("hidden"))) hsa_isa_t triple_to_hsa_isa(const std::string& triple) { const std::string isa{isa_name(std::move(triple))};