diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h --- a/clang/lib/Headers/__clang_hip_libdevice_declares.h +++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -318,7 +318,7 @@ __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 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -14,6 +14,7 @@ #include #include #include +#include #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 +struct __compare_result{}; +template<> +struct __compare_result { + static const bool valid; +}; + +__DEVICE__ +inline void __suppress_unused_warning(bool b) {}; +template +__DEVICE__ +inline void __static_assert_equal_size() { + __suppress_unused_warning(__compare_result::valid); +} + +#define __static_assert_type_size_equal(A, B) \ + __static_assert_equal_size() + +#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 @@ 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 @@ 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 @@ 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(&val); diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h --- 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);