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

 __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';
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__

clang/lib/Headers/__clang_hip_runtime_wrapper.h

 #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