diff --git a/projects/clr/hipamd/rocclr/hip_platform.cpp b/projects/clr/hipamd/rocclr/hip_platform.cpp index e41a8af26a..cb3cdf7f97 100755 --- a/projects/clr/hipamd/rocclr/hip_platform.cpp +++ b/projects/clr/hipamd/rocclr/hip_platform.cpp @@ -561,202 +561,6 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, namespace hip_impl { -struct dl_phdr_info { - ELFIO::Elf64_Addr dlpi_addr; - const char *dlpi_name; - const ELFIO::Elf64_Phdr *dlpi_phdr; - ELFIO::Elf64_Half dlpi_phnum; -}; - -extern "C" int dl_iterate_phdr( - int (*callback) (struct dl_phdr_info *info, size_t size, void *data), void *data -); - -struct Symbol { - std::string name; - ELFIO::Elf64_Addr value = 0; - ELFIO::Elf_Xword size = 0; - ELFIO::Elf_Half sect_idx = 0; - uint8_t bind = 0; - uint8_t type = 0; - uint8_t other = 0; -}; - -inline Symbol read_symbol(const ELFIO::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 ELFIO::section* find_section_if(ELFIO::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; -} - -std::vector> function_names_for(const ELFIO::elfio& reader, - ELFIO::section* symtab) { - std::vector> r; - ELFIO::symbol_section_accessor symbols{reader, symtab}; - - for (auto i = 0u; i != symbols.get_symbols_num(); ++i) { - auto tmp = read_symbol(symbols, i); - - if (tmp.type == STT_FUNC && tmp.sect_idx != SHN_UNDEF && !tmp.name.empty()) { - r.emplace_back(tmp.value, tmp.name); - } - } - - return r; -} - -const std::vector>& function_names_for_process() { - static constexpr const char self[] = "/proc/self/exe"; - - static std::vector> r; - static std::once_flag f; - - std::call_once(f, []() { - ELFIO::elfio reader; - - if (reader.load(self)) { - const auto it = find_section_if( - reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; }); - - if (it) r = function_names_for(reader, it); - } - }); - - return r; -} - - -const std::unordered_map& function_names() -{ - static std::unordered_map r{ - function_names_for_process().cbegin(), - function_names_for_process().cend()}; - static std::once_flag f; - - std::call_once(f, []() { - dl_iterate_phdr([](dl_phdr_info* info, size_t, void*) { - ELFIO::elfio reader; - - if (reader.load(info->dlpi_name)) { - const auto it = find_section_if( - reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; }); - - if (it) { - auto n = function_names_for(reader, it); - - for (auto&& f : n) f.first += info->dlpi_addr; - - r.insert(make_move_iterator(n.begin()), make_move_iterator(n.end())); - } - } - return 0; - }, - nullptr); - }); - - return r; -} - -std::vector bundles_for_process() { - static constexpr const char self[] = "/proc/self/exe"; - static constexpr const char kernel_section[] = ".kernel"; - std::vector r; - - ELFIO::elfio reader; - - if (reader.load(self)) { - auto it = find_section_if( - reader, [](const ELFIO::section* x) { return x->get_name() == kernel_section; }); - - if (it) r.insert(r.end(), it->get_data(), it->get_data() + it->get_size()); - } - - return r; -} - -const std::vector& modules() { - static std::vector r; - static std::once_flag f; - - std::call_once(f, []() { - static std::vector> bundles{bundles_for_process()}; - - dl_iterate_phdr( - [](dl_phdr_info* info, std::size_t, void*) { - ELFIO::elfio tmp; - if (tmp.load(info->dlpi_name)) { - const auto it = find_section_if( - tmp, [](const ELFIO::section* x) { return x->get_name() == ".kernel"; }); - - if (it) bundles.emplace_back(it->get_data(), it->get_data() + it->get_size()); - } - return 0; - }, - nullptr); - - for (auto&& bundle : bundles) { - if (bundle.empty()) { - continue; - } - std::string magic(&bundle[0], sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1); - if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) - continue; - - const auto obheader = reinterpret_cast(&bundle[0]); - const auto* desc = &obheader->desc[0]; - for (uint64_t i = 0; i < obheader->numBundles; ++i, - desc = reinterpret_cast( - reinterpret_cast(&desc->triple[0]) + desc->tripleSize)) { - - std::string triple(desc->triple, sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1); - if (triple.compare(HCC_AMDGCN_AMDHSA_TRIPLE)) - continue; - - std::string target(desc->triple + sizeof(HCC_AMDGCN_AMDHSA_TRIPLE), - desc->tripleSize - sizeof(HCC_AMDGCN_AMDHSA_TRIPLE)); - - if (isCompatibleCodeObject(target, hip::getCurrentDevice()->devices()[0]->info().name_)) { - hipModule_t module; - if (hipSuccess == hipModuleLoadData(&module, reinterpret_cast( - reinterpret_cast(obheader) + desc->offset))) - r.push_back(module); - break; - } - } - } - }); - - return r; -} - -const std::unordered_map& functions() -{ - static std::unordered_map r; - static std::once_flag f; - - std::call_once(f, []() { - for (auto&& function : function_names()) { - for (auto&& module : modules()) { - hipFunction_t f; - if (hipSuccess == hipModuleGetFunction(&f, module, function.second.c_str())) { - r[function.first] = f; - } - } - } - }); - - return r; -} - void hipLaunchKernelGGLImpl( uintptr_t function_address, const dim3& numBlocks, @@ -767,11 +571,19 @@ void hipLaunchKernelGGLImpl( { HIP_INIT(); - const auto it = functions().find(function_address); - if (it == functions().cend()) - assert(0); + hip::Stream* s = reinterpret_cast(stream); + int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); + if (deviceId == -1) { + DevLogPrintfError("Wrong Device Id: %d \n", deviceId); + } - hipModuleLaunchKernel(it->second, + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, reinterpret_cast(function_address), deviceId); + if ((hip_error != hipSuccess) || (func == nullptr)) { + DevLogPrintfError("Cannot find the static function: 0x%x", function_address); + } + + hipModuleLaunchKernel(func, numBlocks.x, numBlocks.y, numBlocks.z, dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, kernarg); @@ -815,16 +627,7 @@ hipError_t ihipLaunchKernel(const void* hostFunction, hipFunction_t func = nullptr; hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId); if ((hip_error != hipSuccess) || (func == nullptr)) { -#ifdef ATI_OS_LINUX - const auto it = hip_impl::functions().find(reinterpret_cast(hostFunction)); - if (it == hip_impl::functions().cend()) { - DevLogPrintfError("Cannot find function: 0x%x \n", hostFunction); - HIP_RETURN(hipErrorInvalidDeviceFunction); - } - func = it->second; -#else HIP_RETURN(hipErrorInvalidDeviceFunction); -#endif } HIP_RETURN(ihipModuleLaunchKernel(func, (gridDim.x * blockDim.x), (gridDim.y * blockDim.y), (gridDim.z * blockDim.z), blockDim.x, blockDim.y, blockDim.z,