From a4a3132c640b01cf93f825c19a8472bb0f7c528f Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 16 May 2019 15:58:54 +0300 Subject: [PATCH] Add HIPRTC, glorious ersatz for NVRTC (#1097) * Add ersatz for NVRTC. * Fix extraneous paren and use correct namespace. * Use lowerCamelCase (yuck, yuck) consistently. * Link against FS when building hiprtc lib. * Correctly mark Manipulators. Fix dual compile. * Add unit tests. Extend HIT to accept linker options. * Make sure the HIPRTC library is installed. * Better logging. Try to auto-detect the target. * Stop specifying the target explicitly. * Add missing flavour of `hipModuleLaunchKernel`. * Program was already destroyed. * Don't use `--genco`. Fix mangled name trimming. * Fix HIPRTC breakage due to upstream noise. * [dtests] Replace RUN -> TEST in hiprtc tests Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3 * [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a [ROCm/clr commit: a538eb705a20ac740fb4a182ec945796ad6a83ee] --- projects/clr/hipamd/CMakeLists.txt | 10 +- .../hipamd/include/hip/hcc_detail/hiprtc.h | 111 ++++ .../include/hip/hcc_detail/program_state.hpp | 2 +- projects/clr/hipamd/include/hip/hiprtc.h | 32 + projects/clr/hipamd/packaging/hip_hcc.txt | 1 + projects/clr/hipamd/src/hip_hcc_internal.h | 2 + projects/clr/hipamd/src/hip_module.cpp | 45 +- projects/clr/hipamd/src/hiprtc.cpp | 601 ++++++++++++++++++ projects/clr/hipamd/src/program_state.inl | 13 +- projects/clr/hipamd/tests/hit/HIT.cmake | 23 +- .../tests/src/hiprtc/hiprtcGetLoweredName.cpp | 152 +++++ .../tests/src/hiprtc/hiprtcGetTypeName.cpp | 138 ++++ .../clr/hipamd/tests/src/hiprtc/saxpy.cpp | 146 +++++ 13 files changed, 1250 insertions(+), 26 deletions(-) create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h create mode 100644 projects/clr/hipamd/include/hip/hiprtc.h create mode 100644 projects/clr/hipamd/src/hiprtc.cpp create mode 100644 projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp create mode 100644 projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp create mode 100644 projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 30aa071515..30a699c7ce 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -253,7 +253,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") @@ -263,6 +263,12 @@ 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) + target_link_libraries(hiprtc PUBLIC stdc++fs) endif() string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS}) @@ -297,7 +303,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/projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h b/projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h new file mode 100644 index 0000000000..d4174e4ce1 --- /dev/null +++ b/projects/clr/hipamd/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 = hip_impl::demangle(typeid(T).name()); + + return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR : + HIPRTC_SUCCESS; + } + #endif +#endif \ No newline at end of file diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp index ca8b6dcd88..19db8f9e0d 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp @@ -75,7 +75,7 @@ public: hipFunction_t kernel_descriptor(std::uintptr_t, hsa_agent_t); - + kernargs_size_align get_kernargs_size_align(std::uintptr_t); hsa_executable_t load_executable(const char*, const size_t, hsa_executable_t, diff --git a/projects/clr/hipamd/include/hip/hiprtc.h b/projects/clr/hipamd/include/hip/hiprtc.h new file mode 100644 index 0000000000..22d78d2656 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/packaging/hip_hcc.txt b/projects/clr/hipamd/packaging/hip_hcc.txt index fe866e47f9..5aebc6c36d 100644 --- a/projects/clr/hipamd/packaging/hip_hcc.txt +++ b/projects/clr/hipamd/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) diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index b40fac93a5..769d6b7914 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index e029ec6e78..e8a8801e98 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -91,6 +91,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 <> @@ -132,6 +133,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; @@ -146,19 +149,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) { @@ -236,10 +246,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); } @@ -461,7 +467,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; } @@ -549,6 +555,11 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { (*module)->executable = get_program_state().load_executable( content.data(), content.size(), (*module)->executable, this_agent()); + istringstream elf{content}; + ELFIO::elfio reader; + 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/projects/clr/hipamd/src/hiprtc.cpp b/projects/clr/hipamd/src/hiprtc.cpp new file mode 100644 index 0000000000..aadf48c3ed --- /dev/null +++ b/projects/clr/hipamd/src/hiprtc.cpp @@ -0,0 +1,601 @@ +/* +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 +#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."}; + }; +} + +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; + static std::mutex mtx; + + // DATA + std::vector> headers; + std::vector> names; + std::vector loweredNames; + 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_first_of("(<")}; + + if (dx == string::npos) return name; + + if (name[dx] == '<') { + auto cnt{1u}; + do { + ++dx; + cnt += (name[dx] == '<') ? 1 : ((name[dx] == '>') ? -1 : 0); + } while (cnt); + + name.erase(++dx); + } + else name.erase(dx); + + 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 + { + return std::find_if(programs.cbegin(), programs.cend(), + [=](const std::unique_ptr<_hiprtcProgram>& x) { + return x.get() == p; + }) != programs.cend(); + } + + // MANIPULATORS + bool compile(const std::vector& args, + const std::experimental::filesystem::path& program_folder) + { + using namespace ELFIO; + using namespace redi; + using namespace std; + + ipstream compile{args.front(), args, pstreambuf::pstderr}; + + constexpr const auto tmp_size{1024u}; + char tmp[tmp_size]{}; + while (!compile.eof()) { + log.append(tmp, tmp + compile.readsome(tmp, tmp_size)); + } + + compile.close(); + + if (compile.rdbuf()->exited() && + compile.rdbuf()->status() != EXIT_SUCCESS) return false; + + elfio reader; + if (!reader.load(args.back())) return false; + + const auto it{find_if(reader.sections.begin(), reader.sections.end(), + [](const section* x) { + return x->get_name() == ".kernel"; + })}; + + if (it == reader.sections.end()) return false; + + hip_impl::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; + + istringstream blob{string{elf.cbegin(), elf.cend()}}; + + elfio reader; + + if (!reader.load(blob)) return false; + + 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}; + + auto n{symbols.get_symbols_num()}; + + if (n < loweredNames.size()) return false; + + 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; + } + + loweredNames[distance(names.cbegin(), it)] = tmp.name; + } + + return true; + } + + // ACCESSORS + std::experimental::filesystem::path writeTemporaryFiles( + const std::experimental::filesystem::path& programFolder) const + { + using namespace std; + + vector> fut{headers.size()}; + transform(headers.cbegin(), headers.cend(), begin(fut), + [&](const pair& x) { + return async([&]() { + ofstream h{programFolder / x.first}; + h.write(x.second.data(), x.second.size()); + }); + }); + + auto tmp{(programFolder / 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); + p->loweredNames.emplace_back(); + + 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 +{ + 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) + { + using namespace std; + + bool hasTarget{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)), "--amdgpu-target"); + hasTarget = true; + + break; + } + if (!hasTarget) args.push_back("--amdgpu-target=" + defaultTarget()); + } +} // 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->writeTemporaryFiles(tmp.path())}; + + vector args{hipcc, "-shared"}; + if (n) args.insert(args.cend(), o, o + n); + + handleTarget(args); + + args.emplace_back(src); + args.emplace_back("-o"); + args.emplace_back(tmp.path() / "hiprtc.out"); + + if (!p->compile(args, tmp.path())) return HIPRTC_ERROR_INTERNAL_ERROR; + if (!p->readLoweredNames()) return HIPRTC_ERROR_INTERNAL_ERROR; + + p->compiled = true; + + return HIPRTC_SUCCESS; +} + +hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, + const char* name, int n, const char** hdrs, + const char** incs) +{ + using namespace std; + + if (!p) return HIPRTC_ERROR_INVALID_PROGRAM; + if (n < 0) return HIPRTC_ERROR_INVALID_INPUT; + if (n && (!hdrs || !incs)) return HIPRTC_ERROR_INVALID_INPUT; + + vector> h; + for (auto i = 0; i != n; ++i) h.emplace_back(incs[i], hdrs[i]); + + *p = _hiprtcProgram::make(src, name ? name : "default_name", 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->loweredNames[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.empty() ? 0 : 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 diff --git a/projects/clr/hipamd/src/program_state.inl b/projects/clr/hipamd/src/program_state.inl index 9729da8115..f1397b3fe9 100644 --- a/projects/clr/hipamd/src/program_state.inl +++ b/projects/clr/hipamd/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< diff --git a/projects/clr/hipamd/tests/hit/HIT.cmake b/projects/clr/hipamd/tests/hit/HIT.cmake index f62ba65657..d4c162f4b9 100644 --- a/projects/clr/hipamd/tests/hit/HIT.cmake +++ b/projects/clr/hipamd/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() @@ -143,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() #------------------------------------------------------------------------------- @@ -165,7 +179,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}) @@ -173,6 +187,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/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetLoweredName.cpp new file mode 100644 index 0000000000..e3fa057a81 --- /dev/null +++ b/projects/clr/hipamd/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 + * TEST: %t + * 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(); +} diff --git a/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp b/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp new file mode 100644 index 0000000000..b0348408f3 --- /dev/null +++ b/projects/clr/hipamd/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 + * TEST: %t + * 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(); +} diff --git a/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp b/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp new file mode 100644 index 0000000000..5f9dc7a125 --- /dev/null +++ b/projects/clr/hipamd/tests/src/hiprtc/saxpy.cpp @@ -0,0 +1,146 @@ +/* +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 + * TEST: %t + * 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 + + 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()); + + 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); + + passed(); +}