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 experimental Explicit SIMD math APIs.
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>
#include <sycl/ext/intel/esimd/detail/types.hpp>
#include <sycl/ext/intel/esimd/math.hpp>
#include <sycl/ext/intel/experimental/esimd/common.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/util.hpp>
 
namespace sycl {
__SYCL_INLINE_VER_NAMESPACE(_V1) {
namespace ext::intel::experimental::esimd {
 
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             __ESIMD_NS::simd<T0, SZ>>
shl(__ESIMD_NS::simd<T1, SZ> src0, U 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, __ESIMD_NS::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, typename T2,
          class Sat = __ESIMD_NS::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 = {}) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result =
      esimd::shl<T0, T1, 1, T2, Sat>(Src0, src1, sat);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             __ESIMD_NS::simd<T0, SZ>>
lsr(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
  using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
  typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
  __ESIMD_NS::simd<ComputationTy, SZ> Src0 = src0;
  __ESIMD_NS::simd<ComputationTy, SZ> Src1 = src1;
  // TODO H/W supports saturation with this op - map to more efficient version.
  __ESIMD_NS::simd<ComputationTy, SZ> Result = Src0.data() >> Src1.data();
 
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T0>(Result);
}
 
template <typename T0, typename T1, typename T2,
          class Sat = __ESIMD_NS::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 = {}) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result =
      esimd::lsr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
 
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             __ESIMD_NS::simd<T0, SZ>>
asr(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
  using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
  typedef typename std::make_signed<IntermedTy>::type ComputationTy;
  __ESIMD_NS::simd<ComputationTy, SZ> Src0 = src0;
  // TODO H/W supports saturation with this op - map to more efficient version.
  __ESIMD_NS::simd<ComputationTy, SZ> Result = Src0 >> src1;
 
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T0>(Result);
}
 
template <typename T0, typename T1, typename T2,
          class Sat = __ESIMD_NS::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 = {}) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result =
      esimd::asr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value &&
                                 std::is_integral<U>::value,
                             __ESIMD_NS::simd<T0, SZ>>
shr(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
  if constexpr (std::is_unsigned<T1>::value) {
    return esimd::lsr<T0, T1, SZ, U, Sat>(src0, src1, sat);
  } else {
    return esimd::asr<T0, T1, SZ, U, Sat>(src0, src1, sat);
  }
}
 
template <typename T0, typename T1, typename T2,
          class Sat = __ESIMD_NS::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 = {}) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result =
      esimd::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<
    __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>(),
    __ESIMD_NS::simd<T0, SZ>>
rol(__ESIMD_NS::simd<T1, SZ> src0, __ESIMD_NS::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<
    __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<U, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>(),
    __ESIMD_NS::simd<T0, SZ>>
rol(__ESIMD_NS::simd<T1, SZ> src0, U src1) {
  __ESIMD_NS::simd<T1, SZ> Src1 = src1;
  return esimd::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 &&
        __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>() &&
        __ESIMD_NS::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) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result = esimd::rol<T0, T1, 1, T2>(Src0, src1);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ>
__ESIMD_API std::enable_if_t<
    __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>(),
    __ESIMD_NS::simd<T0, SZ>>
ror(__ESIMD_NS::simd<T1, SZ> src0, __ESIMD_NS::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<
    __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<U, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>(),
    __ESIMD_NS::simd<T0, SZ>>
ror(__ESIMD_NS::simd<T1, SZ> src0, U src1) {
  __ESIMD_NS::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 &&
        __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>() &&
        __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
                                    int64_t, uint64_t>() &&
        __ESIMD_NS::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) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result = esimd::ror<T0, T1, 1, T2>(Src0, src1);
  return Result[0];
}
 
 
 
// imul
#ifndef ESIMD_HAS_LONG_LONG
// use mulh instruction for high half
template <typename T0, typename T1, typename U, int SZ>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value &&
                                      __ESIMD_DNS::is_dword_type<T1>::value &&
                                      __ESIMD_DNS::is_dword_type<U>::value,
                                  __ESIMD_NS::simd<T0, SZ>>
    imul(__ESIMD_NS::simd<T0, SZ> &rmd, __ESIMD_NS::simd<T1, SZ> src0, U src1) {
  using ComputationTy = __ESIMD_DNS::computation_type_t<decltype(src0), U>;
  ComputationTy Src0 = src0;
  ComputationTy Src1 = src1;
  rmd = Src0 * Src1;
  if constexpr (std::is_unsigned<T0>::value)
    return __esimd_umulh(Src0.data(), Src1.data());
  else
    return __esimd_smulh(Src0.data(), Src1.data());
}
 
#else
// imul bdw+ version: use qw=dw*dw multiply.
// We need to special case SZ==1 to avoid "error: when select size is 1, the
// stride must also be 1" on the selects.
template <typename T0, typename T1, typename U, int SZ>
__ESIMD_API
    std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value &&
                         __ESIMD_DNS::is_dword_type<T1>::value &&
                         __ESIMD_DNS::is_dword_type<U>::value && SZ == 1,
                     __ESIMD_NS::simd<T0, SZ>>
    imul(__ESIMD_NS::simd<T0, SZ> &rmd, __ESIMD_NS::simd<T1, SZ> src0, U src1) {
  using ComputationTy =
      __ESIMD_DNS::computation_type_t<decltype(rmd), long long>;
  ComputationTy Product = convert<long long>(src0);
  Product *= src1;
  rmd = Product.bit_cast_view<T0>().select<1, 1>[0];
  return Product.bit_cast_view<T0>().select<1, 1>[1];
}
 
template <typename T0, typename T1, typename U, int SZ>
__ESIMD_API
    std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value &&
                         __ESIMD_DNS::is_dword_type<T1>::value &&
                         __ESIMD_DNS::is_dword_type<U>::value && SZ != 1,
                     __ESIMD_NS::simd<T0, SZ>>
    imul(__ESIMD_NS::simd<T0, SZ> &rmd, __ESIMD_NS::simd<T1, SZ> src0, U src1) {
  using ComputationTy =
      __ESIMD_DNS::computation_type_t<decltype(rmd), long long>;
  ComputationTy Product = convert<long long>(src0);
  Product *= src1;
  rmd = Product.bit_cast_view<T0>().select<SZ, 2>(0);
  return Product.bit_cast_view<T0>().select<SZ, 2>(1);
}
#endif
 
// TODO: document
template <typename T0, typename T1, typename U, int SZ>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<U>::value,
                             __ESIMD_NS::simd<T0, SZ>>
imul(__ESIMD_NS::simd<T0, SZ> &rmd, U src0, __ESIMD_NS::simd<T1, SZ> src1) {
  return esimd::imul(rmd, src1, src0);
}
 
// TODO: document
template <typename T0, typename T, typename U>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<U>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T0>::value,
                                  T0>
    imul(__ESIMD_NS::simd<T0, 1> &rmd, T src0, U src1) {
  __ESIMD_NS::simd<T, 1> src_0 = src0;
  __ESIMD_NS::simd<U, 1> src_1 = src1;
  __ESIMD_NS::simd<T0, 1> res =
      esimd::imul(rmd, src_0.select_all(), src_1.select_all());
  return res[0];
}
 
template <typename T, int SZ, typename U>
__ESIMD_API
    std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
                     __ESIMD_NS::simd<T, SZ>>
    quot(__ESIMD_NS::simd<T, SZ> src0, U src1) {
  return src0 / src1;
}
 
template <typename T0, typename T1>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                                 std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value,
                             std::remove_const_t<T0>>
quot(T0 src0, T1 src1) {
  return src0 / src1;
}
 
template <typename T, int SZ, typename U>
__ESIMD_API
    std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
                     __ESIMD_NS::simd<T, SZ>>
    mod(__ESIMD_NS::simd<T, SZ> src0, U src1) {
  return src0 % src1;
}
 
template <typename T0, typename T1>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
                                 std::is_integral<T0>::value &&
                                 std::is_integral<T1>::value,
                             std::remove_const_t<T0>>
mod(T0 src0, T1 src1) {
  return src0 % src1;
}
 
template <typename T, int SZ, typename U>
__ESIMD_API
    std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
                     __ESIMD_NS::simd<T, SZ>>
    div(__ESIMD_NS::simd<T, SZ> &remainder, __ESIMD_NS::simd<T, SZ> src0,
        U src1) {
  remainder = src0 % src1;
  return src0 / src1;
}
 
template <typename T, int SZ, typename U>
__ESIMD_API
    std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value &&
                         __ESIMD_DNS::is_esimd_scalar<U>::value,
                     __ESIMD_NS::simd<T, SZ>>
    div(__ESIMD_NS::simd<T, SZ> &remainder, U src0,
        __ESIMD_NS::simd<T, SZ> src1) {
  remainder = src0 % src1;
  return src0 / src1;
}
 
template <typename RT, typename T0, typename T1>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<RT>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T1>::value,
                                  std::remove_const_t<RT>>
    div(__ESIMD_NS::simd<std::remove_const_t<RT>, 1> &remainder, T0 src0,
        T1 src1) {
  remainder[0] = src0 % src1;
  return src0 / src1;
}
 
// Dot product builtins
#if defined(ESIMD_GEN7_5) || defined(ESIMD_GEN8) || defined(ESIMD_GEN8_5) ||   \
    defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<T0, SZ> dp2(__ESIMD_NS::simd<T1, SZ> src0, U src1,
                                         Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
  __ESIMD_NS::simd<float, SZ> Src0 = src0;
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp2(Src0.data(), Src1.data());
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T0>(Result);
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<T0, SZ> dp3(__ESIMD_NS::simd<T1, SZ> src0, U src1,
                                         Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
  __ESIMD_NS::simd<float, SZ> Src0 = src0;
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp3(Src0.data(), Src1.data());
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T0>(Result);
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<T0, SZ> dp4(__ESIMD_NS::simd<T1, SZ> src0, U src1,
                                         Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
  __ESIMD_NS::simd<float, SZ> Src0 = src0;
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp4(Src0.data(), Src1.data());
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T0>(Result);
}
 
template <typename T0, typename T1, typename U, int SZ,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<T0, SZ> dph(__ESIMD_NS::simd<T1, SZ> src0, U src1,
                                         Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
  __ESIMD_NS::simd<float, SZ> Src0 = src0;
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result = __esimd_dph(Src0.data(), Src1.data());
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T0>(Result);
}
 
template <typename RT, typename T1, typename T2, int SZ,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<RT, SZ> line(__ESIMD_NS::simd<T1, 4> src0,
                                          __ESIMD_NS::simd<T2, SZ> src1,
                                          Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
  __ESIMD_NS::simd<float, 4> Src0 = src0;
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result = __esimd_line(Src0.data(), Src1.data());
 
  __ESIMD_NS::simd<RT, SZ> Result;
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<RT>(Result);
}
 
template <typename RT, typename T, int SZ,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<RT, SZ>
line(float P, float Q, __ESIMD_NS::simd<T, SZ> src1, Sat sat = {}) {
  __ESIMD_NS::simd<float, 4> Src0 = P;
  Src0(3) = Q;
  return esimd::line<RT>(Src0, src1, sat);
}
 
#else
// The old implementation is to generate vISA IRs for dp2/dp3/dp4/dph/line.
// Now We change to use direct mul/add, and hope to generate mad instructions
// at the end, to still get the performance as good as HW solution.
// We rely on "pragma unroll" to get better code.
// The only input and return types for these APIs are floats.
// In order to be able to use the old emu code, we keep the template argument
// for the type, although the type "T" can only be float.
// We use enable_if to force the float type only.
// If the gen is not specified we warn the programmer that they are potentially
// using a less efficient implementation if not on GEN10 or above.
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
                         std::is_floating_point<T1>::value &&
                         __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
                         std::is_floating_point<U>::value,
                     __ESIMD_NS::simd<T0, SZ>>
    dp2(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result;
#pragma unroll
  for (int i = 0; i < SZ; i += 4) {
    Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1];
  }
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T1>(Result);
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
                         std::is_floating_point<T1>::value &&
                         __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
                         std::is_floating_point<U>::value,
                     __ESIMD_NS::simd<T0, SZ>>
    dp3(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result;
#pragma unroll
  for (int i = 0; i < SZ; i += 4) {
    Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
                             src0[i + 2] * Src1[i + 2];
  }
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T1>(Result);
}
 
template <typename T0, typename T1, int SZ, typename U,
          class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
                         std::is_floating_point<T1>::value &&
                         __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
                         std::is_floating_point<U>::value,
                     __ESIMD_NS::simd<T0, SZ>>
    dp4(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
  __ESIMD_NS::simd<T1, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result;
#pragma unroll
  for (int i = 0; i < SZ; i += 4) {
    Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
                             src0[i + 2] * Src1[i + 2] +
                             src0[i + 3] * Src1[i + 3];
  }
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T1>(Result);
}
 
template <typename T, typename U, int SZ,
          class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
                         std::is_floating_point<T>::value &&
                         __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
                         std::is_floating_point<U>::value,
                     __ESIMD_NS::simd<T, SZ>>
    dph(__ESIMD_NS::simd<T, SZ> src0, U src1, Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Result;
#pragma unroll
  for (int i = 0; i < SZ; i += 4) {
    Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
                             src0[i + 2] * Src1[i + 2] + 1.0 * Src1[i + 3];
  }
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T>(Result);
}
 
template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
                                      std::is_floating_point<T>::value,
                                  __ESIMD_NS::simd<T, SZ>>
    line(__ESIMD_NS::simd<T, 4> src0, __ESIMD_NS::simd<T, SZ> src1,
         Sat sat = {}) {
  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
 
  __ESIMD_NS::simd<T, SZ> Src1 = src1;
  __ESIMD_NS::simd<T, SZ> Result;
#pragma unroll
  for (int i = 0; i < SZ; i += 4) {
    Result.select<4, 1>(i) = src0[0] * src1[i] + src0[3];
  }
 
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T>(Result);
}
 
template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
                                      std::is_floating_point<T>::value,
                                  __ESIMD_NS::simd<T, SZ>>
    line(float P, float Q, __ESIMD_NS::simd<T, SZ> src1, Sat sat = {}) {
  __ESIMD_NS::simd<T, 4> Src0 = P;
  Src0(3) = Q;
  return esimd::line<T>(Src0, src1, sat);
}
 
#endif
 
template <typename T, int SZ>
__ESIMD_API __ESIMD_NS::simd<T, SZ> frc(__ESIMD_NS::simd<T, SZ> src0) {
  __ESIMD_NS::simd<float, SZ> Src0 = src0;
  return __esimd_frc(Src0.data());
}
 
template <typename T> __ESIMD_API T frc(T src0) {
  __ESIMD_NS::simd<T, 1> Src0 = src0;
  __ESIMD_NS::simd<T, 1> Result = esimd::frc<T>(Src0);
  return Result[0];
}
 
// lzd
template <typename RT, typename T0, int SZ,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<RT, SZ> lzd(__ESIMD_NS::simd<T0, SZ> src0,
                                         Sat sat = {}) {
  // Saturation parameter ignored
  __ESIMD_NS::simd<__ESIMD_NS::uint, SZ> Src0 = src0;
  return __esimd_lzd<__ESIMD_NS::uint, SZ>(Src0.data());
}
 
template <typename RT, typename T0, class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<RT>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T0>::value,
                                  std::remove_const_t<RT>>
    lzd(T0 src0, Sat sat = {}) {
  __ESIMD_NS::simd<T0, 1> Src0 = src0;
  __ESIMD_NS::simd<RT, 1> Result = esimd::lzd<RT>(Src0);
  return Result[0];
}
 
// lrp
#if defined(ESIMD_GEN7_5) || defined(ESIMD_GEN8) || defined(ESIMD_GEN8_5) ||   \
    defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
 
template <int SZ, typename U, typename V,
          class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<float, SZ> lrp(__ESIMD_NS::simd<float, SZ> src0,
                                            U src1, V src2, Sat sat = {}) {
  static_assert(SZ >= 4 && (SZ & 0x3) == 0,
                "vector size must be a multiple of 4");
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Src2 = src2;
  __ESIMD_NS::simd<float, SZ> Result =
      __esimd_lrp<SZ>(src0.data(), Src1.data(), Src2.data());
 
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<float>(Result);
}
 
#else
 
// The old implementation is to generate vISA IRs for lrp.
// Now We change to use direct mul/add, and hope to generate mad instructions
// at the end, to still get the performance as good as HW solution.
// The only input and return types for these APIs are floats.
// In order to be able to use the old emu code, we keep the template argument
// for the type, although the type "T" can only be float.
// We use enable_if to force the float type only.
// If the gen is not specified we warn the programmer that they are potentially
// using less efficient implementation.
template <typename T, int SZ, typename U, typename V,
          class Sat = __ESIMD_NS::saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
                         std::is_floating_point<T>::value &&
                         __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
                         std::is_floating_point<U>::value,
                     __ESIMD_NS::simd<T, SZ>>
    lrp(__ESIMD_NS::simd<T, SZ> src0, U src1, V src2, Sat sat = {}) {
 
  __ESIMD_NS::simd<float, SZ> Src1 = src1;
  __ESIMD_NS::simd<float, SZ> Src2 = src2;
  __ESIMD_NS::simd<float, SZ> Result;
  Result = Src1 * src0 + Src2 * (1.0f - src0);
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return Result;
  else
    return __ESIMD_NS::saturate<T>(Result);
}
#endif
 
 
 
template <typename T0, typename T1, int SZ>
__ESIMD_API __ESIMD_NS::simd<T0, SZ> bf_reverse(__ESIMD_NS::simd<T1, SZ> src0) {
  __ESIMD_NS::simd<unsigned, SZ> Src0 = src0;
  return __esimd_bfrev<unsigned>(Src0.data());
}
 
template <typename T0, typename T1>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T1>::value,
                                  std::remove_const_t<T0>>
    bf_reverse(T1 src0) {
  __ESIMD_NS::simd<T1, 1> Src0 = src0;
  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_reverse<T0>(Src0);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, typename U, typename V, typename W>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<std::is_integral<T1>::value, __ESIMD_NS::simd<T0, SZ>>
    bf_insert(U src0, V src1, W src2, __ESIMD_NS::simd<T1, SZ> src3) {
  typedef typename __ESIMD_DNS::dword_type<T1> DT1;
  static_assert(std::is_integral<DT1>::value && sizeof(DT1) == sizeof(int),
                "operand conversion failed");
  __ESIMD_NS::simd<DT1, SZ> Src0 = src0;
  __ESIMD_NS::simd<DT1, SZ> Src1 = src1;
  __ESIMD_NS::simd<DT1, SZ> Src2 = src2;
  __ESIMD_NS::simd<DT1, SZ> Src3 = src3;
 
  return __esimd_bfi<DT1>(Src0.data(), Src1.data(), Src2.data(), Src3.data());
}
 
template <typename T0, typename T1, typename T2, typename T3, typename T4>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T4>::value,
                                  std::remove_const_t<T0>>
    bf_insert(T1 src0, T2 src1, T3 src2, T4 src3) {
  __ESIMD_NS::simd<T4, 1> Src3 = src3;
  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_insert<T0>(src0, src1, src2, Src3);
  return Result[0];
}
 
template <typename T0, typename T1, int SZ, typename U, typename V>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<std::is_integral<T1>::value, __ESIMD_NS::simd<T0, SZ>>
    bf_extract(U src0, V src1, __ESIMD_NS::simd<T1, SZ> src2) {
  typedef typename __ESIMD_DNS::dword_type<T1> DT1;
  static_assert(std::is_integral<DT1>::value && sizeof(DT1) == sizeof(int),
                "operand conversion failed");
  __ESIMD_NS::simd<DT1, SZ> Src0 = src0;
  __ESIMD_NS::simd<DT1, SZ> Src1 = src1;
  __ESIMD_NS::simd<DT1, SZ> Src2 = src2;
 
  return __esimd_sbfe<DT1>(Src0.data(), Src1.data(), Src2.data());
}
 
template <typename T0, typename T1, typename T2, typename T3>
ESIMD_NODEBUG
    ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
                                      __ESIMD_DNS::is_esimd_scalar<T3>::value,
                                  std::remove_const_t<T0>>
    bf_extract(T1 src0, T2 src1, T3 src2) {
  __ESIMD_NS::simd<T3, 1> Src2 = src2;
  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_extract<T0>(src0, src1, Src2);
  return Result[0];
}
 
 
 
// sincos
template <int SZ, typename U, class Sat = __ESIMD_NS::saturation_off_tag>
__ESIMD_API __ESIMD_NS::simd<float, SZ>
sincos(__ESIMD_NS::simd<float, SZ> &dstcos, U src0, Sat sat = {}) {
  dstcos = __ESIMD_NS::cos(src0, sat);
  return __ESIMD_NS::sin(src0, sat);
}
 
// atan
 
namespace detail {
constexpr double __ESIMD_CONST_PI = 3.1415926535897932384626433832795;
} // namespace detail
 
template <typename T, int SZ>
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<T, SZ>
atan(__ESIMD_NS::simd<T, SZ> src0) {
  static_assert(std::is_floating_point<T>::value,
                "Floating point argument type is expected.");
  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
 
  __ESIMD_NS::simd<T, SZ> OneP((T)1.0);
  __ESIMD_NS::simd<T, SZ> OneN((T)-1.0);
  __ESIMD_NS::simd<T, SZ> sign;
  __ESIMD_NS::simd_mask<SZ> Gt1 = Src0 > T(1.0);
 
  sign.merge(OneN, OneP, src0 < 0);
 
  Src0.merge(__ESIMD_NS::inv(Src0), Gt1);
 
  __ESIMD_NS::simd<T, SZ> Src0P2 = Src0 * Src0;
  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
 
  __ESIMD_NS::simd<T, SZ> Result =
      (Src0P4 * T(0.185696) + ((Src0 * T(0.787997) + T(0.63693)) * Src0P2) +
       Src0) /
      (((((Src0 * -T(0.000121387) + T(0.00202308)) * Src0P2) +
         (Src0 * -T(0.0149145)) + T(0.182569)) *
        Src0P4) +
       ((Src0 * T(0.395889) + T(1.12158)) * Src0P2) + (Src0 * T(0.636918)) +
       T(1.0));
 
  Result.merge(Result - T(detail::__ESIMD_CONST_PI) / T(2.0), Gt1);
 
  return __ESIMD_NS::abs(Result) * sign;
}
 
template <typename T> __ESIMD_API T atan(T src0) {
  static_assert(std::is_floating_point<T>::value,
                "Floating point argument type is expected.");
  __ESIMD_NS::simd<T, 1> Src0 = src0;
  __ESIMD_NS::simd<T, 1> Result = esimd::atan(Src0);
  return Result[0];
}
 
// acos
 
template <typename T, int SZ>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
    acos(__ESIMD_NS::simd<T, SZ> src0) {
  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
 
  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
  __ESIMD_NS::simd_mask<SZ> TooBig = Src0 >= T(0.999998);
 
  // Replace oversized values to ensure no possibility of sqrt of
  // a negative value later
  Src0.merge(T(0.0), TooBig);
 
  __ESIMD_NS::simd<T, SZ> Src01m = T(1.0) - Src0;
 
  __ESIMD_NS::simd<T, SZ> Src0P2 = Src01m * Src01m;
  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
 
  __ESIMD_NS::simd<T, SZ> Result =
      (((Src01m * T(0.015098965761299077) - T(0.005516443930088506)) * Src0P4) +
       ((Src01m * T(0.047654245891495528) + T(0.163910606547823220)) * Src0P2) +
       Src01m * T(2.000291665285952400) - T(0.000007239283986332)) *
      __ESIMD_NS::rsqrt(Src01m * T(2.0));
 
  Result.merge(T(0.0), TooBig);
  Result.merge(T(detail::__ESIMD_CONST_PI) - Result, Neg);
  return Result;
}
 
template <typename T>
__ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> acos(T src0) {
  __ESIMD_NS::simd<T, 1> Src0 = src0;
  __ESIMD_NS::simd<T, 1> Result = esimd::acos(Src0);
  return Result[0];
}
 
// asin
 
template <typename T, int SZ>
ESIMD_NODEBUG ESIMD_INLINE
    std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
    asin(__ESIMD_NS::simd<T, SZ> src0) {
  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
 
  __ESIMD_NS::simd<T, SZ> Result =
      T(detail::__ESIMD_CONST_PI / 2.0) - esimd::acos(__ESIMD_NS::abs(src0));
 
  Result.merge(-Result, Neg);
  return Result;
}
 
template <typename T>
__ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> asin(T src0) {
  __ESIMD_NS::simd<T, 1> Src0 = src0;
  __ESIMD_NS::simd<T, 1> Result = esimd::asin(Src0);
  return Result[0];
}
 
 
/* atan2_fast - a fast atan2 implementation */
/* vector input */
template <int N>
__ESIMD_NS::simd<float, N> atan2_fast(__ESIMD_NS::simd<float, N> y,
                                      __ESIMD_NS::simd<float, N> x);
/* scalar input */
template <typename T> float atan2_fast(T y, T x);
 
/* atan2 - atan2 implementation */
/* For Vector input */
template <int N>
__ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
                                 __ESIMD_NS::simd<float, N> x);
/* scalar Input */
template <typename T> float atan2(T y, T x);
 
/* fmod: */
/* vector input */
template <int N>
__ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
                                __ESIMD_NS::simd<float, N> x);
/* scalar Input */
template <typename T> float fmod(T y, T x);
 
/* sin_emu - EU emulation for sin(x) */
/* For Vector input */
template <int N>
__ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x);
/* scalar Input */
template <typename T> float sin_emu(T x);
 
/* cos_emu - EU emulation for cos(x) */
/* For Vector input */
template <int N>
__ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x);
 
/* scalar Input */
template <typename T> float cos_emu(T x);
 
/* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
/* float input */
float tanh_cody_waite(float x);
/* vector input */
template <int N>
__ESIMD_NS::simd<float, N> tanh_cody_waite(__ESIMD_NS::simd<float, N> x);
/* tanh - opencl like implementation for tanh(x) */
/* float input */
float tanh(float x);
/* vector input */
template <int N> __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x);
 
/* ------------------------- Extended Math Routines
 * -------------------------------------------------*/
 
// For vector input
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N>
atan2_fast(__ESIMD_NS::simd<float, N> y, __ESIMD_NS::simd<float, N> x) {
  /* smallest such that 1.0+CONST_DBL_EPSILON != 1.0 */
  constexpr float CONST_DBL_EPSILON = 0.00001f;
  __ESIMD_NS::simd<float, N> OneP(1.0f);
  __ESIMD_NS::simd<float, N> OneN(-1.0f);
  __ESIMD_NS::simd<float, N> sign;
  __ESIMD_NS::simd<float, N> atan2;
  __ESIMD_NS::simd<float, N> r;
  __ESIMD_NS::simd_mask<N> mask = x < 0;
  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y) + CONST_DBL_EPSILON;
 
  r.merge((x + abs_y) / (abs_y - x), (x - abs_y) / (x + abs_y), mask);
  atan2.merge(float(detail::__ESIMD_CONST_PI) * 0.75f,
              float(detail::__ESIMD_CONST_PI) * 0.25f, mask);
  atan2 += (0.1963f * r * r - 0.9817f) * r;
 
  sign.merge(OneN, OneP, y < 0);
 
  return atan2 * sign;
}
 
//   For Scalar Input
template <> ESIMD_INLINE float atan2_fast(float y, float x) {
  __ESIMD_NS::simd<float, 1> vy = y;
  __ESIMD_NS::simd<float, 1> vx = x;
  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2_fast(vy, vx);
  return atan2[0];
}
 
// atan2
// For Vector input
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
                                              __ESIMD_NS::simd<float, N> x) {
  __ESIMD_NS::simd<float, N> v_distance;
  __ESIMD_NS::simd<float, N> atan2;
  __ESIMD_NS::simd_mask<N> mask;
 
  constexpr float CONST_DBL_EPSILON = 0.00001f;
 
  mask = (x < -CONST_DBL_EPSILON && y < CONST_DBL_EPSILON && y >= 0.f);
  atan2.merge(float(detail::__ESIMD_CONST_PI), 0.f, mask);
  mask = (x < -CONST_DBL_EPSILON && y > -CONST_DBL_EPSILON && y < 0);
  atan2.merge(float(-detail::__ESIMD_CONST_PI), mask);
  mask = (x < CONST_DBL_EPSILON && __ESIMD_NS::abs(y) > CONST_DBL_EPSILON);
  v_distance = __ESIMD_NS::sqrt(x * x + y * y);
  atan2.merge(2.0f * esimd::atan((v_distance - x) / y), mask);
 
  mask = (x > 0.f);
  atan2.merge(2.0f * esimd::atan(y / (v_distance + x)), mask);
 
  return atan2;
}
 
// For Scalar Input
template <> ESIMD_INLINE float atan2(float y, float x) {
  __ESIMD_NS::simd<float, 1> vy = y;
  __ESIMD_NS::simd<float, 1> vx = x;
  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2(vy, vx);
  return atan2[0];
}
 
// fmod:
// For Vector input
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
                                             __ESIMD_NS::simd<float, N> x) {
  __ESIMD_NS::simd<float, N> abs_x = __ESIMD_NS::abs(x);
  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y);
 
  auto fmod_sign_mask = (y.template bit_cast_view<int32_t>()) & 0x80000000;
 
  __ESIMD_NS::simd<float, N> reminder =
      abs_y - abs_x * __ESIMD_NS::trunc<float>(abs_y / abs_x);
 
  abs_x.merge(0.0f, reminder >= 0);
  __ESIMD_NS::simd<float, N> fmod = reminder + abs_x;
  __ESIMD_NS::simd<float, N> fmod_abs = __ESIMD_NS::abs(fmod);
 
  auto fmod_bits =
      (fmod_abs.template bit_cast_view<int32_t>()) | fmod_sign_mask;
  return fmod_bits.template bit_cast_view<float>();
}
 
// For Scalar Input
template <> ESIMD_INLINE float fmod(float y, float x) {
  return fmod(__ESIMD_NS::simd<float, 1>(y), __ESIMD_NS::simd<float, 1>(x))[0];
}
 
// sin_emu - EU emulation for sin(x)
// For Vector input
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x) {
  __ESIMD_NS::simd<float, N> x1;
  __ESIMD_NS::simd<float, N> x2;
  __ESIMD_NS::simd<float, N> t3;
 
  __ESIMD_NS::simd<float, N> sign;
  __ESIMD_NS::simd<float, N> fTrig;
  __ESIMD_NS::simd<float, N> TwoPI(float(detail::__ESIMD_CONST_PI) * 2.0f);
  __ESIMD_NS::simd<float, N> CmpI((float)detail::__ESIMD_CONST_PI);
  __ESIMD_NS::simd<float, N> OneP(1.0f);
  __ESIMD_NS::simd<float, N> OneN(-1.0f);
 
  x = esimd::fmod(x, TwoPI);
  x.merge(TwoPI + x, x < 0);
 
  x1.merge(CmpI - x, x - CmpI, (x <= float(detail::__ESIMD_CONST_PI)));
  x1.merge(x, (x <= float(detail::__ESIMD_CONST_PI) * 0.5f));
  x1.merge(TwoPI - x, (x > float(detail::__ESIMD_CONST_PI) * 1.5f));
 
  sign.merge(OneN, OneP, (x > float(detail::__ESIMD_CONST_PI)));
 
  x2 = x1 * x1;
  t3 = x2 * x1 * 0.1666667f;
 
  fTrig =
      x1 + t3 * (OneN + x2 * 0.05f *
                            (OneP + x2 * 0.0238095f *
                                        (OneN + x2 * 0.0138889f *
                                                    (OneP - x2 * 0.0090909f))));
  fTrig *= sign;
  return fTrig;
}
 
// scalar Input
template <> ESIMD_INLINE float sin_emu(float x0) {
  return esimd::sin_emu(__ESIMD_NS::simd<float, 1>(x0))[0];
}
 
// cos_emu - EU emulation for sin(x)
// For Vector input
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x) {
  return esimd::sin_emu(0.5f * float(detail::__ESIMD_CONST_PI) - x);
}
 
// scalar Input
template <> ESIMD_INLINE float cos_emu(float x0) {
  return esimd::cos_emu(__ESIMD_NS::simd<float, 1>(x0))[0];
}
 
namespace detail {
 
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N>
tanh_cody_waite_impl(__ESIMD_NS::simd<float, N> x) {
  /*
   *      0           x_small             x_medium            x_large
   *  |   x   | rational polynomial | 1 - 2/(1 + exp(2*x)) |  1
   *
   * rational polynomial for single precision = x + x * (g * (p[1] * g + p[0]) /
   * (g + q[0]) g = x^2 p0 = -0.82377 28127 E+00 p1 = -0.38310 10665 E-02 q0 =
   * 0.24713 19654 E+01 q1 = 1.00000 00000 E+00
   *
   */
 
  constexpr float p0 = -0.8237728127E+00f;
  constexpr float p1 = -0.3831010665E-02f;
  constexpr float q0 = 0.2471319654E+01f;
  constexpr float q1 = 1.0000000000E+00f;
  constexpr float xsmall = 4.22863966691620432990E-04f;
  constexpr float xmedium = 0.54930614433405484570f;
  constexpr float xlarge = 8.66433975699931636772f;
 
  using RT = __ESIMD_NS::simd<float, N>;
 
  RT absX = __ESIMD_NS::abs(x);
  RT g = absX * absX;
 
  RT sign;
  sign.merge(-1.f, 1.f, x < 0.f);
 
  auto isLarge = absX > xlarge;
  auto minor = absX <= xlarge;
  auto isGtMed = minor & (absX > xmedium);
  auto isGtSmall = (absX > xsmall) & (absX <= xmedium);
 
  RT res;
  res.merge(sign, x, isLarge);
  auto temp = __ESIMD_NS::exp(absX * 2.0f) + 1.f;
  temp = ((temp - 2.f) / temp) * sign;
  res.merge(temp, isGtMed);
  res.merge((absX + absX * g * (g * p1 + p0) / (g + q0)) * sign, isGtSmall);
 
  return res;
}
 
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N>
tanh_impl(__ESIMD_NS::simd<float, N> x) {
  /*
   *      0                       x_small                          x_large
   * |    x    |  ( exp(x) - exp(-x) ) / ( exp(x) + exp(-x) )  |      1
   *
   */
 
  constexpr float xsmall = 0.000045f; // same as exp(-10.0f)
  constexpr float xlarge = 40.f;
 
  using RT = __ESIMD_NS::simd<float, N>;
 
  RT absX = __ESIMD_NS::abs(x);
 
  RT sign;
  sign.merge(-1.f, 1.f, x < 0.f);
 
  auto isLarge = (absX > xlarge);
  auto isLessE = (absX <= xlarge);
 
  RT res;
  res.merge(sign, x, isLarge);
 
  RT exp;
  exp = __ESIMD_NS::exp(absX * 2.f);
 
  res.merge(((exp - 1.f) / (exp + 1.f)) * sign, (absX > xsmall) & isLessE);
 
  return res;
}
} // namespace detail
 
/* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
/* float input */
ESIMD_INLINE float tanh_cody_waite(float x) {
  return detail::tanh_cody_waite_impl(__ESIMD_NS::simd<float, 1>(x))[0];
}
/* vector input */
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N>
tanh_cody_waite(__ESIMD_NS::simd<float, N> x) {
  return detail::tanh_cody_waite_impl(x);
}
 
/* tanh - opencl like implementation for tanh(x) */
/* float input */
ESIMD_INLINE float tanh(float x) {
  return esimd::detail::tanh_impl(__ESIMD_NS::simd<float, 1>(x))[0];
}
/* vector input */
template <int N>
ESIMD_INLINE __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x) {
  return esimd::detail::tanh_impl(x);
}
 
template <typename T, int N>
__ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,
                           __ESIMD_NS::simd<T, N> v2) {
  auto retv = __esimd_dp4<T, N>(v1.data(), v2.data());
  return retv;
}
 
 
 
template <argument_type src1_precision, argument_type src2_precision,
          typename T, int systolic_depth, int repeat_count, typename T0,
          typename T1, typename T2, int N, int N1, int N2,
          typename Sat = __ESIMD_NS::saturation_off_tag>
__SYCL_DEPRECATED("use sycl::ext::intel::esimd::native::dpas()")
__ESIMD_API __ESIMD_NS::simd<T, N> dpas(
    __ESIMD_NS::simd<T0, N> src0, __ESIMD_NS::simd<T1, N1> src1,
    __ESIMD_NS::simd<T2, N2> src2,
    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
  auto result =
      __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T0, T1, T2,
                            src1_precision, src2_precision>(src0, src1, src2);
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return result;
  else
    return __ESIMD_NS::saturate<T>(result);
}
 
template <argument_type src1_precision, argument_type src2_precision,
          int systolic_depth, int repeat_count, typename T, typename T1,
          typename T2, int N, int N1, int N2,
          typename Sat = __ESIMD_NS::saturation_off_tag>
__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
__ESIMD_API __ESIMD_NS::simd<T, N> dpas(
    __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
    __ESIMD_NS::simd<T2, N2> src2,
    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
  return dpas<src1_precision, src2_precision, T, systolic_depth, repeat_count>(
      src0, src1, src2, sat);
}
 
template <argument_type src1_precision, argument_type src2_precision,
          int systolic_depth, int repeat_count, typename T, typename T1,
          typename T2, int N, int N1, int N2,
          typename Sat = __ESIMD_NS::saturation_off_tag>
__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
__ESIMD_API __ESIMD_NS::simd<T, N> dpas(
    __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
 
  __ESIMD_NS::simd<T, N> result =
      __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T1, T2,
                            src1_precision, src2_precision>(src1, src2);
 
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return result;
  else
    return __ESIMD_NS::saturate<T>(result);
}
 
template <argument_type src1_precision, argument_type src2_precision,
          int systolic_depth, int repeat_count, typename T, typename T1,
          typename T2, int N, int N1, int N2,
          typename Sat = __ESIMD_NS::saturation_off_tag>
__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
__ESIMD_API __ESIMD_NS::simd<T, N> dpasw(
    __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
    __ESIMD_NS::simd<T2, N2> src2,
    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
 
  __ESIMD_NS::simd<T, N> result =
      __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
                             src1_precision, src2_precision>(src0, src1, src2);
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return result;
  else
    return __ESIMD_NS::saturate<T>(result);
}
 
template <argument_type src1_precision, argument_type src2_precision,
          int systolic_depth, int repeat_count, typename T, typename T1,
          typename T2, int N, int N1, int N2,
          typename Sat = __ESIMD_NS::saturation_off_tag>
__SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
__ESIMD_API __ESIMD_NS::simd<T, N> dpasw2(
    __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
    std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
 
  __ESIMD_NS::simd<T, N> result =
      __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
                             src1_precision, src2_precision>(src1, src2);
 
  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
    return result;
  else
    return __ESIMD_NS::saturate<T>(result);
}
 
} // namespace ext::intel::experimental::esimd
} // __SYCL_INLINE_VER_NAMESPACE(_V1)
} // namespace sycl