DPC++ Runtime
Runtime libraries for oneAPI DPC++
Macros | Functions
Hardware-accelerated math.
DPC++ Explicit SIMD API » ESIMD math operations.

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]; \
}
simd
Definition: simd.hpp:1387
sycl::_V1::ext::intel::esimd::src1
__ESIMD_API SZ simd< T, SZ > src1
Definition: math.hpp:187
sycl::_V1::ext::intel::esimd::sat
__ESIMD_API SZ simd< T, SZ > Sat sat
Definition: math.hpp:187
sycl::_V1::ext::intel::esimd::Sat
__ESIMD_API SZ simd< T, SZ > Sat int class Sat
Definition: math.hpp:222
sycl::_V1::ext::intel::esimd::src0
__ESIMD_API SZ src0
Definition: math.hpp:187

Definition at line 479 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 395 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 398 of file math.hpp.

◆ __ESIMD_SCALAR_IMPL

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

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

◆ __ESIMD_VECTOR_IMPL

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

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

Scalar version.


Definition at line 451 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 554 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 554 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 554 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 413 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 ( T  src,
Sat  sat = {} 
)

Scalar version.


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

Scalar version.


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

Scalar version.


Definition at line 409 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 513 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,
U  src1,
Sat  sat = {} 
)

(vector, scalar) version.


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

(scalar, scalar) version.


Definition at line 538 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 443 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 458 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 ( T  src,
Sat  sat = {} 
)

Scalar version.


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

Scalar version.


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

Scalar version.


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

Scalar version.


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

Scalar version.


Definition at line 434 of file math.hpp.