Improve scalar access into vector types. (#1531)

The improvement is based on the ideas here: https://t0rakka.silvrback.com/simd-scalar-accessor. It yields significantly better ISA when the base's .xyzw members are used.

[ROCm/clr commit: 84d5b399f6]

projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h

@@ -47,6 +47,95 @@ THE SOFTWARE.
 #if defined(__cplusplus)
     #include <type_traits>
 
+    namespace hip_impl {
+        template<typename T, typename Vector, unsigned int idx>
+        struct Scalar_accessor {
+            // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor
+            Vector data;
+
+            __host__ __device__
+            operator T() const noexcept { return data[idx]; }
+
+            __host__ __device__
+            Scalar_accessor& operator=(T x) noexcept {
+                data[idx] = x;
+
+                return *this;
+            }
+
+            __host__ __device__
+            Scalar_accessor& operator+=(T x) noexcept {
+                data[idx] += x;
+                return *this;
+            }
+            __host__ __device__
+            Scalar_accessor& operator-=(T x) noexcept {
+                data[idx] -= x;
+                return *this;
+            }
+
+            __host__ __device__
+            Scalar_accessor& operator*=(T x) noexcept {
+                data[idx] *= x;
+                return *this;
+            }
+            __host__ __device__
+            Scalar_accessor& operator/=(T x) noexcept {
+                data[idx] /= x;
+                return *this;
+            }
+            template<
+                typename U = T,
+                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            __host__ __device__
+            Scalar_accessor& operator%=(T x) noexcept {
+                data[idx] %= x;
+                return *this;
+            }
+
+            template<
+                typename U = T,
+                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            __host__ __device__
+            Scalar_accessor& operator>>=(T x) noexcept {
+                data[idx] >>= x;
+                return *this;
+            }
+            template<
+                typename U = T,
+                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            __host__ __device__
+            Scalar_accessor& operator<<=(T x) noexcept {
+                data[idx] <<= x;
+                return *this;
+            }
+            template<
+                typename U = T,
+                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            __host__ __device__
+            Scalar_accessor& operator&=(T x) noexcept {
+                data[idx] &= x;
+                return *this;
+            }
+            template<
+                typename U = T,
+                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            __host__ __device__
+            Scalar_accessor& operator|=(T x) noexcept {
+                data[idx] |= x;
+                return *this;
+            }
+            template<
+                typename U = T,
+                typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
+            __host__ __device__
+            Scalar_accessor& operator^=(T x) noexcept {
+                data[idx] ^= x;
+                return *this;
+            }
+        };
+    } // Namespace hip_impl.
+
     template<typename T, unsigned int n> struct HIP_vector_base;
 
     template<typename T>
@@ -55,9 +144,7 @@ THE SOFTWARE.
 
         union {
             Native_vec_ data;
-            struct {
-                T x;
-            };
+            hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
         };
     };
 
@@ -67,10 +154,8 @@ THE SOFTWARE.
 
         union {
             Native_vec_ data;
-            struct {
-                T x;
-                T y;
-            };
+            hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
+            hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
         };
     };
 
@@ -238,12 +323,10 @@ THE SOFTWARE.
 
         union {
             Native_vec_ data;
-            struct {
-                T x;
-                T y;
-                T z;
-                T w;
-            };
+            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;
         };
     };