From c323c0b74dacbe7b9ec8a3fbf41f3590d9e92a6e Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 27 Apr 2018 21:21:31 -0400 Subject: [PATCH] P4 to Git Change 1547563 by lmoriche@lmoriche_opencl_dev2 on 2018/04/27 21:10:47 SWDEV-145570 - [HIP] - Add support for HCC compiled binaries. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/build/Makefile.hip#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elf_types.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_dump.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_dynamic.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_header.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_note.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_relocation.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_section.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_segment.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_strings.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_symbols.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/elfio/elfio_utils.hpp#1 add ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#7 edit --- hipamd/api/hip/hip_platform.cpp | 305 ++++++++++++++++++++++++-------- 1 file changed, 227 insertions(+), 78 deletions(-) diff --git a/hipamd/api/hip/hip_platform.cpp b/hipamd/api/hip/hip_platform.cpp index a9c52bac2e..9c9c3b1ecf 100644 --- a/hipamd/api/hip/hip_platform.cpp +++ b/hipamd/api/hip/hip_platform.cpp @@ -26,9 +26,10 @@ THE SOFTWARE. #include "platform/program.hpp" #include "platform/runtime.hpp" -constexpr unsigned __cudaFatMAGIC = 0x1ee55a01; +#include +#include "elfio.hpp" + constexpr unsigned __cudaFatMAGIC2 = 0x466243b1; -constexpr unsigned __cudaFatMAGIC3 = 0xba55ed50; struct __CudaFatBinaryWrapper { unsigned int magic; @@ -37,53 +38,9 @@ struct __CudaFatBinaryWrapper { void* dummy1; }; -struct __CudaFatBinaryHeader { - unsigned int magic; - unsigned short version; - unsigned short headerSize; - unsigned long long int fatSize; -}; - -struct __CudaPartHeader { - unsigned short type; - unsigned short dummy1; - unsigned int headerSize; - unsigned long long int partSize; - unsigned long long int dummy2; - unsigned int dummy3; - unsigned int subarch; -}; - -static hipModule_t registerCudaFatBinary(const __CudaFatBinaryHeader* fbheader) -{ - const __CudaPartHeader* pheader = reinterpret_cast( - reinterpret_cast(fbheader) + fbheader->headerSize); - const __CudaPartHeader* end = reinterpret_cast( - reinterpret_cast(pheader) + fbheader->fatSize); - - amd::Program* program = new amd::Program(*g_context); - if (!program) return nullptr; - - while (pheader < end) { - if (true/*pheader->subarch == match a device in the context*/) { - const void *image = reinterpret_cast( - reinterpret_cast(pheader) + pheader->headerSize); - size_t size = pheader->partSize; - if (CL_SUCCESS != program->addDeviceProgram(*g_context->devices()[0], image, size) || - CL_SUCCESS != program->build(g_context->devices(), nullptr, nullptr, nullptr)) { - return nullptr; - } - break; - } - pheader = reinterpret_cast( - reinterpret_cast(pheader) + pheader->headerSize + pheader->partSize); - } - - return reinterpret_cast(as_cl(program)); -} - #define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__" -#define AMDGCN_AMDHSA_TRIPLE "openmp-amdgcn--amdhsa" +#define OPENMP_AMDGCN_AMDHSA_TRIPLE "openmp-amdgcn--amdhsa" +#define HCC_AMDGCN_AMDHSA_TRIPLE "hcc-amdgcn--amdhsa" struct __ClangOffloadBundleDesc { uint64_t offset; @@ -98,23 +55,35 @@ struct __ClangOffloadBundleHeader { __ClangOffloadBundleDesc desc[1]; }; -static hipModule_t registerOffloadBundle(const __ClangOffloadBundleHeader* obheader) +extern "C" hipModule_t __hipRegisterFatBinary(const void* data) { + HIP_INIT(); + + const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast(data); + if (fbwrapper->magic != __cudaFatMAGIC2 || fbwrapper->version != 1) { + return nullptr; + } + std::string magic((char*)fbwrapper->binary, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1); + if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) { + return nullptr; + } + amd::Program* program = new amd::Program(*g_context); if (!program) return nullptr; - const __ClangOffloadBundleDesc* desc = &obheader->desc[0]; + const auto obheader = reinterpret_cast(fbwrapper->binary); + 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(AMDGCN_AMDHSA_TRIPLE) - 1); - if (triple.compare(AMDGCN_AMDHSA_TRIPLE)) + std::string triple(desc->triple, sizeof(OPENMP_AMDGCN_AMDHSA_TRIPLE) - 1); + if (triple.compare(OPENMP_AMDGCN_AMDHSA_TRIPLE)) continue; - std::string target(desc->triple + sizeof(AMDGCN_AMDHSA_TRIPLE), - desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)); + std::string target(desc->triple + sizeof(OPENMP_AMDGCN_AMDHSA_TRIPLE), + desc->tripleSize - sizeof(OPENMP_AMDGCN_AMDHSA_TRIPLE)); if (target.compare(g_context->devices()[0]->info().name_)) continue; @@ -130,28 +99,6 @@ static hipModule_t registerOffloadBundle(const __ClangOffloadBundleHeader* obhea return reinterpret_cast(as_cl(program)); } - -extern "C" hipModule_t __hipRegisterFatBinary(const void* data) -{ - HIP_INIT(); - - const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast(data); - if (fbwrapper->magic != __cudaFatMAGIC2 || fbwrapper->version != 1) { - return nullptr; - } - const __CudaFatBinaryHeader* fbheader = reinterpret_cast(fbwrapper->binary); - if (fbheader->magic == __cudaFatMAGIC3 && fbheader->version == 1) { - return registerCudaFatBinary(fbheader); - } - - std::string magic((char*)fbwrapper->binary, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1); - if (!magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) { - return registerOffloadBundle(reinterpret_cast(fbwrapper->binary)); - } - - return nullptr; -} - std::map g_functions; @@ -265,15 +212,217 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) 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) { + 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 (!target.compare(g_context->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, const dim3& dimBlocks, uint32_t sharedMemBytes, hipStream_t stream, - void** kernarg) { + void** kernarg) +{ + HIP_INIT(); - assert(0 && "Unimplemented"); + const auto it = functions().find(function_address); + if (it == functions().cend()) + return; + + hipModuleLaunchKernel(it->second, + numBlocks.x, numBlocks.y, numBlocks.z, + dimBlocks.x, dimBlocks.y, dimBlocks.z, + sharedMemBytes, stream, nullptr, kernarg); } }