From fb1021cc0a98cd0e501e971284cbbb847941cfd1 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 21 Nov 2017 13:15:13 +0000 Subject: [PATCH] This corrects how addresses are formed for symbols which reside in shared objects. For this case, the .value component of an ELF symbol holds the offset from the base VA where the shared object was loaded. Thus, to correctly obtain the VA of the object refered by the symbol, we must add the offset to the VA where the shared object is loaded. We were already doing this correctly for symbols denoting functions, but we were incorrect for those denoting objects. --- hipamd/src/program_state.cpp | 191 +++++++++++++++++++---------------- 1 file changed, 106 insertions(+), 85 deletions(-) diff --git a/hipamd/src/program_state.cpp b/hipamd/src/program_state.cpp index 2bb115981b..79d692e06f 100644 --- a/hipamd/src/program_state.cpp +++ b/hipamd/src/program_state.cpp @@ -49,6 +49,38 @@ bool operator==(hsa_isa_t x, hsa_isa_t y) namespace { + struct Symbol { + std::string name; + ELFIO::Elf64_Addr value = 0; + ELFIO::Elf_Xword size = 0; + ELFIO::Elf_Half sect_idx = 0; + std::uint8_t bind = 0; + std::uint8_t type = 0; + std::uint8_t other = 0; + }; + + inline + Symbol read_symbol(const symbol_section_accessor& section, unsigned int idx) + { + assert(idx < section.get_symbols_num()); + + Symbol r; + section.get_symbol( + idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other); + + return r; + } + + template + inline + section* find_section_if(elfio& reader, P p) + { + const auto it = find_if( + reader.sections.begin(), reader.sections.end(), move(p)); + + return it != reader.sections.end() ? *it : nullptr; + } + vector copy_names_of_undefined_symbols( const symbol_section_accessor& section) { @@ -57,47 +89,57 @@ namespace for (auto i = 0u; i != section.get_symbols_num(); ++i) { // TODO: this is boyscout code, caching the temporaries // may be of worth. - string name; - Elf64_Addr value = 0; - Elf_Xword size = 0; - Elf_Half sect_idx = 0; - uint8_t bind = 0; - uint8_t type = 0; - uint8_t other = 0; - section.get_symbol( - i, name, value, size, bind, type, sect_idx, other); - - if (sect_idx == SHN_UNDEF && !name.empty()) { - r.push_back(std::move(name)); + auto tmp = read_symbol(section, i); + if (tmp.sect_idx == SHN_UNDEF && !tmp.name.empty()) { + r.push_back(std::move(tmp.name)); } } return r; } - pair find_symbol_address( - const symbol_section_accessor& section, - const string& symbol_name) + const std::unordered_map< + std::string, + std::pair>& symbol_addresses() { - static constexpr pair r{0, 0}; + static unordered_map> r; + static once_flag f; - for (auto i = 0u; i != section.get_symbols_num(); ++i) { - // TODO: this is boyscout code, caching the temporaries - // may be of worth. - string name; - Elf64_Addr value = 0; - Elf_Xword size = 0; - Elf_Half sect_idx = 0; - uint8_t bind = 0; - uint8_t type = 0; - uint8_t other = 0; + call_once(f, []() { + dl_iterate_phdr([](dl_phdr_info* info, size_t, void*) { + static constexpr const char self[] = "/proc/self/exe"; + elfio reader; - section.get_symbol( - i, name, value, size, bind, type, sect_idx, other); + static unsigned int iter = 0u; + if (reader.load(!iter ? self : info->dlpi_name)) { + auto it = find_section_if( + reader, [](const class section* x) { + return x->get_type() == SHT_SYMTAB; + }); - if (name == symbol_name) return make_pair(value, size); - } + if (it) { + const symbol_section_accessor symtab{reader, it}; + + for (auto i = 0u; i != symtab.get_symbols_num(); ++i) { + auto tmp = read_symbol(symtab, i); + + if (tmp.type == STT_OBJECT && + tmp.sect_idx != SHN_UNDEF) { + const auto addr = + tmp.value + (iter ? info->dlpi_addr : 0); + r.emplace( + move(tmp.name), make_pair(addr, tmp.size)); + } + } + } + + ++iter; + } + + return 0; + }, nullptr); + }); return r; } @@ -116,55 +158,43 @@ namespace symbol_section_accessor{reader, code_object_dynsym}); for (auto&& x : undefined_symbols) { - const auto tmp = find_symbol_address( - symbol_section_accessor{self_reader, process_symtab}, x); + using RAII_global = + unique_ptr; - if (!tmp.first) { - throw runtime_error{ - "The global variable: " + x + ", could not be found."}; + static unordered_map globals; + static once_flag f; + call_once(f, [=]() { globals.reserve(symbol_addresses().size()); }); + + if (globals.find(x) != globals.cend()) return; + + const auto it1 = symbol_addresses().find(x); + + if (it1 == symbol_addresses().cend()) { + throw runtime_error{"Global symbol: " + x + " is undefined."}; } - static unordered_map< - Elf64_Addr, - unique_ptr> globals; + static mutex mtx; + lock_guard lck{mtx}; - if (globals.count(tmp.first) == 0) { - void* p = nullptr; - hsa_amd_memory_lock( - reinterpret_cast(tmp.first), - tmp.second, - &agent, - 1, - &p); + if (globals.find(x) != globals.cend()) return; - static mutex mtx; + void* p = nullptr; + hsa_amd_memory_lock( + reinterpret_cast(it1->second.first), + it1->second.second, + nullptr, // All agents. + 0, + &p); - lock_guard lck{mtx}; - globals.emplace( - piecewise_construct, - make_tuple(tmp.first), - make_tuple(p, hsa_amd_memory_unlock)); - } - - const auto it = globals.find(tmp.first); - - assert(it != globals.cend()); + if (!p) { cerr << it1->first << endl; assert(false); } hsa_executable_agent_global_variable_define( - executable, agent, x.c_str(), it->second.get()); + executable, agent, x.c_str(), p); + + globals.emplace(x, RAII_global{p, hsa_amd_memory_unlock}); } } - template - inline - section* find_section_if(elfio& reader, P p) - { - const auto it = find_if( - reader.sections.begin(), reader.sections.end(), std::move(p)); - - return it != reader.sections.end() ? *it : nullptr; - } - vector code_object_blob_for_process() { static constexpr const char self[] = "/proc/self/exe"; @@ -217,8 +247,8 @@ namespace Bundled_code_header tmp{blob}; if (valid(tmp)) { for (auto&& bundle : bundles(tmp)) { - r[triple_to_hsa_isa(bundle.triple)] - .push_back(bundle.blob); + r[triple_to_hsa_isa(bundle.triple)].push_back( + bundle.blob); } } } @@ -233,24 +263,15 @@ namespace vector> r; symbol_section_accessor symbols{reader, symtab}; - auto foo = reader.get_entry(); - for (auto i = 0u; i != symbols.get_symbols_num(); ++i) { // TODO: this is boyscout code, caching the temporaries // may be of worth. - string name; - Elf64_Addr value = 0; - Elf_Xword size = 0; - Elf_Half sect_idx = 0; - uint8_t bind = 0; - uint8_t type = 0; - uint8_t other = 0; + auto tmp = read_symbol(symbols, i); - symbols.get_symbol( - i, name, value, size, bind, type, sect_idx, other); - - if (type == STT_FUNC && sect_idx != SHN_UNDEF && !name.empty()) { - r.emplace_back(value, name); + if (tmp.type == STT_FUNC && + tmp.sect_idx != SHN_UNDEF && + !tmp.name.empty()) { + r.emplace_back(tmp.value, tmp.name); } } @@ -417,7 +438,7 @@ namespace hip_impl for (auto&& acc : accelerators) { auto agent = static_cast(acc.get_hsa_agent()); - if (!agent) continue; + if (!agent || !acc.is_hsa_accelerator()) continue; hsa_agent_iterate_isas(*agent, [](hsa_isa_t x, void* pa) { const auto it = code_object_blobs().find(x);