From 988dcd1e4a39c74ac34d303b55d3729899bfcca9 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Tue, 13 Nov 2018 22:28:00 -0500 Subject: [PATCH] Let hip-clang support --genco --- hipamd/CMakeLists.txt | 1 + hipamd/bin/hipcc | 6 ++- hipamd/src/hip_clang.cpp | 39 ++--------------- hipamd/src/hip_fatbin.cpp | 91 +++++++++++++++++++++++++++++++++++++++ hipamd/src/hip_fatbin.h | 58 +++++++++++++++++++++++++ hipamd/src/hip_module.cpp | 7 +++ 6 files changed, 165 insertions(+), 37 deletions(-) create mode 100644 hipamd/src/hip_fatbin.cpp create mode 100644 hipamd/src/hip_fatbin.h diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index 45d7640643..7730950ad5 100644 --- a/hipamd/CMakeLists.txt +++ b/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 diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index 68e4a96721..50e9004acf 100755 --- a/hipamd/bin/hipcc +++ b/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; diff --git a/hipamd/src/hip_clang.cpp b/hipamd/src/hip_clang.cpp index fe08bbe45c..ef8d456103 100644 --- a/hipamd/src/hip_clang.cpp +++ b/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* __hipRegisterFatBinary(const void* data) { @@ -108,21 +83,13 @@ __hipRegisterFatBinary(const void* data) std::string image{reinterpret_cast( reinterpret_cast(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 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(); diff --git a/hipamd/src/hip_fatbin.cpp b/hipamd/src/hip_fatbin.cpp new file mode 100644 index 0000000000..8fe7740ed7 --- /dev/null +++ b/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 +#include + +#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 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(data); + std::string magic(reinterpret_cast(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( + reinterpret_cast(&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( + reinterpret_cast(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; +} + diff --git a/hipamd/src/hip_fatbin.h b/hipamd/src/hip_fatbin.h new file mode 100644 index 0000000000..7b4a063c68 --- /dev/null +++ b/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 diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 2f1e4ee322..657fb06b5e 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -48,6 +48,7 @@ THE SOFTWARE. #include #include #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);