DPC++ Runtime
Runtime libraries for oneAPI DPC++
Hardware-accelerated math.

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_VECTOR_IMPL(T, name, iname)   return 0;
 
#define __ESIMD_SCALAR_IMPL(T, name, iname)   return 0;
 
#define __ESIMD_UNARY_INTRINSIC_DEF(COND, 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>
 
#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> >>
__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> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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> >>
__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> >>
__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<__ESIMD_EMATH_IEEE_COND>>
__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<__ESIMD_EMATH_IEEE_COND>>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::cos (T src, Sat sat={})
 Scalar version. More...
 
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 > > sycl::_V1::ext::intel::esimd::rsqrt (simd< T, N > src, Sat sat={})
 Square root reciprocal - calculates 1/sqrt(x). More...
 
template<class T , class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< std::is_same_v< T, double >, double > sycl::_V1::ext::intel::esimd::rsqrt (T src, Sat sat={})
 Scalar version. More...
 
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 > 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<__ESIMD_EMATH_SPIRV_COND>>
__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<__ESIMD_EMATH_SPIRV_COND>>
__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: simd.hpp:1387
__ESIMD_API SZ simd< T, SZ > Sat sat
Definition: math.hpp:180
__ESIMD_API SZ simd< T, SZ > Sat int class Sat
Definition: math.hpp:211

Definition at line 464 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 380 of file math.hpp.

◆ __ESIMD_EMATH_SPIRV_COND

#define __ESIMD_EMATH_SPIRV_COND    std::is_same_v<T, float> || std::is_same_v<T, sycl::half>

Definition at line 383 of file math.hpp.

◆ __ESIMD_SCALAR_IMPL

#define __ESIMD_SCALAR_IMPL (   T,
  name,
  iname 
)    return 0;

Definition at line 362 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_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) \
}

Definition at line 365 of file math.hpp.

◆ __ESIMD_VECTOR_IMPL

#define __ESIMD_VECTOR_IMPL (   T,
  name,
  iname 
)    return 0;

Definition at line 361 of file math.hpp.

Function Documentation

◆ cos() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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 436 of file math.hpp.

◆ cos() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::cos ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 436 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 539 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,
src1,
Sat  sat = {} 
)

(vector, scalar) version.


Definition at line 539 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 ( src0,
src1,
Sat  sat = {} 
)

(scalar, scalar) version.


Definition at line 539 of file math.hpp.

◆ exp2() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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 398 of file math.hpp.

◆ exp2() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::exp2 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 398 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> >>
__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::inv ( simd< T, N >  src,
Sat  sat = {} 
)

Inversion - calculates (1/x).

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

Definition at line 388 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> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::inv ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 388 of file math.hpp.

◆ log2() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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 394 of file math.hpp.

◆ log2() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::log2 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 394 of file math.hpp.

◆ pow() [1/3]

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> 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.

Definition at line 498 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<__ESIMD_EMATH_SPIRV_COND>>
__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::pow ( simd< T, N >  src0,
src1,
Sat  sat = {} 
)

(vector, scalar) version.


Definition at line 516 of file math.hpp.

◆ pow() [3/3]

template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
__ESIMD_API T sycl::_V1::ext::intel::esimd::pow ( src0,
src1,
Sat  sat = {} 
)

(scalar, scalar) version.


Definition at line 523 of file math.hpp.

◆ rsqrt() [1/4]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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 428 of file math.hpp.

◆ rsqrt() [2/4]

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> > sycl::_V1::ext::intel::esimd::rsqrt ( simd< T, N >  src,
Sat  sat = {} 
)

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

Supports double. Precision: 4 ULP.

Definition at line 443 of file math.hpp.

◆ rsqrt() [3/4]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::rsqrt ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 428 of file math.hpp.

◆ rsqrt() [4/4]

template<class T , class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t<std::is_same_v<T, double>, double> sycl::_V1::ext::intel::esimd::rsqrt ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 453 of file math.hpp.

◆ sin() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__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 432 of file math.hpp.

◆ sin() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< std::is_same_v<T, float> || std::is_same_v<T, sycl::half> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::sin ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 432 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> >>
__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, float and double. Precision: 4 ULP. Vector version.

Definition at line 402 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> >>
__ESIMD_API T sycl::_V1::ext::intel::esimd::sqrt ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 402 of file math.hpp.

◆ sqrt_ieee() [1/2]

template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t<__ESIMD_EMATH_IEEE_COND>>
__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.

Definition at line 407 of file math.hpp.

◆ sqrt_ieee() [2/2]

template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t<__ESIMD_EMATH_IEEE_COND>>
__ESIMD_API T sycl::_V1::ext::intel::esimd::sqrt_ieee ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 419 of file math.hpp.