diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h index 18ad566121..23ad461065 100644 --- a/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -1126,8 +1126,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __half2_raw{__builtin_convertvector(-r, _Float16_2)}; } inline __device__ @@ -1135,8 +1134,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data != static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1144,8 +1142,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data <= static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1153,8 +1150,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data >= static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1162,8 +1158,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1171,8 +1166,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__