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. 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. 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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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 cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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_COND >>
__ESIMD_API simd< T, N > cl::sycl::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_COND >>
__ESIMD_API T cl::sycl::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 > cl::sycl::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< __ESIMD_EMATH_IEEE_COND >>
__ESIMD_API simd< T, N > cl::sycl::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< __ESIMD_EMATH_IEEE_COND >>
__ESIMD_API T cl::sycl::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>(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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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> cl::sycl::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< __ESIMD_EMATH_IEEE_COND >>
__ESIMD_API simd<T, N> cl::sycl::ext::intel::esimd::div_ieee ( simd< T, N >  src0,
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< __ESIMD_EMATH_IEEE_COND >>
__ESIMD_API T cl::sycl::ext::intel::esimd::div_ieee ( T  src0,
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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::ext::intel::esimd::inv ( T  src,
Sat  sat = {} 
)

Scalar version.


Definition at line 374 of file math.hpp.

Referenced by cl::sycl::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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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> cl::sycl::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.

◆ pow() [2/3]

template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< __ESIMD_EMATH_COND >>
__ESIMD_API simd<T, N> cl::sycl::ext::intel::esimd::pow ( simd< T, N >  src0,
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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::ext::intel::esimd::pow ( T  src0,
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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::ext::intel::esimd::rsqrt ( T  src,
Sat  sat = {} 
)

Scalar version.


Definition at line 396 of file math.hpp.

Referenced by cl::sycl::ext::intel::experimental::esimd::acos().

◆ 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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::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> cl::sycl::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< __ESIMD_EMATH_COND >>
__ESIMD_API T cl::sycl::ext::intel::esimd::sqrt ( T  src,
Sat  sat = {} 
)

Scalar version.


Definition at line 388 of file math.hpp.

Referenced by cl::sycl::ext::intel::experimental::esimd::atan2().

◆ 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> cl::sycl::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< __ESIMD_EMATH_IEEE_COND >>
__ESIMD_API T cl::sycl::ext::intel::esimd::sqrt_ieee ( T  src,
Sat  sat = {} 
)

Scalar version.


Definition at line 391 of file math.hpp.

T
simd
Definition: simd.hpp:1027