Don't use accelerated vector element access for hip-clang. (#1796)

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;