diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp
index 46856ab03e544..fad5dda896ee7 100644
--- a/sycl/include/sycl/detail/type_traits.hpp
+++ b/sycl/include/sycl/detail/type_traits.hpp
@@ -24,10 +24,6 @@ template <class T> struct is_fixed_size_group : std::false_type {};
 
 template <class T>
 inline constexpr bool is_fixed_size_group_v = is_fixed_size_group<T>::value;
-
-template <typename VecT, typename OperationLeftT, typename OperationRightT,
-          template <typename> class OperationCurrentT, int... Indexes>
-class SwizzleOp;
 } // namespace detail
 
 template <int Dimensions> class group;
@@ -154,12 +150,20 @@ template <typename T, size_t N> struct get_elem_type_unqual<marray<T, N>> {
 template <typename T, int N> struct get_elem_type_unqual<vec<T, N>> {
   using type = T;
 };
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
 template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 struct get_elem_type_unqual<SwizzleOp<VecT, OperationLeftT, OperationRightT,
                                OperationCurrentT, Indexes...>> {
   using type = typename get_elem_type_unqual<std::remove_cv_t<VecT>>::type;
 };
+#else
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+struct get_elem_type_unqual<detail::hide_swizzle_from_adl::Swizzle<
+    IsConstVec, DataT, VecSize, Indexes...>> {
+  using type = DataT;
+};
+#endif
 
 template <typename ElementType, access::address_space Space,
           access::decorated DecorateAddress>
diff --git a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp b/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
index 7196a7ca4f868..afc2a676023b8 100644
--- a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
+++ b/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
@@ -32,9 +32,16 @@ template <typename DataT, int NumElements> class __SYCL_EBO vec;
 template <typename DataT, std::size_t N> class marray;
 
 namespace detail {
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
 template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 class SwizzleOp;
+#else
+namespace hide_swizzle_from_adl {
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+class __SYCL_EBO Swizzle;
+}
+#endif
 
 // Utility for converting a swizzle to a vector or preserve the type if it isn't
 // a swizzle.
@@ -42,12 +49,20 @@ template <typename T> struct simplify_if_swizzle {
   using type = T;
 };
 
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
 template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 struct simplify_if_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
                                      OperationCurrentT, Indexes...>> {
   using type = vec<typename VecT::element_type, sizeof...(Indexes)>;
 };
+#else
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+struct simplify_if_swizzle<detail::hide_swizzle_from_adl::Swizzle<
+    IsConstVec, DataT, VecSize, Indexes...>> {
+  using type = vec<DataT, sizeof...(Indexes)>;
+};
+#endif
 
 template <typename T>
 using simplify_if_swizzle_t = typename simplify_if_swizzle<T>::type;
@@ -79,10 +94,17 @@ inline constexpr bool is_valid_type_for_ext_vector_v =
     is_valid_type_for_ext_vector<T>::value;
 
 template <typename> struct is_swizzle : std::false_type {};
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
 template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 struct is_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
                             OperationCurrentT, Indexes...>> : std::true_type {};
+#else
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+struct is_swizzle<detail::hide_swizzle_from_adl::Swizzle<IsConstVec, DataT,
+                                                         VecSize, Indexes...>>
+    : std::true_type {};
+#endif
 template <typename T> constexpr bool is_swizzle_v = is_swizzle<T>::value;
 
 template <typename T>
@@ -108,11 +130,18 @@ struct num_elements<T __attribute__((ext_vector_type(N)))>
     : std::integral_constant<std::size_t, N> {};
 #endif
 #endif
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
 template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 struct num_elements<SwizzleOp<VecT, OperationLeftT, OperationRightT,
                               OperationCurrentT, Indexes...>>
     : std::integral_constant<std::size_t, sizeof...(Indexes)> {};
+#else
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+struct num_elements<detail::hide_swizzle_from_adl::Swizzle<IsConstVec, DataT,
+                                                           VecSize, Indexes...>>
+    : std::integral_constant<std::size_t, sizeof...(Indexes)> {};
+#endif
 
 template <typename T>
 inline constexpr std::size_t num_elements_v = num_elements<T>::value;
diff --git a/sycl/include/sycl/detail/vector_arith.hpp b/sycl/include/sycl/detail/vector_arith.hpp
index 6d92213e825e5..377605a391d67 100644
--- a/sycl/include/sycl/detail/vector_arith.hpp
+++ b/sycl/include/sycl/detail/vector_arith.hpp
@@ -16,11 +16,7 @@
 
 namespace sycl {
 inline namespace _V1 {
-
-template <typename DataT, int NumElem> class __SYCL_EBO vec;
-
 namespace detail {
-
 template <typename T> struct from_incomplete;
 template <typename T>
 struct from_incomplete<const T> : public from_incomplete<T> {};
@@ -31,6 +27,33 @@ struct from_incomplete<vec<DataT, NumElements>> {
   static constexpr size_t size() { return NumElements; }
 };
 
+#if !__SYCL_USE_LIBSYCL8_VEC_IMPL
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+struct from_incomplete<
+    hide_swizzle_from_adl::Swizzle<IsConstVec, DataT, VecSize, Indexes...>> {
+  using element_type = DataT;
+  static constexpr size_t size() { return sizeof...(Indexes); }
+
+  using vec_ty = std::conditional_t<IsConstVec, const vec<DataT, VecSize>,
+                                    vec<DataT, VecSize>>;
+  using result_vec_ty = vec<DataT, size()>;
+  static constexpr int vec_size = VecSize;
+  static constexpr bool is_over_const_vec = IsConstVec;
+  static constexpr bool has_repeating_indexes = []() constexpr {
+    int Idxs[] = {Indexes...};
+    for (std::size_t i = 1; i < sizeof...(Indexes); ++i) {
+      for (std::size_t j = 0; j < i; ++j)
+        if (Idxs[j] == Idxs[i])
+          // Repeating index
+          return true;
+    }
+
+    return false;
+  }();
+  static constexpr bool is_assignable = !IsConstVec && !has_repeating_indexes;
+};
+#endif
+
 template <bool Cond, typename Mixin> struct ApplyIf {};
 template <typename Mixin> struct ApplyIf<true, Mixin> : Mixin {};
 
@@ -477,6 +500,172 @@ template <typename Self> struct VecOperators {
                     OpAssign<ShiftRight>, IncDec> {};
 };
 
+#if !__SYCL_USE_LIBSYCL8_VEC_IMPL
+template <typename Self> struct SwizzleOperators {
+  using element_type = typename from_incomplete<Self>::element_type;
+  using vec_ty = typename from_incomplete<Self>::result_vec_ty;
+  static constexpr int N = from_incomplete<Self>::size();
+
+  template <typename T>
+  static constexpr bool is_compatible_scalar =
+      std::is_convertible_v<T, typename from_incomplete<Self>::element_type> &&
+      !is_swizzle_v<T>;
+
+  // Can't use partial specialization on constexpr variables because it took too
+  // long for gcc to fix https://gcc.gnu.org/bugzilla/show_bug.cgi?id=71954 and
+  // we need to support older versions without the fix.
+  template <typename OtherSwizzle, typename = void>
+  struct is_compatible_swizzle_impl : std::false_type {};
+
+  template <typename OtherSwizzle>
+  struct is_compatible_swizzle_impl<
+      OtherSwizzle, std::enable_if_t<is_swizzle_v<OtherSwizzle>>>
+      : std::bool_constant<
+            std::is_same_v<typename from_incomplete<OtherSwizzle>::element_type,
+                           typename from_incomplete<Self>::element_type> &&
+            from_incomplete<OtherSwizzle>::size() ==
+                from_incomplete<Self>::size()> {};
+
+  template <typename OtherSwizzle>
+  static constexpr bool is_compatible_swizzle =
+      is_compatible_swizzle_impl<OtherSwizzle>::value;
+
+  template <typename OtherSwizzle, typename = void>
+  struct is_compatible_swizzle_opposite_const_impl : std::false_type {};
+
+  template <typename OtherSwizzle>
+  struct is_compatible_swizzle_opposite_const_impl<
+      OtherSwizzle, std::enable_if_t<is_swizzle_v<OtherSwizzle>>>
+      : std::bool_constant<is_compatible_swizzle<OtherSwizzle> &&
+                           from_incomplete<OtherSwizzle>::is_over_const_vec !=
+                               from_incomplete<Self>::is_over_const_vec> {};
+
+  template <typename OtherSwizzle>
+  static constexpr bool is_compatible_swizzle_opposite_const =
+      is_compatible_swizzle_opposite_const_impl<OtherSwizzle>::value;
+
+  template <typename Op>
+  using result_t = std::conditional_t<
+      is_logical<Op>, vec<fixed_width_signed<sizeof(element_type)>, N>, vec_ty>;
+
+  // Uglier than possible due to
+  // https://gcc.gnu.org/bugzilla/show_bug.cgi?id=85282.
+  template <typename Op, typename = void> struct OpMixin;
+
+  template <typename Op>
+  struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, IncDec>>>
+      : public IncDecImpl<const Self> {};
+
+#define __SYCL_SWIZZLE_BINOP_MIXIN(OP, OPERATOR)                               \
+  template <typename Op>                                                       \
+  struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, OP>>> {               \
+    friend result_t<OP> operator OPERATOR(const Self & lhs,                    \
+                                          const vec_ty & rhs) {                \
+      return OP{}(vec_ty{lhs}, rhs);                                           \
+    }                                                                          \
+    friend result_t<OP> operator OPERATOR(const vec_ty & lhs,                  \
+                                          const Self & rhs) {                  \
+      return OP{}(lhs, vec_ty{rhs});                                           \
+    }                                                                          \
+    template <typename T>                                                      \
+    friend std::enable_if_t<is_compatible_scalar<T>, result_t<OP>>             \
+    operator OPERATOR(const Self & lhs, const T & rhs) {                       \
+      return OP{}(vec_ty{lhs}, vec_ty{rhs});                                   \
+    }                                                                          \
+    template <typename T>                                                      \
+    friend std::enable_if_t<is_compatible_scalar<T>, result_t<OP>>             \
+    operator OPERATOR(const T & lhs, const Self & rhs) {                       \
+      return OP{}(vec_ty{lhs}, vec_ty{rhs});                                   \
+    }                                                                          \
+    template <typename OtherSwizzle>                                           \
+    friend std::enable_if_t<is_compatible_swizzle<OtherSwizzle>, result_t<OP>> \
+    operator OPERATOR(const Self & lhs, const OtherSwizzle & rhs) {            \
+      return OP{}(vec_ty{lhs}, vec_ty{rhs});                                   \
+    }                                                                          \
+    template <typename OtherSwizzle>                                           \
+    friend std::enable_if_t<                                                   \
+        is_compatible_swizzle_opposite_const<OtherSwizzle>, result_t<OP>>      \
+    operator OPERATOR(const OtherSwizzle & lhs, const Self & rhs) {            \
+      return OP{}(vec_ty{lhs}, vec_ty{rhs});                                   \
+    }                                                                          \
+  };
+
+#define __SYCL_SWIZZLE_OPASSIGN_MIXIN(OP, OPERATOR)                            \
+  template <typename Op>                                                       \
+  struct OpMixin<OpAssign<Op>, std::enable_if_t<std::is_same_v<Op, OP>>> {     \
+    friend const Self &operator OPERATOR(const Self & lhs,                     \
+                                         const vec_ty & rhs) {                 \
+      lhs = OP{}(vec_ty{lhs}, rhs);                                            \
+      return lhs;                                                              \
+    }                                                                          \
+    template <typename T>                                                      \
+    friend std::enable_if_t<is_compatible_scalar<T>, const Self &>             \
+    operator OPERATOR(const Self & lhs, const T & rhs) {                       \
+      lhs = OP{}(vec_ty{lhs}, vec_ty{rhs});                                    \
+      return lhs;                                                              \
+    }                                                                          \
+    template <typename OtherSwizzle>                                           \
+    friend std::enable_if_t<is_compatible_swizzle<OtherSwizzle>, const Self &> \
+    operator OPERATOR(const Self & lhs, const OtherSwizzle & rhs) {            \
+      lhs = OP{}(vec_ty{lhs}, vec_ty{rhs});                                    \
+      return lhs;                                                              \
+    }                                                                          \
+  };
+
+#define __SYCL_SWIZZLE_UOP_MIXIN(OP, OPERATOR)                                 \
+  template <typename Op>                                                       \
+  struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, OP>>> {               \
+    friend result_t<OP> operator OPERATOR(const Self & v) {                    \
+      return OP{}(vec_ty{v});                                                  \
+    }                                                                          \
+  };
+
+  __SYCL_INSTANTIATE_OPERATORS(__SYCL_SWIZZLE_BINOP_MIXIN,
+                               __SYCL_SWIZZLE_OPASSIGN_MIXIN,
+                               __SYCL_SWIZZLE_UOP_MIXIN)
+
+#undef __SYCL_SWIZZLE_UOP_MIXIN
+#undef __SYCL_SWIZZLE_OPASSIGN_MIXIN
+#undef __SYCL_SWIZZLE_BINOP_MIXIN
+
+  template <typename... Op>
+  struct __SYCL_EBO CombineImpl
+      : ApplyIf<is_op_available_for_type<Op, element_type>, OpMixin<Op>>... {};
+
+  template <typename _Self, typename = void>
+  struct CombinedImpl
+      : CombineImpl<std::plus<void>, std::minus<void>, std::multiplies<void>,
+                    std::divides<void>, std::modulus<void>, std::bit_and<void>,
+                    std::bit_or<void>, std::bit_xor<void>, std::equal_to<void>,
+                    std::not_equal_to<void>, std::less<void>,
+                    std::greater<void>, std::less_equal<void>,
+                    std::greater_equal<void>, std::logical_and<void>,
+                    std::logical_or<void>, ShiftLeft, ShiftRight,
+                    std::negate<void>, std::logical_not<void>,
+                    std::bit_not<void>, UnaryPlus> {};
+
+  template <typename _Self>
+  struct CombinedImpl<_Self,
+                      std::enable_if_t<from_incomplete<_Self>::is_assignable>>
+      : CombineImpl<std::plus<void>, std::minus<void>, std::multiplies<void>,
+                    std::divides<void>, std::modulus<void>, std::bit_and<void>,
+                    std::bit_or<void>, std::bit_xor<void>, std::equal_to<void>,
+                    std::not_equal_to<void>, std::less<void>,
+                    std::greater<void>, std::less_equal<void>,
+                    std::greater_equal<void>, std::logical_and<void>,
+                    std::logical_or<void>, ShiftLeft, ShiftRight,
+                    std::negate<void>, std::logical_not<void>,
+                    std::bit_not<void>, UnaryPlus, OpAssign<std::plus<void>>,
+                    OpAssign<std::minus<void>>, OpAssign<std::multiplies<void>>,
+                    OpAssign<std::divides<void>>, OpAssign<std::modulus<void>>,
+                    OpAssign<std::bit_and<void>>, OpAssign<std::bit_or<void>>,
+                    OpAssign<std::bit_xor<void>>, OpAssign<ShiftLeft>,
+                    OpAssign<ShiftRight>, IncDec> {};
+
+  using Combined = CombinedImpl<Self>;
+};
+#endif
+
 #if __SYCL_USE_LIBSYCL8_VEC_IMPL
 template <typename DataT, int NumElements>
 class vec_arith : public VecOperators<vec<DataT, NumElements>>::Combined {};
diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp
index 556e91e7912ef..0f274566ebd91 100644
--- a/sycl/include/sycl/vector.hpp
+++ b/sycl/include/sycl/vector.hpp
@@ -232,6 +232,7 @@ template <typename DataT, int NumElements> class vec_base {
   template <typename DataT_, typename T> class FlattenVecArg {
     template <std::size_t... Is>
     static constexpr auto helper(const T &V, std::index_sequence<Is...>) {
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
       // FIXME: Swizzle's `operator[]` for expression trees seems to be broken
       // and returns values of the underlying vector of some of the operands. On
       // the other hand, `getValue()` gives correct results. This can be changed
@@ -239,6 +240,7 @@ template <typename DataT, int NumElements> class vec_base {
       if constexpr (is_swizzle_v<T>)
         return std::array{static_cast<DataT_>(V.getValue(Is))...};
       else
+#endif
         return std::array{static_cast<DataT_>(V[Is])...};
     }
 
@@ -300,17 +302,172 @@ template <typename DataT> class vec_base<DataT, 1> {
 
   // Not `explicit` on purpose, differs from NumElements > 1.
   constexpr vec_base(const DataT &arg) : m_Data{{arg}} {}
+};
+
+template <typename Self> class ConversionToVecMixin {
+  using vec_ty = typename from_incomplete<Self>::result_vec_ty;
+
+public:
+  operator vec_ty() const {
+    vec_ty res{*static_cast<const Self *>(this)};
+    return res;
+  }
+};
+
+template <typename Self, typename = void> class SwizzleBase {
+  using VecT = typename from_incomplete<Self>::vec_ty;
+
+public:
+  explicit SwizzleBase(VecT &Vec) : Vec(Vec) {}
+
+  const Self &operator=(const Self &) = delete;
+
+protected:
+  VecT &Vec;
+};
+
+template <typename Self>
+class SwizzleBase<Self,
+                  std::enable_if_t<from_incomplete<Self>::is_assignable>> {
+  using VecT = typename from_incomplete<Self>::vec_ty;
+  using ResultVecT = typename from_incomplete<Self>::result_vec_ty;
+
+  using DataT = typename from_incomplete<Self>::element_type;
+  static constexpr int N = from_incomplete<Self>::size();
+
+public:
+  explicit SwizzleBase(VecT &Vec) : Vec(Vec) {}
+
+  template <access::address_space AddressSpace, access::decorated IsDecorated>
+  void load(size_t offset,
+            multi_ptr<const DataT, AddressSpace, IsDecorated> ptr) const {
+    ResultVecT v;
+    v.load(offset, ptr);
+    *static_cast<Self *>(this) = v;
+  }
+
+  template <bool OtherIsConstVec, int OtherVecSize, int... OtherIndexes>
+  std::enable_if_t<sizeof...(OtherIndexes) == N, const Self &>
+  operator=(const detail::hide_swizzle_from_adl::Swizzle<
+            OtherIsConstVec, DataT, OtherVecSize, OtherIndexes...> &rhs) {
+    return (*this = static_cast<ResultVecT>(rhs));
+  }
+
+  const Self &operator=(const ResultVecT &rhs) const {
+    for (int i = 0; i < N; ++i)
+      (*static_cast<const Self *>(this))[i] = rhs[i];
+
+    return *static_cast<const Self *>(this);
+  }
+
+  template <typename T,
+            typename = std::enable_if_t<std::is_convertible_v<T, DataT> &&
+                                        !is_swizzle_v<T>>>
+  const Self &operator=(const T &rhs) const {
+    for (int i = 0; i < N; ++i)
+      (*static_cast<const Self *>(this))[i] = static_cast<DataT>(rhs);
+
+    return *static_cast<const Self *>(this);
+  }
+
+  // Default copy-assignment. Self's implicitly generated copy-assignment uses
+  // this.
+  //
+  // We're templated on "Self", so each swizzle has its own SwizzleBase and the
+  // following is ok (1-to-1 bidirectional mapping between Self and its
+  // SwizzleBase instantiation) even if a bit counterintuitive.
+  const SwizzleBase &operator=(const SwizzleBase &rhs) const {
+    const Self &self = (*static_cast<const Self *>(this));
+    self = static_cast<ResultVecT>(static_cast<const Self &>(rhs));
+    return self;
+  }
+
+protected:
+  VecT &Vec;
+};
+
+namespace hide_swizzle_from_adl {
+// Can't have sycl::vec anywhere in template parameters because that would bring
+// its hidden friends into ADL. Put it in a dedicated namespace to avoid
+// anything extra via ADL as well.
+template <bool IsConstVec, typename DataT, int VecSize, int... Indexes>
+class __SYCL_EBO Swizzle
+    : public SwizzleBase<Swizzle<IsConstVec, DataT, VecSize, Indexes...>>,
+      public SwizzleOperators<
+          Swizzle<IsConstVec, DataT, VecSize, Indexes...>>::Combined,
+      public ApplyIf<sizeof...(Indexes) == 1,
+                     ScalarConversionOperatorsMixIn<
+                         Swizzle<IsConstVec, DataT, VecSize, Indexes...>>>,
+      public ApplyIf<sizeof...(Indexes) != 1,
+                     ConversionToVecMixin<
+                         Swizzle<IsConstVec, DataT, VecSize, Indexes...>>>,
+      public NamedSwizzlesMixinBoth<
+          Swizzle<IsConstVec, DataT, VecSize, Indexes...>> {
+  using Base = SwizzleBase<Swizzle<IsConstVec, DataT, VecSize, Indexes...>>;
+
+  static constexpr int NumElements = sizeof...(Indexes);
+  using ResultVec = vec<DataT, NumElements>;
+
+  // Get underlying vec index for (*this)[idx] access.
+  static constexpr auto get_vec_idx(int idx) {
+    int counter = 0;
+    int result = -1;
+    ((result = counter++ == idx ? Indexes : result), ...);
+    return result;
+  }
+
+public:
+  using Base::Base;
+  using Base::operator=;
+
+  using element_type = DataT;
+  using value_type = DataT;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  using vector_t = typename vec<DataT, NumElements>::vector_t;
+#endif // __SYCL_DEVICE_ONLY__
+
+  Swizzle() = delete;
+  Swizzle(const Swizzle &) = delete;
+
+  static constexpr size_t byte_size() noexcept {
+    return ResultVec::byte_size();
+  }
+  static constexpr size_t size() noexcept { return ResultVec::size(); }
+
+  __SYCL2020_DEPRECATED(
+      "get_size() is deprecated, please use byte_size() instead")
+  size_t get_size() const { return static_cast<ResultVec>(*this).get_size(); }
+
+  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
+  size_t get_count() const {
+    return static_cast<ResultVec>(*this).get_count();
+  };
+
+  template <typename ConvertT,
+            rounding_mode RoundingMode = rounding_mode::automatic>
+  vec<ConvertT, NumElements> convert() const {
+    return static_cast<ResultVec>(*this)
+        .template convert<ConvertT, RoundingMode>();
+  }
+
+  template <typename asT> asT as() const {
+    return static_cast<ResultVec>(*this).template as<asT>();
+  }
+
+  template <access::address_space AddressSpace, access::decorated IsDecorated>
+  void store(size_t offset,
+             multi_ptr<DataT, AddressSpace, IsDecorated> ptr) const {
+    return static_cast<ResultVec>(*this).store(offset, ptr);
+  }
 
-  // FIXME: Temporary workaround because swizzle's `operator DataT` is a
-  // template.
-  template <typename Swizzle,
-            typename = std::enable_if_t<is_swizzle_v<Swizzle>>,
-            typename = std::enable_if_t<Swizzle::size() == 1>,
-            typename = std::enable_if<
-                std::is_convertible_v<typename Swizzle::element_type, DataT>>>
-  constexpr vec_base(const Swizzle &other)
-      : vec_base(static_cast<DataT>(other)) {}
+  template <int... swizzleIndexes> auto swizzle() const {
+    return this->Vec.template swizzle<get_vec_idx(swizzleIndexes)...>();
+  }
+
+  auto &operator[](int index) const { return this->Vec[get_vec_idx(index)]; }
 };
+} // namespace hide_swizzle_from_adl
 #endif
 } // namespace detail
 
@@ -385,6 +542,7 @@ class __SYCL_EBO vec :
 private:
 #endif // __SYCL_DEVICE_ONLY__
 
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
   template <int... Indexes>
   using Swizzle =
       detail::SwizzleOp<vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
@@ -394,6 +552,17 @@ class __SYCL_EBO vec :
   using ConstSwizzle =
       detail::SwizzleOp<const vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
                         detail::GetOp, Indexes...>;
+#else
+  template <int... Indexes>
+  using Swizzle =
+      detail::hide_swizzle_from_adl::Swizzle<false, DataT, NumElements,
+                                             Indexes...>;
+
+  template <int... Indexes>
+  using ConstSwizzle =
+      detail::hide_swizzle_from_adl::Swizzle<true, DataT, NumElements,
+                                             Indexes...>;
+#endif
 
   // Element type for relational operator return value.
   using rel_t = detail::fixed_width_signed<sizeof(DataT)>;
@@ -479,12 +648,20 @@ class __SYCL_EBO vec :
   template <typename asT> asT as() const { return sycl::bit_cast<asT>(*this); }
 
   template <int... SwizzleIndexes> Swizzle<SwizzleIndexes...> swizzle() {
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
     return this;
+#else
+    return Swizzle<SwizzleIndexes...>{*this};
+#endif
   }
 
   template <int... SwizzleIndexes>
   ConstSwizzle<SwizzleIndexes...> swizzle() const {
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
     return this;
+#else
+    return ConstSwizzle<SwizzleIndexes...>{*this};
+#endif
   }
 
   const DataT &operator[](int i) const { return this->m_Data[i]; }
@@ -563,6 +740,7 @@ template <class T, class... U,
 vec(T, U...) -> vec<T, sizeof...(U) + 1>;
 #endif
 
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
 namespace detail {
 
 // Special type for working SwizzleOp with scalars, stores a scalar and gives
@@ -1398,5 +1576,6 @@ class SwizzleOp : public detail::NamedSwizzlesMixinBoth<
 };
 ///////////////////////// class SwizzleOp /////////////////////////
 } // namespace detail
+#endif
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/test/basic_tests/vectors/assign.cpp b/sycl/test/basic_tests/vectors/assign.cpp
index 2045eca0303e3..115bb9af9f2dd 100644
--- a/sycl/test/basic_tests/vectors/assign.cpp
+++ b/sycl/test/basic_tests/vectors/assign.cpp
@@ -33,8 +33,8 @@ static_assert(                  std::is_assignable_v<vec<half, 1>, vec<half, 1>>
 static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<half, 1>, vec<float, 1>>);
 static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<half, 1>, vec<double, 1>>);
 static_assert(                  std::is_assignable_v<vec<half, 1>, sw_half_1>);
-static_assert(                  std::is_assignable_v<vec<half, 1>, sw_float_1>);
-static_assert(                  std::is_assignable_v<vec<half, 1>, sw_double_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<half, 1>, sw_float_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<half, 1>, sw_double_1>);
 static_assert(                 !std::is_assignable_v<vec<half, 1>, sw_half_2>);
 static_assert(                 !std::is_assignable_v<vec<half, 1>, sw_float_2>);
 static_assert(                 !std::is_assignable_v<vec<half, 1>, sw_double_2>);
@@ -46,8 +46,8 @@ static_assert(                  std::is_assignable_v<vec<half, 2>, vec<half, 1>>
 static_assert(                 !std::is_assignable_v<vec<half, 2>, vec<float, 1>>);
 static_assert(                 !std::is_assignable_v<vec<half, 2>, vec<double, 1>>);
 static_assert(                  std::is_assignable_v<vec<half, 2>, sw_half_1>);
-static_assert(                  std::is_assignable_v<vec<half, 2>, sw_float_1>);
-static_assert(                  std::is_assignable_v<vec<half, 2>, sw_double_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<half, 2>, sw_float_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<half, 2>, sw_double_1>);
 static_assert(                  std::is_assignable_v<vec<half, 2>, sw_half_2>);
 static_assert(                 !std::is_assignable_v<vec<half, 2>, sw_float_2>);
 static_assert(                 !std::is_assignable_v<vec<half, 2>, sw_double_2>);
@@ -62,7 +62,7 @@ static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<float, 1>, vec<half, 1>
 #endif
 static_assert(                  std::is_assignable_v<vec<float, 1>, vec<float, 1>>);
 static_assert(                  std::is_assignable_v<vec<float, 1>, vec<double, 1>>);
-static_assert(                  std::is_assignable_v<vec<float, 1>, sw_half_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<float, 1>, sw_half_1>);
 static_assert(                  std::is_assignable_v<vec<float, 1>, sw_float_1>);
 static_assert(                  std::is_assignable_v<vec<float, 1>, sw_double_1>);
 static_assert(                 !std::is_assignable_v<vec<float, 1>, sw_half_2>);
@@ -79,7 +79,7 @@ static_assert(                 !std::is_assignable_v<vec<float, 2>, vec<half, 1>
 #endif
 static_assert(                  std::is_assignable_v<vec<float, 2>, vec<float, 1>>);
 static_assert(                  std::is_assignable_v<vec<float, 2>, vec<double, 1>>);
-static_assert(                  std::is_assignable_v<vec<float, 2>, sw_half_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<float, 2>, sw_half_1>);
 static_assert(                  std::is_assignable_v<vec<float, 2>, sw_float_1>);
 static_assert(                  std::is_assignable_v<vec<float, 2>, sw_double_1>);
 static_assert(                 !std::is_assignable_v<vec<float, 2>, sw_half_2>);
diff --git a/sycl/test/basic_tests/vectors/cxx_conversions.cpp b/sycl/test/basic_tests/vectors/cxx_conversions.cpp
index b59adc718c31f..4973b07cd00db 100644
--- a/sycl/test/basic_tests/vectors/cxx_conversions.cpp
+++ b/sycl/test/basic_tests/vectors/cxx_conversions.cpp
@@ -53,9 +53,9 @@ using sw_double_2 = decltype(std::declval<vec<double, 4>>().swizzle<1, 2>());
 static_assert(                  std::is_invocable_v<decltype(f_half_v1), half>);
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), float>);
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), double>);
-static_assert(                  std::is_invocable_v<decltype(f_half_v1), sw_half_1>);
-static_assert(                  std::is_invocable_v<decltype(f_half_v1), sw_float_1>);
-static_assert(                  std::is_invocable_v<decltype(f_half_v1), sw_double_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), sw_half_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), sw_float_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), sw_double_1>);
 static_assert(                  std::is_invocable_v<decltype(f_half_v1), vec<half, 1>>);
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), vec<float, 1>>);
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), vec<double, 1>>);
@@ -63,9 +63,9 @@ static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_half_v1), vec<dou
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_float_v1), half>);
 static_assert(                  std::is_invocable_v<decltype(f_float_v1), float>);
 static_assert(                  std::is_invocable_v<decltype(f_float_v1), double>);
-static_assert(                  std::is_invocable_v<decltype(f_float_v1), sw_half_1>);
-static_assert(                  std::is_invocable_v<decltype(f_float_v1), sw_float_1>);
-static_assert(                  std::is_invocable_v<decltype(f_float_v1), sw_double_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_float_v1), sw_half_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_float_v1), sw_float_1>);
+static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_float_v1), sw_double_1>);
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_float_v1), vec<half, 1>>);
 static_assert(                  std::is_invocable_v<decltype(f_float_v1), vec<float, 1>>);
 static_assert(EXCEPT_IN_PREVIEW std::is_invocable_v<decltype(f_float_v1), vec<double, 1>>);
diff --git a/sycl/test/basic_tests/vectors/swizzle.cpp b/sycl/test/basic_tests/vectors/swizzle.cpp
index 2c6ce60331dc8..4de5fde499c5c 100644
--- a/sycl/test/basic_tests/vectors/swizzle.cpp
+++ b/sycl/test/basic_tests/vectors/swizzle.cpp
@@ -9,8 +9,12 @@ int main() {
   assert(sw.lo()[0] == 2);
   assert(sw.hi()[0] == 3);
 
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
   // FIXME: Should be "4":
   assert((sw + sw).lo()[0] == 2);
+#else
+  assert((sw + sw).lo()[0] == 4);
+#endif
 
   assert(sw.swizzle<0>()[0] == 2);
   assert(sw.swizzle<1>()[0] == 3);
@@ -24,9 +28,14 @@ int main() {
   {
     auto tmp = (sw + sw).swizzle<1, 0>();
 
+#if __SYCL_USE_LIBSYCL8_VEC_IMPL
     // FIXME: Should be "6" and "4", respectively.
     assert(tmp[0] == 3);
     assert(tmp[1] == 2);
+#else
+    assert(tmp[0] == 6);
+    assert(tmp[1] == 4);
+#endif
   }
 
   return 0;