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

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

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

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__

hipamd/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>

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

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

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