From 3bfff0a23db52a869709b32aa588bdeee77347b0 Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Mon, 5 Aug 2019 15:21:36 +0530 Subject: [PATCH] Enabled gcc for hip host code (#1214) * Enabled gcc for hip host code * Adding tests for hip code + (gcc & g++), without kernels * Excluding nvcc platforms for gcc and g++ tests + Addressing review comments * minor code clean-up * Add rocm include path * Added relative path for library * Hiding non supported functions for gcc * Incorporating review comments --- hip_prof_gen.py | 7 +++-- include/hip/hcc_detail/hip_runtime_api.h | 18 ++++++++++++- include/hip/hcc_detail/hip_vector_types.h | 21 +++++++++++++++ include/hip/hip_runtime_api.h | 2 ++ .../complex_loading_behavior.cpp | 2 +- tests/src/g++/hipMalloc.cpp | 27 +++++++++++++++++++ tests/src/gcc/hipMalloc.c | 27 +++++++++++++++++++ 7 files changed, 100 insertions(+), 4 deletions(-) diff --git a/hip_prof_gen.py b/hip_prof_gen.py index 587b2d545a..e2876d87bc 100755 --- a/hip_prof_gen.py +++ b/hip_prof_gen.py @@ -333,7 +333,7 @@ def generate_prof_header(f, api_map, opts_map): # Generating the callbacks ID enumaration f.write('\n// Return HIP API string\n') - f.write('static inline const char* hip_api_name(const uint32_t& id) {\n') + f.write('static inline const char* hip_api_name(const uint32_t id) {\n') f.write(' switch(id) {\n') for name in api_map.keys(): f.write(' case HIP_API_ID_' + name + ': return "' + name + '";\n') @@ -353,7 +353,10 @@ def generate_prof_header(f, api_map, opts_map): if len(args) != 0: f.write(' struct {\n') for arg_tuple in args: - f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + if arg_tuple[0] == "hipLimit_t": + f.write(' enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + else: + f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') f.write(' } ' + name + ';\n') f.write( ' } args;\n' + diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 755d15d8aa..e8871912fa 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -45,7 +45,7 @@ THE SOFTWARE. #include #include -#if !__HIP_VDI__ +#if !__HIP_VDI__ && defined(__cplusplus) #include #include #endif @@ -79,9 +79,11 @@ THE SOFTWARE. #define __dparm(x) #endif +#ifdef __cplusplus namespace hip_impl { hipError_t hip_init(); } // namespace hip_impl +#endif // Structure definitions: #ifdef __cplusplus @@ -1459,6 +1461,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, #else hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*); +#ifdef __cplusplus //Start : Not supported in gcc namespace hip_impl { inline __attribute__((visibility("hidden"))) @@ -1466,6 +1469,7 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, const char* name); } // Namespace hip_impl. + /** * @brief Copies the memory address of symbol @p symbolName to @p devPtr * @@ -1504,15 +1508,18 @@ hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { void* devPtr = nullptr; return hip_impl::read_agent_global_from_process(&devPtr, size, (const char*)symbolName); } +#endif // End : Not supported in gcc #if defined(__cplusplus) } // extern "C" #endif +#ifdef __cplusplus namespace hip_impl { hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind, const char*); } // Namespace hip_impl. +#endif #if defined(__cplusplus) extern "C" { @@ -1541,6 +1548,7 @@ extern "C" { * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync */ +#ifdef __cplusplus inline __attribute__((visibility("hidden"))) hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, @@ -1554,11 +1562,13 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, return hip_impl::hipMemcpyToSymbol(dst, src, sizeBytes, offset, kind, (const char*)symbolName); } +#endif #if defined(__cplusplus) } // extern "C" #endif +#ifdef __cplusplus namespace hip_impl { hipError_t hipMemcpyToSymbolAsync(void*, const void*, size_t, size_t, hipMemcpyKind, hipStream_t, const char*); @@ -1567,6 +1577,7 @@ hipError_t hipMemcpyFromSymbol(void*, const void*, size_t, size_t, hipError_t hipMemcpyFromSymbolAsync(void*, const void*, size_t, size_t, hipMemcpyKind, hipStream_t, const char*); } // Namespace hip_impl. +#endif #if defined(__cplusplus) extern "C" { @@ -1597,6 +1608,8 @@ extern "C" { * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync */ + +#ifdef __cplusplus //Start : Not supported in gcc inline __attribute__((visibility("hidden"))) hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, @@ -1641,6 +1654,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, stream, (const char*)symbolName); } +#endif // End : Not supported in gcc #endif // __HIP_VDI__ /** @@ -2604,6 +2618,7 @@ hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func } // extern "C" #endif +#ifdef __cplusplus namespace hip_impl { class agent_globals_impl; class agent_globals { @@ -2635,6 +2650,7 @@ namespace hip_impl { return get_agent_globals().read_agent_global_from_process(dptr, bytes, name); } } // Namespace hip_impl. +#endif #if defined(__cplusplus) extern "C" { diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 3407b8a752..a03a46b8cf 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -849,6 +849,7 @@ __MAKE_VECTOR_TYPE__(longlong, long long); __MAKE_VECTOR_TYPE__(float, float); __MAKE_VECTOR_TYPE__(double, double); +#ifdef __cplusplus #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ static inline __device__ __host__ \ type make_##type(comp x) { type r{x}; return r; } @@ -867,6 +868,26 @@ __MAKE_VECTOR_TYPE__(double, double); type r{x, y, z, w}; \ return r; \ } +#else + #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x) { type r; r.x =x; return r; } + + #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; } + + #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; } + + #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x, comp y, comp z, comp w) { \ + type r; r.x=x; r.y=y; r.z=z; r.w=w; \ + return r; \ + } +#endif DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1); DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2); diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index c067f5bf54..1c8e971545 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -255,6 +255,8 @@ typedef enum __HIP_NODISCARD hipError_t { 1071, ///< Produced when the IPC memory attach failed from ROCr. hipErrorAssert = 1081, ///< Produced when the kernel calls assert. + hipErrorNotSupported = + 1082, ///< Produced when the hip API is not supported/implemented hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; diff --git a/tests/src/dynamicLoading/complex_loading_behavior.cpp b/tests/src/dynamicLoading/complex_loading_behavior.cpp index 6ee21a9cad..b8ed0f16e0 100644 --- a/tests/src/dynamicLoading/complex_loading_behavior.cpp +++ b/tests/src/dynamicLoading/complex_loading_behavior.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc * BUILD_CMD: libfoo_nvidia %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM hcc * BUILD_CMD: %t %hc %S/%s -o %T/%t -ldl * TEST: %t diff --git a/tests/src/g++/hipMalloc.cpp b/tests/src/g++/hipMalloc.cpp index 9c1f695a85..3aab48aad5 100644 --- a/tests/src/g++/hipMalloc.cpp +++ b/tests/src/g++/hipMalloc.cpp @@ -1,7 +1,34 @@ +/* + * Copyright (c) 2019-2020 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * */ + +/* HIT_START + * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_HCC__ -I%hip-path/include -I/opt/rocm/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ + + #include #include int main() { int* Ad; hipMalloc((void**)&Ad, 1024); + std::cout<<"PASSED!"< #include @@ -5,4 +31,5 @@ int main() { int *Ad; hipMalloc((void**)&Ad, 1024); + printf("PASSED!\n"); }