From 2d617bba3bbe148d3ba7066d2dc2965e4029abac Mon Sep 17 00:00:00 2001
From: kpyzhov <konstantin.pyzhov@gmail.com>
Date: Wed, 15 Jan 2020 21:17:08 -0500
Subject: [PATCH] Don't use accelerated vector element access for hip-clang.
 (#1796)

[ROCm/clr commit: 1d6a2fb0646f83064435bfbf5ea8536597c81d7d]
---
 .../include/hip/hcc_detail/hip_vector_types.h | 30 ++++++++++++++-----
 1 file changed, 22 insertions(+), 8 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
index 453e92ebfc..9ef8e32516 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
@@ -105,14 +105,6 @@ THE SOFTWARE.
             // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor
             Vector data;
 
-            __host__ __device__
-            Scalar_accessor() = default;
-
-            __host__ __device__
-            Scalar_accessor(T x) {
-                data[idx] = x;
-            }
-
             __host__ __device__
             operator T() const noexcept { return data[idx]; }
             __host__ __device__
@@ -302,7 +294,13 @@ THE SOFTWARE.
 
         union {
             Native_vec_ data;
+#if __HIP_CLANG_ONLY__
+            struct {
+                T x;
+            };
+#else
             hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
+#endif
         };
 
         using value_type = T;
@@ -325,8 +323,15 @@ THE SOFTWARE.
 
         union {
             Native_vec_ data;
+#if __HIP_CLANG_ONLY__
+            struct {
+                T x;
+                T y;
+            };
+#else
             hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
             hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
+#endif
         };
 
         using value_type = T;
@@ -506,10 +511,19 @@ THE SOFTWARE.
 
         union {
             Native_vec_ data;
+#if __HIP_CLANG_ONLY__
+            struct {
+                T x;
+                T y;
+                T z;
+                T w;
+            };
+#else
             hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
             hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
             hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
             hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
+#endif            
         };
 
         using value_type = T;