From 55fd1363e25f33d05158ed6895fe48befdedfa1f Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Wed, 6 Nov 2019 02:43:04 +0200
Subject: [PATCH] __half2 should walk like CUDA and talk like CUDA

---
 hipamd/include/hip/hcc_detail/hip_fp16.h | 18 ++++++------------
 1 file changed, 6 insertions(+), 12 deletions(-)

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__