From 0a513d8a02bc0dcc029dbe951f4a8b65fcd5598c Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 4 Jun 2020 13:50:16 -0400 Subject: [PATCH] Fix include path and wrapper header Currently std::complex and some other std functions require uses to include hip_runtime.h before any other headers to work, which is not reliable. changes are made in clang to fix this issue: https://reviews.llvm.org/D81176 which requires hipcc and HIP headers to make corresponding changes. This patch will make sure the clang change will not break HIP/ROCclr during this transition. After the transition is done, we can remove explicitly setting include path for HIP-Clang and HIP header in hipcc and hip config cmake files and rely on clang driver to set it automatically. Change-Id: I5d226861c2560ffa6c5ab17343a43cc378048061 --- hipamd/bin/hipcc | 5 ++++- hipamd/hip-config.cmake.in | 2 +- hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h | 3 ++- hipamd/include/hip/hcc_detail/hip_runtime.h | 6 ++++-- hipamd/include/hip/hcc_detail/host_defines.h | 2 ++ hipamd/include/hip/hcc_detail/math_functions.h | 5 ++++- hipamd/include/hip/hcc_detail/math_fwd.h | 4 +++- 7 files changed, 20 insertions(+), 7 deletions(-) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index c1e144300b..ad46e8e322 100755 --- a/hipamd/bin/hipcc +++ b/hipamd/bin/hipcc @@ -214,6 +214,7 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { $HIP_LIB_PATH = "$HIP_PATH/lib"; } if ($verbose & 0x2) { + print ("ROCM_PATH=$ROCM_PATH\n"); if (defined $HIP_ROCCLR_HOME) { print ("HIP_ROCCLR_HOME=$HIP_ROCCLR_HOME\n"); } @@ -823,7 +824,9 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { $HIPLDFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; } } - $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; + if ($DEVICE_LIB_PATH ne "$ROCM_PATH/amdgcn/bitcode") { + $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; + } if ($hasHIP) { $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS"; if ($HIP_RUNTIME ne "HCC") { diff --git a/hipamd/hip-config.cmake.in b/hipamd/hip-config.cmake.in index cc80e1ff4a..6ac076e279 100644 --- a/hipamd/hip-config.cmake.in +++ b/hipamd/hip-config.cmake.in @@ -150,7 +150,7 @@ if(HIP_COMPILER STREQUAL "clang") if (EXISTS ${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode) set_property(TARGET hip::device APPEND PROPERTY - INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode + INTERFACE_COMPILE_OPTIONS -x hip ) else() # This path is to support an older build of the device library diff --git a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h index 95403e6ca8..16d834cca8 100644 --- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h +++ b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h @@ -27,7 +27,7 @@ THE SOFTWARE. // */ #include "host_defines.h" - +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ extern "C" { __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); @@ -82,3 +82,4 @@ extern "C" __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); } +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index a166935823..117af3e1fa 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -313,6 +313,7 @@ static constexpr Coordinates threadIdx{}; #endif // defined __HCC__ #if __HCC_OR_HIP_CLANG__ +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #if __HIP_ENABLE_DEVICE_MALLOC__ extern "C" __device__ void* __hip_malloc(size_t); extern "C" __device__ void* __hip_free(void* ptr); @@ -322,7 +323,7 @@ static inline __device__ void* free(void* ptr) { return __hip_free(ptr); } static inline __device__ void* malloc(size_t size) { __builtin_trap(); return nullptr; } static inline __device__ void* free(void* ptr) { __builtin_trap(); return nullptr; } #endif - +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #endif //__HCC_OR_HIP_CLANG__ #ifdef __HCC__ @@ -550,6 +551,7 @@ hc_get_workitem_absolute_id(int dim) #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // Support std::complex. #if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ #pragma push_macro("__CUDA__") @@ -567,7 +569,7 @@ hc_get_workitem_absolute_id(int dim) #undef __CUDA__ #pragma pop_macro("__CUDA__") #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ - +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #endif // defined(__clang__) && defined(__HIP__) #include diff --git a/hipamd/include/hip/hcc_detail/host_defines.h b/hipamd/include/hip/hcc_detail/host_defines.h index ad28cc7626..72f3932aff 100644 --- a/hipamd/include/hip/hcc_detail/host_defines.h +++ b/hipamd/include/hip/hcc_detail/host_defines.h @@ -64,11 +64,13 @@ THE SOFTWARE. #elif defined(__clang__) && defined(__HIP__) +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #define __host__ __attribute__((host)) #define __device__ __attribute__((device)) #define __global__ __attribute__((global)) #define __shared__ __attribute__((shared)) #define __constant__ __attribute__((constant)) +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #define __noinline__ __attribute__((noinline)) #define __forceinline__ inline __attribute__((always_inline)) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index 494685e261..876a14eaa9 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -71,6 +71,7 @@ struct __numeric_type<_Float16> #define __RETURN_TYPE bool #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ __DEVICE__ inline uint64_t __make_mantissa_base8(const char* tagp) @@ -139,6 +140,7 @@ uint64_t __make_mantissa(const char* tagp) return __make_mantissa_base10(tagp); } +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // DOT FUNCTIONS #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__ @@ -174,6 +176,7 @@ uint amd_mixed_dot(uint a, uint b, uint c, bool saturate) { } #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // BEGIN FLOAT __DEVICE__ inline @@ -1507,7 +1510,7 @@ __host__ inline static int min(int arg1, int arg2) { __host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); } - +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #pragma pop_macro("__DEF_FLOAT_FUN") #pragma pop_macro("__DEF_FLOAT_FUN2") diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h index c25b5e90b4..4c0fde591c 100644 --- a/hipamd/include/hip/hcc_detail/math_fwd.h +++ b/hipamd/include/hip/hcc_detail/math_fwd.h @@ -23,7 +23,6 @@ THE SOFTWARE. #pragma once #include "host_defines.h" - #if defined(__cplusplus) extern "C" { #endif @@ -67,6 +66,7 @@ __attribute__((const)) unsigned int __ockl_udot8(unsigned int, unsigned int, unsigned int, bool); #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // BEGIN FLOAT __device__ __attribute__((const)) @@ -701,6 +701,8 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); // END INTRINSICS // END DOUBLE +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ + #if defined(__cplusplus) } // extern "C" #endif