[SYCL] vec abi unification and trivially copyable (#9492)

adapted from #6816   

sycl::vec now uses std::array on host in all cases, but device access is
via `ext_vector_type` so as to not lose performance. Alignment is set
the same as the size to a maximum of 64 bytes, which aligns with an
upcoming spec change proposal. In the cases where the size is greater
than 64 we use std::array on the device. We no longer need to emit a
pragma on Windows and we can use `alignas()` on all platforms.

Author: cperkinsintel

sycl/include/sycl/detail/vector_traits.hpp

@@ -17,12 +17,20 @@ inline namespace _V1 {
 namespace detail {
 
 // 4.10.2.6 Memory layout and alignment
-template <typename T, int N>
+// due to MSVC the maximum alignment for sycl::vec is 64 and this proposed
+// change is being brought to the spec committee.
+constexpr size_t MaxVecAlignment = 64;
+template <typename T, size_t N>
 struct vector_alignment_impl
-    : std::conditional_t<N == 3, std::integral_constant<int, sizeof(T) * 4>,
-                         std::integral_constant<int, sizeof(T) * N>> {};
+    : std::conditional_t<
+          N == 3,
+          std::integral_constant<size_t,
+                                 (std::min)(sizeof(T) * 4, MaxVecAlignment)>,
+          std::integral_constant<size_t,
+                                 (std::min)(sizeof(T) * N, MaxVecAlignment)>> {
+};
 
-template <typename T, int N>
+template <typename T, size_t N>
 struct vector_alignment
     : vector_alignment_impl<std::remove_cv_t<std::remove_reference_t<T>>, N> {};
 } // namespace detail

sycl/include/sycl/half_type.hpp

@@ -253,28 +253,12 @@ using BIsRepresentationT = half;
 // for vec because they are actually defined as an integer type under the
 // hood. As a result half values will be converted to the integer and passed
 // as a kernel argument which is expected to be floating point number.
-template <int NumElements> struct half_vec {
-  alignas(
-      vector_alignment<StorageT, NumElements>::value) StorageT s[NumElements];
 
-  __SYCL_CONSTEXPR_HALF half_vec() : s{0.0f} { initialize_data(); }
-  template <typename... Ts,
-            typename = std::enable_if_t<(sizeof...(Ts) == NumElements) &&
-                                        (std::is_same_v<half, Ts> && ...)>>
-  __SYCL_CONSTEXPR_HALF half_vec(const Ts &...hs) : s{hs...} {}
-
-  constexpr void initialize_data() {
-    for (size_t i = 0; i < NumElements; ++i) {
-      s[i] = StorageT(0.0f);
-    }
-  }
-};
-
-using Vec2StorageT = half_vec<2>;
-using Vec3StorageT = half_vec<3>;
-using Vec4StorageT = half_vec<4>;
-using Vec8StorageT = half_vec<8>;
-using Vec16StorageT = half_vec<16>;
+using Vec2StorageT = std::array<StorageT, 2>;
+using Vec3StorageT = std::array<StorageT, 3>;
+using Vec4StorageT = std::array<StorageT, 4>;
+using Vec8StorageT = std::array<StorageT, 8>;
+using Vec16StorageT = std::array<StorageT, 16>;
 #endif
 
 #ifndef __SYCL_DEVICE_ONLY__