[SYCL] Switch to preview implementation of sycl::vec (#14317)

I'm keeping the `vector_preview.hpp` file name to avoid misleading diff
in this PR.

Also, a minor fix in swizzles implementation is included in this change
fixing a bug that would have affected `SubGroup/load_store.cpp` E2E test
otherwise.

Author: aelovikov-intel

libdevice/nativecpu_utils.cpp

@@ -77,9 +77,9 @@ __spirv_ControlBarrier(uint32_t Execution, uint32_t Memory,
 
 namespace ncpu_types {
 template <class T> struct vtypes {
-  using v2 = typename sycl::detail::VecStorage<T, 2>::DataType;
-  using v4 = typename sycl::detail::VecStorage<T, 4>::DataType;
-  using v8 = typename sycl::detail::VecStorage<T, 8>::DataType;
+  using v2 = typename sycl::vec<T, 2>::vector_t;
+  using v4 = typename sycl::vec<T, 4>::vector_t;
+  using v8 = typename sycl::vec<T, 8>::vector_t;
 };
 } // namespace ncpu_types
 
@@ -252,8 +252,8 @@ DefShuffleINTEL_All(double, f64, double)
 DefShuffleINTEL_All(float, f32, float)
 
 #define DefineShuffleVec(T, N, Sfx, MuxType)                                   \
-  using vt##T##N = sycl::detail::VecStorage<T, N>::DataType;                   \
-  using vt##MuxType##N = sycl::detail::VecStorage<MuxType, N>::DataType;       \
+  using vt##T##N = sycl::vec<T, N>::vector_t;                                  \
+  using vt##MuxType##N = sycl::vec<MuxType, N>::vector_t;                      \
   DefShuffleINTEL_All(vt##T##N, v##N##Sfx, vt##MuxType##N)
 
 #define DefineShuffleVec2to16(Type, Sfx, MuxType)                              \

sycl/include/sycl/detail/generic_type_traits.hpp

@@ -344,20 +344,7 @@ template <typename T> auto convertToOpenCLType(T &&x) {
                                                    std::declval<ElemTy>()))>,
                             no_ref::size()>;
 #ifdef __SYCL_DEVICE_ONLY__
-
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-    // TODO: for some mysterious reasons on NonUniformGroups E2E tests fail if
-    // we use the "else" version only. I suspect that's an issues with
-    // non-uniform groups implementation.
-    if constexpr (std::is_same_v<MatchingVec, no_ref>)
-      return static_cast<typename MatchingVec::vector_t>(x);
-    else
-      return static_cast<typename MatchingVec::vector_t>(
-          x.template as<MatchingVec>());
-#else  // __INTEL_PREVIEW_BREAKING_CHANGES
     return sycl::bit_cast<typename MatchingVec::vector_t>(x);
-#endif // __INTEL_PREVIEW_BREAKING_CHANGES
-
 #else
     return x.template as<MatchingVec>();
 #endif

sycl/include/sycl/ext/oneapi/bfloat16.hpp

@@ -78,25 +78,6 @@ template <int N> void BF16VecToFloatVec(const bfloat16 src[N], float dst[N]) {
   }
 #endif
 }
-
-// sycl::vec support
-namespace bf16 {
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-#ifdef __SYCL_DEVICE_ONLY__
-using Vec2StorageT = Bfloat16StorageT __attribute__((ext_vector_type(2)));
-using Vec3StorageT = Bfloat16StorageT __attribute__((ext_vector_type(3)));
-using Vec4StorageT = Bfloat16StorageT __attribute__((ext_vector_type(4)));
-using Vec8StorageT = Bfloat16StorageT __attribute__((ext_vector_type(8)));
-using Vec16StorageT = Bfloat16StorageT __attribute__((ext_vector_type(16)));
-#else
-using Vec2StorageT = std::array<Bfloat16StorageT, 2>;
-using Vec3StorageT = std::array<Bfloat16StorageT, 3>;
-using Vec4StorageT = std::array<Bfloat16StorageT, 4>;
-using Vec8StorageT = std::array<Bfloat16StorageT, 8>;
-using Vec16StorageT = std::array<Bfloat16StorageT, 16>;
-#endif
-#endif // __INTEL_PREVIEW_BREAKING_CHANGES
-} // namespace bf16
 } // namespace detail
 
 class bfloat16 {

sycl/include/sycl/half_type.hpp

@@ -169,34 +169,12 @@ class half;
 using StorageT = _Float16;
 using BIsRepresentationT = _Float16;
 using VecElemT = _Float16;
-
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-using Vec2StorageT = VecElemT __attribute__((ext_vector_type(2)));
-using Vec3StorageT = VecElemT __attribute__((ext_vector_type(3)));
-using Vec4StorageT = VecElemT __attribute__((ext_vector_type(4)));
-using Vec8StorageT = VecElemT __attribute__((ext_vector_type(8)));
-using Vec16StorageT = VecElemT __attribute__((ext_vector_type(16)));
-#endif // __INTEL_PREVIEW_BREAKING_CHANGES
-
 #else // SYCL_DEVICE_ONLY
 using StorageT = uint16_t;
 // No need to extract underlying data type for built-in functions operating on
 // host
 using BIsRepresentationT = half;
 using VecElemT = half;
-
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-// On the host side we cannot use OpenCL cl_half# types as an underlying type
-// 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.
-using Vec2StorageT = std::array<VecElemT, 2>;
-using Vec3StorageT = std::array<VecElemT, 3>;
-using Vec4StorageT = std::array<VecElemT, 4>;
-using Vec8StorageT = std::array<VecElemT, 8>;
-using Vec16StorageT = std::array<VecElemT, 16>;
-#endif // __INTEL_PREVIEW_BREAKING_CHANGES
-
 #endif // SYCL_DEVICE_ONLY
 
 // Creation token to disambiguate constructors.

sycl/include/sycl/types.hpp

@@ -21,10 +21,6 @@
 #include <sycl/half_type.hpp>                  // for StorageT, half, Vec16...
 #include <sycl/marray.hpp>                     // for __SYCL_BINOP, __SYCL_...
 #include <sycl/multi_ptr.hpp>                  // for multi_ptr
-#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
-#include <sycl/vector_preview.hpp> // for sycl::vec and swizzles
-#else
-#include <sycl/vector.hpp> // for sycl::vec and swizzles
-#endif
+#include <sycl/vector_preview.hpp>             // for sycl::vec and swizzles
 
 #include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16