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
2018-04-27 21:21:31 -04:00
@@ -26,9 +26,10 @@ THE SOFTWARE.
 #include "platform/program.hpp"
 #include "platform/runtime.hpp"

-constexpr unsigned __cudaFatMAGIC  = 0x1ee55a01;
+#include <unordered_map>
+#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<const __CudaPartHeader*>(
-      reinterpret_cast<uintptr_t>(fbheader) + fbheader->headerSize);
-  const __CudaPartHeader* end = reinterpret_cast<const __CudaPartHeader*>(
-      reinterpret_cast<uintptr_t>(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<void*>(
-          reinterpret_cast<uintptr_t>(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<const __CudaPartHeader*>(
-        reinterpret_cast<uintptr_t>(pheader) + pheader->headerSize + pheader->partSize);
-  }
-
-  return reinterpret_cast<hipModule_t>(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<const __CudaFatBinaryWrapper*>(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<const __ClangOffloadBundleHeader*>(fbwrapper->binary);
+  const auto* desc = &obheader->desc[0];
  for (uint64_t i = 0; i < obheader->numBundles; ++i,
       desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
           reinterpret_cast<uintptr_t>(&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<hipModule_t>(as_cl(program));
 }

-
-extern "C" hipModule_t __hipRegisterFatBinary(const void* data)
-{
-  HIP_INIT();
-
-  const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);
-  if (fbwrapper->magic != __cudaFatMAGIC2 || fbwrapper->version != 1) {
-    return nullptr;
-  }
-  const __CudaFatBinaryHeader* fbheader = reinterpret_cast<const __CudaFatBinaryHeader*>(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<const __ClangOffloadBundleHeader*>(fbwrapper->binary));
-  }
-
-  return nullptr;
-}
-
 std::map<const void*, hipFunction_t> 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 <typename P>
+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<std::pair<uintptr_t, std::string>> function_names_for(const ELFIO::elfio& reader,
+                                                                  ELFIO::section* symtab) {
+  std::vector<std::pair<uintptr_t, std::string>> 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<std::pair<uintptr_t, std::string>>& function_names_for_process() {
+  static constexpr const char self[] = "/proc/self/exe";
+
+  static std::vector<std::pair<uintptr_t, std::string>> 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<uintptr_t, std::string>& function_names()
+{
+  static std::unordered_map<uintptr_t, std::string> 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<char> bundles_for_process() {
+  static constexpr const char self[] = "/proc/self/exe";
+  static constexpr const char kernel_section[] = ".kernel";
+  std::vector<char> 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<hipModule_t>& modules() {
+    static std::vector<hipModule_t> r;
+    static std::once_flag f;
+
+    std::call_once(f, []() {
+      static std::vector<std::vector<char>> 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<const __ClangOffloadBundleHeader*>(&bundle[0]);
+        const auto* desc = &obheader->desc[0];
+        for (uint64_t i = 0; i < obheader->numBundles; ++i,
+             desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
+                 reinterpret_cast<uintptr_t>(&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<const void*>(
+                reinterpret_cast<uintptr_t>(obheader) + desc->offset)))
+              r.push_back(module);
+              break;
+          }
+        }
+      }
+    });
+
+    return r;
+}
+
+const std::unordered_map<uintptr_t, hipFunction_t>& functions()
+{
+  static std::unordered_map<uintptr_t, hipFunction_t> 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);
 }

 }