From f423584145d55afc3602db590169d7a6dae3f1b1 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Thu, 20 Jul 2023 15:39:46 +0000 Subject: [PATCH] SWDEV-379212 - Handle template kernels with hiprtc lowered names Change-Id: Ib8e6493a1f342f92a35031d5ee39b2e22132b56a [ROCm/clr commit: dc8f66b86f26d41355883c12ee9f63cad638825d] --- .../hipamd/src/hiprtc/hiprtcComgrHelper.cpp | 34 ++++---------- .../hipamd/src/hiprtc/hiprtcComgrHelper.hpp | 4 +- .../clr/hipamd/src/hiprtc/hiprtcInternal.cpp | 44 +++++++------------ .../clr/hipamd/src/hiprtc/hiprtcInternal.hpp | 3 +- projects/clr/rocclr/device/comgrctx.cpp | 2 + projects/clr/rocclr/device/comgrctx.hpp | 10 +++++ 6 files changed, 40 insertions(+), 57 deletions(-) diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp index 485c44516a..debecbe910 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp @@ -895,7 +895,7 @@ std::string handleMangledName(std::string loweredName) { return loweredName; } -bool fillMangledNames(std::vector& dataVec, std::vector& mangledNames, +bool fillMangledNames(std::vector& dataVec, std::map& mangledNames, bool isBitcode) { amd_comgr_data_t dataObject; if (auto res = amd::Comgr::create_data( @@ -910,48 +910,32 @@ bool fillMangledNames(std::vector& dataVec, std::vector& mang } size_t Count; - if (auto res = amd::Comgr::populate_mangled_names(dataObject, &Count)) { + if (auto res = amd::Comgr::populate_name_expression_map(dataObject, &Count)) { amd::Comgr::release_data(dataObject); return false; } - for (size_t i = 0; i < Count; i++) { + for (auto &it : mangledNames) { size_t Size; - if (auto res = amd::Comgr::get_mangled_name(dataObject, i, &Size, NULL)) { + char *data = const_cast(it.first.data()); + + if (auto res = amd::Comgr::map_name_expression_to_symbol_name(dataObject, &Size, data, NULL)) { amd::Comgr::release_data(dataObject); return false; } - char* mName = new char[Size](); - if (auto res = amd::Comgr::get_mangled_name(dataObject, i, &Size, mName)) { + std::unique_ptr mName(new char[Size]()); + if (auto res = amd::Comgr::map_name_expression_to_symbol_name(dataObject, &Size, data, mName.get())) { amd::Comgr::release_data(dataObject); return false; } - mangledNames.push_back(std::string(mName)); - delete [] mName; + it.second = std::string(mName.get()); } amd::Comgr::release_data(dataObject); return true; } -bool getDemangledNames(const std::vector& mangledNames, - std::map& demangledNames) { - for (auto& i : mangledNames) { - std::string demangledName; - if (!demangleName(i, demangledName)) return false; - demangledName = handleMangledName(demangledName); - - demangledName.erase(std::remove_if(demangledName.begin(), demangledName.end(), - [](unsigned char c) { return std::isspace(c); }), - demangledName.end()); - - if (auto dres = demangledNames.find(demangledName); dres != demangledNames.end()) { - dres->second = i; - } - } - return true; -} } // namespace helpers } // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp index fc9d0612fa..d34c9264c6 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp @@ -54,10 +54,8 @@ bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa, std::vector& exeOptions, std::string name, std::string& buildLog); bool demangleName(const std::string& mangledName, std::string& demangledName); std::string handleMangledName(std::string loweredName); -bool fillMangledNames(std::vector& executable, std::vector& mangledNames, +bool fillMangledNames(std::vector& executable, std::map& mangledNames, bool isBitcode); -bool getDemangledNames(const std::vector& mangledNames, - std::map& demangledNames); void GenerateUniqueFileName(std::string& name); } // namespace helpers } // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp index f7de2315b2..9dfbb2caff 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -313,17 +313,12 @@ bool RTCCompileProgram::compile(const std::vector& options, bool fg return false; } - if (fgpu_rdc_) { - std::vector mangledNames; - if (!fillMangledNames(LLVMBitcode_, mangledNames, true)) { + if (fgpu_rdc_ && !mangled_names_.empty()) { + if (!fillMangledNames(LLVMBitcode_, mangled_names_, true)) { LogError("Error in hiprtc: unable to fill mangled names"); return false; } - if (!getDemangledNames(mangledNames, demangled_names_)) { - LogError("Error in hiprtc: unable to get demangled names"); - return false; - } return true; } @@ -373,15 +368,11 @@ bool RTCCompileProgram::compile(const std::vector& options, bool fg return false; } - std::vector mangledNames; - if (!fillMangledNames(executable_, mangledNames, false)) { - LogError("Error in hiprtc: unable to fill mangled names"); - return false; - } - - if (!getDemangledNames(mangledNames, demangled_names_)) { - LogError("Error in hiprtc: unable to get demangled names"); - return false; + if (!mangled_names_.empty()) { + if (!fillMangledNames(executable_, mangled_names_, false)) { + LogError("Error in hiprtc: unable to fill mangled names"); + return false; + } } return true; @@ -395,10 +386,7 @@ void RTCCompileProgram::stripNamedExpression(std::string& strippedName) { if (strippedName.front() == '&') { strippedName.erase(0, 1); } - // Removes the spaces from strippedName if present - strippedName.erase(std::remove_if(strippedName.begin(), strippedName.end(), - [](unsigned char c) { return std::isspace(c); }), - strippedName.end()); + } bool RTCCompileProgram::trackMangledName(std::string& name) { @@ -406,14 +394,16 @@ bool RTCCompileProgram::trackMangledName(std::string& name) { if (name.size() == 0) return false; - std::string strippedNameNoSpace = name; - stripNamedExpression(strippedNameNoSpace); + std::string strippedName = name; + stripNamedExpression(strippedName); - stripped_names_.insert(std::pair(name, strippedNameNoSpace)); - demangled_names_.insert(std::pair(strippedNameNoSpace, "")); + mangled_names_.insert(std::pair(strippedName, "")); - const auto var{"__hiprtc_" + std::to_string(stripped_names_.size())}; - const auto code{"\nextern \"C\" constexpr auto " + var + " = " + name + ";\n"}; + std::string gcn_expr = "__amdgcn_name_expr_"; + std::string size = std::to_string(mangled_names_.size()); + const auto var1{"\n static __device__ const void* " + gcn_expr + size + "[]= {\"" + strippedName + "\", (void*)&" + strippedName + "};"}; + const auto var2{"\n static auto __amdgcn_name_expr_stub_" + size + " = " + gcn_expr + size + ";\n"}; + const auto code{var1 + var2}; source_code_ += code; return true; @@ -423,7 +413,7 @@ bool RTCCompileProgram::getMangledName(const char* name_expression, const char** std::string strippedName = name_expression; stripNamedExpression(strippedName); - if (auto dres = demangled_names_.find(strippedName); dres != demangled_names_.end()) { + if (auto dres = mangled_names_.find(strippedName); dres != mangled_names_.end()) { if (dres->second.size() != 0) { *loweredName = dres->second.c_str(); return true; diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp index 1461b70519..a1965d1b19 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp @@ -140,8 +140,7 @@ class RTCCompileProgram : public RTCProgram { std::string source_code_; std::string source_name_; - std::map stripped_names_; - std::map demangled_names_; + std::map mangled_names_; std::vector compile_options_; std::vector link_options_; diff --git a/projects/clr/rocclr/device/comgrctx.cpp b/projects/clr/rocclr/device/comgrctx.cpp index a09b241c18..67184ca944 100644 --- a/projects/clr/rocclr/device/comgrctx.cpp +++ b/projects/clr/rocclr/device/comgrctx.cpp @@ -115,6 +115,8 @@ bool Comgr::LoadLib(bool is_versioned) { GET_COMGR_OPTIONAL_SYMBOL(amd_comgr_demangle_symbol_name) GET_COMGR_SYMBOL(amd_comgr_populate_mangled_names) GET_COMGR_SYMBOL(amd_comgr_get_mangled_name) + GET_COMGR_SYMBOL(amd_comgr_populate_name_expression_map) + GET_COMGR_SYMBOL(amd_comgr_map_name_expression_to_symbol_name) is_ready_ = true; return true; } diff --git a/projects/clr/rocclr/device/comgrctx.hpp b/projects/clr/rocclr/device/comgrctx.hpp index 4d7a23a38d..d8dd5e17f5 100644 --- a/projects/clr/rocclr/device/comgrctx.hpp +++ b/projects/clr/rocclr/device/comgrctx.hpp @@ -74,6 +74,8 @@ typedef amd_comgr_status_t (*t_amd_comgr_symbol_get_info)(amd_comgr_symbol_t sym typedef amd_comgr_status_t (*t_amd_comgr_demangle_symbol_name)(amd_comgr_data_t MangledSymbolName, amd_comgr_data_t* DemangledSymbolName); typedef amd_comgr_status_t (*t_amd_comgr_populate_mangled_names)(amd_comgr_data_t data, size_t *count); typedef amd_comgr_status_t (*t_amd_comgr_get_mangled_name)(amd_comgr_data_t data, size_t index, size_t *size, char *mangled_name); +typedef amd_comgr_status_t (*t_amd_comgr_populate_name_expression_map)(amd_comgr_data_t data, size_t *count); +typedef amd_comgr_status_t (*t_amd_comgr_map_name_expression_to_symbol_name)(amd_comgr_data_t data, size_t *size, char *name_expression, char* symbol_name); struct ComgrEntryPoints { void* handle; @@ -125,6 +127,8 @@ struct ComgrEntryPoints { t_amd_comgr_demangle_symbol_name amd_comgr_demangle_symbol_name; t_amd_comgr_populate_mangled_names amd_comgr_populate_mangled_names; t_amd_comgr_get_mangled_name amd_comgr_get_mangled_name; + t_amd_comgr_populate_name_expression_map amd_comgr_populate_name_expression_map; + t_amd_comgr_map_name_expression_to_symbol_name amd_comgr_map_name_expression_to_symbol_name; }; #ifdef COMGR_DYN_DLL @@ -300,6 +304,12 @@ public: static amd_comgr_status_t get_mangled_name(amd_comgr_data_t data, size_t index, size_t *size, char *mangled_name) { return COMGR_DYN(amd_comgr_get_mangled_name)(data, index, size, mangled_name); } + static amd_comgr_status_t populate_name_expression_map(amd_comgr_data_t data, size_t *count) { + return COMGR_DYN(amd_comgr_populate_name_expression_map)(data, count); + } + static amd_comgr_status_t map_name_expression_to_symbol_name(amd_comgr_data_t data, size_t *size, char *name_expression, char* symbol_name) { + return COMGR_DYN(amd_comgr_map_name_expression_to_symbol_name)(data, size, name_expression, symbol_name); + } private: