From 21712e6fb853adb69cebb67d362824a137b282d9 Mon Sep 17 00:00:00 2001 From: saleelk Date: Fri, 28 Feb 2020 03:17:29 -0800 Subject: [PATCH] Fix HIPRTC headers to export C style symbols (#1879) [ROCm/clr commit: 40a28e767e46e0dfaafd5b3f422bc91d6eb0379e] --- .../hipamd/include/hip/hcc_detail/hiprtc.h | 66 +++----- projects/clr/hipamd/src/hiprtc.cpp | 64 ++++---- .../tests/src/hiprtc/hiprtcGetTypeName.cpp | 147 ------------------ 3 files changed, 58 insertions(+), 219 deletions(-) delete mode 100644 projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h b/projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h index 26d3129dbc..624f1ea157 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hiprtc.h @@ -19,10 +19,14 @@ 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 +#ifndef HIPRTC_H +#define HIPRTC_H -#include -#include +#ifdef __cplusplus +extern "C" { +#endif /* __cplusplus */ + +#include enum hiprtcResult { HIPRTC_SUCCESS = 0, @@ -41,29 +45,22 @@ enum hiprtcResult { 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; +hiprtcResult hiprtcVersion(int* major, int* minor); - return HIPRTC_SUCCESS; -} - -struct _hiprtcProgram; -using hiprtcProgram = _hiprtcProgram*; +typedef struct _hiprtcProgram* hiprtcProgram; hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression); -hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, +hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, + int numOptions, const char** options); -hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, - const char* name, int numHeaders, +hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, + const char* src, + const char* name, + int numHeaders, const char** headers, const char** includeNames); @@ -76,37 +73,14 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log); hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, - std::size_t* logSizeRet); + size_t* logSizeRet); hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code); -hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, std::size_t* codeSizeRet); +hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet); -namespace hip_impl -{ - char* demangle(const char* mangled_expression); +#ifdef __cplusplus } +#endif /* __cplusplus */ -#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; - - char * res= hip_impl::demangle(typeid(T).name()); - result->assign(res == nullptr ? "" : res); - std::free(res); - return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR : - HIPRTC_SUCCESS; - } - #endif -#endif +#endif //HIPRTC_H diff --git a/projects/clr/hipamd/src/hiprtc.cpp b/projects/clr/hipamd/src/hiprtc.cpp index 5198bf0cbb..3c7fe6e78c 100644 --- a/projects/clr/hipamd/src/hiprtc.cpp +++ b/projects/clr/hipamd/src/hiprtc.cpp @@ -50,7 +50,7 @@ THE SOFTWARE. #include #include -const char* hiprtcGetErrorString(hiprtcResult x) +extern "C" const char* hiprtcGetErrorString(hiprtcResult x) { switch (x) { case HIPRTC_SUCCESS: @@ -95,6 +95,21 @@ inline bool fileExists (const std::string& name) { } } // namespace hip_impl +namespace +{ + char* demangle(const char* x) + { + if (!x) return nullptr; + + int s{}; + char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s); + + if (s != 0) return nullptr; + + return tmp; + } +} // Unnamed namespace. + namespace { struct Symbol { @@ -158,7 +173,7 @@ struct _hiprtcProgram { { using namespace std; - char* demangled = hip_impl::demangle(name.c_str()); + char* demangled = demangle(name.c_str()); name.assign(demangled == nullptr ? "" : demangled); free(demangled); @@ -352,7 +367,7 @@ namespace } } // Unnamed namespace. -hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) +extern "C" hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n) { if (!n) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -413,21 +428,6 @@ namespace }; } // Unnamed namespace. -namespace hip_impl -{ - char* demangle(const char* x) - { - if (!x) return nullptr; - - int s{}; - char* tmp = abi::__cxa_demangle(x, nullptr, nullptr, &s); - - if (s != 0) return nullptr; - - return tmp; - } -} // Namespace hip_impl. - namespace { const std::string& defaultTarget() @@ -492,7 +492,7 @@ namespace } } // Unnamed namespace. -hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) +extern "C" hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) { using namespace std; @@ -530,7 +530,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, +extern "C" hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, const char* name, int n, const char** hdrs, const char** incs) { @@ -548,14 +548,14 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src, return HIPRTC_SUCCESS; } -hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p) +extern "C" hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p) { if (!p) return HIPRTC_SUCCESS; return _hiprtcProgram::destroy(*p); } -hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, +extern "C" hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, const char** ln) { using namespace std; @@ -576,7 +576,7 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n, return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) +extern "C" hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) { if (!l) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -588,7 +588,7 @@ hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) +extern "C" hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) { if (!sz) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -599,7 +599,7 @@ hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) +extern "C" hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) { if (!c) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -610,7 +610,7 @@ hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c) return HIPRTC_SUCCESS; } -hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) +extern "C" hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) { if (!sz) return HIPRTC_ERROR_INVALID_INPUT; if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM; @@ -620,3 +620,15 @@ hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz) return HIPRTC_SUCCESS; } + +extern "C" hiprtcResult hiprtcVersion(int* major, int* minor) +{ + if (major == nullptr || minor == nullptr) { + return HIPRTC_ERROR_INVALID_INPUT; + } + + *major = 9; + *minor = 0; + + return HIPRTC_SUCCESS; +} diff --git a/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp b/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp deleted file mode 100644 index 812229f81f..0000000000 --- a/projects/clr/hipamd/tests/src/hiprtc/hiprtcGetTypeName.cpp +++ /dev/null @@ -1,147 +0,0 @@ -/* -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()); - - hipDeviceProp_t props; - int device = 0; - hipGetDeviceProperties(&props, device); - std::string gfxName = "gfx" + std::to_string(props.gcnArch); - std::string sarg = "--gpu-architecture=" + gfxName; - const char* options[] = { - sarg.c_str() - }; - - hiprtcResult compileResult = hiprtcCompileProgram(prog, 1, options); - - 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(); -}