(Maybe) Match alignment between Clang and GCC. (#1789)

Should fix #1740 and the related internal bug.

hipamd/.gitignore

@@ -15,4 +15,4 @@ include/hip/hip_version.h
 tags
 samples/1_Utils/hipInfo/hipInfo
 samples/1_Utils/hipBusBandwidth/hipBusBandwidth
-samples/1_Utils/hipDispatchLatency/hipDispatchLatency
+samples/1_Utils/hipDispatchLatency/hipDispatchLatency

hipamd/include/hip/hcc_detail/hip_vector_types.h

@@ -34,11 +34,11 @@ THE SOFTWARE.
 
 #include "hip/hcc_detail/host_defines.h"
 
-#if !defined(_MSC_VER)|| __clang__
+#if !defined(_MSC_VER) || __clang__
     #if __has_attribute(ext_vector_type)
-        #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
+        #define __NATIVE_VECTOR__(n, T) T __attribute__((ext_vector_type(n)))
     #else
-        #define __NATIVE_VECTOR__(n, ...) [n]
+        #define __NATIVE_VECTOR__(n, T) T[n]
     #endif
 
 #if defined(__cplusplus)
@@ -49,16 +49,16 @@ THE SOFTWARE.
     namespace hip_impl {
         template<typename, typename, unsigned int> struct Scalar_accessor;
     } // Namespace hip_impl.
-    
+
     namespace std {
         template<typename T, typename U, unsigned int n>
         struct is_integral<hip_impl::Scalar_accessor<T, U, n>>
             : is_integral<T> {};
         template<typename T, typename U, unsigned int n>
-        struct is_floating_point<hip_impl::Scalar_accessor<T, U, n>> 
+        struct is_floating_point<hip_impl::Scalar_accessor<T, U, n>>
             : is_floating_point<T> {};
     } // Namespace std.
-    
+
     namespace hip_impl {
         template<typename T, typename Vector, unsigned int idx>
         struct Scalar_accessor {
@@ -146,7 +146,7 @@ THE SOFTWARE.
 
             __host__ __device__
             Address operator&() const noexcept { return Address{this}; }
-            
+
             __host__ __device__
             Scalar_accessor& operator=(const Scalar_accessor& x) noexcept {
                 data[idx] = x.data[idx];
@@ -284,13 +284,20 @@ THE SOFTWARE.
                 return *this;
             }
         };
+
+        inline
+        constexpr
+        unsigned int next_pot(unsigned int x) {
+            // Precondition: x > 1.
+	        return 1u << (32u - __builtin_clz(x - 1u));
+        }
     } // Namespace hip_impl.
 
     template<typename T, unsigned int n> struct HIP_vector_base;
 
     template<typename T>
     struct HIP_vector_base<T, 1> {
-        using Native_vec_ = T __NATIVE_VECTOR__(1, T);
+        using Native_vec_ = __NATIVE_VECTOR__(1, T);
 
         union {
             Native_vec_ data;
@@ -304,7 +311,7 @@ THE SOFTWARE.
         };
 
         using value_type = T;
-        
+
         __host__ __device__
         HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
             #if __has_attribute(ext_vector_type)
@@ -319,9 +326,13 @@ THE SOFTWARE.
 
     template<typename T>
     struct HIP_vector_base<T, 2> {
-        using Native_vec_ = T __NATIVE_VECTOR__(2, T);
+        using Native_vec_ = __NATIVE_VECTOR__(2, T);
 
-        union {
+        union
+        #if !__has_attribute(ext_vector_type)
+            alignas(hip_impl::next_pot(2 * sizeof(T)))
+        #endif
+        {
             Native_vec_ data;
 #if __HIP_CLANG_ONLY__
             struct {
@@ -335,7 +346,7 @@ THE SOFTWARE.
         };
 
         using value_type = T;
-        
+
         __host__ __device__
         HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
             #if __has_attribute(ext_vector_type)
@@ -507,9 +518,13 @@ THE SOFTWARE.
 
     template<typename T>
     struct HIP_vector_base<T, 4> {
-        using Native_vec_ = T __NATIVE_VECTOR__(4, T);
+        using Native_vec_ = __NATIVE_VECTOR__(4, T);
 
-        union {
+        union
+        #if !__has_attribute(ext_vector_type)
+            alignas(hip_impl::next_pot(4 * sizeof(T)))
+        #endif
+        {
             Native_vec_ data;
 #if __HIP_CLANG_ONLY__
             struct {
@@ -527,7 +542,7 @@ THE SOFTWARE.
         };
 
         using value_type = T;
-        
+
         __host__ __device__
         HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
             #if __has_attribute(ext_vector_type)