From: Yaxun (Sam) Liu <yaxun.liu@amd.com>
Date: Fri, 7 Aug 2020 17:50:22 +0000 (-0400)
Subject: Make clang HIP headers compatible with C++98
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=ac3e720dc1af9555c5591017bd2b6c7f6641a34d;p=platform%2Fupstream%2Fllvm.git

Make clang HIP headers compatible with C++98

Automation to detect compiler features, such as CMake's target_compile_features,
would attempt to detect compiler features by explicitly using langugage flags.
This change ensures that the HIP headers would still work with C++98.

Patch by Siu Chi Chan

Differential Revision: https://reviews.llvm.org/D85471

Change-Id: I304e964b18a525b0fde55efd841da74b6c4dc8ed
---

diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h
index 2cf9cc7..de2f82c 100644
--- a/clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -318,7 +318,7 @@ __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
 __device__ inline __2f16
 __llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL.
 {
-  return (__2f16){__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y)};
+  return (__2f16)(__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y));
 }
 __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index f9ca9bf..0c27ef6 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -14,6 +14,7 @@
 #include <limits.h>
 #include <limits>
 #include <stdint.h>
+#include <assert.h>
 
 #pragma push_macro("__DEVICE__")
 #pragma push_macro("__RETURN_TYPE")
@@ -22,6 +23,34 @@
 #define __DEVICE__ static __device__
 #define __RETURN_TYPE bool
 
+#if defined (__cplusplus) && __cplusplus < 201103L
+//emulate static_assert on type sizes
+template<bool>
+struct __compare_result{};
+template<>
+struct __compare_result<true> {
+  static const bool valid;
+};
+
+__DEVICE__
+inline void __suppress_unused_warning(bool b) {};
+template<unsigned int S, unsigned int T>
+__DEVICE__
+inline void __static_assert_equal_size() {
+    __suppress_unused_warning(__compare_result<S==T>::valid);
+}
+
+#define __static_assert_type_size_equal(A, B) \
+  __static_assert_equal_size<A,B>()
+
+#else
+
+#define __static_assert_type_size_equal(A,B) \
+  static_assert((A) == (B), "")
+
+#endif
+
+
 __DEVICE__
 inline uint64_t __make_mantissa_base8(const char *__tagp) {
   uint64_t __r = 0;
@@ -252,9 +281,8 @@ inline float nanf(const char *__tagp) {
       uint32_t exponent : 8;
       uint32_t sign : 1;
     } bits;
-
-    static_assert(sizeof(float) == sizeof(struct ieee_float), "");
   } __tmp;
+  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
 
   __tmp.bits.sign = 0u;
   __tmp.bits.exponent = ~0u;
@@ -716,8 +744,8 @@ inline double nan(const char *__tagp) {
       uint32_t exponent : 11;
       uint32_t sign : 1;
     } bits;
-    static_assert(sizeof(double) == sizeof(struct ieee_double), "");
   } __tmp;
+  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
 
   __tmp.bits.sign = 0u;
   __tmp.bits.exponent = ~0u;
@@ -726,7 +754,7 @@ inline double nan(const char *__tagp) {
 
   return __tmp.val;
 #else
-  static_assert(sizeof(uint64_t) == sizeof(double));
+  __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double));
   uint64_t val = __make_mantissa(__tagp);
   val |= 0xFFF << 51;
   return *reinterpret_cast<double *>(&val);
diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
index addae56..007b5f1 100644
--- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -28,6 +28,10 @@
 #define __shared__ __attribute__((shared))
 #define __constant__ __attribute__((constant))
 
+#if !defined(__cplusplus) || __cplusplus < 201103L
+  #define nullptr NULL;
+#endif
+
 #if __HIP_ENABLE_DEVICE_MALLOC__
 extern "C" __device__ void *__hip_malloc(size_t __size);
 extern "C" __device__ void *__hip_free(void *__ptr);