This actually (tries) to do the right thing all the way, by using memcpy for bitcasting, and not rely on undefined behaviour of a different flavour as a substitute for the original undefined behaviour. Note that the compiler will (should) optimise down to the same emitted code, since this is a pattern it understands.

hipamd/src/device_functions.cpp

@@ -43,11 +43,21 @@ __device__ float __double2float_rz(double x)
 
 __device__ int __double2hiint(double x)
 {
-  return reinterpret_cast<int(&)[2]>(x)[1];
+  static_assert(sizeof(double) == 2 * sizeof(int), "");
+
+  int tmp[2];
+  __builtin_memcpy(tmp, &x, sizeof(tmp));
+
+  return tmp[1];
 }
 __device__ int __double2loint(double x)
 {
-  return reinterpret_cast<int(&)[2]>(x)[0];
+  static_assert(sizeof(double) == 2 * sizeof(int), "");
+
+  int tmp[2];
+  __builtin_memcpy(tmp, &x, sizeof(tmp));
+
+  return tmp[0];
 }
 
 
@@ -122,7 +132,12 @@ __device__ unsigned long long int __double2ull_rz(double x)
 
 __device__ long long int __double_as_longlong(double x)
 {
-  return reinterpret_cast<long long&>(x);
+  static_assert(sizeof(long long) == sizeof(double), "");
+
+  long long tmp;
+  __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+  return tmp;
 }
 
 __device__ int __float2int_rd(float x)
@@ -195,17 +210,32 @@ __device__ unsigned long long int __float2ull_rz(float x)
 
 __device__ int __float_as_int(float x)
 {
-  return reinterpret_cast<int&>(x);
+  static_assert(sizeof(int) == sizeof(float), "");
+
+  int tmp;
+  __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+  return tmp;
 }
 __device__ unsigned int __float_as_uint(float x)
 {
-  return reinterpret_cast<unsigned int&>(x);
+  static_assert(sizeof(unsigned int) == sizeof(float), "");
+
+  unsigned int tmp;
+  __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+  return tmp;
 }
-__device__ double __hiloint2double(int hi, int lo)
-{ // TODO: this matches the original in not considering endianness, is that
-  //       correct though?
-  int tmp[] = {lo, hi};
-  return reinterpret_cast<double&>(tmp);
+__device__ double __hiloint2double(int32_t hi, int32_t lo)
+{
+  static_assert(sizeof(double) == sizeof(uint64_t), "");
+
+  uint64_t tmp0 =
+    (static_cast<uint64_t>(hi) << 32ull) | static_cast<uint32_t>(lo);
+  double tmp1;
+  __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
+
+  return tmp1;
 }
 __device__ double __int2double_rn(int x)
 {
@@ -231,7 +261,12 @@ __device__ float __int2float_rz(int x)
 
 __device__ float __int_as_float(int x)
 {
-  return reinterpret_cast<float&>(x);
+  static_assert(sizeof(float) == sizeof(int), "");
+
+  float tmp;
+  __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+  return tmp;
 }
 
 __device__ double __ll2double_rd(long long int x)
@@ -270,7 +305,12 @@ __device__ float __ll2float_rz(long long int x)
 
 __device__ double __longlong_as_double(long long int x)
 {
-  return reinterpret_cast<double&>(x);
+  static_assert(sizeof(double) == sizeof(long long), "");
+
+  double tmp;
+  __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+  return x;
 }
 
 __device__ double __uint2double_rn(int x)
@@ -297,7 +337,12 @@ __device__ float __uint2float_rz(unsigned int x)
 
 __device__ float __uint_as_float(unsigned int x)
 {
-  return reinterpret_cast<float&>(x);
+  static_assert(sizeof(float) == sizeof(unsigned int), "");
+
+  float tmp;
+  __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+  return tmp;
 }
 
 __device__ double __ull2double_rd(unsigned long long int x)