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

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__

include/hip/hcc_detail/hip_runtime.h

@@ -313,6 +313,7 @@ static constexpr Coordinates<hip_impl::WorkitemId> 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 <hip/hcc_detail/hip_memory.h>

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))

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")

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