SWDEV-379212 - Handle template kernels with hiprtc lowered names

Change-Id: Ib8e6493a1f342f92a35031d5ee39b2e22132b56a

hipamd/src/hiprtc/hiprtcComgrHelper.cpp

@@ -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

hipamd/src/hiprtc/hiprtcComgrHelper.hpp

@@ -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

hipamd/src/hiprtc/hiprtcInternal.cpp

@@ -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;

hipamd/src/hiprtc/hiprtcInternal.hpp

@@ -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_;