Merge pull request #761 from yxsamliu/genco

Let hip-clang support --genco

hipamd/CMakeLists.txt

@@ -213,6 +213,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/hip_device.cpp
         src/hip_error.cpp
         src/hip_event.cpp
+        src/hip_fatbin.cpp
         src/hip_memory.cpp
         src/hip_peer.cpp
         src/hip_stream.cpp

hipamd/bin/hipcc

@@ -334,7 +334,7 @@ foreach $arg (@ARGV)
     $trimarg = $arg;
     $trimarg =~ s/^\s+|\s+$//g;  # Remive whitespace
     my $swallowArg = 0;
-    if ($arg eq '-c') {
+    if ($arg eq '-c' or $arg eq '--genco') {
         $compileOnly = 1;
         $needCXXFLAGS = 1;
         $needLDFLAGS  = 0;
@@ -386,6 +386,10 @@ foreach $arg (@ARGV)
         $swallowArg = 1;
     }
 
+    if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) {
+        $arg = "--cuda-device-only";
+    }
+
     if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
     {
         $HIPCXXFLAGS .= $HCC_WA_FLAGS;

hipamd/src/hip_clang.cpp

@@ -26,34 +26,9 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 #include "hip_hcc_internal.h"
+#include "hip_fatbin.h"
 #include "trace_helper.h"
 
-constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF"
-
-#define CLANG_OFFLOAD_BUNDLER_MAGIC "__CLANG_OFFLOAD_BUNDLE__"
-#define AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
-
-struct __ClangOffloadBundleDesc {
-  uint64_t offset;
-  uint64_t size;
-  uint64_t tripleSize;
-  const char triple[1];
-};
-
-struct __ClangOffloadBundleHeader {
-  const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC) - 1];
-  uint64_t numBundles;
-  __ClangOffloadBundleDesc desc[1];
-};
-
-struct __CudaFatBinaryWrapper {
-  unsigned int                magic;
-  unsigned int                version;
-  __ClangOffloadBundleHeader* binary;
-  void*                       unused;
-};
-
-
 extern "C" std::vector<hipModule_t>*
 __hipRegisterFatBinary(const void* data)
 {
@@ -108,21 +83,13 @@ __hipRegisterFatBinary(const void* data)
 
       std::string image{reinterpret_cast<const char*>(
           reinterpret_cast<uintptr_t>(header) + desc->offset), desc->size};
+      if (HIP_DUMP_CODE_OBJECT)
+        __hipDumpCodeObject(image);
       module->executable = hip_impl::load_executable(image, module->executable, agent);
 
       if (module->executable.handle) {
         modules->at(deviceId) = module;
         tprintf(DB_FB, "Loaded code object for %s\n", name);
-        if (HIP_DUMP_CODE_OBJECT) {
-          char fname[30];
-          static std::atomic<int> index;
-          sprintf(fname, "__hip_dump_code_object%04d.o", index++);
-          tprintf(DB_FB, "Dump code object %s\n", fname);
-          std::ofstream ofs;
-          ofs.open(fname, std::ios::binary);
-          ofs << image;
-          ofs.close();
-        }
       } else {
         fprintf(stderr, "Failed to load code object for %s\n", name);
         abort();

hipamd/src/hip_fatbin.cpp

@@ -0,0 +1,91 @@
+/*
+Copyright (c) 2018 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <string>
+#include <fstream>
+
+#include "hip_fatbin.h"
+#include "hip/hip_runtime.h"
+#include "hip_hcc_internal.h"
+#include "trace_helper.h"
+
+void __hipDumpCodeObject(const std::string& image) {
+    char fname[30];
+    static std::atomic<int> index;
+    sprintf(fname, "__hip_dump_code_object%04d.o", index++);
+    tprintf(DB_FB, "Dump code object %s\n", fname);
+    std::ofstream ofs;
+    ofs.open(fname, std::ios::binary);
+    ofs << image;
+    ofs.close();
+}
+
+// Returns a pointer to the code object in the fatbin. The pointer should not
+// be freed.
+const void* __hipExtractCodeObjectFromFatBinary(const void* data,
+    const char* agent_name)
+{
+  HIP_INIT();
+
+  tprintf(DB_FB, "Enter __hipExtractCodeObjectFromFatBinary(%p, \"%s\")\n",
+      data, agent_name);
+
+  const __ClangOffloadBundleHeader* header
+      = reinterpret_cast<const __ClangOffloadBundleHeader*>(data);
+  std::string magic(reinterpret_cast<const char*>(header),
+      sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC) - 1);
+  if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC)) {
+    return nullptr;
+  }
+
+  const __ClangOffloadBundleDesc* desc = &header->desc[0];
+  for (uint64_t i = 0; i < header->numBundles; ++i,
+       desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
+           reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
+
+    std::string triple{&desc->triple[0], sizeof(AMDGCN_AMDHSA_TRIPLE) - 1};
+    if (triple.compare(AMDGCN_AMDHSA_TRIPLE))
+      continue;
+
+    std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
+      desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)};
+    tprintf(DB_FB, "Found hip-clang bundle for %s\n", target.c_str());
+    if (target.compare(agent_name)) {
+       continue;
+    }
+
+    auto *codeobj = reinterpret_cast<const char*>(
+        reinterpret_cast<uintptr_t>(header) + desc->offset);
+    if (HIP_DUMP_CODE_OBJECT)
+      __hipDumpCodeObject(std::string{codeobj, desc->size});
+
+    tprintf(DB_FB, "__hipExtractCodeObjectFromFatBinary succeeds and returns %p\n",
+        codeobj);
+    return codeobj;
+  }
+
+  // hipcc --genco for HCC generates fat binaries with different triple strings.
+  // It will reach here and return a null pointer. The fat binary itself will
+  // be handled in a different place.
+  tprintf(DB_FB, "No hip-clang device code bundle for %s\n", agent_name);
+  return nullptr;
+}
+

hipamd/src/hip_fatbin.h

@@ -0,0 +1,58 @@
+/*
+Copyright (c) 2018 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#ifndef HIP_SRC_HIP_FATBIN_H
+#define HIP_SRC_HIP_FATBIN_H
+
+#include "hip/hip_runtime.h"
+#include "hip_hcc_internal.h"
+
+// hip-clang fatbin format
+constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF"
+
+#define CLANG_OFFLOAD_BUNDLER_MAGIC "__CLANG_OFFLOAD_BUNDLE__"
+#define AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
+
+struct __ClangOffloadBundleDesc {
+  uint64_t offset;
+  uint64_t size;
+  uint64_t tripleSize;
+  const char triple[1];
+};
+
+struct __ClangOffloadBundleHeader {
+  const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC) - 1];
+  uint64_t numBundles;
+  __ClangOffloadBundleDesc desc[1];
+};
+
+struct __CudaFatBinaryWrapper {
+  unsigned int                magic;
+  unsigned int                version;
+  __ClangOffloadBundleHeader* binary;
+  void*                       unused;
+};
+
+const void* __hipExtractCodeObjectFromFatBinary(const void* data,
+    const char* agent_name);
+void __hipDumpCodeObject(const std::string& image);
+
+#endif // HIP_SRC_HIP_FATBIN_H

hipamd/src/hip_module.cpp

@@ -48,6 +48,7 @@ THE SOFTWARE.
 #include <utility>
 #include <vector>
 #include "../include/hip/hcc_detail/code_object_bundle.hpp"
+#include "hip_fatbin.h"
 // TODO Use Pool APIs from HCC to get memory regions.
 
 using namespace ELFIO;
@@ -556,6 +557,12 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
     auto ctx = ihipGetTlsDefaultCtx();
     if (!ctx) return hipErrorInvalidContext;
 
+    // try extracting code object from image as fatbin.
+    char name[64] = {};
+    hsa_agent_get_info(this_agent(), HSA_AGENT_INFO_NAME, name);
+    if (auto *code_obj = __hipExtractCodeObjectFromFatBinary(image, name))
+      image = code_obj;
+
     hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr,
                               &(*module)->executable);