From e5bd00d06b4bdcbc6912a78ef913cef7ae51ed95 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 6 Nov 2019 02:43:04 +0200 Subject: [PATCH 1/2] __half2 should walk like CUDA and talk like CUDA --- include/hip/hcc_detail/hip_fp16.h | 18 ++++++------------ 1 file changed, 6 insertions(+), 12 deletions(-) diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index 18ad566121..23ad461065 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/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__ From b9faa9f8ae72962cd0ea48bbccf65c81134f34ea Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 6 Nov 2019 02:46:21 +0200 Subject: [PATCH 2/2] Remove leftover noise. --- include/hip/hcc_detail/hip_fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index 23ad461065..52abc1a004 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/include/hip/hcc_detail/hip_fp16.h @@ -1126,7 +1126,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data; - return __half2_raw{__builtin_convertvector(-r, _Float16_2)}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__