AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-379212 - Handle template kernels with hiprtc lowered names

Browse Source 
Change-Id: Ib8e6493a1f342f92a35031d5ee39b2e22132b56a


[ROCm/clr commit: dc8f66b86f]
This commit is contained in:
Satyanvesh Dittakavi
2023-07-20 15:39:46 +00:00
parent 5f13c16039
commit f423584145
6 changed files with 40 additions and 57 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp
+9 -25
View File
@@ -895,7 +895,7 @@ std::string handleMangledName(std::string loweredName) {
return loweredName;
}
bool fillMangledNames(std::vector<char>& dataVec, std::vector<std::string>& mangledNames,
bool fillMangledNames(std::vector<char>& dataVec, std::map<std::string, std::string>& mangledNames,
bool isBitcode) {
amd_comgr_data_t dataObject;
if (auto res = amd::Comgr::create_data(
@@ -910,48 +910,32 @@ bool fillMangledNames(std::vector<char>& dataVec, std::vector<std::string>& 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<char*>(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<char[]> 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<std::string>& mangledNames,
std::map<std::string, std::string>& 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
projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp
+1 -3
View File
@@ -54,10 +54,8 @@ bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa,
std::vector<std::string>& 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<char>& executable, std::vector<std::string>& mangledNames,
bool fillMangledNames(std::vector<char>& executable, std::map<std::string, std::string>& mangledNames,
bool isBitcode);
bool getDemangledNames(const std::vector<std::string>& mangledNames,
std::map<std::string, std::string>& demangledNames);
void GenerateUniqueFileName(std::string& name);
} // namespace helpers
} // namespace hiprtc
projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp
+17 -27
View File
@@ -313,17 +313,12 @@ bool RTCCompileProgram::compile(const std::vector<std::string>& options, bool fg
return false;
}
if (fgpu_rdc_) {
std::vector<std::string> 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<std::string>& options, bool fg
return false;
}
std::vector<std::string> 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<std::string, std::string>(name, strippedNameNoSpace));
demangled_names_.insert(std::pair<std::string, std::string>(strippedNameNoSpace, ""));
mangled_names_.insert(std::pair<std::string, std::string>(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;
projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp
+1 -2
View File
@@ -140,8 +140,7 @@ class RTCCompileProgram : public RTCProgram {
std::string source_code_;
std::string source_name_;
std::map<std::string, std::string> stripped_names_;
std::map<std::string, std::string> demangled_names_;
std::map<std::string, std::string> mangled_names_;
std::vector<std::string> compile_options_;
std::vector<std::string> link_options_;
projects/clr/rocclr/device/comgrctx.cpp
+2
View File
@@ -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;
}
projects/clr/rocclr/device/comgrctx.hpp
+10
View File
@@ -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: