18 #if defined(__SPIR__) || defined(__SPIRV__)
32 #include <type_traits>
64 inline namespace _V1 {
65 namespace ext::intel::math {
68 "sycl::half is not compatible with _iml_half_internal.");
70 template <
typename Tp>
71 std::enable_if_t<std::is_same_v<Tp, float>,
float>
saturate(Tp x) {
75 template <
typename Tp>
76 std::enable_if_t<std::is_same_v<Tp, float>,
float>
copysign(Tp x, Tp y) {
80 template <
typename Tp>
81 std::enable_if_t<std::is_same_v<Tp, double>,
double>
copysign(Tp x, Tp y) {
85 template <
typename Tp>
93 template <
typename Tp>
94 std::enable_if_t<std::is_same_v<Tp, float>,
float>
ceil(Tp x) {
98 template <
typename Tp>
99 std::enable_if_t<std::is_same_v<Tp, double>,
double>
ceil(Tp x) {
103 template <
typename Tp>
109 template <
typename Tp>
110 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
ceil(Tp x) {
111 return sycl::half2{
ceil(
x.s0()),
ceil(
x.s1())};
114 template <
typename Tp>
115 std::enable_if_t<std::is_same_v<Tp, float>,
float>
floor(Tp x) {
119 template <
typename Tp>
120 std::enable_if_t<std::is_same_v<Tp, double>,
double>
floor(Tp x) {
124 template <
typename Tp>
130 template <
typename Tp>
131 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
floor(Tp x) {
135 template <
typename Tp>
136 std::enable_if_t<std::is_same_v<Tp, float>,
float>
inv(Tp x) {
140 template <
typename Tp>
141 std::enable_if_t<std::is_same_v<Tp, double>,
double>
inv(Tp x) {
145 template <
typename Tp>
151 template <
typename Tp>
152 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
inv(Tp x) {
153 return sycl::half2{
inv(
x.s0()),
inv(
x.s1())};
156 template <
typename Tp>
157 std::enable_if_t<std::is_same_v<Tp, float>,
float>
rint(Tp x) {
161 template <
typename Tp>
162 std::enable_if_t<std::is_same_v<Tp, double>,
double>
rint(Tp x) {
166 template <
typename Tp>
172 template <
typename Tp>
173 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
rint(Tp x) {
174 return sycl::half2{
rint(
x.s0()),
rint(
x.s1())};
177 template <
typename Tp>
178 std::enable_if_t<std::is_same_v<Tp, float>,
float>
sqrt(Tp x) {
182 template <
typename Tp>
183 std::enable_if_t<std::is_same_v<Tp, double>,
double>
sqrt(Tp x) {
187 template <
typename Tp>
193 template <
typename Tp>
194 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
sqrt(Tp x) {
195 return sycl::half2{
sqrt(
x.s0()),
sqrt(
x.s1())};
198 template <
typename Tp>
199 std::enable_if_t<std::is_same_v<Tp, float>,
float>
rsqrt(Tp x) {
203 template <
typename Tp>
204 std::enable_if_t<std::is_same_v<Tp, double>,
double>
rsqrt(Tp x) {
208 template <
typename Tp>
214 template <
typename Tp>
215 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
rsqrt(Tp x) {
219 template <
typename Tp>
220 std::enable_if_t<std::is_same_v<Tp, float>,
float>
trunc(Tp x) {
224 template <
typename Tp>
225 std::enable_if_t<std::is_same_v<Tp, double>,
double>
trunc(Tp x) {
229 template <
typename Tp>
235 template <
typename Tp>
236 std::enable_if_t<std::is_same_v<Tp, sycl::half2>, sycl::half2>
trunc(Tp x) {
240 template <
typename Tp>
241 std::enable_if_t<std::is_same_v<Tp, double>,
double>
rcp64h(Tp x) {
uint16_t _iml_half_internal
float __imf_truncf(float)
float __imf_saturatef(float)
float __imf_rsqrtf(float)
double __imf_rint(double)
double __imf_ceil(double)
_iml_half_internal __imf_invf16(_iml_half_internal)
double __imf_sqrt(double)
double __imf_trunc(double)
_iml_half_internal __imf_rintf16(_iml_half_internal)
double __imf_rsqrt(double)
float __imf_copysignf(float, float)
_iml_half_internal __imf_floorf16(_iml_half_internal)
double __imf_floor(double)
_iml_half_internal __imf_sqrtf16(_iml_half_internal)
_iml_half_internal __imf_truncf16(_iml_half_internal)
double __imf_rcp64h(double)
float __imf_floorf(float)
double __imf_copysign(double, double)
_iml_half_internal __imf_copysignf16(_iml_half_internal, _iml_half_internal)
_iml_half_internal __imf_rsqrtf16(_iml_half_internal)
_iml_half_internal __imf_ceilf16(_iml_half_internal)
std::enable_if_t< std::is_same_v< Tp, float >, float > sqrt(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > ceil(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > trunc(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > copysign(Tp x, Tp y)
std::enable_if_t< std::is_same_v< Tp, float >, float > rint(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > floor(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > saturate(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > inv(Tp x)
std::enable_if_t< std::is_same_v< Tp, double >, double > rcp64h(Tp x)
std::enable_if_t< std::is_same_v< Tp, float >, float > rsqrt(Tp x)