[SYCL] vec abi changes behind -fpreview-breaking-changes flag

sycl/include/sycl/buffer.hpp

@@ -168,8 +168,13 @@ template <typename T, int dimensions = 1,
               typename std::enable_if_t<(dimensions > 0) && (dimensions <= 3)>>
 class buffer : public detail::buffer_plain,
                public detail::OwnerLessBase<buffer<T, dimensions, AllocatorT>> {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
   static_assert(is_device_copyable_v<T>,
                 "Underlying type of a buffer must be device copyable!");
+#else
+  static_assert(!std::is_same_v<T, std::string>,
+                "'std::string' is not a device copyable type");
+#endif
 
 public:
   using value_type = T;

sycl/include/sycl/detail/vector_traits.hpp

@@ -16,6 +16,7 @@ inline namespace _V1 {
 namespace detail {
 
 // 4.10.2.6 Memory layout and alignment
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
 // 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;
@@ -28,6 +29,12 @@ struct vector_alignment_impl
           std::integral_constant<size_t,
                                  (std::min)(sizeof(T) * N, MaxVecAlignment)>> {
 };
+#else
+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>> {};
+#endif
 
 template <typename T, size_t N>
 struct vector_alignment

sycl/include/sycl/half_type.hpp

@@ -249,7 +249,7 @@ using Vec3StorageT = StorageT __attribute__((ext_vector_type(3)));
 using Vec4StorageT = StorageT __attribute__((ext_vector_type(4)));
 using Vec8StorageT = StorageT __attribute__((ext_vector_type(8)));
 using Vec16StorageT = StorageT __attribute__((ext_vector_type(16)));
-#else
+#else // SYCL_DEVICE_ONLY
 using StorageT = detail::host_half_impl::half;
 // No need to extract underlying data type for built-in functions operating on
 // host
@@ -259,13 +259,37 @@ 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.
-
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
 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
+#else  // __INTEL_PREVIEW_BREAKING_CHANGES
+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>;
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
+#endif // SYCL_DEVICE_ONLY
 
 #ifndef __SYCL_DEVICE_ONLY__
 class half {