diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 7b97503c16..89b341afa7 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -33,7 +33,48 @@ THE SOFTWARE. #include namespace hip_impl { -hsa_isa_t triple_to_hsa_isa(const std::string& triple); + +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); + + return triple; +} + +inline +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/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 7b7ffb0638..b011911a9d 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2506,7 +2506,41 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); struct Agent_global { - std::string name; + + 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; }; @@ -2518,13 +2552,13 @@ 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( 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 ? @@ -2543,7 +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; - auto key = hash_for(hmod); + std::string key(hash_for(hmod)); if (agent_globals.count(key) == 0) { static std::mutex mtx; 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_[]; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 52d1fb38bd..aec9c58e7f 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() { @@ -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,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) { - static_cast(out)->push_back(Agent_global{name(x), address(x), size(x)}); + 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); }