This fixes some outright quaint choices made when implementing HIP's bitwise conversion functions, by using simple reinterpret_casts, as is idiomatic. These functions are supposed to be re-entrant, correct and efficient. Sadly, they were neither: they hid a massive race condition against a value stored in global memory, which means that they were also unreasonably slow if they ever managed to be correct, and relied on union based type punning which is in a grey area of the standard. It is difficult to ascertain what may have been the reason for coming up with this quirky solution.

[ROCm/hip commit: 153878e368]
This commit is contained in:
Alex Voicu
2017-11-17 16:00:28 +00:00
parent 0fd5239b14
commit d9ecde1b60
+12 -41
View File
@@ -23,27 +23,6 @@ THE SOFTWARE.
#include <hc_math.hpp> #include <hc_math.hpp>
#include "device_util.h" #include "device_util.h"
struct holder64Bit{
union{
double d;
unsigned long int uli;
signed long int sli;
signed int si[2];
unsigned int ui[2];
};
} __attribute__((aligned(8)));
struct holder32Bit {
union {
float f;
unsigned int ui;
signed int si;
};
} __attribute__((aligned(4)));
__device__ struct holder64Bit hold64;
__device__ struct holder32Bit hold32;
__device__ float __double2float_rd(double x) __device__ float __double2float_rd(double x)
{ {
return (double)x; return (double)x;
@@ -64,13 +43,11 @@ __device__ float __double2float_rz(double x)
__device__ int __double2hiint(double x) __device__ int __double2hiint(double x)
{ {
hold64.d = x; return reinterpret_cast<int(&)[2]>(x)[1];
return hold64.si[1];
} }
__device__ int __double2loint(double x) __device__ int __double2loint(double x)
{ {
hold64.d = x; return reinterpret_cast<int(&)[2]>(x)[0];
return hold64.si[0];
} }
@@ -145,8 +122,7 @@ __device__ unsigned long long int __double2ull_rz(double x)
__device__ long long int __double_as_longlong(double x) __device__ long long int __double_as_longlong(double x)
{ {
hold64.d = x; return reinterpret_cast<long long&>(x);
return hold64.sli;
} }
__device__ int __float2int_rd(float x) __device__ int __float2int_rd(float x)
@@ -219,19 +195,17 @@ __device__ unsigned long long int __float2ull_rz(float x)
__device__ int __float_as_int(float x) __device__ int __float_as_int(float x)
{ {
hold32.f = x; return reinterpret_cast<int&>(x);
return hold32.si;
} }
__device__ unsigned int __float_as_uint(float x) __device__ unsigned int __float_as_uint(float x)
{ {
hold32.f = x; return reinterpret_cast<unsigned int&>(x);
return hold32.ui;
} }
__device__ double __hiloint2double(int hi, int lo) __device__ double __hiloint2double(int hi, int lo)
{ { // TODO: this matches the original in not considering endianness, is that
hold64.si[1] = hi; // correct though?
hold64.si[0] = lo; int tmp[] = {lo, hi};
return hold64.d; return reinterpret_cast<double&>(tmp);
} }
__device__ double __int2double_rn(int x) __device__ double __int2double_rn(int x)
{ {
@@ -257,8 +231,7 @@ __device__ float __int2float_rz(int x)
__device__ float __int_as_float(int x) __device__ float __int_as_float(int x)
{ {
hold32.si = x; return reinterpret_cast<float&>(x);
return hold32.f;
} }
__device__ double __ll2double_rd(long long int x) __device__ double __ll2double_rd(long long int x)
@@ -297,8 +270,7 @@ __device__ float __ll2float_rz(long long int x)
__device__ double __longlong_as_double(long long int x) __device__ double __longlong_as_double(long long int x)
{ {
hold64.sli = x; return reinterpret_cast<double&>(x);
return hold64.d;
} }
__device__ double __uint2double_rn(int x) __device__ double __uint2double_rn(int x)
@@ -325,8 +297,7 @@ __device__ float __uint2float_rz(unsigned int x)
__device__ float __uint_as_float(unsigned int x) __device__ float __uint_as_float(unsigned int x)
{ {
hold32.ui = x; return reinterpret_cast<float&>(x);
return hold32.f;
} }
__device__ double __ull2double_rd(unsigned long long int x) __device__ double __ull2double_rd(unsigned long long int x)