Merge pull request #958 from aaronenyeshi/cxxabi-mismatch-workaround

CXX11 ABI Mismatch Workaround

include/hip/hcc_detail/code_object_bundle.hpp

@@ -33,7 +33,48 @@ THE SOFTWARE.
 #include <vector>
 
 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 {

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<typename ForwardIterator>
 std::pair<hipDeviceptr_t, std::size_t> 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_global>> agent_globals;
-    auto key = hash_for(hmod);
+    std::string key(hash_for(hmod));
 
     if (agent_globals.count(key) == 0) {
         static std::mutex mtx;

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_[];
 

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<Container*>(out)->push_back(Agent_global{name(x), address(x), size(x)});
+        Agent_global tmp(name(x).c_str(), address(x), size(x));
+        static_cast<Container*>(out)->push_back(std::move(tmp));
 
         track(static_cast<Container*>(out)->back(),agent);
     }