From c8ad49e2d1aedc8c56034e6a1df2ede0bcd01f2e Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 30 Apr 2019 03:48:44 +0300 Subject: [PATCH 01/15] Add ersatz for NVRTC. --- CMakeLists.txt | 9 +- include/hip/hcc_detail/hiprtc.h | 111 +++++++ include/hip/hiprtc.h | 32 ++ src/hiprtc.cpp | 514 ++++++++++++++++++++++++++++++++ 4 files changed, 664 insertions(+), 2 deletions(-) create mode 100644 include/hip/hcc_detail/hiprtc.h create mode 100644 include/hip/hiprtc.h create mode 100644 src/hiprtc.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 78d2bc7a30..dc8d8188e3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -252,7 +252,7 @@ if(HIP_PLATFORM STREQUAL "hcc") src/h2f.cpp) execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) - set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS}") set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906") if(COMPILE_HIP_ATP_MARKER) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger") @@ -262,6 +262,11 @@ if(HIP_PLATFORM STREQUAL "hcc") if(HIP_COMPILER STREQUAL "hcc") target_link_libraries(hip_hcc PRIVATE hc_am) target_link_libraries(hip_hcc_static PRIVATE hc_am) + + add_library(hiprtc SHARED src/hiprtc.cpp) + target_include_directories( + hiprtc SYSTEM + PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include) endif() string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS}) @@ -296,7 +301,7 @@ endif() ############################# # Install hip_hcc if platform is hcc if(HIP_PLATFORM STREQUAL "hcc") - install(TARGETS hip_hcc_static hip_hcc DESTINATION lib) + install(TARGETS hip_hcc_static hip_hcc hiprtc DESTINATION lib) # Install .hipInfo install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib) diff --git a/include/hip/hcc_detail/hiprtc.h b/include/hip/hcc_detail/hiprtc.h new file mode 100644 index 0000000000..89d1fe57a4 --- /dev/null +++ b/include/hip/hcc_detail/hiprtc.h @@ -0,0 +1,111 @@ +/* +Copyright (c) 2015 - 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. +*/ +#pragma once + +#include +#include + +enum hiprtcResult { + HIPRTC_SUCCESS = 0, + HIPRTC_ERROR_OUT_OF_MEMORY = 1, + HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, + HIPRTC_ERROR_INVALID_INPUT = 3, + HIPRTC_ERROR_INVALID_PROGRAM = 4, + HIPRTC_ERROR_INVALID_OPTION = 5, + HIPRTC_ERROR_COMPILATION = 6, + HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7, + HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8, + HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9, + HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10, + HIPRTC_ERROR_INTERNAL_ERROR = 11 +}; + +const char* hiprtcGetErrorString(hiprtcResult result); + +inline +hiprtcResult hiprtcVersion(int* major, int* minor) noexcept +{ // TODO: NVRTC versioning is somewhat unclear. + if (!major || !minor) return HIPRTC_ERROR_INVALID_INPUT; + + // TODO: this should be generic / set by the build infrastructure. + *major = 9; + *minor = 0; + + return HIPRTC_SUCCESS; +} + +struct _hiprtcProgram; +using hiprtcProgram = _hiprtcProgram*; + +hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, + const char* name_expression); + +hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, + const char** options); + +hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, + const char* name, int numHeaders, + const char** headers, + const char** includeNames); + +hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog); + +hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, + const char* name_expression, + const char** lowered_name); + +hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log); + +hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, + std::size_t* logSizeRet); + +hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code); + +hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, std::size_t* codeSizeRet); + +namespace hip_impl +{ + std::string demangle(const char* mangled_expression); +} + +#if defined(HIPRTC_GET_TYPE_NAME) + #include + + #if defined(_WIN32) + #include + + template + hiprtcResult hiprtcGetTypeName(std::string*) = delete; + #else + template + inline + hiprtcResult hiprtcGetTypeName(std::string* result) + { + if (!result) return HIPRTC_ERROR_INVALID_INPUT; + + *result = demangle(typeid(T).name()); + + return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR : + HIPRTC_SUCCESS; + } + #endif +#endif \ No newline at end of file diff --git a/include/hip/hiprtc.h b/include/hip/hiprtc.h new file mode 100644 index 0000000000..22d78d2656 --- /dev/null +++ b/include/hip/hiprtc.h @@ -0,0 +1,32 @@ +/* +Copyright (c) 2015 - 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. +*/ +#pragma once + +#include + +#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__) + #include +#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__) + #include +#else + #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif \ No newline at end of file diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp new file mode 100644 index 0000000000..81a7fda82a --- /dev/null +++ b/src/hiprtc.cpp @@ -0,0 +1,514 @@ +/* +Copyright (c) 2015 - 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/hip/hiprtc.h" +#include "../include/hip/hcc_detail/code_object_bundle.hpp" +#include "../include/hip/hcc_detail/elfio/elfio.hpp" +#include "../include/hip/hcc_detail/program_state.hpp" + +#include "../lpl_ca/pstreams/pstream.h" + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +const char* hiprtcGetErrorString(hiprtcResult x) +{ + switch (x) { + case HIPRTC_SUCCESS: + return "HIPRTC_SUCCESS"; + case HIPRTC_ERROR_OUT_OF_MEMORY: + return "HIPRTC_ERROR_OUT_OF_MEMORY"; + case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE: + return "HIPRTC_ERROR_PROGRAM_CREATION_FAILURE"; + case HIPRTC_ERROR_INVALID_INPUT: + return "HIPRTC_ERROR_INVALID_INPUT"; + case HIPRTC_ERROR_INVALID_PROGRAM: + return "HIPRTC_ERROR_INVALID_PROGRAM"; + case HIPRTC_ERROR_INVALID_OPTION: + return "HIPRTC_ERROR_INVALID_OPTION"; + case HIPRTC_ERROR_COMPILATION: + return "HIPRTC_ERROR_COMPILATION"; + case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE: + return "HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE"; + case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION: + return "HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION"; + case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION: + return "HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION"; + case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID: + return "HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID"; + case HIPRTC_ERROR_INTERNAL_ERROR: + return "HIPRTC_ERROR_INTERNAL_ERROR"; + default: throw std::logic_error{"Invalid HIPRTC result."}; + }; +} + +struct _hiprtcProgram { + // DATA - STATICS + static std::vector> programs; + static std::mutex mtx; + + // DATA + std::vector> headers; + std::vector> names; + std::vector lowered_names; + std::vector elf; + std::string source; + std::string name; + std::string log; + bool compiled; + + // STATICS + static + hiprtcResult destroy(_hiprtcProgram* p) + { + using namespace std; + + lock_guard lck{mtx}; + + const auto it{find_if(programs.cbegin(), programs.cend(), + [=](const unique_ptr<_hiprtcProgram>& x) { + return x.get() == p; + })}; + + if (it == programs.cend()) return HIPRTC_ERROR_INVALID_PROGRAM; + + return HIPRTC_SUCCESS; + } + + static + std::string handleMangledName(std::string name) + { + using namespace std; + + name = hip_impl::demangle(name.c_str()); + + if (name.empty()) return name; + + if (name.find("void ") == 0) name.erase(0, strlen("void ")); + + auto dx{name.find('<')}; + if (dx != string::npos) { + auto cnt{1u}; + do { + ++dx; + cnt += (name[dx] == '<') ? 1 : ((name[dx] == '>') ? -1 : 0); + } while (cnt); + + name.erase(++dx); + } + else if (name.find('(') != string::npos) name.erase(name.find('(')); + + return name; + } + + static + _hiprtcProgram* make(std::string s, std::string n, + std::vector> h) + { + using namespace std; + + unique_ptr<_hiprtcProgram> tmp{new _hiprtcProgram{move(h), {}, {}, {}, + move(s), move(n), {}, + false}}; + + lock_guard lck{mtx}; + + programs.push_back(move(tmp)); + + return programs.back().get(); + } + + static + bool isValid(_hiprtcProgram* p) noexcept + { + using namespace std; + + return find_if(programs.cbegin(), programs.cend(), + [=](const unique_ptr<_hiprtcProgram>& x) { + return x.get() == p; + }) != programs.cend(); + } + + // ACCESSORS + bool compile(const std::vector& args, + const std::experimental::filesystem::path& program_folder) + { + using namespace std; + + redi::pstream compile{args.front(), args}; + + compile.close(); + + ostringstream{log} << compile.rdbuf(); + + if (compile.rdbuf()->status() != EXIT_SUCCESS) return false; + + cerr << log << endl; + + ifstream in{args.back()}; + elf.resize(experimental::filesystem::file_size(args.back())); + in.read(elf.data(), elf.size()); + + return true; + } + + bool read_lowered_names() + { + using namespace hip_impl; + using namespace std; + + if (names.empty()) return true; + + Bundled_code_header h{elf.data()}; + istringstream blob{string{bundles(h).back().blob.cbegin(), + bundles(h).back().blob.cend()}}; + + ELFIO::elfio reader; + + if (!reader.load(blob)) return false; + + const auto it{find_section_if(reader, [](const ELFIO::section* x) { + return x->get_type() == SHT_SYMTAB; + })}; + + ELFIO::symbol_section_accessor symbols{reader, it}; + + lowered_names.resize(names.size()); + + auto n{symbols.get_symbols_num()}; + while (n--) { + const auto tmp{read_symbol(symbols, n)}; + + auto it{find_if(names.cbegin(), names.cend(), + [&](const pair& x) { + return x.second == tmp.name; + })}; + + if (it == names.cend()) { + const auto name{handleMangledName(tmp.name)}; + + if (name.empty()) continue; + + it = find_if(names.cbegin(), names.cend(), + [&](const pair& x) { + return x.second == name; + }); + + if (it == names.cend()) continue; + } + + lowered_names[distance(names.cbegin(), it)] = tmp.name; + } + + return true; + } + + std::experimental::filesystem::path write_temporary_files( + const std::experimental::filesystem::path& program_folder) + { + using namespace std; + + vector> fut{headers.size()}; + transform(headers.cbegin(), headers.cend(), begin(fut), + [&](const pair& x) { + return async([&]() { + ofstream h{program_folder / x.first}; + h.write(x.second.data(), x.second.size()); + }); + }); + + auto tmp{(program_folder / name).replace_extension(".cpp")}; + ofstream{tmp}.write(source.data(), source.size()); + + return tmp; + } + + +}; +std::vector> _hiprtcProgram::programs{}; +std::mutex _hiprtcProgram::mtx{}; + +namespace +{ + inline + bool isValidProgram(const hiprtcProgram p) + { + if (!p) return false; + + std::lock_guard lck{_hiprtcProgram::mtx}; + + return _hiprtcProgram::isValid(p); + } +} // Unnamed namespace. + +hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) +{ + if (!n) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (p->compiled) return HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION; + + const auto id{p->names.size()}; + + p->names.emplace_back(n, n); + + if (p->names.back().second.back() == ')') { + p->names.back().second.pop_back(); + p->names.back().second.erase(0, p->names.back().second.find('(')); + } + if (p->names.back().second.front() == '&') p->names.back().second.erase(0, 1); + + const auto var{"__hiprtc_" + std::to_string(id)}; + p->source.append("\nextern \"C\" constexpr auto " + var + " = " + n + ';'); + + return HIPRTC_SUCCESS; +} + +namespace +{ + class Unique_temporary_path { + // DATA + std::experimental::filesystem::path path_{}; + public: + // CREATORS + Unique_temporary_path() : path_{std::tmpnam(nullptr)} + { + while (std::experimental::filesystem::exists(path_)) { + path_ = std::tmpnam(nullptr); + } + } + Unique_temporary_path(const std::string& extension) + : Unique_temporary_path{} + { + path_.replace_extension(extension); + } + + Unique_temporary_path(const Unique_temporary_path&) = default; + Unique_temporary_path(Unique_temporary_path&&) = default; + + ~Unique_temporary_path() noexcept + { + std::experimental::filesystem::remove_all(path_); + } + + // MANIPULATORS + Unique_temporary_path& operator=( + const Unique_temporary_path&) = default; + Unique_temporary_path& operator=(Unique_temporary_path&&) = default; + + // ACCESSORS + const std::experimental::filesystem::path& path() const noexcept + { + return path_; + } + }; +} // Unnamed namespace. + +namespace hip_impl +{ + inline + std::string demangle(const char* x) + { + if (!x) return {}; + + int s{}; + std::unique_ptr tmp{ + abi::__cxa_demangle(x, nullptr, nullptr, &s), std::free}; + + if (s != 0) return {}; + + return tmp.get(); + } +} // Namespace hip_impl. + +namespace +{ + inline + void handle_target(std::vector& args) + { + using namespace std; + + bool has_target{false}; + for (auto&& x : args) { + const auto dx{x.find("--gpu-architecture")}; + const auto dy{(dx == string::npos) ? x.find("-arch") + : string::npos}; + + if (dx == dy) continue; + + x.replace(0, x.find('=', min(dx, dy)), "--targets"); + has_target = true; + + break; + } + if (!has_target) args.emplace_back("--targets=gfx900"); + } +} // Unnamed namespace. + +hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) +{ + using namespace std; + + if (n && !o) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (p->compiled) return HIPRTC_ERROR_COMPILATION; + + static const string hipcc{ + getenv("HIP_PATH") ? (getenv("HIP_PATH") + string{"/bin/hipcc"}) + : "/opt/rocm/bin/hipcc"}; + + if (!experimental::filesystem::exists(hipcc)) { + return HIPRTC_ERROR_INTERNAL_ERROR; + } + + Unique_temporary_path tmp{}; + experimental::filesystem::create_directory(tmp.path()); + + const auto src{p->write_temporary_files(tmp.path())}; + + vector args{hipcc, "--genco"}; + if (n) args.insert(args.cend(), o, o + n); + + handle_target(args); + + args.emplace_back(src); + args.emplace_back("-o"); + args.emplace_back(tmp.path() / "hiprtc.out"); + + const auto compile{p->compile(args, tmp.path())}; + + if (!p->compile(args, tmp.path())) return HIPRTC_ERROR_INTERNAL_ERROR; + if (!p->read_lowered_names()) return HIPRTC_ERROR_INTERNAL_ERROR; + + p->compiled = true; + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, + const char* name, int n_hdr, const char** hdrs, + const char** inc_names) +{ + using namespace std; + + if (!p) return HIPRTC_ERROR_INVALID_PROGRAM; + if (n_hdr < 0) return HIPRTC_ERROR_INVALID_INPUT; + if (n_hdr && (!hdrs || !inc_names)) return HIPRTC_ERROR_INVALID_INPUT; + + string s{src}; + string n{name ? name : "default_name"}; + vector> h; + + for (auto i = 0; i != n_hdr; ++i) h.emplace_back(inc_names[i], hdrs[i]); + + *p = _hiprtcProgram::make(move(s), move(n), move(h)); + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p) +{ + if (!p) return HIPRTC_SUCCESS; + + return _hiprtcProgram::destroy(*p); +} + +hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, + const char** ln) +{ + using namespace std; + + if (!n || !ln) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (!p->compiled) return HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION; + + const auto it{find_if(p->names.cbegin(), p->names.cend(), + [=](const pair& x) { + return x.first == n; + })}; + + if (it == p->names.cend()) return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID; + + *ln = p->lowered_names[distance(p->names.cbegin(), it)].c_str(); + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) +{ + if (!l) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM; + + l = std::copy_n(p->log.data(), p->log.size(), l); + *l = '\0'; + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) +{ + if (!sz) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM; + + *sz = p->log.size() + 1; + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) +{ + if (!c) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM; + + std::copy_n(p->elf.data(), p->elf.size(), c); + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) +{ + if (!sz) return HIPRTC_ERROR_INVALID_INPUT; + if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; + if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM; + + *sz = p->elf.size(); + + return HIPRTC_SUCCESS; +} \ No newline at end of file From 14ff2f71ed82f773bd2256ebb7015f51b2d1cd7b Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 30 Apr 2019 13:45:44 +0300 Subject: [PATCH 02/15] Fix extraneous paren and use correct namespace. --- include/hip/hcc_detail/hiprtc.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hiprtc.h b/include/hip/hcc_detail/hiprtc.h index 89d1fe57a4..d4174e4ce1 100644 --- a/include/hip/hcc_detail/hiprtc.h +++ b/include/hip/hcc_detail/hiprtc.h @@ -102,7 +102,7 @@ namespace hip_impl { if (!result) return HIPRTC_ERROR_INVALID_INPUT; - *result = demangle(typeid(T).name()); + *result = hip_impl::demangle(typeid(T).name()); return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR : HIPRTC_SUCCESS; From a13dcd380141aec8ac12c292e57f0f65f3b45d84 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 30 Apr 2019 13:46:32 +0300 Subject: [PATCH 03/15] Use lowerCamelCase (yuck, yuck) consistently. --- src/hiprtc.cpp | 70 +++++++++++++++++++++++++++----------------------- 1 file changed, 38 insertions(+), 32 deletions(-) diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp index 81a7fda82a..a5f733d458 100644 --- a/src/hiprtc.cpp +++ b/src/hiprtc.cpp @@ -85,7 +85,7 @@ struct _hiprtcProgram { // DATA std::vector> headers; std::vector> names; - std::vector lowered_names; + std::vector loweredNames; std::vector elf; std::string source; std::string name; @@ -131,7 +131,11 @@ struct _hiprtcProgram { name.erase(++dx); } - else if (name.find('(') != string::npos) name.erase(name.find('(')); + else { + const auto dy{name.find('(')}; + + if (dy != string::npos) name.erase(dy); + } return name; } @@ -156,10 +160,8 @@ struct _hiprtcProgram { static bool isValid(_hiprtcProgram* p) noexcept { - using namespace std; - - return find_if(programs.cbegin(), programs.cend(), - [=](const unique_ptr<_hiprtcProgram>& x) { + return std::find_if(programs.cbegin(), programs.cend(), + [=](const std::unique_ptr<_hiprtcProgram>& x) { return x.get() == p; }) != programs.cend(); } @@ -187,7 +189,7 @@ struct _hiprtcProgram { return true; } - bool read_lowered_names() + bool readLoweredNames() { using namespace hip_impl; using namespace std; @@ -206,11 +208,14 @@ struct _hiprtcProgram { return x->get_type() == SHT_SYMTAB; })}; - ELFIO::symbol_section_accessor symbols{reader, it}; + loweredNames.resize(names.size()); - lowered_names.resize(names.size()); + ELFIO::symbol_section_accessor symbols{reader, it}; auto n{symbols.get_symbols_num()}; + + if (n < loweredNames.size()) return false; + while (n--) { const auto tmp{read_symbol(symbols, n)}; @@ -232,14 +237,14 @@ struct _hiprtcProgram { if (it == names.cend()) continue; } - lowered_names[distance(names.cbegin(), it)] = tmp.name; + loweredNames[distance(names.cbegin(), it)] = tmp.name; } return true; } - std::experimental::filesystem::path write_temporary_files( - const std::experimental::filesystem::path& program_folder) + std::experimental::filesystem::path writeTemporaryFiles( + const std::experimental::filesystem::path& programFolder) { using namespace std; @@ -247,12 +252,12 @@ struct _hiprtcProgram { transform(headers.cbegin(), headers.cend(), begin(fut), [&](const pair& x) { return async([&]() { - ofstream h{program_folder / x.first}; + ofstream h{programFolder / x.first}; h.write(x.second.data(), x.second.size()); }); }); - auto tmp{(program_folder / name).replace_extension(".cpp")}; + auto tmp{(programFolder / name).replace_extension(".cpp")}; ofstream{tmp}.write(source.data(), source.size()); return tmp; @@ -290,7 +295,9 @@ hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) p->names.back().second.pop_back(); p->names.back().second.erase(0, p->names.back().second.find('(')); } - if (p->names.back().second.front() == '&') p->names.back().second.erase(0, 1); + if (p->names.back().second.front() == '&') { + p->names.back().second.erase(0, 1); + } const auto var{"__hiprtc_" + std::to_string(id)}; p->source.append("\nextern \"C\" constexpr auto " + var + " = " + n + ';'); @@ -357,12 +364,14 @@ namespace hip_impl namespace { + constexpr const char defaultTarget[]{"--targets=gfx900"}; + inline - void handle_target(std::vector& args) + void handleTarget(std::vector& args) { using namespace std; - bool has_target{false}; + bool hasTarget{false}; for (auto&& x : args) { const auto dx{x.find("--gpu-architecture")}; const auto dy{(dx == string::npos) ? x.find("-arch") @@ -371,11 +380,11 @@ namespace if (dx == dy) continue; x.replace(0, x.find('=', min(dx, dy)), "--targets"); - has_target = true; + hasTarget = true; break; } - if (!has_target) args.emplace_back("--targets=gfx900"); + if (!hasTarget) args.emplace_back(defaultTarget); } } // Unnamed namespace. @@ -398,12 +407,12 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) Unique_temporary_path tmp{}; experimental::filesystem::create_directory(tmp.path()); - const auto src{p->write_temporary_files(tmp.path())}; + const auto src{p->writeTemporaryFiles(tmp.path())}; vector args{hipcc, "--genco"}; if (n) args.insert(args.cend(), o, o + n); - handle_target(args); + handleTarget(args); args.emplace_back(src); args.emplace_back("-o"); @@ -412,7 +421,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) const auto compile{p->compile(args, tmp.path())}; if (!p->compile(args, tmp.path())) return HIPRTC_ERROR_INTERNAL_ERROR; - if (!p->read_lowered_names()) return HIPRTC_ERROR_INTERNAL_ERROR; + if (!p->readLoweredNames()) return HIPRTC_ERROR_INTERNAL_ERROR; p->compiled = true; @@ -420,22 +429,19 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) } hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, - const char* name, int n_hdr, const char** hdrs, - const char** inc_names) + const char* name, int n, const char** hdrs, + const char** incs) { using namespace std; if (!p) return HIPRTC_ERROR_INVALID_PROGRAM; - if (n_hdr < 0) return HIPRTC_ERROR_INVALID_INPUT; - if (n_hdr && (!hdrs || !inc_names)) return HIPRTC_ERROR_INVALID_INPUT; + if (n < 0) return HIPRTC_ERROR_INVALID_INPUT; + if (n && (!hdrs || !incs)) return HIPRTC_ERROR_INVALID_INPUT; - string s{src}; - string n{name ? name : "default_name"}; vector> h; + for (auto i = 0; i != n; ++i) h.emplace_back(incs[i], hdrs[i]); - for (auto i = 0; i != n_hdr; ++i) h.emplace_back(inc_names[i], hdrs[i]); - - *p = _hiprtcProgram::make(move(s), move(n), move(h)); + *p = _hiprtcProgram::make(src, name ? name : "default_name", move(h)); return HIPRTC_SUCCESS; } @@ -463,7 +469,7 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, if (it == p->names.cend()) return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID; - *ln = p->lowered_names[distance(p->names.cbegin(), it)].c_str(); + *ln = p->loweredNames[distance(p->names.cbegin(), it)].c_str(); return HIPRTC_SUCCESS; } From a26f81ea43006bacb524df1154a124f54f8cb895 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 30 Apr 2019 17:04:15 +0300 Subject: [PATCH 04/15] Link against FS when building hiprtc lib. --- CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index dc8d8188e3..21774b7423 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -267,6 +267,7 @@ if(HIP_PLATFORM STREQUAL "hcc") target_include_directories( hiprtc SYSTEM PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include) + target_link_libraries(hiprtc PUBLIC stdc++fs) endif() string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS}) From 6d92e958b7e87083597a2314027191e5f0ff290c Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 30 Apr 2019 18:18:50 +0300 Subject: [PATCH 05/15] Correctly mark Manipulators. Fix dual compile. --- src/hiprtc.cpp | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp index a5f733d458..6d9048f500 100644 --- a/src/hiprtc.cpp +++ b/src/hiprtc.cpp @@ -166,7 +166,7 @@ struct _hiprtcProgram { }) != programs.cend(); } - // ACCESSORS + // MANIPULATORS bool compile(const std::vector& args, const std::experimental::filesystem::path& program_folder) { @@ -208,8 +208,6 @@ struct _hiprtcProgram { return x->get_type() == SHT_SYMTAB; })}; - loweredNames.resize(names.size()); - ELFIO::symbol_section_accessor symbols{reader, it}; auto n{symbols.get_symbols_num()}; @@ -243,8 +241,9 @@ struct _hiprtcProgram { return true; } + // ACCESSORS std::experimental::filesystem::path writeTemporaryFiles( - const std::experimental::filesystem::path& programFolder) + const std::experimental::filesystem::path& programFolder) const { using namespace std; @@ -290,6 +289,7 @@ hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) const auto id{p->names.size()}; p->names.emplace_back(n, n); + p->loweredNames.emplace_back(); if (p->names.back().second.back() == ')') { p->names.back().second.pop_back(); @@ -418,8 +418,6 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) args.emplace_back("-o"); args.emplace_back(tmp.path() / "hiprtc.out"); - const auto compile{p->compile(args, tmp.path())}; - if (!p->compile(args, tmp.path())) return HIPRTC_ERROR_INTERNAL_ERROR; if (!p->readLoweredNames()) return HIPRTC_ERROR_INTERNAL_ERROR; From 4413c5c85bb714dc6cf383f7a277d300d773e3f6 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 30 Apr 2019 19:17:13 +0300 Subject: [PATCH 06/15] Add unit tests. Extend HIT to accept linker options. --- tests/hit/HIT.cmake | 19 ++- tests/src/hiprtc/hiprtcGetLoweredName.cpp | 152 ++++++++++++++++++++++ tests/src/hiprtc/hiprtcGetTypeName.cpp | 138 ++++++++++++++++++++ tests/src/hiprtc/saxpy.cpp | 151 +++++++++++++++++++++ 4 files changed, 458 insertions(+), 2 deletions(-) create mode 100644 tests/src/hiprtc/hiprtcGetLoweredName.cpp create mode 100644 tests/src/hiprtc/hiprtcGetTypeName.cpp create mode 100644 tests/src/hiprtc/saxpy.cpp diff --git a/tests/hit/HIT.cmake b/tests/hit/HIT.cmake index 82e8508dcd..15c259238e 100644 --- a/tests/hit/HIT.cmake +++ b/tests/hit/HIT.cmake @@ -3,17 +3,19 @@ find_package(HIP REQUIRED) #------------------------------------------------------------------------------- # Helper macro to parse BUILD instructions -macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _exclude_platforms _dir) +macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _dir) set(${_target}) set(${_sources}) set(${_hipcc_options}) set(${_hcc_options}) set(${_nvcc_options}) + set(${_link_options}) set(${_exclude_platforms}) set(_target_found FALSE) set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found FALSE) + set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) foreach(arg ${ARGN}) if(NOT _target_found) @@ -23,21 +25,31 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt set(_hipcc_options_found TRUE) set(_hcc_options_found FALSE) set(_nvcc_options_found FALSE) + set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) elseif("x${arg}" STREQUAL "xHCC_OPTIONS") set(_hipcc_options_found FALSE) set(_hcc_options_found TRUE) set(_nvcc_options_found FALSE) + set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) elseif("x${arg}" STREQUAL "xNVCC_OPTIONS") set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found TRUE) + set(_link_options_found FALSE) + set(_exclude_platforms_found FALSE) + elseif("x${arg}" STREQUAL "xLINK_OPTIONS") + set(_hipcc_options_found FALSE) + set(_hcc_options_found FALSE) + set(_nvcc_options_found FALSE) + set(_link_options_found TRUE) set(_exclude_platforms_found FALSE) elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM") set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found FALSE) + set(_link_options_found FALSE) set(_exclude_platforms_found TRUE) else() if(_hipcc_options_found) @@ -46,6 +58,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt list(APPEND ${_hcc_options} ${arg}) elseif(_nvcc_options_found) list(APPEND ${_nvcc_options} ${arg}) + elseif(_link_options_found) + list(APPEND ${_link_options} ${arg}) elseif(_exclude_platforms_found) set(${_exclude_platforms} ${arg}) else() @@ -142,7 +156,7 @@ macro(HIT_ADD_FILES _dir _label _parent) string(REGEX REPLACE "\n" ";" _contents "${_contents}") foreach(_cmd ${_contents}) string(REGEX REPLACE " " ";" _cmd "${_cmd}") - parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _exclude_platforms ${_dir} ${_cmd}) + parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms ${_dir} ${_cmd}) string(REGEX REPLACE "/" "." target ${_label}/${_target}) insert_into_map("_exclude" "${target}" "${_exclude_platforms}") if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM}) @@ -150,6 +164,7 @@ macro(HIT_ADD_FILES _dir _label _parent) set_source_files_properties(${_sources} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) hip_reset_flags() hip_add_executable(${target} ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options} EXCLUDE_FROM_ALL) + target_link_libraries(${target} PRIVATE ${_link_options}) set_target_properties(${target} PROPERTIES OUTPUT_NAME ${_target} RUNTIME_OUTPUT_DIRECTORY ${_label} LINK_DEPENDS "${HIP_LIB_FILES}") add_dependencies(${_parent} ${target}) endif() diff --git a/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/tests/src/hiprtc/hiprtcGetLoweredName.cpp new file mode 100644 index 0000000000..2e4f4e9dde --- /dev/null +++ b/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -0,0 +1,152 @@ +/* +Copyright (c) 2015 - 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. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ +#include + +#include +#include + +#include +#include +#include + + +static constexpr const char gpu_program[]{ +R"( +#include + +__device__ int V1; // set from host code +static __global__ void f1(int *result) { *result = V1 + 10; } +namespace N1 { +namespace N2 { +__constant__ int V2; // set from host code +__global__ void f2(int *result) { *result = V2 + 20; } +} +} +template +__global__ void f3(int *result) { *result = sizeof(T); } +)"}; + +int main() +{ + using namespace std; + + hiprtcProgram prog; + hiprtcCreateProgram(&prog, gpu_program, "prog.cu", 0, nullptr, nullptr); + + vector kernel_name_vec; + vector variable_name_vec; + vector variable_initial_value; + vector expected_result; + + kernel_name_vec.push_back("&f1"); + expected_result.push_back(10 + 100); + kernel_name_vec.push_back("N1::N2::f2"); + expected_result.push_back(20 + 200); + kernel_name_vec.push_back("f3"); + expected_result.push_back(sizeof(int)); + kernel_name_vec.push_back("f3"); + expected_result.push_back(sizeof(double)); + + for (auto&& x : kernel_name_vec) hiprtcAddNameExpression(prog, x.c_str()); + + variable_name_vec.push_back("&V1"); + variable_initial_value.push_back(100); + variable_name_vec.push_back("&N1::N2::V2"); + variable_initial_value.push_back(200); + + for (auto&& x : variable_name_vec) hiprtcAddNameExpression(prog, x.c_str()); + + hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr); + + // Obtain compilation log from the program. + size_t logSize; + hiprtcGetProgramLogSize(prog, &logSize); + + if (logSize) { + string log(logSize, '\0'); + hiprtcGetProgramLog(prog, &log[0]); + + cout << log << '\n'; + } + + if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); } + + size_t codeSize; + hiprtcGetCodeSize(prog, &codeSize); + + vector code(codeSize); + hiprtcGetCode(prog, code.data()); + + hipModule_t module; + hipModuleLoadData(&module, code.data()); + + hipDeviceptr_t dResult; + int hResult = 0; + hipMalloc(&dResult, sizeof(hResult)); + hipMemcpyHtoD(dResult, &hResult, sizeof(hResult)); + + for (decltype(variable_name_vec.size()) i = 0; i != variable_name_vec.size(); ++i) { + const char* name; + hiprtcGetLoweredName(prog, variable_name_vec[i].c_str(), &name); + + int initial_value = variable_initial_value[i]; + + hipDeviceptr_t variable_addr; + size_t bytes{}; + hipModuleGetGlobal(&variable_addr, &bytes, module, name); + hipMemcpyHtoD(variable_addr, &initial_value, sizeof(initial_value)); + } + + for (decltype(kernel_name_vec.size()) i = 0; i != kernel_name_vec.size(); ++i) { + const char* name; + hiprtcGetLoweredName(prog, kernel_name_vec[i].c_str(), &name); + + hipFunction_t kernel; + hipModuleGetFunction(&kernel, module, name); + + struct { hipDeviceptr_t a_; } args{dResult}; + + auto size = sizeof(args); + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + + hipModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, + config); + + hipMemcpyDtoH(&hResult, dResult, sizeof(hResult)); + + if (expected_result[i] != hResult) { failed("Validation failed."); } + } + + hipFree(dResult); + hipModuleUnload(module); + + hiprtcDestroyProgram(&prog); + + passed(); +} \ No newline at end of file diff --git a/tests/src/hiprtc/hiprtcGetTypeName.cpp b/tests/src/hiprtc/hiprtcGetTypeName.cpp new file mode 100644 index 0000000000..da94c926ea --- /dev/null +++ b/tests/src/hiprtc/hiprtcGetTypeName.cpp @@ -0,0 +1,138 @@ +/* +Copyright (c) 2015 - 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. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ + +#include + +#define HIPRTC_GET_TYPE_NAME +#include +#include + +#include +#include +#include + +static constexpr auto gpu_program{ +R"( +#include + +namespace N1 { struct S1_t { int i; double d; }; } +template +__global__ void f3(int *result) { *result = sizeof(T); } +)"}; + +// note: this structure is also defined in GPU code string. Should ideally +// be in a header file included by both GPU code string and by CPU code. +namespace N1 { struct S1_t { int i; double d; }; }; + +template +std::string getKernelNameForType(void) +{ + std::string type_name; + hiprtcGetTypeName(&type_name); + return std::string{"f3<"} + type_name + '>'; +} + +int main() +{ + using namespace std; + + hiprtcProgram prog; + hiprtcCreateProgram(&prog, gpu_program, "gpu_program.cu", 0, nullptr, + nullptr); + + vector name_vec; + vector expected_result; + + name_vec.push_back(getKernelNameForType()); + expected_result.push_back(sizeof(int)); + name_vec.push_back(getKernelNameForType()); + expected_result.push_back(sizeof(double)); + name_vec.push_back(getKernelNameForType()); + expected_result.push_back(sizeof(N1::S1_t)); + + for (auto&& x : name_vec) hiprtcAddNameExpression(prog, x.c_str()); + + hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr); + + size_t logSize; + hiprtcGetProgramLogSize(prog, &logSize); + + if (logSize) { + string log(logSize, '\0'); + hiprtcGetProgramLog(prog, &log[0]); + + cout << log << '\n'; + } + + if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); } + + size_t codeSize; + hiprtcGetCodeSize(prog, &codeSize); + + vector code(codeSize); + hiprtcGetCode(prog, code.data()); + + hipModule_t module; + hipModuleLoadDataEx(&module, code.data(), 0, nullptr, nullptr); + + hipDeviceptr_t dResult; + int hResult = 0; + hipMalloc(&dResult, sizeof(hResult)); + hipMemcpyHtoD(dResult, &hResult, sizeof(hResult)); + + for (size_t i = 0; i < name_vec.size(); ++i) { + const char *name; + hiprtcGetLoweredName(prog, name_vec[i].c_str(), &name); + + hipFunction_t kernel; + hipModuleGetFunction(&kernel, module, name); + + struct { hipDeviceptr_t a_; } args{dResult}; + + auto size = sizeof(args); + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + + hipModuleLaunchKernel(kernel, + 1, 1, 1, + 1, 1, 1, + 0, nullptr, + nullptr, config); + + hipMemcpyDtoH(&hResult, dResult, sizeof(hResult)); + + if (expected_result[i] != hResult) { failed("Validation failed."); } + } + + hipFree(dResult); + hipModuleUnload(module); + + hiprtcDestroyProgram(&prog); + + passed(); +} \ No newline at end of file diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp new file mode 100644 index 0000000000..de95839ef4 --- /dev/null +++ b/tests/src/hiprtc/saxpy.cpp @@ -0,0 +1,151 @@ +/* +Copyright (c) 2015 - 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. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ + +#include + +#include +#include + +#include +#include +#include +#include +#include + +static constexpr auto NUM_THREADS{128}; +static constexpr auto NUM_BLOCKS{32}; + +static constexpr auto saxpy{ +R"( +#include + +extern "C" +__global__ +void saxpy(float a, float* x, float* y, float* out, size_t n) +{ + size_t tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid < n) { + out[tid] = a * x[tid] + y[tid]; + } +} +)"}; + +int main() +{ + using namespace std; + + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + saxpy, // buffer + "saxpy.cu", // name + 0, // numHeaders + nullptr, // headers + nullptr); // includeNames + + static const char* opts[]{"--gpu-architecture=gfx900"}; + + hiprtcResult compileResult{hiprtcCompileProgram(prog, + sizeof(opts) / sizeof(opts[0]), + opts)}; + + size_t logSize; + hiprtcGetProgramLogSize(prog, &logSize); + + if (logSize) { + string log(logSize, '\0'); + hiprtcGetProgramLog(prog, &log[0]); + + cout << log << '\n'; + } + + if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); } + + size_t codeSize; + hiprtcGetCodeSize(prog, &codeSize); + + vector code(codeSize); + hiprtcGetCode(prog, code.data()); + + hiprtcDestroyProgram(&prog); + + hipModule_t module; + hipFunction_t kernel; + + hipModuleLoadData(&module, code.data()); + hipModuleGetFunction(&kernel, module, "saxpy"); + + size_t n = NUM_THREADS * NUM_BLOCKS; + size_t bufferSize = n * sizeof(float); + + float a = 5.1f; + unique_ptr hX{new float[n]}; + unique_ptr hY{new float[n]}; + unique_ptr hOut{new float[n]}; + + for (size_t i = 0; i < n; ++i) { + hX[i] = static_cast(i); + hY[i] = static_cast(i * 2); + } + + hipDeviceptr_t dX, dY, dOut; + hipMalloc(&dX, bufferSize); + hipMalloc(&dY, bufferSize); + hipMalloc(&dOut, bufferSize); + hipMemcpyHtoD(dX, hX.get(), bufferSize); + hipMemcpyHtoD(dY, hY.get(), bufferSize); + + struct { + float a_; + hipDeviceptr_t b_; + hipDeviceptr_t c_; + hipDeviceptr_t d_; + size_t e_; + } args{a, dX, dY, dOut, n}; + + auto size = sizeof(args); + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + + hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1, + 0, nullptr, nullptr, config); + + hipMemcpyDtoH(hOut.get(), dOut, bufferSize); + + for (size_t i = 0; i < n; ++i) { + if (a * hX[i] + hY[i] != hOut[i]) { failed("Validation failed."); } + } + + hipFree(dX); + hipFree(dY); + hipFree(dOut); + + hipModuleUnload(module); + hiprtcDestroyProgram(&prog); + + passed(); +} \ No newline at end of file From 6f321612499675d2c0738b3db3d1ca1d84a07833 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 7 May 2019 06:18:27 +0300 Subject: [PATCH 07/15] Make sure the HIPRTC library is installed. --- packaging/hip_hcc.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt index fe866e47f9..5aebc6c36d 100644 --- a/packaging/hip_hcc.txt +++ b/packaging/hip_hcc.txt @@ -3,6 +3,7 @@ project(hip_hcc) install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib) +install(FILES @PROJECT_BINARY_DIR@/libhiprtc.so DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip) install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip) From 5f0bb143bd396aa86892fb9498b0e06f6733a3a8 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 7 May 2019 08:24:26 +0300 Subject: [PATCH 08/15] Better logging. Try to auto-detect the target. --- src/hiprtc.cpp | 62 ++++++++++++++++++++++++++++++++++++++++++-------- 1 file changed, 53 insertions(+), 9 deletions(-) diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp index 6d9048f500..60f64dee0b 100644 --- a/src/hiprtc.cpp +++ b/src/hiprtc.cpp @@ -27,6 +27,8 @@ THE SOFTWARE. #include "../lpl_ca/pstreams/pstream.h" +#include + #include #include @@ -170,17 +172,21 @@ struct _hiprtcProgram { bool compile(const std::vector& args, const std::experimental::filesystem::path& program_folder) { + using namespace redi; using namespace std; - redi::pstream compile{args.front(), args}; - - compile.close(); + ipstream compile{args.front(), args, pstreambuf::pstderr}; - ostringstream{log} << compile.rdbuf(); + constexpr const auto tmp_size{1024u}; + char tmp[tmp_size]{}; + while (!compile.eof()) { + log.append(tmp, tmp + compile.readsome(tmp, tmp_size)); + } - if (compile.rdbuf()->status() != EXIT_SUCCESS) return false; + compile.close(); - cerr << log << endl; + if (compile.rdbuf()->exited() && + compile.rdbuf()->status() != EXIT_SUCCESS) return false; ifstream in{args.back()}; elf.resize(experimental::filesystem::file_size(args.back())); @@ -364,7 +370,45 @@ namespace hip_impl namespace { - constexpr const char defaultTarget[]{"--targets=gfx900"}; + const std::string& defaultTarget() + { + using namespace std; + + static string r{"gfx900"}; + static once_flag f{}; + + call_once(f, []() { + static hsa_agent_t a{}; + hsa_iterate_agents([](hsa_agent_t x, void*) { + hsa_device_type_t t{}; + hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t); + + if (t != HSA_DEVICE_TYPE_GPU) return HSA_STATUS_SUCCESS; + + a = x; + + return HSA_STATUS_INFO_BREAK; + }, nullptr); + + if (!a.handle) return; + + hsa_agent_iterate_isas(a, [](hsa_isa_t x, void*){ + uint32_t n{}; + hsa_isa_get_info_alt(x, HSA_ISA_INFO_NAME_LENGTH, &n); + + if (n == 0) return HSA_STATUS_SUCCESS; + + r.resize(n); + hsa_isa_get_info_alt(x, HSA_ISA_INFO_NAME, &r[0]); + + r.erase(0, r.find("gfx")); + + return HSA_STATUS_INFO_BREAK; + }, nullptr); + }); + + return r; + } inline void handleTarget(std::vector& args) @@ -384,7 +428,7 @@ namespace break; } - if (!hasTarget) args.emplace_back(defaultTarget); + if (!hasTarget) args.push_back("--targets=" + defaultTarget()); } } // Unnamed namespace. @@ -490,7 +534,7 @@ hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM; - *sz = p->log.size() + 1; + *sz = p->log.empty() ? 0 : p->log.size() + 1; return HIPRTC_SUCCESS; } From 7f70758d9c8335d95964f593ae81516ba5a924ba Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 7 May 2019 08:24:54 +0300 Subject: [PATCH 09/15] Stop specifying the target explicitly. --- tests/src/hiprtc/saxpy.cpp | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp index de95839ef4..b289c84c9c 100644 --- a/tests/src/hiprtc/saxpy.cpp +++ b/tests/src/hiprtc/saxpy.cpp @@ -64,13 +64,9 @@ int main() "saxpy.cu", // name 0, // numHeaders nullptr, // headers - nullptr); // includeNames + nullptr); // includeNames - static const char* opts[]{"--gpu-architecture=gfx900"}; - - hiprtcResult compileResult{hiprtcCompileProgram(prog, - sizeof(opts) / sizeof(opts[0]), - opts)}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 0, nullptr)}; size_t logSize; hiprtcGetProgramLogSize(prog, &logSize); From 88e4605405137de7c32cdfa4a9d4c6f5e710fb49 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 9 May 2019 11:30:54 +0300 Subject: [PATCH 10/15] Add missing flavour of `hipModuleLaunchKernel`. --- include/hip/hcc_detail/program_state.hpp | 11 ++++-- src/hip_hcc_internal.h | 2 ++ src/hip_module.cpp | 43 ++++++++++++++---------- 3 files changed, 37 insertions(+), 19 deletions(-) diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index f05f41d3f5..c22580b847 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -90,10 +90,17 @@ class Kernel_descriptor { std::uint64_t kernel_object_{}; amd_kernel_code_t const* kernel_header_{nullptr}; std::string name_{}; + std::vector> kernarg_layout_{}; public: Kernel_descriptor() = default; - Kernel_descriptor(std::uint64_t kernel_object, const std::string& name) - : kernel_object_{kernel_object}, name_{name} + Kernel_descriptor( + std::uint64_t kernel_object, + const std::string& name, + std::vector> kernarg_layout = {}) + : + kernel_object_{kernel_object}, + name_{name}, + kernarg_layout_{std::move(kernarg_layout)} { bool supported{false}; std::uint16_t min_v{UINT16_MAX}; diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index b40fac93a5..769d6b7914 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -371,6 +371,8 @@ struct ihipModule_t { hsa_executable_t executable = {}; hsa_code_object_reader_t coReader = {}; std::string hash; + std::unordered_map< + std::string, std::vector>> kernargs; ~ihipModule_t() { if (executable.handle) hsa_executable_destroy(executable); diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e1a3b1ff65..fa6ca36e1d 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -90,6 +90,7 @@ struct ihipModuleSymbol_t { uint64_t _object{}; // The kernel object. amd_kernel_code_t const* _header{}; string _name; // TODO - review for performance cost. Name is just used for debug. + vector> _kernarg_layout{}; }; template <> @@ -131,6 +132,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void** kernelParams, void** extra, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) { + using namespace hip_impl; + auto ctx = ihipGetTlsDefaultCtx(); hipError_t ret = hipSuccess; @@ -145,19 +148,26 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, void* config[5] = {0}; size_t kernArgSize; - if (kernelParams != NULL) { - std::string name = f->_name; - struct ihipKernArgInfo pl = kernelArguments[name]; - char* argBuf = (char*)malloc(pl.totalSize); - memset(argBuf, 0, pl.totalSize); - int index = 0; - for (int i = 0; i < pl.Size.size(); i++) { - memcpy(argBuf + index, kernelParams[i], pl.Size[i]); - index += pl.Align[i]; + std::vector tmp{}; + if (kernelParams) { + if (extra) return hipErrorInvalidValue; + + for (auto&& x : f->_kernarg_layout) { + const auto p{static_cast(*kernelParams)}; + + tmp.insert( + tmp.cend(), + round_up_to_next_multiple_nonnegative( + tmp.size(), x.second) - tmp.size(), + '\0'); + tmp.insert(tmp.cend(), p, p + x.first); + + ++kernelParams; } - config[1] = (void*)argBuf; - kernArgSize = pl.totalSize; - } else if (extra != NULL) { + config[1] = static_cast(tmp.data()); + + kernArgSize = tmp.size(); + } else if (extra) { memcpy(config, extra, sizeof(size_t) * 5); if (config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END) { @@ -235,10 +245,6 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, stopEvent->attachToCompletionFuture(&cf, hStream, hipEventTypeStopCommand); } - - if (kernelParams != NULL) { - free(config[1]); - } ihipPostLaunchKernel(f->_name.c_str(), hStream, lp); } @@ -460,7 +466,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch // TODO: refactor the whole ihipThisThat, which is a mess and yields the // below, due to hipFunction_t being a pointer to ihipModuleSymbol_t. func[0][0] = *static_cast( - Kernel_descriptor{kernel_object(kernel), name}); + Kernel_descriptor{kernel_object(kernel), name, hmod->kernargs[name]}); return hipSuccess; } @@ -550,6 +556,9 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { (*module)->executable = load_executable(content, (*module)->executable, this_agent()); + istringstream elf{content}; + ELFIO::elfio reader; + if (reader.load(elf)) read_kernarg_metadata(reader, (*module)->kernargs); // compute the hash of the code object (*module)->hash = checksum(content.length(), content.data()); From bc69066e95a7e51935723df36a157f9bedba1fb3 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 9 May 2019 07:22:20 -0400 Subject: [PATCH 11/15] Program was already destroyed. --- tests/src/hiprtc/saxpy.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp index b289c84c9c..a0c20feb91 100644 --- a/tests/src/hiprtc/saxpy.cpp +++ b/tests/src/hiprtc/saxpy.cpp @@ -141,7 +141,6 @@ int main() hipFree(dOut); hipModuleUnload(module); - hiprtcDestroyProgram(&prog); passed(); } \ No newline at end of file From f0304546944f3ef496c2586222d4c8c87d00a29a Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 10 May 2019 13:05:48 +0300 Subject: [PATCH 12/15] Don't use `--genco`. Fix mangled name trimming. --- src/hiprtc.cpp | 48 ++++++++++++++++++++++++++++++------------------ 1 file changed, 30 insertions(+), 18 deletions(-) diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp index 60f64dee0b..1e5d9e0af9 100644 --- a/src/hiprtc.cpp +++ b/src/hiprtc.cpp @@ -123,8 +123,11 @@ struct _hiprtcProgram { if (name.find("void ") == 0) name.erase(0, strlen("void ")); - auto dx{name.find('<')}; - if (dx != string::npos) { + auto dx{name.find_first_of("(<")}; + + if (dx == string::npos) return name; + + if (name[dx] == '<') { auto cnt{1u}; do { ++dx; @@ -133,11 +136,7 @@ struct _hiprtcProgram { name.erase(++dx); } - else { - const auto dy{name.find('(')}; - - if (dy != string::npos) name.erase(dy); - } + else name.erase(dx); return name; } @@ -172,6 +171,8 @@ struct _hiprtcProgram { bool compile(const std::vector& args, const std::experimental::filesystem::path& program_folder) { + using namespace ELFIO; + using namespace hip_impl; using namespace redi; using namespace std; @@ -188,29 +189,40 @@ struct _hiprtcProgram { if (compile.rdbuf()->exited() && compile.rdbuf()->status() != EXIT_SUCCESS) return false; - ifstream in{args.back()}; - elf.resize(experimental::filesystem::file_size(args.back())); - in.read(elf.data(), elf.size()); + elfio reader; + if (!reader.load(args.back())) return false; + + const auto it{find_section_if(reader, [](const section* x) { + return x->get_name() == ".kernel"; + })}; + + if (!it) return false; + + Bundled_code_header h{it->get_data()}; + + if (bundles(h).empty()) return false; + + elf.assign(bundles(h).back().blob.cbegin(), + bundles(h).back().blob.cend()); return true; } bool readLoweredNames() { + using namespace ELFIO; using namespace hip_impl; using namespace std; if (names.empty()) return true; - Bundled_code_header h{elf.data()}; - istringstream blob{string{bundles(h).back().blob.cbegin(), - bundles(h).back().blob.cend()}}; + istringstream blob{string{elf.cbegin(), elf.cend()}}; - ELFIO::elfio reader; + elfio reader; if (!reader.load(blob)) return false; - const auto it{find_section_if(reader, [](const ELFIO::section* x) { + const auto it{find_section_if(reader, [](const section* x) { return x->get_type() == SHT_SYMTAB; })}; @@ -423,12 +435,12 @@ namespace if (dx == dy) continue; - x.replace(0, x.find('=', min(dx, dy)), "--targets"); + x.replace(0, x.find('=', min(dx, dy)), "--amdgpu-target"); hasTarget = true; break; } - if (!hasTarget) args.push_back("--targets=" + defaultTarget()); + if (!hasTarget) args.push_back("--amdgpu-target=" + defaultTarget()); } } // Unnamed namespace. @@ -453,7 +465,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) const auto src{p->writeTemporaryFiles(tmp.path())}; - vector args{hipcc, "--genco"}; + vector args{hipcc, "-shared"}; if (n) args.insert(args.cend(), o, o + n); handleTarget(args); From 800673e38f46fa6cd501f75d807aa9b491a8d0f0 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 14 May 2019 12:35:27 +0300 Subject: [PATCH 13/15] Fix HIPRTC breakage due to upstream noise. --- src/hip_module.cpp | 4 +++- src/hiprtc.cpp | 39 +++++++++++++++++++++++++++++++++------ src/program_state.inl | 13 +++++++++++-- 3 files changed, 47 insertions(+), 9 deletions(-) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 686dc45a71..e8a8801e98 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -557,7 +557,9 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { this_agent()); istringstream elf{content}; ELFIO::elfio reader; - if (reader.load(elf)) read_kernarg_metadata(reader, (*module)->kernargs); + if (reader.load(elf)) { + program_state_impl::read_kernarg_metadata(reader, (*module)->kernargs); + } // compute the hash of the code object (*module)->hash = checksum(content.length(), content.data()); diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp index 1e5d9e0af9..aadf48c3ed 100644 --- a/src/hiprtc.cpp +++ b/src/hiprtc.cpp @@ -32,6 +32,7 @@ THE SOFTWARE. #include #include +#include #include #include #include @@ -79,6 +80,31 @@ const char* hiprtcGetErrorString(hiprtcResult x) }; } +namespace +{ + struct Symbol { + std::string name; + ELFIO::Elf64_Addr value = 0; + ELFIO::Elf_Xword size = 0; + ELFIO::Elf_Half sect_idx = 0; + std::uint8_t bind = 0; + std::uint8_t type = 0; + std::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; + } +} // Unnamed namespace. + struct _hiprtcProgram { // DATA - STATICS static std::vector> programs; @@ -172,7 +198,6 @@ struct _hiprtcProgram { const std::experimental::filesystem::path& program_folder) { using namespace ELFIO; - using namespace hip_impl; using namespace redi; using namespace std; @@ -192,13 +217,14 @@ struct _hiprtcProgram { elfio reader; if (!reader.load(args.back())) return false; - const auto it{find_section_if(reader, [](const section* x) { + const auto it{find_if(reader.sections.begin(), reader.sections.end(), + [](const section* x) { return x->get_name() == ".kernel"; })}; - if (!it) return false; + if (it == reader.sections.end()) return false; - Bundled_code_header h{it->get_data()}; + hip_impl::Bundled_code_header h{(*it)->get_data()}; if (bundles(h).empty()) return false; @@ -222,11 +248,12 @@ struct _hiprtcProgram { if (!reader.load(blob)) return false; - const auto it{find_section_if(reader, [](const section* x) { + const auto it{find_if(reader.sections.begin(), reader.sections.end(), + [](const section* x) { return x->get_type() == SHT_SYMTAB; })}; - ELFIO::symbol_section_accessor symbols{reader, it}; + ELFIO::symbol_section_accessor symbols{reader, *it}; auto n{symbols.get_symbols_num()}; diff --git a/src/program_state.inl b/src/program_state.inl index 9729da8115..f1397b3fe9 100644 --- a/src/program_state.inl +++ b/src/program_state.inl @@ -66,10 +66,17 @@ class Kernel_descriptor { std::uint64_t kernel_object_{}; amd_kernel_code_t const* kernel_header_{nullptr}; std::string name_{}; + std::vector> kernarg_layout_{}; public: Kernel_descriptor() = default; - Kernel_descriptor(std::uint64_t kernel_object, const std::string& name) - : kernel_object_{kernel_object}, name_{name} + Kernel_descriptor( + std::uint64_t kernel_object, + const std::string& name, + std::vector> kernarg_layout = {}) + : + kernel_object_{kernel_object}, + name_{name}, + kernarg_layout_{std::move(kernarg_layout)} { bool supported{false}; std::uint16_t min_v{UINT16_MAX}; @@ -548,6 +555,7 @@ public: return functions[agent].second; } + static std::size_t parse_args( const std::string& metadata, std::size_t f, @@ -576,6 +584,7 @@ public: } while (true); } + static void read_kernarg_metadata( ELFIO::elfio& reader, std::unordered_map< From 2225942869bcdfbcb017cca50df0dd726c2b9be9 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 15 May 2019 14:07:30 +0530 Subject: [PATCH 14/15] [dtests] Replace RUN -> TEST in hiprtc tests Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3 --- tests/src/hiprtc/hiprtcGetLoweredName.cpp | 4 ++-- tests/src/hiprtc/hiprtcGetTypeName.cpp | 4 ++-- tests/src/hiprtc/saxpy.cpp | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/tests/src/hiprtc/hiprtcGetLoweredName.cpp index 2e4f4e9dde..e3fa057a81 100644 --- a/tests/src/hiprtc/hiprtcGetLoweredName.cpp +++ b/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t * HIT_END */ #include @@ -149,4 +149,4 @@ int main() hiprtcDestroyProgram(&prog); passed(); -} \ No newline at end of file +} diff --git a/tests/src/hiprtc/hiprtcGetTypeName.cpp b/tests/src/hiprtc/hiprtcGetTypeName.cpp index da94c926ea..b0348408f3 100644 --- a/tests/src/hiprtc/hiprtcGetTypeName.cpp +++ b/tests/src/hiprtc/hiprtcGetTypeName.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t * HIT_END */ @@ -135,4 +135,4 @@ int main() hiprtcDestroyProgram(&prog); passed(); -} \ No newline at end of file +} diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp index a0c20feb91..5f9dc7a125 100644 --- a/tests/src/hiprtc/saxpy.cpp +++ b/tests/src/hiprtc/saxpy.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc - * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t * HIT_END */ @@ -143,4 +143,4 @@ int main() hipModuleUnload(module); passed(); -} \ No newline at end of file +} From 49951190074bf3e7e5369ca34bdc459ae2a9b22e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 16 May 2019 16:16:49 +0530 Subject: [PATCH 15/15] [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a --- tests/hit/HIT.cmake | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/hit/HIT.cmake b/tests/hit/HIT.cmake index 173d3a8dca..d4c162f4b9 100644 --- a/tests/hit/HIT.cmake +++ b/tests/hit/HIT.cmake @@ -157,12 +157,12 @@ macro(MAKE_TEST exe) string(REPLACE " " "" smush_args ${ARGN}) set(testname ${exe}${smush_args}.tst) add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}) - set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED") + set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" ENVIRONMENT HIP_PATH=${HIP_ROOT_DIR}) endmacro() macro(MAKE_NAMED_TEST exe testname) add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}) - set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED") + set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" ENVIRONMENT HIP_PATH=${HIP_ROOT_DIR}) endmacro() #-------------------------------------------------------------------------------