Merge pull request #469 from yxsamliu/fun

Add math function declaration for hip-clang

[ROCm/hip commit: a1fea02c6e]

projects/hip/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") {

projects/hip/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>
@@ -514,7 +537,9 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
  *   @}
  */
 
-
+//
+// hip-clang functions
+//
 #elif defined(__clang__) && defined(__HIP__)
 
 #define HIP_KERNEL_NAME(...) __VA_ARGS__
@@ -626,8 +651,13 @@ __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>
+
 #endif
 
 #endif  // HIP_HCC_DETAIL_RUNTIME_H