diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index f81cbecf41..7b6ace6113 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -66,11 +66,11 @@ THE SOFTWARE. union { static_assert(sizeof(_Float16_2) == sizeof(unsigned short[2]), ""); - _Float16_2 data; struct { unsigned short x; unsigned short y; }; + _Float16_2 data; }; }; @@ -353,11 +353,11 @@ THE SOFTWARE. static_assert( sizeof(_Float16_2) == sizeof(unsigned short[2]), ""); - _Float16_2 data; struct { unsigned short x; unsigned short y; }; + _Float16_2 data; }; // CREATORS @@ -443,7 +443,11 @@ THE SOFTWARE. __HOST_DEVICE__ operator decltype(data)() const { return data; } __HOST_DEVICE__ - operator __half2_raw() const { return __half2_raw{data}; } + operator __half2_raw() const { + __half2_raw r; + r.data = data; + return r; + } // ACCESSORS - DEVICE ONLY #if !defined(__HIP_NO_HALF_OPERATORS__) @@ -587,7 +591,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __high2half2(__half2 x) { - return __half2_raw{ + return __half2{ _Float16_2{ static_cast<__half2_raw>(x).data.y, static_cast<__half2_raw>(x).data.y}}; @@ -597,7 +601,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __lows2half2(__half2 x, __half2 y) { - return __half2_raw{ + return __half2{ _Float16_2{ static_cast<__half2_raw>(x).data.x, static_cast<__half2_raw>(y).data.x}}; @@ -607,7 +611,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __highs2half2(__half2 x, __half2 y) { - return __half2_raw{ + return __half2{ _Float16_2{ static_cast<__half2_raw>(x).data.y, static_cast<__half2_raw>(y).data.y}}; @@ -617,7 +621,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __lowhigh2highlow(__half2 x) { - return __half2_raw{ + return __half2{ _Float16_2{ static_cast<__half2_raw>(x).data.y, static_cast<__half2_raw>(x).data.x}}; @@ -710,7 +714,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __float2half2_rn(float x) { - return __half2_raw{ + return __half2{ _Float16_2{ static_cast<_Float16>(x), static_cast<_Float16>(x)}}; } @@ -718,7 +722,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __floats2half2_rn(float x, float y) { - return __half2_raw{_Float16_2{ + return __half2{_Float16_2{ static_cast<_Float16>(x), static_cast<_Float16>(y)}}; } inline @@ -1406,7 +1410,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __hadd2(__half2 x, __half2 y) { - return __half2_raw{ + return __half2{ static_cast<__half2_raw>(x).data + static_cast<__half2_raw>(y).data}; } @@ -1414,14 +1418,14 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __habs2(__half2 x) { - return __half2_raw{ + return __half2{ __ocml_fabs_2f16(static_cast<__half2_raw>(x).data)}; } inline __HOST_DEVICE__ __half2 __hsub2(__half2 x, __half2 y) { - return __half2_raw{ + return __half2{ static_cast<__half2_raw>(x).data - static_cast<__half2_raw>(y).data}; } @@ -1429,7 +1433,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __hmul2(__half2 x, __half2 y) { - return __half2_raw{ + return __half2{ static_cast<__half2_raw>(x).data * static_cast<__half2_raw>(y).data}; } @@ -1464,7 +1468,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __hfma2(__half2 x, __half2 y, __half2 z) { - return __half2_raw{__ocml_fma_2f16(x, y, z)}; + return __half2{__ocml_fma_2f16(x, y, z)}; } inline __HOST_DEVICE__ @@ -1479,7 +1483,7 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 __h2div(__half2 x, __half2 y) { - return __half2_raw{ + return __half2{ static_cast<__half2_raw>(x).data / static_cast<__half2_raw>(y).data}; } @@ -1622,61 +1626,61 @@ THE SOFTWARE. __HOST_DEVICE__ __half2 h2trunc(__half2 x) { - return __half2_raw{__ocml_trunc_2f16(x)}; + return __half2{__ocml_trunc_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2ceil(__half2 x) { - return __half2_raw{__ocml_ceil_2f16(x)}; + return __half2{__ocml_ceil_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2floor(__half2 x) { - return __half2_raw{__ocml_floor_2f16(x)}; + return __half2{__ocml_floor_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2rint(__half2 x) { - return __half2_raw{__ocml_rint_2f16(x)}; + return __half2{__ocml_rint_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2sin(__half2 x) { - return __half2_raw{__ocml_sin_2f16(x)}; + return __half2{__ocml_sin_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2cos(__half2 x) { - return __half2_raw{__ocml_cos_2f16(x)}; + return __half2{__ocml_cos_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2exp(__half2 x) { - return __half2_raw{__ocml_exp_2f16(x)}; + return __half2{__ocml_exp_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2exp2(__half2 x) { - return __half2_raw{__ocml_exp2_2f16(x)}; + return __half2{__ocml_exp2_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2exp10(__half2 x) { - return __half2_raw{__ocml_exp10_2f16(x)}; + return __half2{__ocml_exp10_2f16(x)}; } inline __HOST_DEVICE__ __half2 h2log2(__half2 x) { - return __half2_raw{__ocml_log2_2f16(x)}; + return __half2{__ocml_log2_2f16(x)}; } inline __HOST_DEVICE__ @@ -1701,7 +1705,7 @@ THE SOFTWARE. __half2 __hisinf2(__half2 x) { auto r = __ocml_isinf_2f16(x); - return __half2_raw{_Float16_2{ + return __half2{_Float16_2{ static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; } inline @@ -1709,14 +1713,14 @@ THE SOFTWARE. __half2 __hisnan2(__half2 x) { auto r = __ocml_isnan_2f16(x); - return __half2_raw{_Float16_2{ + return __half2{_Float16_2{ static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; } inline __HOST_DEVICE__ __half2 __hneg2(__half2 x) { - return __half2_raw{-static_cast<__half2_raw>(x).data}; + return __half2{-static_cast<__half2_raw>(x).data}; } } // Anonymous namespace.