From 9c43970a27df175efea7247618a3ca3da596ec34 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Fri, 25 May 2018 16:18:46 -0400
Subject: [PATCH] Add more function declarations to hip-clang

---
 hipamd/bin/hipcc                            |  2 +-
 hipamd/include/hip/hcc_detail/hip_runtime.h | 30 +++++++++++++++++++--
 2 files changed, 29 insertions(+), 3 deletions(-)

diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc
index b162de2dce..c1446d0f7d 100755
--- a/hipamd/bin/hipcc
+++ b/hipamd/bin/hipcc
@@ -90,7 +90,7 @@ if ($HIP_PLATFORM eq "clang") {
     }
     $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
     $HIPCC="$HIP_CLANG_PATH/clang++";
-    $HIPCXXFLAGS .= " -I$HIP_PATH/include";
+    $HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include";
 
     $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
 } elsif ($HIP_PLATFORM eq "hcc") {
diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index d682d21dbd..1e469f5d03 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -29,6 +29,20 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
 
+#if defined(__HCC__)
+#define __HCC_OR_HIP_CLANG__ 1
+#define __HCC_ONLY__ 1
+#define __HIP_CLANG_ONLY__ 0
+#elif defined(__clang__) && defined(__HIP__)
+#define __HCC_OR_HIP_CLANG__ 1
+#define __HCC_ONLY__ 0
+#define __HIP_CLANG_ONLY__ 1
+#else
+#define __HCC_OR_HIP_CLANG__ 0
+#define __HCC_ONLY__ 0
+#define __HIP_CLANG_ONLY__ 0
+#endif
+
 //---
 // Top part of file can be compiled with any compiler
 
@@ -41,15 +55,16 @@ THE SOFTWARE.
 #include <stddef.h>
 #endif  //__cplusplus
 
-#if __HCC__
+#if __HCC_OR_HIP_CLANG__
 
 // Define NVCC_COMPAT for CUDA compatibility
 #define NVCC_COMPAT
 #define CUDA_SUCCESS hipSuccess
 
 #include <hip/hip_runtime_api.h>
+#endif  // __HCC_OR_HIP_CLANG__
 
-
+#if __HCC__
 // define HIP_ENABLE_PRINTF to enable printf
 #ifdef HIP_ENABLE_PRINTF
 #define HCC_ENABLE_ACCELERATOR_PRINTF 1
@@ -164,6 +179,10 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif
 
+#endif  // defined __HCC__
+
+#if __HCC_OR_HIP_CLANG__
+
 // TODO - hipify-clang - change to use the function call.
 //#define warpSize hc::__wavesize()
 static constexpr int warpSize = 64;
@@ -371,6 +390,10 @@ __device__ void __threadfence_system(void);
  * @}
  */
 
+#endif  // __HCC_OR_CLANG__
+
+#if defined __HCC__
+
 template <
     typename std::common_type<decltype(hc_get_group_id), decltype(hc_get_group_size),
                               decltype(hc_get_num_groups), decltype(hc_get_workitem_id)>::type f>
@@ -628,6 +651,9 @@ __DEVICE__ void inline __assert_fail(const char * __assertion,
     // Ignore all the args for now.
     __device_trap();
 }
+
+__DEVICE__ void __syncthreads();
+
 #pragma push_macro("__DEVICE__")
 
 #include <hip/hcc_detail/math_functions.h>