math.hpp

Go to the documentation of this file.

//==-------------- math.hpp - DPC++ Explicit SIMD API   --------------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// Implement Explicit SIMD math APIs.
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/ext/intel/esimd/common.hpp>
#include <sycl/ext/intel/esimd/detail/math_intrin.hpp>
#include <sycl/ext/intel/esimd/detail/operators.hpp>
#include <sycl/ext/intel/esimd/detail/types.hpp>
#include <sycl/ext/intel/esimd/detail/util.hpp>
#include <sycl/ext/intel/esimd/simd.hpp>
#include <sycl/ext/intel/esimd/simd_view.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp>
 
#include <cstdint>
 
namespace sycl {
inline namespace _V1 {
namespace ext::intel::esimd {
 
 
 
template <typename T0, typename T1, int SZ>
__ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v<T0> ||
                                 std::is_same_v<T1, T0>,
                             simd<T0, SZ>>
saturate(simd<T1, SZ> src) {
  if constexpr (detail::is_generic_floating_point_v<T0>)
    return __esimd_sat<T0, T1, SZ>(src.data());
  else if constexpr (detail::is_generic_floating_point_v<T1>) {
    if constexpr (std::is_unsigned_v<T0>)
      return __esimd_fptoui_sat<T0, T1, SZ>(src.data());
    else
      return __esimd_fptosi_sat<T0, T1, SZ>(src.data());
  } else if constexpr (std::is_unsigned_v<T0>) {
    if constexpr (std::is_unsigned_v<T1>)
      return __esimd_uutrunc_sat<T0, T1, SZ>(src.data());
    else
      return __esimd_ustrunc_sat<T0, T1, SZ>(src.data());
  } else {
    if constexpr (std::is_signed_v<T1>)
      return __esimd_sstrunc_sat<T0, T1, SZ>(src.data());
    else
      return __esimd_sutrunc_sat<T0, T1, SZ>(src.data());
  }
}
 
// abs
namespace detail {
 
template <typename TRes, typename TArg, int SZ>
ESIMD_NODEBUG ESIMD_INLINE simd<TRes, SZ>
__esimd_abs_common_internal(simd<TArg, SZ> src0) {
  simd<TArg, SZ> Result;
  if constexpr (detail::is_generic_floating_point_v<TArg>) {
    using CppT = __ESIMD_DNS::element_type_traits<TArg>::EnclosingCppT;
    Result =
        __ESIMD_DNS::convert_vector<TArg, CppT, SZ>(__spirv_ocl_fabs<CppT, SZ>(
            __ESIMD_DNS::convert_vector<CppT, TArg, SZ>(src0.data())));
  } else
    Result = simd<TArg, SZ>(__spirv_ocl_s_abs<TArg, SZ>(src0.data()));
  return convert<TRes>(Result);
}
 
template <typename TRes, typename TArg>
 
__ESIMD_API std::enable_if_t<detail::is_esimd_scalar<TRes>::value &&
                                 detail::is_esimd_scalar<TArg>::value,
                             TRes>
__esimd_abs_common_internal(TArg src0) {
  simd<TArg, 1> Src0 = src0;
  simd<TArg, 1> Result = __esimd_abs_common_internal<TArg>(Src0);
  return convert<TRes>(Result)[0];
}
} // namespace detail
 
template <typename TRes, typename TArg, int SZ>
__ESIMD_API std::enable_if_t<
    !std::is_same<std::remove_const_t<TRes>, std::remove_const_t<TArg>>::value,
    simd<TRes, SZ>>
abs(simd<TArg, SZ> src0) {
  return detail::__esimd_abs_common_internal<TRes, TArg, SZ>(src0.data());
}
 
template <typename TRes, typename TArg>
__ESIMD_API std::enable_if_t<!std::is_same<std::remove_const_t<TRes>,
                                           std::remove_const_t<TArg>>::value &&
                                 detail::is_esimd_scalar<TRes>::value &&
                                 detail::is_esimd_scalar<TArg>::value,
                             std::remove_const_t<TRes>>
abs(TArg src0) {
  return detail::__esimd_abs_common_internal<TRes, TArg>(src0);
}
 
template <typename T1, int SZ> __ESIMD_API simd<T1, SZ> abs(simd<T1, SZ> src0) {
  return detail::__esimd_abs_common_internal<T1, T1, SZ>(src0.data());
}
 
template <typename T1>
__ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T1>::value,
                             std::remove_const_t<T1>>
abs(T1 src0) {
  return detail::__esimd_abs_common_internal<T1, T1>(src0);
}
 
template <typename T, int SZ, class Sat = saturation_off_tag>
__ESIMD_API simd<T, SZ>(max)(simd<T, SZ> src0, simd<T, SZ> src1, Sat sat = {}) {
  constexpr bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
 
  if constexpr (detail::is_generic_floating_point_v<T>) {
    using CppT = __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
    auto Result =
        __ESIMD_DNS::convert_vector<T, CppT, SZ>(__spirv_ocl_fmax<CppT, SZ>(
            __ESIMD_DNS::convert_vector<CppT, T, SZ>(src0.data()),
            __ESIMD_DNS::convert_vector<CppT, T, SZ>(src1.data())));
    if constexpr (is_sat)
      Result = __esimd_sat<T, T, SZ>(Result);
    return simd<T, SZ>(Result);
  } else if constexpr (std::is_unsigned<T>::value) {
    auto Result = __esimd_umax<T, SZ>(src0.data(), src1.data());
    if constexpr (is_sat)
      Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
    return simd<T, SZ>(Result);
  } else {
    auto Result = __esimd_smax<T, SZ>(src0.data(), src1.data());
    if constexpr (is_sat)
      Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
    return simd<T, SZ>(Result);
  }
}
 
template <typename T, int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
    max)(simd<T, SZ> src0, T src1, Sat sat = {}) {
  simd<T, SZ> Src1 = src1;
  simd<T, SZ> Result = (esimd::max)(src0, Src1, sat);
  return Result;
}
 
template <typename T, int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
    max)(T src0, simd<T, SZ> src1, Sat sat = {}) {
  simd<T, SZ> Src0 = src0;
  simd<T, SZ> Result = (esimd::max)(Src0, src1, sat);
  return Result;
}
 
template <typename T, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(max)(T src0, T src1,
                                                            Sat sat = {}) {
  simd<T, 1> Src0 = src0;
  simd<T, 1> Src1 = src1;
  simd<T, 1> Result = (esimd::max)(Src0, Src1, sat);
  return Result[0];
}
 
template <typename T, int SZ, class Sat = saturation_off_tag>
__ESIMD_API simd<T, SZ>(min)(simd<T, SZ> src0, simd<T, SZ> src1, Sat sat = {}) {
  constexpr bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
 
  if constexpr (detail::is_generic_floating_point_v<T>) {
    using CppT = __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
    auto Result =
        __ESIMD_DNS::convert_vector<T, CppT, SZ>(__spirv_ocl_fmin<CppT, SZ>(
            __ESIMD_DNS::convert_vector<CppT, T, SZ>(src0.data()),
            __ESIMD_DNS::convert_vector<CppT, T, SZ>(src1.data())));
    if constexpr (is_sat)
      Result = __esimd_sat<T, T, SZ>(Result);
    return simd<T, SZ>(Result);
  } else if constexpr (std::is_unsigned<T>::value) {
    auto Result = __esimd_umin<T, SZ>(src0.data(), src1.data());
    if constexpr (is_sat)
      Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
    return simd<T, SZ>(Result);
  } else {
    auto Result = __esimd_smin<T, SZ>(src0.data(), src1.data());
    if constexpr (is_sat)
      Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
    return simd<T, SZ>(Result);
  }
}
 
template <typename T, int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
    min)(simd<T, SZ> src0, T src1, Sat sat = {}) {
  simd<T, SZ> Src1 = src1;
  simd<T, SZ> Result = (esimd::min)(src0, Src1, sat);
  return Result;
}
 
template <typename T, int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
    min)(T src0, simd<T, SZ> src1, Sat sat = {}) {
  simd<T, SZ> Src0 = src0;
  simd<T, SZ> Result = (esimd::min)(Src0, src1, sat);
  return Result;
}
 
template <typename T, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(min)(T src0, T src1,
                                                            Sat sat = {}) {
  simd<T, 1> Src0 = src0;
  simd<T, 1> Src1 = src1;
  simd<T, 1> Result = (esimd::min)(Src0, Src1, sat);
  return Result[0];
}
 
 
 
#if defined(__SYCL_DEVICE_ONLY__)
#define __ESIMD_VECTOR_IMPL(T, name, iname)                                    \
  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res =                 \
      __spirv_ocl_native_##iname<__ESIMD_DNS::__raw_t<T>, N>(src.data());      \
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)                       \
    return res;                                                                \
  else                                                                         \
    return esimd::saturate<T>(simd<T, N>(res));
#define __ESIMD_SCALAR_IMPL(T, name, iname)                                    \
  __ESIMD_DNS::__raw_t<T> res =                                                \
      __spirv_ocl_native_##iname<__ESIMD_DNS::__raw_t<T>>(src);                \
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)                       \
    return res;                                                                \
  else                                                                         \
    return esimd::saturate<T>(simd<T, 1>(res))[0];
#else
#define __ESIMD_VECTOR_IMPL(T, name, iname) return 0;
#define __ESIMD_SCALAR_IMPL(T, name, iname) return 0;
#endif // __SYCL_DEVICE_ONLY__
 
#define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname)                         \
 \
  template <class T, int N, class Sat = saturation_off_tag,                    \
            class = std::enable_if_t<COND>>                                    \
  __ESIMD_API simd<T, N> name(simd<T, N> src, Sat sat = {}) {                  \
    __ESIMD_VECTOR_IMPL(T, name, iname)                                        \
  }                                                                            \
                                                                               \
 \
  template <typename T, class Sat = saturation_off_tag,                        \
            class = std::enable_if_t<COND>>                                    \
  __ESIMD_API T name(T src, Sat sat = {}) {                                    \
    __ESIMD_SCALAR_IMPL(T, name, iname)                                        \
  }
 
#define __ESIMD_EMATH_IEEE_COND                                                \
  detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4)
 
#define __ESIMD_EMATH_SPIRV_COND                                               \
  std::is_same_v<T, float> || std::is_same_v<T, sycl::half>
 
__ESIMD_UNARY_INTRINSIC_DEF(detail::is_generic_floating_point_v<T>, inv, recip)
 
__ESIMD_UNARY_INTRINSIC_DEF(__ESIMD_EMATH_SPIRV_COND, log2, log2)
 
__ESIMD_UNARY_INTRINSIC_DEF(__ESIMD_EMATH_SPIRV_COND, exp2, exp2)
 
__ESIMD_UNARY_INTRINSIC_DEF(detail::is_generic_floating_point_v<T>, sqrt, sqrt)
 
template <class T, int N, class Sat = saturation_off_tag,
          class = std::enable_if_t<__ESIMD_EMATH_IEEE_COND>>
__ESIMD_API simd<T, N> sqrt_ieee(simd<T, N> src, Sat sat = {}) {
  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res =
      __esimd_ieee_sqrt<T, N>(src.data());
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return res;
  else
    return esimd::saturate<T>(simd<T, N>(res));
}
 
template <typename T, class Sat = saturation_off_tag,
          class = std::enable_if_t<__ESIMD_EMATH_IEEE_COND>>
__ESIMD_API T sqrt_ieee(T src, Sat sat = {}) {
  simd<T, 1> src_vec = src;
  simd<T, 1> res = sqrt_ieee<T, 1>(src_vec, sat);
  return res[0];
}
 
__ESIMD_UNARY_INTRINSIC_DEF(__ESIMD_EMATH_SPIRV_COND, rsqrt, rsqrt)
 
__ESIMD_UNARY_INTRINSIC_DEF(__ESIMD_EMATH_SPIRV_COND, sin, sin)
 
__ESIMD_UNARY_INTRINSIC_DEF(__ESIMD_EMATH_SPIRV_COND, cos, cos)
 
template <class T, int N, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_same_v<T, double>, simd<double, N>>
rsqrt(simd<T, N> src, Sat sat = {}) {
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return inv(sqrt(src));
  else
    return esimd::saturate<double>(inv(sqrt(src)));
}
 
template <class T, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_same_v<T, double>, double>
rsqrt(T src, Sat sat = {}) {
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return inv(sqrt(src));
  else
    return esimd::saturate<double>(inv(sqrt(src)));
}
 
#undef __ESIMD_UNARY_INTRINSIC_DEF
#undef __ESIMD_VECTOR_IMPL
#undef __ESIMD_SCALAR_IMPL
 
#define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname)                        \
 \
  template <class T, int N, class U, class Sat = saturation_off_tag,           \
            class = std::enable_if_t<COND>>                                    \
  __ESIMD_API simd<T, N> name(simd<T, N> src0, simd<U, N> src1,                \
                              Sat sat = {}) {                                  \
    using RawVecT = __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N>;    \
    RawVecT src1_raw_conv = detail::convert_vector<T, U, N>(src1.data());      \
    RawVecT res_raw = __esimd_##iname<T, N>(src0.data(), src1_raw_conv);       \
    if constexpr (std::is_same_v<Sat, saturation_off_tag>)                     \
      return res_raw;                                                          \
    else                                                                       \
      return esimd::saturate<T>(simd<T, N>(res_raw));                          \
  }                                                                            \
                                                                               \
 \
  template <class T, int N, class U, class Sat = saturation_off_tag,           \
            class = std::enable_if_t<COND>>                                    \
  __ESIMD_API simd<T, N> name(simd<T, N> src0, U src1, Sat sat = {}) {         \
    return name<T, N, U>(src0, simd<U, N>(src1), sat);                         \
  }                                                                            \
                                                                               \
 \
  template <class T, class U, class Sat = saturation_off_tag,                  \
            class = std::enable_if_t<COND>>                                    \
  __ESIMD_API T name(T src0, U src1, Sat sat = {}) {                           \
    simd<T, 1> res = name<T, 1, U>(simd<T, 1>(src0), simd<U, 1>(src1), sat);   \
    return res[0];                                                             \
  }
 
template <class T, int N, class U, class Sat = saturation_off_tag,
          class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
__ESIMD_API simd<T, N> pow(simd<T, N> src0, simd<U, N> src1, Sat sat = {}) {
#if defined(__SYCL_DEVICE_ONLY__)
  using RawVecT = __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N>;
  RawVecT src1_raw_conv = detail::convert_vector<T, U, N>(src1.data());
  RawVecT res_raw = __spirv_ocl_native_powr<__ESIMD_DNS::__raw_t<T>, N>(
      src0.data(), src1_raw_conv);
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return res_raw;
  else
    return esimd::saturate<T>(simd<T, N>(res_raw));
#else
  return 0;
#endif // __SYCL_DEVICE_ONLY__
}
 
template <class T, int N, class U, class Sat = saturation_off_tag,
          class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
__ESIMD_API simd<T, N> pow(simd<T, N> src0, U src1, Sat sat = {}) {
  return pow<T, N, U>(src0, simd<U, N>(src1), sat);
}
 
template <class T, class U, class Sat = saturation_off_tag,
          class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
__ESIMD_API T pow(T src0, U src1, Sat sat = {}) {
#if defined(__SYCL_DEVICE_ONLY__)
  using ResT = __ESIMD_DNS::__raw_t<T>;
  ResT src1_raw_conv = detail::convert_scalar<T, U>(src1);
  ResT res_raw =
      __spirv_ocl_native_powr<__ESIMD_DNS::__raw_t<T>>(src0, src1_raw_conv);
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return res_raw;
  else
    return esimd::saturate<T>(simd<T, 1>(res_raw))[0];
#else
  return 0;
#endif // __SYCL_DEVICE_ONLY__
}
 
__ESIMD_BINARY_INTRINSIC_DEF(__ESIMD_EMATH_IEEE_COND, div_ieee, ieee_div)
 
#undef __ESIMD_BINARY_INTRINSIC_DEF
#undef __ESIMD_EMATH_IEEE_COND
#undef __ESIMD_EMATH_SPIRV_COND
 
 
 
namespace detail {
// std::numbers::ln2_v<float> in c++20
constexpr float ln2 = 0.69314718f;
// std::numbers::log2e_v<float> in c++20
constexpr float log2e = 1.442695f;
} // namespace detail
 
template <class T, int SZ, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> log(simd<T, SZ> src0, Sat sat = {}) {
  using CppT = __ESIMD_DNS::__cpp_t<T>;
  simd<T, SZ> Result =
      esimd::log2<T, SZ, saturation_off_tag>(src0) * detail::ln2;
 
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return Result;
  else
    return esimd::saturate<T>(Result);
}
 
template <class T, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE T log(T src0, Sat sat = {}) {
  return esimd::log<T, 1>(src0, sat)[0];
}
 
template <class T, int SZ, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> exp(simd<T, SZ> src0, Sat sat = {}) {
  using CppT = __ESIMD_DNS::__cpp_t<T>;
  return esimd::exp2<T, SZ>(src0 * detail::log2e, sat);
}
 
template <class T, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE T exp(T src0, Sat sat = {}) {
  return esimd::exp<T, 1>(src0, sat)[0];
}
 
 
 
// Rounding intrinsics.
 
#define __ESIMD_INTRINSIC_DEF(name)                                            \
 \
 \
                                    \
 \
 \
 \
  template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>    \
  __ESIMD_API __ESIMD_NS::simd<T, SZ> name(__ESIMD_NS::simd<float, SZ> src0,   \
                                           Sat sat = {}) {                     \
    __ESIMD_NS::simd<float, SZ> Result = __esimd_##name<SZ>(src0.data());      \
    if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)         \
      return Result;                                                           \
    else if constexpr (!std::is_same_v<float, T>) {                            \
      auto RawRes = __ESIMD_NS::saturate<float>(Result).data();                \
      return __ESIMD_DNS::convert_vector<T, float, SZ>(std::move(RawRes));     \
    } else {                                                                   \
      return __ESIMD_NS::saturate<T>(Result);                                  \
    }                                                                          \
  }                                                                            \
 \
  template <typename T, class Sat = __ESIMD_NS::saturation_off_tag>            \
  __ESIMD_API T name(float src0, Sat sat = {}) {                               \
    __ESIMD_NS::simd<float, 1> Src0 = src0;                                    \
    __ESIMD_NS::simd<T, 1> Result = name<T>(Src0, sat);                        \
    return Result[0];                                                          \
  }
 
__ESIMD_INTRINSIC_DEF(rndd)
 
__ESIMD_INTRINSIC_DEF(rndu)
 
__ESIMD_INTRINSIC_DEF(rnde)
 
__ESIMD_INTRINSIC_DEF(rndz)
 
#undef __ESIMD_INTRINSIC_DEF
 
 
template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
floor(const __ESIMD_NS::simd<float, SZ> src0, Sat sat = {}) {
  return esimd::rndd<RT, SZ>(src0, sat);
}
 
template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_INLINE RT floor(float src0, Sat sat = {}) {
  return esimd::rndd<RT, 1U>(src0, sat)[0];
}
 
template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
ceil(const __ESIMD_NS::simd<float, SZ> src0, Sat sat = {}) {
  return esimd::rndu<RT, SZ>(src0, sat);
}
 
template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_INLINE RT ceil(float src0, Sat sat = {}) {
  return esimd::rndu<RT, 1U>(src0, sat);
}
 
template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<RT, SZ>
trunc(const __ESIMD_NS::simd<float, SZ> &src0, Sat sat = {}) {
  return esimd::rndz<RT, SZ>(src0, sat);
}
 
template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API RT trunc(float src0, Sat sat = {}) {
  return esimd::rndz<RT, 1U>(src0, sat)[0];
}
 
 
 
template <int N>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<(N == 8 || N == 16 || N == 32), uint>
    pack_mask(simd_mask<N> src0) {
  return __esimd_pack_mask<N>(src0.data());
}
 
template <int N>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<(N == 8 || N == 16 || N == 32), simd_mask<N>>
    unpack_mask(uint src0) {
  return __esimd_unpack_mask<N>(src0);
}
 
template <int N>
__ESIMD_API std::enable_if_t<(N != 8 && N != 16 && N < 32), uint>
pack_mask(simd_mask<N> src0) {
  simd_mask<(N < 8 ? 8 : N < 16 ? 16 : 32)> src_0 = 0;
  src_0.template select<N, 1>() = src0.template bit_cast_view<ushort>();
  return esimd::pack_mask(src_0);
}
 
template <typename T, int N>
__ESIMD_API
    std::enable_if_t<(std::is_same_v<T, ushort> || std::is_same_v<T, uint>) &&
                         (N > 0 && N <= 32),
                     uint>
    ballot(simd<T, N> mask) {
  simd_mask<N> cmp = (mask != 0);
  if constexpr (N == 8 || N == 16 || N == 32) {
    return __esimd_pack_mask<N>(cmp.data());
  } else {
    constexpr int N1 = (N <= 8 ? 8 : N <= 16 ? 16 : 32);
    simd<uint16_t, N1> res = 0;
    res.template select<N, 1>() = cmp.data();
    return __esimd_pack_mask<N1>(res.data());
  }
}
 
template <typename T, int N>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
    std::is_integral<T>::value && (sizeof(T) <= 4), simd<uint32_t, N>>
cbit(simd<T, N> src) {
  return __esimd_cbit<T, N>(src.data());
}
 
template <typename T>
__ESIMD_API
    std::enable_if_t<std::is_integral<T>::value && (sizeof(T) <= 4), uint32_t>
    cbit(T src) {
  simd<T, 1> Src = src;
  simd<uint32_t, 1> Result = esimd::cbit(Src);
  return Result[0];
}
 
template <typename BaseTy, typename RegionTy>
__ESIMD_API std::enable_if_t<
    std::is_integral<
        typename simd_view<BaseTy, RegionTy>::element_type>::value &&
        (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) <= 4) &&
        (simd_view<BaseTy, RegionTy>::length == 1),
    uint32_t>
cbit(simd_view<BaseTy, RegionTy> src) {
  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
  simd<Ty, 1> Src = src;
  simd<uint32_t, 1> Result = esimd::cbit(Src);
  return Result[0];
}
 
template <typename T, int N>
__ESIMD_API
    std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), simd<T, N>>
    fbl(simd<T, N> src) {
  return __esimd_fbl<T, N>(src.data());
}
 
template <typename T>
__ESIMD_API std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), T>
fbl(T src) {
  simd<T, 1> Src = src;
  simd<T, 1> Result = esimd::fbl(Src);
  return Result[0];
}
 
template <typename BaseTy, typename RegionTy>
__ESIMD_API std::enable_if_t<
    std::is_integral<
        typename simd_view<BaseTy, RegionTy>::element_type>::value &&
        (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) == 4) &&
        (simd_view<BaseTy, RegionTy>::length == 1),
    typename simd_view<BaseTy, RegionTy>::element_type>
fbl(simd_view<BaseTy, RegionTy> src) {
  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
  simd<Ty, 1> Src = src;
  simd<Ty, 1> Result = esimd::fbl(Src);
  return Result[0];
}
 
template <typename T, int N>
__ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
                                 std::is_signed<T>::value && (sizeof(T) == 4),
                             simd<T, N>>
fbh(simd<T, N> src) {
  return __esimd_sfbh<T, N>(src.data());
}
 
template <typename T, int N>
__ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
                                 !std::is_signed<T>::value && (sizeof(T) == 4),
                             simd<T, N>>
fbh(simd<T, N> src) {
  return __esimd_ufbh<T, N>(src.data());
}
 
template <typename T>
__ESIMD_API std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), T>
fbh(T src) {
  simd<T, 1> Src = src;
  simd<T, 1> Result = esimd::fbh(Src);
  return Result[0];
}
 
template <typename BaseTy, typename RegionTy>
__ESIMD_API std::enable_if_t<
    std::is_integral<
        typename simd_view<BaseTy, RegionTy>::element_type>::value &&
        (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) == 4) &&
        (simd_view<BaseTy, RegionTy>::length == 1),
    typename simd_view<BaseTy, RegionTy>::element_type>
fbh(simd_view<BaseTy, RegionTy> src) {
  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
  simd<Ty, 1> Src = src;
  simd<Ty, 1> Result = esimd::fbh(Src);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, class Sat = saturation_off_tag>
__ESIMD_API
    std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
                     simd<T0, SZ>>
    shl(simd<T1, SZ> src0, simd<T1, SZ> src1, Sat sat = {}) {
  using ComputationTy =
      __ESIMD_DNS::computation_type_t<decltype(src0), int32_t>;
  ComputationTy Src0 = src0;
  ComputationTy Src1 = src1;
 
  if constexpr (std::is_same_v<Sat, saturation_on_tag>) {
    if constexpr (std::is_unsigned<T0>::value) {
      if constexpr (std::is_unsigned<
                        typename ComputationTy::element_type>::value)
        return __esimd_uushl_sat<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
      else
        return __esimd_usshl_sat<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
    } else {
      if constexpr (std::is_signed<typename ComputationTy::element_type>::value)
        return __esimd_sushl_sat<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
      else
        return __esimd_ssshl_sat<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
    }
  } else {
    if constexpr (std::is_unsigned<T0>::value) {
      if constexpr (std::is_unsigned<
                        typename ComputationTy::element_type>::value)
        return __esimd_uushl<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
      else
        return __esimd_usshl<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
    } else {
      if constexpr (std::is_signed<typename ComputationTy::element_type>::value)
        return __esimd_sushl<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
      else
        return __esimd_ssshl<T0, typename ComputationTy::element_type, SZ>(
            Src0.data(), Src1.data());
    }
  }
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             simd<T0, SZ>>
shl(simd<T1, SZ> src0, U src1, Sat sat = {}) {
  simd<U, SZ> Src1 = src1;
  return shl<T0, T1, SZ>(src0, Src1, sat);
}
 
template <typename T0, typename T1, typename T2, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
                                 std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<T2>::value,
                             std::remove_const_t<T0>>
shl(T1 src0, T2 src1, Sat sat = {}) {
  simd<T1, 1> Src0 = src0;
  simd<T0, 1> Result = shl<T0, T1, 1, T2, Sat>(Src0, src1, sat);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, class Sat = saturation_off_tag>
__ESIMD_API
    std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
                     simd<T0, SZ>>
    lsr(simd<T1, SZ> src0, simd<T1, SZ> src1, Sat sat = {}) {
  using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
  typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
  simd<ComputationTy, SZ> Src0 = src0;
  simd<ComputationTy, SZ> Src1 = src1;
  simd<ComputationTy, SZ> Result = Src0.data() >> Src1.data();
 
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return Result;
  else
    return saturate<T0>(Result);
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             simd<T0, SZ>>
lsr(simd<T1, SZ> src0, U src1, Sat sat = {}) {
  simd<T1, SZ> Src1 = src1;
  return lsr<T0, T1, SZ>(src0, Src1, sat);
}
 
template <typename T0, typename T1, typename T2, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
                                 std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<T2>::value,
                             std::remove_const_t<T0>>
lsr(T1 src0, T2 src1, Sat sat = {}) {
  simd<T1, 1> Src0 = src0;
  simd<T0, 1> Result = lsr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
 
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, class Sat = saturation_off_tag>
__ESIMD_API
    std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
                     simd<T0, SZ>>
    asr(simd<T1, SZ> src0, simd<T1, SZ> src1, Sat sat = {}) {
  using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
  typedef typename std::make_signed<IntermedTy>::type ComputationTy;
  simd<ComputationTy, SZ> Src0 = src0;
  simd<ComputationTy, SZ> Src1 = src1;
  simd<ComputationTy, SZ> Result = Src0 >> Src1;
  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
    return Result;
  else
    return saturate<T0>(Result);
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             simd<T0, SZ>>
asr(simd<T1, SZ> src0, U src1, Sat sat = {}) {
  simd<U, SZ> Src1 = src1;
  return asr<T0, T1, SZ>(src0, Src1, sat);
}
 
template <typename T0, typename T1, typename T2, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
                                 std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<T2>::value,
                             std::remove_const_t<T0>>
asr(T1 src0, T2 src1, Sat sat = {}) {
  simd<T1, 1> Src0 = src0;
  simd<T0, 1> Result = esimd::asr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, class Sat = saturation_off_tag>
__ESIMD_API
    std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
                     simd<T0, SZ>>
    shr(simd<T1, SZ> src0, simd<T1, SZ> src1, Sat sat = {}) {
  if constexpr (std::is_unsigned<T1>::value) {
    return esimd::lsr<T0, T1, SZ>(src0, src1, sat);
  } else {
    return esimd::asr<T0, T1, SZ>(src0, src1, sat);
  }
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             simd<T0, SZ>>
shr(simd<T1, SZ> src0, U src1, Sat sat = {}) {
  simd<U, SZ> Src1 = src1;
  return shr<T0, T1, SZ>(src0, Src1, sat);
}
 
template <typename T0, typename T1, typename T2, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
                                 std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<T2>::value,
                             std::remove_const_t<T0>>
shr(T1 src0, T2 src1, Sat sat = {}) {
  simd<T1, 1> Src0 = src0;
  simd<T0, 1> Result = shr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ>
__ESIMD_API std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t,
                                             uint32_t, int64_t, uint64_t>() &&
                                 detail::is_type<T1, int16_t, uint16_t, int32_t,
                                                 uint32_t, int64_t, uint64_t>(),
                             simd<T0, SZ>>
rol(simd<T1, SZ> src0, simd<T1, SZ> src1) {
  return __esimd_rol<T0, T1, SZ>(src0.data(), src1.data());
}
 
template <typename T0, typename T1, int SZ, typename U>
__ESIMD_API
    std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                     int64_t, uint64_t>() &&
                         detail::is_type<T1, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>() &&
                         detail::is_type<U, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>(),
                     simd<T0, SZ>>
    rol(simd<T1, SZ> src0, U src1) {
  simd<T1, SZ> Src1 = src1;
  return rol<T0>(src0, Src1);
}
 
template <typename T0, typename T1, typename T2>
__ESIMD_API
    std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                         __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                         __ESIMD_DNS::is_esimd_scalar<T2>::value &&
                         detail::is_type<T0, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>() &&
                         detail::is_type<T1, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>() &&
                         detail::is_type<T2, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>(),
                     std::remove_const_t<T0>>
    rol(T1 src0, T2 src1) {
  simd<T1, 1> Src0 = src0;
  simd<T0, 1> Result = rol<T0, T1, 1, T2>(Src0, src1);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ>
__ESIMD_API std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t,
                                             uint32_t, int64_t, uint64_t>() &&
                                 detail::is_type<T1, int16_t, uint16_t, int32_t,
                                                 uint32_t, int64_t, uint64_t>(),
                             simd<T0, SZ>>
ror(simd<T1, SZ> src0, simd<T1, SZ> src1) {
  return __esimd_ror<T0, T1, SZ>(src0.data(), src1.data());
}
 
template <typename T0, typename T1, int SZ, typename U>
__ESIMD_API
    std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                     int64_t, uint64_t>() &&
                         detail::is_type<T1, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>() &&
                         detail::is_type<U, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>(),
                     simd<T0, SZ>>
    ror(simd<T1, SZ> src0, U src1) {
  simd<T1, SZ> Src1 = src1;
  return esimd::ror<T0>(src0, Src1);
}
 
template <typename T0, typename T1, typename T2>
__ESIMD_API
    std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                         __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                         __ESIMD_DNS::is_esimd_scalar<T2>::value &&
                         detail::is_type<T0, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>() &&
                         detail::is_type<T1, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>() &&
                         detail::is_type<T2, int16_t, uint16_t, int32_t,
                                         uint32_t, int64_t, uint64_t>(),
                     std::remove_const_t<T0>>
    ror(T1 src0, T2 src1) {
  simd<T1, 1> Src0 = src0;
  simd<T0, 1> Result = esimd::ror<T0, T1, 1, T2>(Src0, src1);
  return Result[0];
}
 
 
 
template <typename T1, typename T2, typename T3, typename T4, int N,
          class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<
    detail::is_dword_type<T1>::value && detail::is_dword_type<T2>::value &&
        detail::is_dword_type<T3>::value && detail::is_dword_type<T4>::value,
    simd<T1, N>>
dp4a(simd<T2, N> src0, simd<T3, N> src1, simd<T4, N> src2, Sat sat = {}) {
#if defined(__SYCL_DEVICE_ONLY__)
  simd<T1, N> Result;
  simd<T2, N> Src0 = src0;
  simd<T3, N> Src1 = src1;
  simd<T4, N> Src2 = src2;
  if constexpr (std::is_same_v<Sat, saturation_off_tag>) {
    if constexpr (std::is_unsigned<T1>::value) {
      if constexpr (std::is_unsigned<T2>::value) {
        Result = __esimd_uudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                   Src2.data());
      } else {
        Result = __esimd_usdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                   Src2.data());
      }
    } else {
      if constexpr (std::is_unsigned<T2>::value) {
        Result = __esimd_sudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                   Src2.data());
      } else {
        Result = __esimd_ssdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                   Src2.data());
      }
    }
  } else {
    if constexpr (std::is_unsigned<T1>::value) {
      if constexpr (std::is_unsigned<T2>::value) {
        Result = __esimd_uudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                       Src2.data());
      } else {
        Result = __esimd_usdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                       Src2.data());
      }
    } else {
      if constexpr (std::is_unsigned<T2>::value) {
        Result = __esimd_sudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                       Src2.data());
      } else {
        Result = __esimd_ssdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
                                                       Src2.data());
      }
    }
  }
  return Result;
#else
  __ESIMD_UNSUPPORTED_ON_HOST;
#endif // __SYCL_DEVICE_ONLY__
}
 
// reduction functions
namespace detail {
template <typename T0, typename T1, int SZ> struct esimd_apply_sum {
  template <typename... T>
  simd<T0, SZ> operator()(simd<T1, SZ> v1, simd<T1, SZ> v2) {
    return v1 + v2;
  }
};
 
template <typename T0, typename T1, int SZ> struct esimd_apply_prod {
  template <typename... T>
  simd<T0, SZ> operator()(simd<T1, SZ> v1, simd<T1, SZ> v2) {
    return v1 * v2;
  }
};
 
template <typename T0, typename T1, int SZ> struct esimd_apply_reduced_max {
  template <typename... T>
  simd<T0, SZ> operator()(simd<T1, SZ> v1, simd<T1, SZ> v2) {
    return __ESIMD_DNS::convert_vector<T0, T1, SZ>(
        __ESIMD_NS::max(v1, v2).data());
  }
};
 
template <typename T0, typename T1, int SZ> struct esimd_apply_reduced_min {
  template <typename... T>
  simd<T0, SZ> operator()(simd<T1, SZ> v1, simd<T1, SZ> v2) {
    return __ESIMD_DNS::convert_vector<T0, T1, SZ>(
        __ESIMD_NS::min(v1, v2).data());
  }
};
 
template <typename T0, typename T1, int SZ,
          template <typename RT, typename T, int N> class OpType>
T0 reduce_single(simd<T1, SZ> v) {
  if constexpr (SZ == 1) {
    return v[0];
  } else {
    static_assert(detail::isPowerOf2(SZ),
                  "Invaid input for reduce_single - the vector size must "
                  "be power of two.");
    constexpr int N = SZ / 2;
    simd<T0, N> tmp = OpType<T0, T1, N>()(v.template select<N, 1>(0),
                                          v.template select<N, 1>(N));
    return reduce_single<T0, T0, N, OpType>(tmp);
  }
}
 
template <typename T0, typename T1, int N1, int N2,
          template <typename RT, typename T, int N> class OpType>
T0 reduce_pair(simd<T1, N1> v1, simd<T1, N2> v2) {
  if constexpr (N1 == N2) {
    simd<T0, N1> tmp = OpType<T0, T1, N1>()(v1, v2);
    return reduce_single<T0, T0, N1, OpType>(tmp);
  } else if constexpr (N1 < N2) {
    simd<T0, N1> tmp1 = OpType<T0, T1, N1>()(v1, v2.template select<N1, 1>(0));
    constexpr int N = N2 - N1;
    using NT = simd<T0, N>;
    NT tmp2 = convert<T0>(v2.template select<N, 1>(N1).read());
    return reduce_pair<T0, T0, N1, N, OpType>(tmp1, tmp2);
  } else {
    static_assert(detail::isPowerOf2(N1),
                  "Invaid input for reduce_pair - N1 must be power of two.");
    constexpr int N = N1 / 2;
    simd<T0, N> tmp = OpType<T0, T1, N>()(v1.template select<N, 1>(0),
                                          v1.template select<N, 1>(N));
    using NT = simd<T0, N2>;
    NT tmp2 = convert<T0>(v2);
    return reduce_pair<T0, T0, N, N2, OpType>(tmp, tmp2);
  }
}
 
template <typename T0, typename T1, int SZ,
          template <typename RT, typename T, int N> class OpType>
T0 reduce(simd<T1, SZ> v) {
  constexpr bool isPowerOf2 = detail::isPowerOf2(SZ);
  if constexpr (isPowerOf2) {
    return reduce_single<T0, T1, SZ, OpType>(v);
  } else {
    constexpr unsigned N1 = 1u << detail::log2<SZ>();
    constexpr unsigned N2 = SZ - N1;
 
    simd<T1, N1> v1 = v.template select<N1, 1>(0);
    simd<T1, N2> v2 = v.template select<N2, 1>(N1);
    return reduce_pair<T0, T1, N1, N2, OpType>(v1, v2);
  }
};
 
template <typename T0, typename T1, int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 sum(simd<T1, SZ> v) {
  using TT = detail::computation_type_t<simd<T1, SZ>>;
  using RT = typename TT::element_type;
  T0 retv = reduce<RT, T1, SZ, esimd_apply_sum>(v);
  return retv;
}
 
template <typename T0, typename T1, int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 prod(simd<T1, SZ> v) {
  using TT = detail::computation_type_t<simd<T1, SZ>>;
  using RT = typename TT::element_type;
  T0 retv = reduce<RT, T1, SZ, esimd_apply_prod>(v);
  return retv;
}
} // namespace detail
 
template <typename T0, typename T1, int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 hmax(simd<T1, SZ> v) {
  T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_max>(v);
  return retv;
}
 
template <typename T0, typename T1, int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {
  T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_min>(v);
  return retv;
}
 
// TODO 1) enforce BinaryOperation constraints 2) support std::minimum/maximum
template <typename T0, typename T1, int SZ, typename BinaryOperation>
ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
  if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
                             std::plus<>>::value) {
    T0 retv = detail::sum<T0>(v);
    return retv;
  } else if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
                                    std::multiplies<>>::value) {
    T0 retv = detail::prod<T0>(v);
    return retv;
  }
}
 
 
enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };
 
static constexpr bfn_t operator~(bfn_t x) {
  uint8_t val = static_cast<uint8_t>(x);
  uint8_t res = ~val;
  return static_cast<bfn_t>(res);
}
 
static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
  uint8_t arg0 = static_cast<uint8_t>(x);
  uint8_t arg1 = static_cast<uint8_t>(y);
  uint8_t res = arg0 | arg1;
  return static_cast<bfn_t>(res);
}
 
static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
  uint8_t arg0 = static_cast<uint8_t>(x);
  uint8_t arg1 = static_cast<uint8_t>(y);
  uint8_t res = arg0 & arg1;
  return static_cast<bfn_t>(res);
}
 
static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
  uint8_t arg0 = static_cast<uint8_t>(x);
  uint8_t arg1 = static_cast<uint8_t>(y);
  uint8_t res = arg0 ^ arg1;
  return static_cast<bfn_t>(res);
}
 
template <bfn_t FuncControl, typename T, int N>
__ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
    __ESIMD_NS::simd<T, N> src2) {
  if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
                ((sizeof(T) == 2) && (N % 2 == 0))) {
    // Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
    // Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
    // Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
    auto Result = __ESIMD_NS::bfn<FuncControl>(
        src0.template bit_cast_view<int32_t>().read(),
        src1.template bit_cast_view<int32_t>().read(),
        src2.template bit_cast_view<int32_t>().read());
    return Result.template bit_cast_view<T>();
  } else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
    constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
    return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
  } else if constexpr (N % 2 == 0) {
    // Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
    auto Result = __ESIMD_NS::bfn<FuncControl>(
        src0.template bit_cast_view<int16_t>().read(),
        src1.template bit_cast_view<int16_t>().read(),
        src2.template bit_cast_view<int16_t>().read());
    return Result.template bit_cast_view<T>();
  } else {
    // Odd number of 1-byte elements.
    __ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
    Src0.template select<N, 1>() = src0;
    Src1.template select<N, 1>() = src1;
    Src2.template select<N, 1>() = src2;
    auto Result = __ESIMD_NS::bfn<FuncControl>(Src0, Src1, Src2);
    return Result.template select<N, 1>();
  }
}
 
template <bfn_t FuncControl, typename T>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
    __ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
bfn(T src0, T src1, T src2) {
  __ESIMD_NS::simd<T, 1> Src0 = src0;
  __ESIMD_NS::simd<T, 1> Src1 = src1;
  __ESIMD_NS::simd<T, 1> Src2 = src2;
  __ESIMD_NS::simd<T, 1> Result =
      esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
  return Result[0];
}
 
 
template <int N>
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
     __ESIMD_NS::simd<uint32_t, N> src1) {
  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
            __ESIMD_DNS::vector_type_t<uint32_t, N>>
      Result = __esimd_addc<uint32_t, N>(src0.data(), src1.data());
 
  carry = Result.first;
  return Result.second;
}
 
template <int N>
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
     uint32_t src1) {
  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
  return addc(carry, src0, Src1V);
}
 
template <int N>
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
     __ESIMD_NS::simd<uint32_t, N> src1) {
  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
  return addc(carry, Src0V, src1);
}
 
__ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1) {
  __ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
  __ESIMD_NS::simd<uint32_t, 1> Res = addc(CarryV, Src0V, Src1V);
  carry = CarryV[0];
  return Res[0];
}
 
template <int N>
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
     __ESIMD_NS::simd<uint32_t, N> src1) {
  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
            __ESIMD_DNS::vector_type_t<uint32_t, N>>
      Result = __esimd_subb<uint32_t, N>(src0.data(), src1.data());
 
  borrow = Result.first;
  return Result.second;
}
 
template <int N>
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
     uint32_t src1) {
  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
  return subb(borrow, src0, Src1V);
}
 
template <int N>
__ESIMD_API __ESIMD_NS::simd<uint32_t, N>
subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
     __ESIMD_NS::simd<uint32_t, N> src1) {
  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
  return subb(borrow, Src0V, src1);
}
 
__ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
  __ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
  __ESIMD_NS::simd<uint32_t, 1> Res = subb(BorrowV, Src0V, Src1V);
  borrow = BorrowV[0];
  return Res[0];
}
 
__ESIMD_API uint64_t rdtsc() {
#ifdef __SYCL_DEVICE_ONLY__
  return __spirv_ReadClockKHR(0);
#else
  __ESIMD_UNSUPPORTED_ON_HOST;
#endif
}
 
template <typename T, int N>
__ESIMD_API __ESIMD_NS::simd<T, N> clamp(__ESIMD_NS::simd<T, N> src,
                                         __ESIMD_NS::simd<T, N> min_val,
                                         __ESIMD_NS::simd<T, N> max_val) {
  __ESIMD_NS::simd<T, N> Result = src;
  Result.merge(min_val, src < min_val);
  Result.merge(max_val, src > max_val);
  return Result;
}
 
template <typename T, int N>
__ESIMD_API __ESIMD_NS::simd<T, N> clamp(__ESIMD_NS::simd<T, N> src, T min_val,
                                         T max_val) {
  __ESIMD_NS::simd<T, N> MinVal = min_val;
  __ESIMD_NS::simd<T, N> MaxVal = max_val;
  return clamp(src, MinVal, MaxVal);
}
 
 
} // namespace ext::intel::esimd
} // namespace _V1
} // namespace sycl