// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Mehdi Goli    Codeplay Software Ltd.
// Ralph Potter  Codeplay Software Ltd.
// Luke Iwanski  Codeplay Software Ltd.
// Contact: <eigen@codeplay.com>
//
// This Source Code Form is subject to the terms of the Mozilla
// Public License v. 2.0. If a copy of the MPL was not distributed
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.

/*****************************************************************
 * PacketMath.h
 *
 * \brief:
 *  PacketMath
 *
 *****************************************************************/

#ifndef EIGEN_PACKET_MATH_SYCL_H
#define EIGEN_PACKET_MATH_SYCL_H
#include <type_traits>

// IWYU pragma: private
#include "../../InternalHeaderCheck.h"

namespace Eigen {

namespace internal {
#ifdef SYCL_DEVICE_ONLY
#define SYCL_PLOAD(packet_type, AlignedType)                          \
  template <>                                                         \
  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type                   \
      pload##AlignedType<packet_type>(                                \
          const typename unpacket_traits<packet_type>::type* from) {  \
   auto ptr = cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>(from);\
    packet_type res{};                                                \
    res.load(0, ptr);                                     \
    return res;                                                       \
  }

SYCL_PLOAD(cl::sycl::cl_float4, u)
SYCL_PLOAD(cl::sycl::cl_float4, )
SYCL_PLOAD(cl::sycl::cl_double2, u)
SYCL_PLOAD(cl::sycl::cl_double2, )
#undef SYCL_PLOAD

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8
    pload<cl::sycl::cl_half8>(
        const typename unpacket_traits<cl::sycl::cl_half8>::type* from) {
  auto ptr = cl::sycl::address_space_cast<
      cl::sycl::access::address_space::generic_space,
      cl::sycl::access::decorated::no>(
      reinterpret_cast<const cl::sycl::cl_half*>(from));
  cl::sycl::cl_half8 res{};
  res.load(0, ptr);
  return res;
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8
ploadu<cl::sycl::cl_half8>(
    const typename unpacket_traits<cl::sycl::cl_half8>::type* from) {
  auto ptr = cl::sycl::address_space_cast<
      cl::sycl::access::address_space::generic_space,
      cl::sycl::access::decorated::no>(
      reinterpret_cast<const cl::sycl::cl_half*>(from));
  cl::sycl::cl_half8 res{};
  res.load(0, ptr);
  return res;
}

#define SYCL_PSTORE(scalar, packet_type, alignment)             \
  template <>                                                   \
  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstore##alignment( \
      scalar* to, const packet_type& from) {                    \
    auto ptr = cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>(to);\
    from.store(0, ptr);                               \
  }

SYCL_PSTORE(float, cl::sycl::cl_float4, )
SYCL_PSTORE(float, cl::sycl::cl_float4, u)
SYCL_PSTORE(double, cl::sycl::cl_double2, )
SYCL_PSTORE(double, cl::sycl::cl_double2, u)
#undef SYCL_PSTORE

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstoreu(
    Eigen::half* to, const cl::sycl::cl_half8& from) {
  auto ptr = cl::sycl::address_space_cast<
      cl::sycl::access::address_space::generic_space,
      cl::sycl::access::decorated::no>(
      reinterpret_cast<cl::sycl::cl_half*>(to));
  from.store(0, ptr);
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstore(
    Eigen::half* to, const cl::sycl::cl_half8& from) {
  auto ptr = cl::sycl::address_space_cast<
      cl::sycl::access::address_space::generic_space,
      cl::sycl::access::decorated::no>(
      reinterpret_cast<cl::sycl::cl_half*>(to));
  from.store(0, ptr);
}

#define SYCL_PSET1(packet_type)                                         \
  template <>                                                           \
  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pset1<packet_type>( \
      const typename unpacket_traits<packet_type>::type& from) {        \
    return packet_type(from);                                           \
  }

// global space
SYCL_PSET1(cl::sycl::cl_half8)
SYCL_PSET1(cl::sycl::cl_float4)
SYCL_PSET1(cl::sycl::cl_double2)

#undef SYCL_PSET1

template <typename packet_type>
struct get_base_packet {
  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type
  get_ploaddup(sycl_multi_pointer) {}

  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type
  get_pgather(sycl_multi_pointer, Index) {}
};

template <>
struct get_base_packet<cl::sycl::cl_half8> {
  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_ploaddup(
      sycl_multi_pointer from) {
    return cl::sycl::cl_half8(static_cast<cl::sycl::half>(from[0]),
                              static_cast<cl::sycl::half>(from[0]),
                              static_cast<cl::sycl::half>(from[1]),
                              static_cast<cl::sycl::half>(from[1]),
                              static_cast<cl::sycl::half>(from[2]),
                              static_cast<cl::sycl::half>(from[2]),
                              static_cast<cl::sycl::half>(from[3]),
                              static_cast<cl::sycl::half>(from[3]));
  }
  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_pgather(
      sycl_multi_pointer from, Index stride) {
    return cl::sycl::cl_half8(static_cast<cl::sycl::half>(from[0 * stride]),
                              static_cast<cl::sycl::half>(from[1 * stride]),
                              static_cast<cl::sycl::half>(from[2 * stride]),
                              static_cast<cl::sycl::half>(from[3 * stride]),
                              static_cast<cl::sycl::half>(from[4 * stride]),
                              static_cast<cl::sycl::half>(from[5 * stride]),
                              static_cast<cl::sycl::half>(from[6 * stride]),
                              static_cast<cl::sycl::half>(from[7 * stride]));
  }

  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(
      sycl_multi_pointer to, const cl::sycl::cl_half8& from, Index stride) {
    auto tmp = stride;
    to[0] = Eigen::half(from.s0());
    to[tmp] = Eigen::half(from.s1());
    to[tmp += stride] = Eigen::half(from.s2());
    to[tmp += stride] = Eigen::half(from.s3());
    to[tmp += stride] = Eigen::half(from.s4());
    to[tmp += stride] = Eigen::half(from.s5());
    to[tmp += stride] = Eigen::half(from.s6());
    to[tmp += stride] = Eigen::half(from.s7());
  }
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 set_plset(
      const cl::sycl::half& a) {
    return cl::sycl::cl_half8(static_cast<cl::sycl::half>(a), static_cast<cl::sycl::half>(a + 1),
                              static_cast<cl::sycl::half>(a + 2),
                              static_cast<cl::sycl::half>(a + 3),
                              static_cast<cl::sycl::half>(a + 4),
                              static_cast<cl::sycl::half>(a + 5),
                              static_cast<cl::sycl::half>(a + 6),
                              static_cast<cl::sycl::half>(a + 7));
  }
};

template <>
struct get_base_packet<cl::sycl::cl_float4> {
  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_ploaddup(
      sycl_multi_pointer from) {
    return cl::sycl::cl_float4(from[0], from[0], from[1], from[1]);
  }
  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_pgather(
      sycl_multi_pointer from, Index stride) {
    return cl::sycl::cl_float4(from[0 * stride], from[1 * stride],
                               from[2 * stride], from[3 * stride]);
  }

  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(
      sycl_multi_pointer to, const cl::sycl::cl_float4& from, Index stride) {
    auto tmp = stride;
    to[0] = from.x();
    to[tmp] = from.y();
    to[tmp += stride] = from.z();
    to[tmp += stride] = from.w();
  }
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 set_plset(
      const float& a) {
    return cl::sycl::cl_float4(static_cast<float>(a), static_cast<float>(a + 1),
                               static_cast<float>(a + 2),
                               static_cast<float>(a + 3));
  }
};

template <>
struct get_base_packet<cl::sycl::cl_double2> {
  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2
  get_ploaddup(const sycl_multi_pointer from) {
    return cl::sycl::cl_double2(from[0], from[0]);
  }

  template <typename sycl_multi_pointer, typename Index>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 get_pgather(
      const sycl_multi_pointer from, Index stride) {
    return cl::sycl::cl_double2(from[0 * stride], from[1 * stride]);
  }

  template <typename sycl_multi_pointer>
  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(
      sycl_multi_pointer to, const cl::sycl::cl_double2& from, Index stride) {
    to[0] = from.x();
    to[stride] = from.y();
  }

  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 set_plset(
      const double& a) {
    return cl::sycl::cl_double2(static_cast<double>(a),
                                static_cast<double>(a + 1));
  }
};

#define SYCL_PLOAD_DUP_SPECILIZE(packet_type)                              \
  template <>                                                              \
  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ploaddup<packet_type>( \
      const typename unpacket_traits<packet_type>::type* from) {           \
    return get_base_packet<packet_type>::get_ploaddup(from);               \
  }

SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_half8)
SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_float4)
SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_double2)

#undef SYCL_PLOAD_DUP_SPECILIZE

#define SYCL_PLSET(packet_type)                                         \
  template <>                                                           \
  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type plset<packet_type>( \
      const typename unpacket_traits<packet_type>::type& a) {           \
    return get_base_packet<packet_type>::set_plset(a);                  \
  }
SYCL_PLSET(cl::sycl::cl_float4)
SYCL_PLSET(cl::sycl::cl_double2)
#undef SYCL_PLSET

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 plset<cl::sycl::cl_half8>(
    const typename unpacket_traits<cl::sycl::cl_half8>::type& a) {
  return get_base_packet<cl::sycl::cl_half8>::set_plset((const cl::sycl::half &) a);
}

#define SYCL_PGATHER_SPECILIZE(scalar, packet_type)                            \
  template <>                                                                  \
  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type                            \
  pgather<scalar, packet_type>(                                                \
      const typename unpacket_traits<packet_type>::type* from, Index stride) { \
    return get_base_packet<packet_type>::get_pgather(from, stride);            \
  }

SYCL_PGATHER_SPECILIZE(Eigen::half, cl::sycl::cl_half8)
SYCL_PGATHER_SPECILIZE(float, cl::sycl::cl_float4)
SYCL_PGATHER_SPECILIZE(double, cl::sycl::cl_double2)
#undef SYCL_PGATHER_SPECILIZE

#define SYCL_PSCATTER_SPECILIZE(scalar, packet_type)                        \
  template <>                                                               \
  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter<scalar, packet_type>( \
      typename unpacket_traits<packet_type>::type * to,                     \
      const packet_type& from, Index stride) {                              \
    get_base_packet<packet_type>::set_pscatter(to, from, stride);           \
  }

SYCL_PSCATTER_SPECILIZE(Eigen::half, cl::sycl::cl_half8)
SYCL_PSCATTER_SPECILIZE(float, cl::sycl::cl_float4)
SYCL_PSCATTER_SPECILIZE(double, cl::sycl::cl_double2)

#undef SYCL_PSCATTER_SPECILIZE

#define SYCL_PMAD(packet_type)                                            \
  template <>                                                             \
  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pmadd(                \
      const packet_type& a, const packet_type& b, const packet_type& c) { \
    return cl::sycl::mad(a, b, c);                                        \
  }

SYCL_PMAD(cl::sycl::cl_half8)
SYCL_PMAD(cl::sycl::cl_float4)
SYCL_PMAD(cl::sycl::cl_double2)
#undef SYCL_PMAD

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half pfirst<cl::sycl::cl_half8>(
    const cl::sycl::cl_half8& a) {
  return Eigen::half(a.s0());
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float pfirst<cl::sycl::cl_float4>(
    const cl::sycl::cl_float4& a) {
  return a.x();
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double pfirst<cl::sycl::cl_double2>(
    const cl::sycl::cl_double2& a) {
  return a.x();
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux<cl::sycl::cl_half8>(
    const cl::sycl::cl_half8& a) {
  return Eigen::half(a.s0() + a.s1() + a.s2() + a.s3() + a.s4() + a.s5()
                     + a.s6() + a.s7());
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux<cl::sycl::cl_float4>(
    const cl::sycl::cl_float4& a) {
  return a.x() + a.y() + a.z() + a.w();
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux<cl::sycl::cl_double2>(
    const cl::sycl::cl_double2& a) {
  return a.x() + a.y();
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_max<cl::sycl::cl_half8>(
    const cl::sycl::cl_half8& a) {
  return Eigen::half(cl::sycl::fmax(
          cl::sycl::fmax(
            cl::sycl::fmax(a.s0(), a.s1()),
            cl::sycl::fmax(a.s2(), a.s3())),
          cl::sycl::fmax(
            cl::sycl::fmax(a.s4(), a.s5()),
            cl::sycl::fmax(a.s6(), a.s7()))));
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_max<cl::sycl::cl_float4>(
    const cl::sycl::cl_float4& a) {
  return cl::sycl::fmax(cl::sycl::fmax(a.x(), a.y()),
                        cl::sycl::fmax(a.z(), a.w()));
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_max<cl::sycl::cl_double2>(
    const cl::sycl::cl_double2& a) {
  return cl::sycl::fmax(a.x(), a.y());
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_min<cl::sycl::cl_half8>(
    const cl::sycl::cl_half8& a) {
  return Eigen::half(cl::sycl::fmin(
      cl::sycl::fmin(
          cl::sycl::fmin(a.s0(), a.s1()),
          cl::sycl::fmin(a.s2(), a.s3())),
      cl::sycl::fmin(
          cl::sycl::fmin(a.s4(), a.s5()),
          cl::sycl::fmin(a.s6(), a.s7()))));
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_min<cl::sycl::cl_float4>(
    const cl::sycl::cl_float4& a) {
  return cl::sycl::fmin(cl::sycl::fmin(a.x(), a.y()),
                        cl::sycl::fmin(a.z(), a.w()));
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_min<cl::sycl::cl_double2>(
    const cl::sycl::cl_double2& a) {
  return cl::sycl::fmin(a.x(), a.y());
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_mul<cl::sycl::cl_half8>(
    const cl::sycl::cl_half8& a) {
  return Eigen::half(a.s0() * a.s1() * a.s2() * a.s3() * a.s4() * a.s5() *
                     a.s6() * a.s7());
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_mul<cl::sycl::cl_float4>(
    const cl::sycl::cl_float4& a) {
  return a.x() * a.y() * a.z() * a.w();
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_mul<cl::sycl::cl_double2>(
    const cl::sycl::cl_double2& a) {
  return a.x() * a.y();
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8
pabs<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
  return cl::sycl::cl_half8(cl::sycl::fabs(a.s0()), cl::sycl::fabs(a.s1()),
                            cl::sycl::fabs(a.s2()), cl::sycl::fabs(a.s3()),
                            cl::sycl::fabs(a.s4()), cl::sycl::fabs(a.s5()),
                            cl::sycl::fabs(a.s6()), cl::sycl::fabs(a.s7()));
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_float4
pabs<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
  return cl::sycl::cl_float4(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()),
                             cl::sycl::fabs(a.z()), cl::sycl::fabs(a.w()));
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_double2
pabs<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
  return cl::sycl::cl_double2(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()));
}

template <typename Packet>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_le(const Packet &a,
                                                          const Packet &b) {
  return (a <= b).template as<Packet>();
}

template <typename Packet>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_lt(const Packet &a,
                                                          const Packet &b) {
  return (a < b).template as<Packet>();
}

template <typename Packet>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_eq(const Packet &a,
                                                          const Packet &b) {
  return (a == b).template as<Packet>();
}

#define SYCL_PCMP(OP, TYPE)                                                    \
  template <>                                                                  \
  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE TYPE pcmp_##OP<TYPE>(const TYPE &a,    \
                                                             const TYPE &b) {  \
    return sycl_pcmp_##OP<TYPE>(a, b);                                         \
  }

SYCL_PCMP(le, cl::sycl::cl_half8)
SYCL_PCMP(lt, cl::sycl::cl_half8)
SYCL_PCMP(eq, cl::sycl::cl_half8)
SYCL_PCMP(le, cl::sycl::cl_float4)
SYCL_PCMP(lt, cl::sycl::cl_float4)
SYCL_PCMP(eq, cl::sycl::cl_float4)
SYCL_PCMP(le, cl::sycl::cl_double2)
SYCL_PCMP(lt, cl::sycl::cl_double2)
SYCL_PCMP(eq, cl::sycl::cl_double2)
#undef SYCL_PCMP

EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(
    PacketBlock<cl::sycl::cl_half8, 8>& kernel) {
  cl::sycl::cl_half tmp = kernel.packet[0].s1();
  kernel.packet[0].s1() = kernel.packet[1].s0();
  kernel.packet[1].s0() = tmp;

  tmp = kernel.packet[0].s2();
  kernel.packet[0].s2() = kernel.packet[2].s0();
  kernel.packet[2].s0() = tmp;

  tmp = kernel.packet[0].s3();
  kernel.packet[0].s3() = kernel.packet[3].s0();
  kernel.packet[3].s0() = tmp;

  tmp = kernel.packet[0].s4();
  kernel.packet[0].s4() = kernel.packet[4].s0();
  kernel.packet[4].s0() = tmp;

  tmp = kernel.packet[0].s5();
  kernel.packet[0].s5() = kernel.packet[5].s0();
  kernel.packet[5].s0() = tmp;

  tmp = kernel.packet[0].s6();
  kernel.packet[0].s6() = kernel.packet[6].s0();
  kernel.packet[6].s0() = tmp;

  tmp = kernel.packet[0].s7();
  kernel.packet[0].s7() = kernel.packet[7].s0();
  kernel.packet[7].s0() = tmp;

  tmp = kernel.packet[1].s2();
  kernel.packet[1].s2() = kernel.packet[2].s1();
  kernel.packet[2].s1() = tmp;

  tmp = kernel.packet[1].s3();
  kernel.packet[1].s3() = kernel.packet[3].s1();
  kernel.packet[3].s1() = tmp;

  tmp = kernel.packet[1].s4();
  kernel.packet[1].s4() = kernel.packet[4].s1();
  kernel.packet[4].s1() = tmp;

  tmp = kernel.packet[1].s5();
  kernel.packet[1].s5() = kernel.packet[5].s1();
  kernel.packet[5].s1() = tmp;

  tmp = kernel.packet[1].s6();
  kernel.packet[1].s6() = kernel.packet[6].s1();
  kernel.packet[6].s1() = tmp;

  tmp = kernel.packet[1].s7();
  kernel.packet[1].s7() = kernel.packet[7].s1();
  kernel.packet[7].s1() = tmp;

  tmp = kernel.packet[2].s3();
  kernel.packet[2].s3() = kernel.packet[3].s2();
  kernel.packet[3].s2() = tmp;

  tmp = kernel.packet[2].s4();
  kernel.packet[2].s4() = kernel.packet[4].s2();
  kernel.packet[4].s2() = tmp;

  tmp = kernel.packet[2].s5();
  kernel.packet[2].s5() = kernel.packet[5].s2();
  kernel.packet[5].s2() = tmp;

  tmp = kernel.packet[2].s6();
  kernel.packet[2].s6() = kernel.packet[6].s2();
  kernel.packet[6].s2() = tmp;

  tmp = kernel.packet[2].s7();
  kernel.packet[2].s7() = kernel.packet[7].s2();
  kernel.packet[7].s2() = tmp;

  tmp = kernel.packet[3].s4();
  kernel.packet[3].s4() = kernel.packet[4].s3();
  kernel.packet[4].s3() = tmp;

  tmp = kernel.packet[3].s5();
  kernel.packet[3].s5() = kernel.packet[5].s3();
  kernel.packet[5].s3() = tmp;

  tmp = kernel.packet[3].s6();
  kernel.packet[3].s6() = kernel.packet[6].s3();
  kernel.packet[6].s3() = tmp;

  tmp = kernel.packet[3].s7();
  kernel.packet[3].s7() = kernel.packet[7].s3();
  kernel.packet[7].s3() = tmp;

  tmp = kernel.packet[4].s5();
  kernel.packet[4].s5() = kernel.packet[5].s4();
  kernel.packet[5].s4() = tmp;

  tmp = kernel.packet[4].s6();
  kernel.packet[4].s6() = kernel.packet[6].s4();
  kernel.packet[6].s4() = tmp;

  tmp = kernel.packet[4].s7();
  kernel.packet[4].s7() = kernel.packet[7].s4();
  kernel.packet[7].s4() = tmp;

  tmp = kernel.packet[5].s6();
  kernel.packet[5].s6() = kernel.packet[6].s5();
  kernel.packet[6].s5() = tmp;

  tmp = kernel.packet[5].s7();
  kernel.packet[5].s7() = kernel.packet[7].s5();
  kernel.packet[7].s5() = tmp;

  tmp = kernel.packet[6].s7();
  kernel.packet[6].s7() = kernel.packet[7].s6();
  kernel.packet[7].s6() = tmp;
}

EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(
    PacketBlock<cl::sycl::cl_float4, 4>& kernel) {
  float tmp = kernel.packet[0].y();
  kernel.packet[0].y() = kernel.packet[1].x();
  kernel.packet[1].x() = tmp;

  tmp = kernel.packet[0].z();
  kernel.packet[0].z() = kernel.packet[2].x();
  kernel.packet[2].x() = tmp;

  tmp = kernel.packet[0].w();
  kernel.packet[0].w() = kernel.packet[3].x();
  kernel.packet[3].x() = tmp;

  tmp = kernel.packet[1].z();
  kernel.packet[1].z() = kernel.packet[2].y();
  kernel.packet[2].y() = tmp;

  tmp = kernel.packet[1].w();
  kernel.packet[1].w() = kernel.packet[3].y();
  kernel.packet[3].y() = tmp;

  tmp = kernel.packet[2].w();
  kernel.packet[2].w() = kernel.packet[3].z();
  kernel.packet[3].z() = tmp;
}

EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(
    PacketBlock<cl::sycl::cl_double2, 2>& kernel) {
  double tmp = kernel.packet[0].y();
  kernel.packet[0].y() = kernel.packet[1].x();
  kernel.packet[1].x() = tmp;
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 pblend(
    const Selector<unpacket_traits<cl::sycl::cl_half8>::size>& ifPacket,
    const cl::sycl::cl_half8& thenPacket,
    const cl::sycl::cl_half8& elsePacket) {
  cl::sycl::cl_short8 condition(
      ifPacket.select[0] ? 0 : -1, ifPacket.select[1] ? 0 : -1,
      ifPacket.select[2] ? 0 : -1, ifPacket.select[3] ? 0 : -1,
      ifPacket.select[4] ? 0 : -1, ifPacket.select[5] ? 0 : -1,
      ifPacket.select[6] ? 0 : -1, ifPacket.select[7] ? 0 : -1);
  return cl::sycl::select(thenPacket, elsePacket, condition);
}

template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_float4 pblend(
    const Selector<unpacket_traits<cl::sycl::cl_float4>::size>& ifPacket,
    const cl::sycl::cl_float4& thenPacket,
    const cl::sycl::cl_float4& elsePacket) {
  cl::sycl::cl_int4 condition(
      ifPacket.select[0] ? 0 : -1, ifPacket.select[1] ? 0 : -1,
      ifPacket.select[2] ? 0 : -1, ifPacket.select[3] ? 0 : -1);
  return cl::sycl::select(thenPacket, elsePacket, condition);
}

template <>
inline cl::sycl::cl_double2 pblend(
    const Selector<unpacket_traits<cl::sycl::cl_double2>::size>& ifPacket,
    const cl::sycl::cl_double2& thenPacket,
    const cl::sycl::cl_double2& elsePacket) {
  cl::sycl::cl_long2 condition(ifPacket.select[0] ? 0 : -1,
                               ifPacket.select[1] ? 0 : -1);
  return cl::sycl::select(thenPacket, elsePacket, condition);
}
#endif  // SYCL_DEVICE_ONLY

}  // end namespace internal

}  // end namespace Eigen

#endif  // EIGEN_PACKET_MATH_SYCL_H