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_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: simd.hpp:1384

Definition at line 409 of file math.hpp.

◆ __ESIMD_EMATH_COND

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

Definition at line 367 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 370 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>(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 345 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 405 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 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 405 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 445 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 445 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 445 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 385 of file math.hpp.

Referenced by __host_std::sycl_host_exp2().

◆ 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 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 385 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 375 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 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 375 of file math.hpp.

Referenced by sycl::_V1::ext::intel::experimental::esimd::atan().

◆ 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 381 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 ( src,
Sat  sat = {} 
)

Scalar version.


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

Referenced by __host_std::sycl_host_exp10().

◆ 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,
src1,
Sat  sat = {} 
)

(vector, scalar) version.


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

(scalar, scalar) version.


Definition at line 442 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 397 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 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 397 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 401 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 ( src,
Sat  sat = {} 
)

Scalar version.


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

Referenced by __host_std::sycl_host_fast_normalize().

◆ 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 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 389 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 392 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 ( src,
Sat  sat = {} 
)

Scalar version.


Definition at line 392 of file math.hpp.