__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);
#include <limits.h>
#include <limits>
#include <stdint.h>
+#include <assert.h>
#pragma push_macro("__DEVICE__")
#pragma push_macro("__RETURN_TYPE")
#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;
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;
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;
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);