Make clang HIP headers compatible with C++98

[scchan]

Summary

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.

Change-Id: I381f76fdaf42196d9272efeaa5a3a159edca0ab1

[yaxunl]

LGTM. thanks

[arsenm]

I don't understand why this is necessary. What does cmake have to do with it? Is HIP just not defaulting to the correct default C++ standard?

[arsenm]

Can you either

1. Add test coverage
2. Delete this entirely

clang/lib/Headers/__clang_hip_libdevice_declares.h

 /*===---- __clang_hip_libdevice_declares.h - HIP device library decls -------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_HIP_LIBDEVICE_DECLARES_H__
 #define __CLANG_HIP_LIBDEVICE_DECLARES_H__
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 // BEGIN FLOAT
 __device__ __attribute__((const)) float __ocml_acos_f32(float);
 __device__ __attribute__((pure)) float __ocml_acosh_f32(float);
 __device__ __attribute__((const)) float __ocml_asin_f32(float);
 __device__ __attribute__((pure)) float __ocml_asinh_f32(float);
 __device__ __attribute__((const)) float __ocml_atan2_f32(float, float);
 __device__ __attribute__((const)) float __ocml_atan_f32(float);
 __device__ __attribute__((pure)) float __ocml_atanh_f32(float);
 __device__ __attribute__((pure)) float __ocml_cbrt_f32(float);
 __device__ __attribute__((const)) float __ocml_ceil_f32(float);
 __device__ __attribute__((const)) __device__ float __ocml_copysign_f32(float,
                                                                        float);
 __device__ float __ocml_cos_f32(float);
 __device__ float __ocml_native_cos_f32(float);
 __device__ __attribute__((pure)) __device__ float __ocml_cosh_f32(float);
 __device__ float __ocml_cospi_f32(float);
 __device__ float __ocml_i0_f32(float);
 __device__ float __ocml_i1_f32(float);
 __device__ __attribute__((pure)) float __ocml_erfc_f32(float);
 __device__ __attribute__((pure)) float __ocml_erfcinv_f32(float);
 __device__ __attribute__((pure)) float __ocml_erfcx_f32(float);
 __device__ __attribute__((pure)) float __ocml_erf_f32(float);
 __device__ __attribute__((pure)) float __ocml_erfinv_f32(float);
 __device__ __attribute__((pure)) float __ocml_exp10_f32(float);
 __device__ __attribute__((pure)) float __ocml_native_exp10_f32(float);
 __device__ __attribute__((pure)) float __ocml_exp2_f32(float);
 __device__ __attribute__((pure)) float __ocml_exp_f32(float);

 __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
 __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
 __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);
 __device__ __2f16 __ocml_sin_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
 
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
 #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__

clang/lib/Headers/__clang_hip_math.h

 /*===---- __clang_hip_math.h - HIP math decls -------------------------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_HIP_MATH_H__
 #define __CLANG_HIP_MATH_H__
 
 #include <algorithm>
 #include <limits.h>
 #include <limits>
 #include <stdint.h>
+#include <assert.h>
 
 #pragma push_macro("__DEVICE__")
 #pragma push_macro("__RETURN_TYPE")
 
 // to be consistent with __clang_cuda_math_forward_declares
 #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;
   while (__tagp) {
     char __tmp = *__tagp;
 
     if (__tmp >= '0' && __tmp <= '7')
       __r = (__r * 8u) + __tmp - '0';
     else
       return 0;
inline float nanf(const char *__tagp) {
   union {
     float val;
     struct ieee_float {
       uint32_t mantissa : 22;
       uint32_t quiet : 1;
       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;
   __tmp.bits.quiet = 1u;
   __tmp.bits.mantissa = __make_mantissa(__tagp);
 
   return __tmp.val;
 }
#if !_WIN32
   union {
     double val;
     struct ieee_double {
       uint64_t mantissa : 51;
       uint32_t quiet : 1;
       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;
   __tmp.bits.quiet = 1u;
   __tmp.bits.mantissa = __make_mantissa(__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);
 #endif
 }
 __DEVICE__
 inline double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); }
 __DEVICE__

 #pragma push_macro("__DEF_FUNI")
 #pragma push_macro("__DEF_FLOAT_FUN2I")
 #pragma push_macro("__HIP_OVERLOAD1")
 #pragma push_macro("__HIP_OVERLOAD2")
 
 // __hip_enable_if::type is a type function which returns __T if __B is true.
 template <bool __B, class __T = void> struct __hip_enable_if {};
 
 template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
 
 // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
 // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
 // floor(double).
 #define __HIP_OVERLOAD1(__retty, __fn)                                         \
   template <typename __T>                                                      \
   __DEVICE__ typename __hip_enable_if<std::numeric_limits<__T>::is_integer,    \
                                       __retty>::type                           \
#define __DEF_FUNI(__retty, __func)                                            \
   __HIP_OVERLOAD1(__retty, __func)
 
 // define cmath functions with two float arguments.
 #define __DEF_FUN2(__retty, __func)                                            \
   __DEVICE__                                                                   \
   inline float __func(float __x, float __y) { return __func##f(__x, __y); }    \
   __HIP_OVERLOAD2(__retty, __func)
 
 __DEF_FUN1(double, acos)
 __DEF_FUN1(double, acosh)
 __DEF_FUN1(double, asin)
 __DEF_FUN1(double, asinh)
 __DEF_FUN1(double, atan)
 __DEF_FUN2(double, atan2);
 __DEF_FUN1(double, atanh)
 __DEF_FUN1(double, cbrt)
 __DEF_FUN1(double, ceil)
 __DEF_FUN2(double, copysign);
 __DEF_FUN1(double, cos)

clang/lib/Headers/__clang_hip_runtime_wrapper.h

 /*===---- __clang_hip_runtime_wrapper.h - HIP runtime support ---------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */

 #include <stdlib.h>
 
 #define __host__ __attribute__((host))
 #define __device__ __attribute__((device))
 #define __global__ __attribute__((global))
 #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);
 static inline __device__ void *malloc(size_t __size) {
   return __hip_malloc(__size);
 }
 static inline __device__ void *free(void *__ptr) { return __hip_free(__ptr); }
 #else