This is a group of APIs implementing standard math operations which are also directly supported by the hardware. More...

Collaboration diagram for Hardware-accelerated math.:

Macros
#define	__ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname)

#define	__ESIMD_EMATH_COND detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4)

#define	__ESIMD_EMATH_IEEE_COND detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4)

#define	__ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname)

Functions
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::inv (simd< T, N > src, Sat sat={})
	Inversion - calculates (1/x). More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::inv (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::log2 (simd< T, N > src, Sat sat={})
	Logarithm base 2. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::log2 (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::exp2 (simd< T, N > src, Sat sat={})
	Exponent base 2. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::exp2 (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::sqrt (simd< T, N > src, Sat sat={})
	Square root. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::sqrt (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::sqrt_ieee (simd< T, N > src, Sat sat={})
	IEEE754-compliant square root. Supports `float` and `double`. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::sqrt_ieee (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::rsqrt (simd< T, N > src, Sat sat={})
	Square root reciprocal - calculates `1/sqrt(x)`. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::rsqrt (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::sin (simd< T, N > src, Sat sat={})
	Sine. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::sin (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::cos (simd< T, N > src, Sat sat={})
	Cosine. More...

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::cos (T src, Sat sat={})
	Scalar version. More...

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::pow (simd< T, N > src0, simd< U, N > src1, Sat sat={})
	Power - calculates `src0` in power of `src1`. More...

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::pow (simd< T, N > src0, U src1, Sat sat={})
	(vector, scalar) version. More...

template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::pow (T src0, U src1, Sat sat={})
	(scalar, scalar) version. More...

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::div_ieee (simd< T, N > src0, simd< U, N > src1, Sat sat={})
	IEEE754-compliant floating-point division. Supports `float` and `double`. More...

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N >	sycl::_V1::ext::intel::esimd::div_ieee (simd< T, N > src0, U src1, Sat sat={})
	(vector, scalar) version. More...

template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API T	sycl::_V1::ext::intel::esimd::div_ieee (T src0, U src1, Sat sat={})
	(scalar, scalar) version. More...

Detailed Description

This is a group of APIs implementing standard math operations which are also directly supported by the hardware.

Usually the hardware support is a specific message to the "extended math" GPU "shared function" unit, sent via the math instruction. Most of the operations do not conform to OpenCL requirements for accuracy, so should be used with care.

TODO Provide detailed spec of each operation.

Macro Definition Documentation

◆ __ESIMD_BINARY_INTRINSIC_DEF

#define __ESIMD_BINARY_INTRINSIC_DEF	(	COND,
		name,
		iname
	)

Value:

 \
  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];                                                             \
  }

Definition at line 408 of file math.hpp.

◆ __ESIMD_EMATH_COND

#define __ESIMD_EMATH_COND detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4)

Definition at line 366 of file math.hpp.

◆ __ESIMD_EMATH_IEEE_COND

#define __ESIMD_EMATH_IEEE_COND detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4)

Definition at line 369 of file math.hpp.

◆ __ESIMD_UNARY_INTRINSIC_DEF

#define __ESIMD_UNARY_INTRINSIC_DEF	(	COND,
		name,
		iname
	)

Value:

 \
  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_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res =               \
        __esimd_##iname<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<COND>>                                    \
  __ESIMD_API T name(T src, Sat sat = {}) {                                    \
    simd<T, 1> src_vec = src;                                                  \
    simd<T, 1> res = name<T, 1>(src_vec, sat);                                 \
    return res[0];                                                             \
  }

Definition at line 344 of file math.hpp.

Function Documentation

◆ cos() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::cos	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Cosine.

Supports half and float. Absolute error: 0.0008 or less for the range [-32767*pi, 32767*pi]. Vector version.

Definition at line 404 of file math.hpp.

◆ cos() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::cos	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 404 of file math.hpp.

◆ div_ieee() [1/3]

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::div_ieee	(	simd< T, N >	src0,
		simd< U, N >	src1,
		Sat	sat = `{}`
	)

IEEE754-compliant floating-point division. Supports float and double.

(vector, vector) version.

Definition at line 444 of file math.hpp.

◆ div_ieee() [2/3]

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::div_ieee	(	simd< T, N >	src0,
		U	src1,
		Sat	sat = `{}`
	)

(vector, scalar) version.

Definition at line 444 of file math.hpp.

◆ div_ieee() [3/3]

template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::div_ieee	(	T	src0,
		U	src1,
		Sat	sat = `{}`
	)

(scalar, scalar) version.

Definition at line 444 of file math.hpp.

◆ exp2() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::exp2	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Exponent base 2.

Supports half and float. Precision: 4 ULP. Vector version.

Definition at line 384 of file math.hpp.

◆ exp2() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::exp2	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 384 of file math.hpp.

◆ inv() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::inv	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Inversion - calculates (1/x).

Supports half and float. Precision: 1 ULP. Vector version.

Definition at line 374 of file math.hpp.

◆ inv() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::inv	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 374 of file math.hpp.

◆ log2() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::log2	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Logarithm base 2.

Supports half and float. Precision depending on argument range:

[0.5..2]: absolute error is 2^-21 or less
(0..0.5) or (2..+INF]: relative error is 2^-21 or less Vector version.

Definition at line 380 of file math.hpp.

◆ log2() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::log2	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 380 of file math.hpp.

◆ pow() [1/3]

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::pow	(	simd< T, N >	src0,
		simd< U, N >	src1,
		Sat	sat = `{}`
	)

Power - calculates src0 in power of src1.

Note available in DG2, PVC. Supports half and float. TODO document accuracy etc. (vector, vector) version.

Definition at line 441 of file math.hpp.

Referenced by sycl::_V1::half_precision::BUILTIN_HALF_CUSTOM(), and sycl::_V1::native::BUILTIN_NATIVE_CUSTOM().

◆ pow() [2/3]

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::pow	(	simd< T, N >	src0,
		U	src1,
		Sat	sat = `{}`
	)

(vector, scalar) version.

Definition at line 441 of file math.hpp.

◆ pow() [3/3]

template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::pow	(	T	src0,
		U	src1,
		Sat	sat = `{}`
	)

(scalar, scalar) version.

Definition at line 441 of file math.hpp.

◆ rsqrt() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::rsqrt	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Square root reciprocal - calculates 1/sqrt(x).

Supports half and float. Precision: 4 ULP. Vector version.

Definition at line 396 of file math.hpp.

◆ rsqrt() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::rsqrt	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 396 of file math.hpp.

◆ sin() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::sin	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Sine.

Supports half and float. Absolute error: 0.0008 or less for the range [-32767*pi, 32767*pi]. Vector version.

Definition at line 400 of file math.hpp.

◆ sin() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::sin	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 400 of file math.hpp.

◆ sqrt() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::sqrt	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

Square root.

Is not IEEE754-compatible. Supports half and float. Precision: 4 ULP. Vector version.

Definition at line 388 of file math.hpp.

◆ sqrt() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::sqrt	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 388 of file math.hpp.

◆ sqrt_ieee() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>

__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::sqrt_ieee	(	simd< T, N >	src,
		Sat	sat = `{}`
	)

IEEE754-compliant square root. Supports float and double.

Vector version.

Definition at line 391 of file math.hpp.

◆ sqrt_ieee() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>

__ESIMD_API T sycl::_V1::ext::intel::esimd::sqrt_ieee	(	T	src,
		Sat	sat = `{}`
	)

Scalar version.

Definition at line 391 of file math.hpp.

Macros

Functions

Detailed Description

Macro Definition Documentation

◆ __ESIMD_BINARY_INTRINSIC_DEF

◆ __ESIMD_EMATH_COND

◆ __ESIMD_EMATH_IEEE_COND

◆ __ESIMD_UNARY_INTRINSIC_DEF

Function Documentation

◆ cos() [1/2]

◆ cos() [2/2]

◆ div_ieee() [1/3]

◆ div_ieee() [2/3]

◆ div_ieee() [3/3]

◆ exp2() [1/2]

◆ exp2() [2/2]

◆ inv() [1/2]

◆ inv() [2/2]

◆ log2() [1/2]

◆ log2() [2/2]

◆ pow() [1/3]

◆ pow() [2/3]

◆ pow() [3/3]

◆ rsqrt() [1/2]

◆ rsqrt() [2/2]

◆ sin() [1/2]

◆ sin() [2/2]

◆ sqrt() [1/2]

◆ sqrt() [2/2]

◆ sqrt_ieee() [1/2]

◆ sqrt_ieee() [2/2]