DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins.hpp
Go to the documentation of this file.
1 //==----------- builtins.hpp - SYCL built-in functions ---------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/detail/boolean.hpp>
12 #include <sycl/detail/builtins.hpp>
13 #include <sycl/detail/common.hpp>
14 #include <sycl/detail/generic_type_traits.hpp>
15 #include <sycl/types.hpp>
16 
17 // TODO Decide whether to mark functions with this attribute.
18 #define __NOEXC /*noexcept*/
19 
20 namespace sycl {
21 
22 __SYCL_INLINE_VER_NAMESPACE(_V1) {
23 
24 namespace detail {
25 template <class T, size_t N> vec<T, 2> to_vec2(marray<T, N> x, size_t start) {
26  return {x[start], x[start + 1]};
27 }
28 } // namespace detail
29 
30 #ifdef __SYCL_DEVICE_ONLY__
31 #define __sycl_std
32 #else
33 namespace __sycl_std = __host_std;
34 #endif
35 
36 #ifdef __FAST_MATH__
37 #define __FAST_MATH_GENFLOAT(T) \
38  (detail::is_svgenfloatd<T>::value || detail::is_svgenfloath<T>::value)
39 #define __FAST_MATH_SGENFLOAT(T) \
40  (std::is_same_v<T, double> || std::is_same_v<T, half>)
41 #else
42 #define __FAST_MATH_GENFLOAT(T) (detail::is_svgenfloat<T>::value)
43 #define __FAST_MATH_SGENFLOAT(T) (detail::is_sgenfloat<T>::value)
44 #endif
45 
46 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
47 
48 // These macros for marray math function implementations use vectorizations of
49 // size two as a simple general optimization. A more complex implementation
50 // using larger vectorizations for large marray sizes is possible; however more
51 // testing is required in order to ascertain the performance implications for
52 // all backends.
53 #define __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
54  marray<T, N> res; \
55  for (size_t i = 0; i < N / 2; i++) { \
56  vec<T, 2> partial_res = \
57  __sycl_std::__invoke_##NAME<vec<T, 2>>(detail::to_vec2(x, i * 2)); \
58  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
59  } \
60  if (N % 2) { \
61  res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1]); \
62  } \
63  return res;
64 
65 #define __SYCL_MATH_FUNCTION_OVERLOAD(NAME) \
66  template <typename T, size_t N> \
67  inline __SYCL_ALWAYS_INLINE \
68  std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
69  NAME(marray<T, N> x) __NOEXC { \
70  __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
71  }
72 
73 __SYCL_MATH_FUNCTION_OVERLOAD(cospi)
74 __SYCL_MATH_FUNCTION_OVERLOAD(sinpi)
75 __SYCL_MATH_FUNCTION_OVERLOAD(tanpi)
76 __SYCL_MATH_FUNCTION_OVERLOAD(sinh)
77 __SYCL_MATH_FUNCTION_OVERLOAD(cosh)
78 __SYCL_MATH_FUNCTION_OVERLOAD(tanh)
79 __SYCL_MATH_FUNCTION_OVERLOAD(asin)
80 __SYCL_MATH_FUNCTION_OVERLOAD(acos)
81 __SYCL_MATH_FUNCTION_OVERLOAD(atan)
82 __SYCL_MATH_FUNCTION_OVERLOAD(asinpi)
83 __SYCL_MATH_FUNCTION_OVERLOAD(acospi)
84 __SYCL_MATH_FUNCTION_OVERLOAD(atanpi)
85 __SYCL_MATH_FUNCTION_OVERLOAD(asinh)
86 __SYCL_MATH_FUNCTION_OVERLOAD(acosh)
87 __SYCL_MATH_FUNCTION_OVERLOAD(atanh)
88 __SYCL_MATH_FUNCTION_OVERLOAD(cbrt)
89 __SYCL_MATH_FUNCTION_OVERLOAD(ceil)
90 __SYCL_MATH_FUNCTION_OVERLOAD(floor)
91 __SYCL_MATH_FUNCTION_OVERLOAD(erfc)
92 __SYCL_MATH_FUNCTION_OVERLOAD(erf)
93 __SYCL_MATH_FUNCTION_OVERLOAD(expm1)
94 __SYCL_MATH_FUNCTION_OVERLOAD(tgamma)
95 __SYCL_MATH_FUNCTION_OVERLOAD(lgamma)
96 __SYCL_MATH_FUNCTION_OVERLOAD(log1p)
97 __SYCL_MATH_FUNCTION_OVERLOAD(logb)
98 __SYCL_MATH_FUNCTION_OVERLOAD(rint)
99 __SYCL_MATH_FUNCTION_OVERLOAD(round)
100 __SYCL_MATH_FUNCTION_OVERLOAD(trunc)
101 
102 #undef __SYCL_MATH_FUNCTION_OVERLOAD
103 
104 // __SYCL_MATH_FUNCTION_OVERLOAD_FM cases are replaced by corresponding native
105 // implementations when the -ffast-math flag is used with float.
106 #define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME) \
107  template <typename T, size_t N> \
108  inline __SYCL_ALWAYS_INLINE \
109  std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray<T, N>> \
110  NAME(marray<T, N> x) __NOEXC { \
111  __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
112  }
113 
114 __SYCL_MATH_FUNCTION_OVERLOAD_FM(sin)
115 __SYCL_MATH_FUNCTION_OVERLOAD_FM(cos)
116 __SYCL_MATH_FUNCTION_OVERLOAD_FM(tan)
117 __SYCL_MATH_FUNCTION_OVERLOAD_FM(exp)
118 __SYCL_MATH_FUNCTION_OVERLOAD_FM(exp2)
119 __SYCL_MATH_FUNCTION_OVERLOAD_FM(exp10)
120 __SYCL_MATH_FUNCTION_OVERLOAD_FM(log)
121 __SYCL_MATH_FUNCTION_OVERLOAD_FM(log2)
122 __SYCL_MATH_FUNCTION_OVERLOAD_FM(log10)
123 __SYCL_MATH_FUNCTION_OVERLOAD_FM(sqrt)
124 __SYCL_MATH_FUNCTION_OVERLOAD_FM(rsqrt)
125 
126 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
127 #undef __SYCL_MATH_FUNCTION_OVERLOAD_IMPL
128 
129 #define __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \
130  marray<T, N> res; \
131  for (size_t i = 0; i < N / 2; i++) { \
132  auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
133  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
134  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
135  } \
136  if (N % 2) { \
137  res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1], y[N - 1]); \
138  } \
139  return res;
140 
141 #define __SYCL_MATH_FUNCTION_2_OVERLOAD(NAME) \
142  template <typename T, size_t N> \
143  inline __SYCL_ALWAYS_INLINE \
144  std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
145  NAME(marray<T, N> x, marray<T, N> y) __NOEXC { \
146  __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \
147  }
148 
149 __SYCL_MATH_FUNCTION_2_OVERLOAD(atan2)
150 __SYCL_MATH_FUNCTION_2_OVERLOAD(atan2pi)
151 __SYCL_MATH_FUNCTION_2_OVERLOAD(copysign)
152 __SYCL_MATH_FUNCTION_2_OVERLOAD(fdim)
153 __SYCL_MATH_FUNCTION_2_OVERLOAD(fmin)
154 __SYCL_MATH_FUNCTION_2_OVERLOAD(fmax)
155 __SYCL_MATH_FUNCTION_2_OVERLOAD(fmod)
156 __SYCL_MATH_FUNCTION_2_OVERLOAD(hypot)
157 __SYCL_MATH_FUNCTION_2_OVERLOAD(maxmag)
158 __SYCL_MATH_FUNCTION_2_OVERLOAD(minmag)
159 __SYCL_MATH_FUNCTION_2_OVERLOAD(nextafter)
160 __SYCL_MATH_FUNCTION_2_OVERLOAD(pow)
161 __SYCL_MATH_FUNCTION_2_OVERLOAD(remainder)
162 
163 #undef __SYCL_MATH_FUNCTION_2_OVERLOAD
164 
165 template <typename T, size_t N>
166 inline __SYCL_ALWAYS_INLINE
167  std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray<T, N>>
168  powr(marray<T, N> x,
169  marray<T, N> y) __NOEXC{__SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(powr)}
170 
171 #undef __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL
172 
173 #define __SYCL_MATH_FUNCTION_3_OVERLOAD(NAME) \
174  template <typename T, size_t N> \
175  inline __SYCL_ALWAYS_INLINE \
176  std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
177  NAME(marray<T, N> x, marray<T, N> y, marray<T, N> z) __NOEXC { \
178  marray<T, N> res; \
179  for (size_t i = 0; i < N / 2; i++) { \
180  auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
181  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2), \
182  detail::to_vec2(z, i * 2)); \
183  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
184  } \
185  if (N % 2) { \
186  res[N - 1] = \
187  __sycl_std::__invoke_##NAME<T>(x[N - 1], y[N - 1], z[N - 1]); \
188  } \
189  return res; \
190  }
191 
192 __SYCL_MATH_FUNCTION_3_OVERLOAD(mad) __SYCL_MATH_FUNCTION_3_OVERLOAD(mix)
193  __SYCL_MATH_FUNCTION_3_OVERLOAD(fma)
194 
195 #undef __SYCL_MATH_FUNCTION_3_OVERLOAD
196 
197  // svgenfloat acos (svgenfloat x)
198  template <typename T>
199  detail::enable_if_t<detail::is_svgenfloat<T>::value, T> acos(T x) __NOEXC {
200  return __sycl_std::__invoke_acos<T>(x);
201 }
202 
203 // svgenfloat acosh (svgenfloat x)
204 template <typename T>
205 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> acosh(T x) __NOEXC {
206  return __sycl_std::__invoke_acosh<T>(x);
207 }
208 
209 // svgenfloat acospi (svgenfloat x)
210 template <typename T>
211 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> acospi(T x) __NOEXC {
212  return __sycl_std::__invoke_acospi<T>(x);
213 }
214 
215 // svgenfloat asin (svgenfloat x)
216 template <typename T>
217 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> asin(T x) __NOEXC {
218  return __sycl_std::__invoke_asin<T>(x);
219 }
220 
221 // svgenfloat asinh (svgenfloat x)
222 template <typename T>
223 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> asinh(T x) __NOEXC {
224  return __sycl_std::__invoke_asinh<T>(x);
225 }
226 
227 // svgenfloat asinpi (svgenfloat x)
228 template <typename T>
229 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> asinpi(T x) __NOEXC {
230  return __sycl_std::__invoke_asinpi<T>(x);
231 }
232 
233 // svgenfloat atan (svgenfloat y_over_x)
234 template <typename T>
235 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>
236 atan(T y_over_x) __NOEXC {
237  return __sycl_std::__invoke_atan<T>(y_over_x);
238 }
239 
240 // svgenfloat atan2 (svgenfloat y, svgenfloat x)
241 template <typename T>
242 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> atan2(T y,
243  T x) __NOEXC {
244  return __sycl_std::__invoke_atan2<T>(y, x);
245 }
246 
247 // svgenfloat atanh (svgenfloat x)
248 template <typename T>
249 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> atanh(T x) __NOEXC {
250  return __sycl_std::__invoke_atanh<T>(x);
251 }
252 
253 // svgenfloat atanpi (svgenfloat x)
254 template <typename T>
255 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> atanpi(T x) __NOEXC {
256  return __sycl_std::__invoke_atanpi<T>(x);
257 }
258 
259 // svgenfloat atan2pi (svgenfloat y, svgenfloat x)
260 template <typename T>
261 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> atan2pi(T y,
262  T x) __NOEXC {
263  return __sycl_std::__invoke_atan2pi<T>(y, x);
264 }
265 
266 // svgenfloat cbrt (svgenfloat x)
267 template <typename T>
268 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> cbrt(T x) __NOEXC {
269  return __sycl_std::__invoke_cbrt<T>(x);
270 }
271 
272 // svgenfloat ceil (svgenfloat x)
273 template <typename T>
274 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> ceil(T x) __NOEXC {
275  return __sycl_std::__invoke_ceil<T>(x);
276 }
277 
278 // svgenfloat copysign (svgenfloat x, svgenfloat y)
279 template <typename T>
280 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> copysign(T x,
281  T y) __NOEXC {
282  return __sycl_std::__invoke_copysign<T>(x, y);
283 }
284 
285 // svgenfloat cos (svgenfloat x)
286 template <typename T>
287 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> cos(T x) __NOEXC {
288  return __sycl_std::__invoke_cos<T>(x);
289 }
290 
291 // svgenfloat cosh (svgenfloat x)
292 template <typename T>
293 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> cosh(T x) __NOEXC {
294  return __sycl_std::__invoke_cosh<T>(x);
295 }
296 
297 // svgenfloat cospi (svgenfloat x)
298 template <typename T>
299 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> cospi(T x) __NOEXC {
300  return __sycl_std::__invoke_cospi<T>(x);
301 }
302 
303 // svgenfloat erfc (svgenfloat x)
304 template <typename T>
305 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> erfc(T x) __NOEXC {
306  return __sycl_std::__invoke_erfc<T>(x);
307 }
308 
309 // svgenfloat erf (svgenfloat x)
310 template <typename T>
311 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> erf(T x) __NOEXC {
312  return __sycl_std::__invoke_erf<T>(x);
313 }
314 
315 // svgenfloat exp (svgenfloat x )
316 template <typename T>
317 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp(T x) __NOEXC {
318  return __sycl_std::__invoke_exp<T>(x);
319 }
320 
321 // svgenfloat exp2 (svgenfloat x)
322 template <typename T>
323 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp2(T x) __NOEXC {
324  return __sycl_std::__invoke_exp2<T>(x);
325 }
326 
327 // svgenfloat exp10 (svgenfloat x)
328 template <typename T>
329 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp10(T x) __NOEXC {
330  return __sycl_std::__invoke_exp10<T>(x);
331 }
332 
333 // svgenfloat expm1 (svgenfloat x)
334 template <typename T>
335 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> expm1(T x) __NOEXC {
336  return __sycl_std::__invoke_expm1<T>(x);
337 }
338 
339 // svgenfloat fabs (svgenfloat x)
340 template <typename T>
341 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> fabs(T x) __NOEXC {
342  return __sycl_std::__invoke_fabs<T>(x);
343 }
344 
345 // svgenfloat fdim (svgenfloat x, svgenfloat y)
346 template <typename T>
347 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> fdim(T x, T y) __NOEXC {
348  return __sycl_std::__invoke_fdim<T>(x, y);
349 }
350 
351 // svgenfloat floor (svgenfloat x)
352 template <typename T>
353 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> floor(T x) __NOEXC {
354  return __sycl_std::__invoke_floor<T>(x);
355 }
356 
357 // svgenfloat fma (svgenfloat a, svgenfloat b, svgenfloat c)
358 template <typename T>
359 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> fma(T a, T b,
360  T c) __NOEXC {
361  return __sycl_std::__invoke_fma<T>(a, b, c);
362 }
363 
364 // svgenfloat fmax (svgenfloat x, svgenfloat y)
365 template <typename T>
366 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> fmax(T x, T y) __NOEXC {
367  return __sycl_std::__invoke_fmax<T>(x, y);
368 }
369 
370 // svgenfloat fmax (svgenfloat x, sgenfloat y)
371 template <typename T>
372 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
373 fmax(T x, typename T::element_type y) __NOEXC {
374  return __sycl_std::__invoke_fmax<T>(x, T(y));
375 }
376 
377 // svgenfloat fmin (svgenfloat x, svgenfloat y)
378 template <typename T>
379 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> fmin(T x, T y) __NOEXC {
380  return __sycl_std::__invoke_fmin<T>(x, y);
381 }
382 
383 // svgenfloat fmin (svgenfloat x, sgenfloat y)
384 template <typename T>
385 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
386 fmin(T x, typename T::element_type y) __NOEXC {
387  return __sycl_std::__invoke_fmin<T>(x, T(y));
388 }
389 
390 // svgenfloat fmod (svgenfloat x, svgenfloat y)
391 template <typename T>
392 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> fmod(T x, T y) __NOEXC {
393  return __sycl_std::__invoke_fmod<T>(x, y);
394 }
395 
396 // svgenfloat fract (svgenfloat x, genfloatptr iptr)
397 template <typename T, typename T2>
398 detail::enable_if_t<
399  detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
400 fract(T x, T2 iptr) __NOEXC {
401  detail::check_vector_size<T, T2>();
402  return __sycl_std::__invoke_fract<T>(x, iptr);
403 }
404 
405 // svgenfloat frexp (svgenfloat x, genintptr exp)
406 template <typename T, typename T2>
407 detail::enable_if_t<
408  detail::is_svgenfloat<T>::value && detail::is_genintptr<T2>::value, T>
409 frexp(T x, T2 exp) __NOEXC {
410  detail::check_vector_size<T, T2>();
411  return __sycl_std::__invoke_frexp<T>(x, exp);
412 }
413 
414 // svgenfloat hypot (svgenfloat x, svgenfloat y)
415 template <typename T>
416 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> hypot(T x,
417  T y) __NOEXC {
418  return __sycl_std::__invoke_hypot<T>(x, y);
419 }
420 
421 // genint ilogb (svgenfloat x)
422 template <typename T,
423  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
424 detail::change_base_type_t<T, int> ilogb(T x) __NOEXC {
425  return __sycl_std::__invoke_ilogb<detail::change_base_type_t<T, int>>(x);
426 }
427 
428 // float ldexp (float x, int k)
429 // double ldexp (double x, int k)
430 // half ldexp (half x, int k)
431 template <typename T>
432 detail::enable_if_t<detail::is_sgenfloat<T>::value, T> ldexp(T x,
433  int k) __NOEXC {
434  return __sycl_std::__invoke_ldexp<T>(x, k);
435 }
436 
437 // vgenfloat ldexp (vgenfloat x, int k)
438 template <typename T>
439 detail::enable_if_t<detail::is_vgenfloat<T>::value, T> ldexp(T x,
440  int k) __NOEXC {
441  return __sycl_std::__invoke_ldexp<T>(x, vec<int, T::size()>(k));
442 }
443 
444 // vgenfloat ldexp (vgenfloat x, genint k)
445 template <typename T, typename T2>
446 detail::enable_if_t<
447  detail::is_vgenfloat<T>::value && detail::is_intn<T2>::value, T>
448 ldexp(T x, T2 k) __NOEXC {
449  detail::check_vector_size<T, T2>();
450  return __sycl_std::__invoke_ldexp<T>(x, k);
451 }
452 
453 // svgenfloat lgamma (svgenfloat x)
454 template <typename T>
455 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> lgamma(T x) __NOEXC {
456  return __sycl_std::__invoke_lgamma<T>(x);
457 }
458 
459 // svgenfloat lgamma_r (svgenfloat x, genintptr signp)
460 template <typename T, typename T2>
461 detail::enable_if_t<
462  detail::is_svgenfloat<T>::value && detail::is_genintptr<T2>::value, T>
463 lgamma_r(T x, T2 signp) __NOEXC {
464  detail::check_vector_size<T, T2>();
465  return __sycl_std::__invoke_lgamma_r<T>(x, signp);
466 }
467 
468 // svgenfloat log (svgenfloat x)
469 template <typename T>
470 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log(T x) __NOEXC {
471  return __sycl_std::__invoke_log<T>(x);
472 }
473 
474 // svgenfloat log2 (svgenfloat x)
475 template <typename T>
476 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log2(T x) __NOEXC {
477  return __sycl_std::__invoke_log2<T>(x);
478 }
479 
480 // svgenfloat log10 (svgenfloat x)
481 template <typename T>
482 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log10(T x) __NOEXC {
483  return __sycl_std::__invoke_log10<T>(x);
484 }
485 
486 // svgenfloat log1p (svgenfloat x)
487 template <typename T>
488 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> log1p(T x) __NOEXC {
489  return __sycl_std::__invoke_log1p<T>(x);
490 }
491 
492 // svgenfloat logb (svgenfloat x)
493 template <typename T>
494 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> logb(T x) __NOEXC {
495  return __sycl_std::__invoke_logb<T>(x);
496 }
497 
498 // svgenfloat mad (svgenfloat a, svgenfloat b, svgenfloat c)
499 template <typename T>
500 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> mad(T a, T b,
501  T c) __NOEXC {
502  return __sycl_std::__invoke_mad<T>(a, b, c);
503 }
504 
505 // svgenfloat maxmag (svgenfloat x, svgenfloat y)
506 template <typename T>
507 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> maxmag(T x,
508  T y) __NOEXC {
509  return __sycl_std::__invoke_maxmag<T>(x, y);
510 }
511 
512 // svgenfloat minmag (svgenfloat x, svgenfloat y)
513 template <typename T>
514 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> minmag(T x,
515  T y) __NOEXC {
516  return __sycl_std::__invoke_minmag<T>(x, y);
517 }
518 
519 // svgenfloat modf (svgenfloat x, genfloatptr iptr)
520 template <typename T, typename T2>
521 detail::enable_if_t<
522  detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
523 modf(T x, T2 iptr) __NOEXC {
524  detail::check_vector_size<T, T2>();
525  return __sycl_std::__invoke_modf<T>(x, iptr);
526 }
527 
528 template <typename T,
529  typename = detail::enable_if_t<detail::is_nan_type<T>::value, T>>
530 detail::nan_return_t<T> nan(T nancode) __NOEXC {
531  return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
532  detail::convert_data_type<T, detail::nan_argument_base_t<T>>()(nancode));
533 }
534 
535 // svgenfloat nextafter (svgenfloat x, svgenfloat y)
536 template <typename T>
537 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> nextafter(T x,
538  T y) __NOEXC {
539  return __sycl_std::__invoke_nextafter<T>(x, y);
540 }
541 
542 // svgenfloat pow (svgenfloat x, svgenfloat y)
543 template <typename T>
544 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> pow(T x, T y) __NOEXC {
545  return __sycl_std::__invoke_pow<T>(x, y);
546 }
547 
548 // svgenfloat pown (svgenfloat x, genint y)
549 template <typename T, typename T2>
550 detail::enable_if_t<
551  detail::is_svgenfloat<T>::value && detail::is_genint<T2>::value, T>
552 pown(T x, T2 y) __NOEXC {
553  detail::check_vector_size<T, T2>();
554  return __sycl_std::__invoke_pown<T>(x, y);
555 }
556 
557 // svgenfloat powr (svgenfloat x, svgenfloat y)
558 template <typename T>
559 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> powr(T x, T y) __NOEXC {
560  return __sycl_std::__invoke_powr<T>(x, y);
561 }
562 
563 // svgenfloat remainder (svgenfloat x, svgenfloat y)
564 template <typename T>
565 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> remainder(T x,
566  T y) __NOEXC {
567  return __sycl_std::__invoke_remainder<T>(x, y);
568 }
569 
570 // svgenfloat remquo (svgenfloat x, svgenfloat y, genintptr quo)
571 template <typename T, typename T2>
572 detail::enable_if_t<
573  detail::is_svgenfloat<T>::value && detail::is_genintptr<T2>::value, T>
574 remquo(T x, T y, T2 quo) __NOEXC {
575  detail::check_vector_size<T, T2>();
576  return __sycl_std::__invoke_remquo<T>(x, y, quo);
577 }
578 
579 // svgenfloat rint (svgenfloat x)
580 template <typename T>
581 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> rint(T x) __NOEXC {
582  return __sycl_std::__invoke_rint<T>(x);
583 }
584 
585 // svgenfloat rootn (svgenfloat x, genint y)
586 template <typename T, typename T2>
587 detail::enable_if_t<
588  detail::is_svgenfloat<T>::value && detail::is_genint<T2>::value, T>
589 rootn(T x, T2 y) __NOEXC {
590  detail::check_vector_size<T, T2>();
591  return __sycl_std::__invoke_rootn<T>(x, y);
592 }
593 
594 // svgenfloat round (svgenfloat x)
595 template <typename T>
596 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> round(T x) __NOEXC {
597  return __sycl_std::__invoke_round<T>(x);
598 }
599 
600 // svgenfloat rsqrt (svgenfloat x)
601 template <typename T>
602 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> rsqrt(T x) __NOEXC {
603  return __sycl_std::__invoke_rsqrt<T>(x);
604 }
605 
606 // svgenfloat sin (svgenfloat x)
607 template <typename T>
608 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sin(T x) __NOEXC {
609  return __sycl_std::__invoke_sin<T>(x);
610 }
611 
612 // svgenfloat sincos (svgenfloat x, genfloatptr cosval)
613 template <typename T, typename T2>
614 detail::enable_if_t<
615  detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
616 sincos(T x, T2 cosval) __NOEXC {
617  detail::check_vector_size<T, T2>();
618  return __sycl_std::__invoke_sincos<T>(x, cosval);
619 }
620 
621 // svgenfloat sinh (svgenfloat x)
622 template <typename T>
623 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> sinh(T x) __NOEXC {
624  return __sycl_std::__invoke_sinh<T>(x);
625 }
626 
627 // svgenfloat sinpi (svgenfloat x)
628 template <typename T>
629 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> sinpi(T x) __NOEXC {
630  return __sycl_std::__invoke_sinpi<T>(x);
631 }
632 
633 // svgenfloat sqrt (svgenfloat x)
634 template <typename T>
635 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sqrt(T x) __NOEXC {
636  return __sycl_std::__invoke_sqrt<T>(x);
637 }
638 
639 // svgenfloat tan (svgenfloat x)
640 template <typename T>
641 detail::enable_if_t<__FAST_MATH_GENFLOAT(T), T> tan(T x) __NOEXC {
642  return __sycl_std::__invoke_tan<T>(x);
643 }
644 
645 // svgenfloat tanh (svgenfloat x)
646 template <typename T>
647 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> tanh(T x) __NOEXC {
648  return __sycl_std::__invoke_tanh<T>(x);
649 }
650 
651 // svgenfloat tanpi (svgenfloat x)
652 template <typename T>
653 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> tanpi(T x) __NOEXC {
654  return __sycl_std::__invoke_tanpi<T>(x);
655 }
656 
657 // svgenfloat tgamma (svgenfloat x)
658 template <typename T>
659 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> tgamma(T x) __NOEXC {
660  return __sycl_std::__invoke_tgamma<T>(x);
661 }
662 
663 // svgenfloat trunc (svgenfloat x)
664 template <typename T>
665 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> trunc(T x) __NOEXC {
666  return __sycl_std::__invoke_trunc<T>(x);
667 }
668 
669 /* --------------- 4.13.5 Common functions. ---------------------------------*/
670 // svgenfloat clamp (svgenfloat x, svgenfloat minval, svgenfloat maxval)
671 template <typename T>
672 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>
673 clamp(T x, T minval, T maxval) __NOEXC {
674  return __sycl_std::__invoke_fclamp<T>(x, minval, maxval);
675 }
676 
677 // vgenfloath clamp (vgenfloath x, half minval, half maxval)
678 // vgenfloatf clamp (vgenfloatf x, float minval, float maxval)
679 // vgenfloatd clamp (vgenfloatd x, double minval, double maxval)
680 template <typename T>
681 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
682 clamp(T x, typename T::element_type minval,
683  typename T::element_type maxval) __NOEXC {
684  return __sycl_std::__invoke_fclamp<T>(x, T(minval), T(maxval));
685 }
686 
687 // svgenfloat degrees (svgenfloat radians)
688 template <typename T>
689 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>
690 degrees(T radians) __NOEXC {
691  return __sycl_std::__invoke_degrees<T>(radians);
692 }
693 
694 // svgenfloat abs (svgenfloat x)
695 template <typename T>
696 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> abs(T x) __NOEXC {
697  return __sycl_std::__invoke_fabs<T>(x);
698 }
699 
700 // svgenfloat max (svgenfloat x, svgenfloat y)
701 template <typename T>
702 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>(max)(T x, T y) __NOEXC {
703  return __sycl_std::__invoke_fmax_common<T>(x, y);
704 }
705 
706 // vgenfloatf max (vgenfloatf x, float y)
707 // vgenfloatd max (vgenfloatd x, double y)
708 // vgenfloath max (vgenfloath x, half y)
709 template <typename T>
710 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(max)(
711  T x, typename T::element_type y) __NOEXC {
712  return __sycl_std::__invoke_fmax_common<T>(x, T(y));
713 }
714 
715 // svgenfloat min (svgenfloat x, svgenfloat y)
716 template <typename T>
717 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>(min)(T x, T y) __NOEXC {
718  return __sycl_std::__invoke_fmin_common<T>(x, y);
719 }
720 
721 // vgenfloatf min (vgenfloatf x, float y)
722 // vgenfloatd min (vgenfloatd x, double y)
723 // vgenfloath min (vgenfloath x, half y)
724 template <typename T>
725 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(min)(
726  T x, typename T::element_type y) __NOEXC {
727  return __sycl_std::__invoke_fmin_common<T>(x, T(y));
728 }
729 
730 // svgenfloat mix (svgenfloat x, svgenfloat y, svgenfloat a)
731 template <typename T>
732 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> mix(T x, T y,
733  T a) __NOEXC {
734  return __sycl_std::__invoke_mix<T>(x, y, a);
735 }
736 
737 // vgenfloatf mix (vgenfloatf x, vgenfloatf y, float a)
738 // vgenfloatd mix (vgenfloatd x, vgenfloatd y, double a)
739 // vgenfloatd mix (vgenfloath x, vgenfloath y, half a)
740 template <typename T>
741 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
742 mix(T x, T y, typename T::element_type a) __NOEXC {
743  return __sycl_std::__invoke_mix<T>(x, y, T(a));
744 }
745 
746 // svgenfloat radians (svgenfloat degrees)
747 template <typename T>
748 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>
749 radians(T degrees) __NOEXC {
750  return __sycl_std::__invoke_radians<T>(degrees);
751 }
752 
753 // svgenfloat step (svgenfloat edge, svgenfloat x)
754 template <typename T>
755 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> step(T edge,
756  T x) __NOEXC {
757  return __sycl_std::__invoke_step<T>(edge, x);
758 }
759 
760 // vgenfloatf step (float edge, vgenfloatf x)
761 // vgenfloatd step (double edge, vgenfloatd x)
762 // vgenfloatd step (half edge, vgenfloath x)
763 template <typename T>
764 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
765 step(typename T::element_type edge, T x) __NOEXC {
766  return __sycl_std::__invoke_step<T>(T(edge), x);
767 }
768 
769 // svgenfloat smoothstep (svgenfloat edge0, svgenfloat edge1, svgenfloat x)
770 template <typename T>
771 detail::enable_if_t<detail::is_svgenfloat<T>::value, T>
772 smoothstep(T edge0, T edge1, T x) __NOEXC {
773  return __sycl_std::__invoke_smoothstep<T>(edge0, edge1, x);
774 }
775 
776 // vgenfloatf smoothstep (float edge0, float edge1, vgenfloatf x)
777 // vgenfloatd smoothstep (double edge0, double edge1, vgenfloatd x)
778 // vgenfloath smoothstep (half edge0, half edge1, vgenfloath x)
779 template <typename T>
780 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
781 smoothstep(typename T::element_type edge0, typename T::element_type edge1,
782  T x) __NOEXC {
783  return __sycl_std::__invoke_smoothstep<T>(T(edge0), T(edge1), x);
784 }
785 
786 // svgenfloat sign (svgenfloat x)
787 template <typename T>
788 detail::enable_if_t<detail::is_svgenfloat<T>::value, T> sign(T x) __NOEXC {
789  return __sycl_std::__invoke_sign<T>(x);
790 }
791 
792 /* --------------- 4.13.4 Integer functions. --------------------------------*/
793 // ugeninteger abs (geninteger x)
794 template <typename T>
795 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> abs(T x) __NOEXC {
796  return __sycl_std::__invoke_u_abs<T>(x);
797 }
798 
799 // ugeninteger abs (geninteger x)
800 template <typename T>
801 detail::enable_if_t<detail::is_igeninteger<T>::value,
802  detail::make_unsigned_t<T>>
803 abs(T x) __NOEXC {
804  return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
805 }
806 
807 // ugeninteger abs_diff (geninteger x, geninteger y)
808 template <typename T>
809 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> abs_diff(T x,
810  T y) __NOEXC {
811  return __sycl_std::__invoke_u_abs_diff<T>(x, y);
812 }
813 
814 // ugeninteger abs_diff (geninteger x, geninteger y)
815 template <typename T>
816 detail::enable_if_t<detail::is_igeninteger<T>::value,
817  detail::make_unsigned_t<T>>
818 abs_diff(T x, T y) __NOEXC {
819  return __sycl_std::__invoke_s_abs_diff<detail::make_unsigned_t<T>>(x, y);
820 }
821 
822 // geninteger add_sat (geninteger x, geninteger y)
823 template <typename T>
824 detail::enable_if_t<detail::is_igeninteger<T>::value, T> add_sat(T x,
825  T y) __NOEXC {
826  return __sycl_std::__invoke_s_add_sat<T>(x, y);
827 }
828 
829 // geninteger add_sat (geninteger x, geninteger y)
830 template <typename T>
831 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> add_sat(T x,
832  T y) __NOEXC {
833  return __sycl_std::__invoke_u_add_sat<T>(x, y);
834 }
835 
836 // geninteger hadd (geninteger x, geninteger y)
837 template <typename T>
838 detail::enable_if_t<detail::is_igeninteger<T>::value, T> hadd(T x,
839  T y) __NOEXC {
840  return __sycl_std::__invoke_s_hadd<T>(x, y);
841 }
842 
843 // geninteger hadd (geninteger x, geninteger y)
844 template <typename T>
845 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> hadd(T x,
846  T y) __NOEXC {
847  return __sycl_std::__invoke_u_hadd<T>(x, y);
848 }
849 
850 // geninteger rhadd (geninteger x, geninteger y)
851 template <typename T>
852 detail::enable_if_t<detail::is_igeninteger<T>::value, T> rhadd(T x,
853  T y) __NOEXC {
854  return __sycl_std::__invoke_s_rhadd<T>(x, y);
855 }
856 
857 // geninteger rhadd (geninteger x, geninteger y)
858 template <typename T>
859 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> rhadd(T x,
860  T y) __NOEXC {
861  return __sycl_std::__invoke_u_rhadd<T>(x, y);
862 }
863 
864 // geninteger clamp (geninteger x, geninteger minval, geninteger maxval)
865 template <typename T>
866 detail::enable_if_t<detail::is_igeninteger<T>::value, T>
867 clamp(T x, T minval, T maxval) __NOEXC {
868  return __sycl_std::__invoke_s_clamp<T>(x, minval, maxval);
869 }
870 
871 // geninteger clamp (geninteger x, geninteger minval, geninteger maxval)
872 template <typename T>
873 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>
874 clamp(T x, T minval, T maxval) __NOEXC {
875  return __sycl_std::__invoke_u_clamp<T>(x, minval, maxval);
876 }
877 
878 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
879 template <typename T>
880 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>
881 clamp(T x, typename T::element_type minval,
882  typename T::element_type maxval) __NOEXC {
883  return __sycl_std::__invoke_s_clamp<T>(x, T(minval), T(maxval));
884 }
885 
886 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
887 template <typename T>
888 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>
889 clamp(T x, typename T::element_type minval,
890  typename T::element_type maxval) __NOEXC {
891  return __sycl_std::__invoke_u_clamp<T>(x, T(minval), T(maxval));
892 }
893 
894 // geninteger clz (geninteger x)
895 template <typename T>
896 detail::enable_if_t<detail::is_geninteger<T>::value, T> clz(T x) __NOEXC {
897  return __sycl_std::__invoke_clz<T>(x);
898 }
899 
900 // geninteger ctz (geninteger x)
901 template <typename T>
902 detail::enable_if_t<detail::is_geninteger<T>::value, T> ctz(T x) __NOEXC {
903  return __sycl_std::__invoke_ctz<T>(x);
904 }
905 
906 // geninteger ctz (geninteger x) for calls with deprecated namespace
907 namespace ext::intel {
908 template <typename T>
909 __SYCL_DEPRECATED(
910  "'sycl::ext::intel::ctz' is deprecated, use 'sycl::ctz' instead")
911 sycl::detail::enable_if_t<sycl::detail::is_geninteger<T>::value, T> ctz(
912  T x) __NOEXC {
913  return sycl::ctz(x);
914 }
915 } // namespace ext::intel
916 
917 namespace __SYCL2020_DEPRECATED("use 'ext::intel' instead") intel {
918 using namespace ext::intel;
919 }
920 
921 // geninteger mad_hi (geninteger a, geninteger b, geninteger c)
922 template <typename T>
923 detail::enable_if_t<detail::is_igeninteger<T>::value, T> mad_hi(T x, T y,
924  T z) __NOEXC {
925  return __sycl_std::__invoke_s_mad_hi<T>(x, y, z);
926 }
927 
928 // geninteger mad_hi (geninteger a, geninteger b, geninteger c)
929 template <typename T>
930 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> mad_hi(T x, T y,
931  T z) __NOEXC {
932  return __sycl_std::__invoke_u_mad_hi<T>(x, y, z);
933 }
934 
935 // geninteger mad_sat (geninteger a, geninteger b, geninteger c)
936 template <typename T>
937 detail::enable_if_t<detail::is_igeninteger<T>::value, T> mad_sat(T a, T b,
938  T c) __NOEXC {
939  return __sycl_std::__invoke_s_mad_sat<T>(a, b, c);
940 }
941 
942 // geninteger mad_sat (geninteger a, geninteger b, geninteger c)
943 template <typename T>
944 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> mad_sat(T a, T b,
945  T c) __NOEXC {
946  return __sycl_std::__invoke_u_mad_sat<T>(a, b, c);
947 }
948 
949 // igeninteger max (igeninteger x, igeninteger y)
950 template <typename T>
951 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(max)(T x,
952  T y) __NOEXC {
953  return __sycl_std::__invoke_s_max<T>(x, y);
954 }
955 
956 // ugeninteger max (ugeninteger x, ugeninteger y)
957 template <typename T>
958 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>(max)(T x,
959  T y) __NOEXC {
960  return __sycl_std::__invoke_u_max<T>(x, y);
961 }
962 
963 // igeninteger max (vigeninteger x, sigeninteger y)
964 template <typename T>
965 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>(max)(
966  T x, typename T::element_type y) __NOEXC {
967  return __sycl_std::__invoke_s_max<T>(x, T(y));
968 }
969 
970 // vugeninteger max (vugeninteger x, sugeninteger y)
971 template <typename T>
972 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>(max)(
973  T x, typename T::element_type y) __NOEXC {
974  return __sycl_std::__invoke_u_max<T>(x, T(y));
975 }
976 
977 // igeninteger min (igeninteger x, igeninteger y)
978 template <typename T>
979 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(min)(T x,
980  T y) __NOEXC {
981  return __sycl_std::__invoke_s_min<T>(x, y);
982 }
983 
984 // ugeninteger min (ugeninteger x, ugeninteger y)
985 template <typename T>
986 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>(min)(T x,
987  T y) __NOEXC {
988  return __sycl_std::__invoke_u_min<T>(x, y);
989 }
990 
991 // vigeninteger min (vigeninteger x, sigeninteger y)
992 template <typename T>
993 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>(min)(
994  T x, typename T::element_type y) __NOEXC {
995  return __sycl_std::__invoke_s_min<T>(x, T(y));
996 }
997 
998 // vugeninteger min (vugeninteger x, sugeninteger y)
999 template <typename T>
1000 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>(min)(
1001  T x, typename T::element_type y) __NOEXC {
1002  return __sycl_std::__invoke_u_min<T>(x, T(y));
1003 }
1004 
1005 // geninteger mul_hi (geninteger x, geninteger y)
1006 template <typename T>
1007 detail::enable_if_t<detail::is_igeninteger<T>::value, T> mul_hi(T x,
1008  T y) __NOEXC {
1009  return __sycl_std::__invoke_s_mul_hi<T>(x, y);
1010 }
1011 
1012 // geninteger mul_hi (geninteger x, geninteger y)
1013 template <typename T>
1014 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> mul_hi(T x,
1015  T y) __NOEXC {
1016  return __sycl_std::__invoke_u_mul_hi<T>(x, y);
1017 }
1018 
1019 // geninteger rotate (geninteger v, geninteger i)
1020 template <typename T>
1021 detail::enable_if_t<detail::is_geninteger<T>::value, T> rotate(T v,
1022  T i) __NOEXC {
1023  return __sycl_std::__invoke_rotate<T>(v, i);
1024 }
1025 
1026 // geninteger sub_sat (geninteger x, geninteger y)
1027 template <typename T>
1028 detail::enable_if_t<detail::is_igeninteger<T>::value, T> sub_sat(T x,
1029  T y) __NOEXC {
1030  return __sycl_std::__invoke_s_sub_sat<T>(x, y);
1031 }
1032 
1033 // geninteger sub_sat (geninteger x, geninteger y)
1034 template <typename T>
1035 detail::enable_if_t<detail::is_ugeninteger<T>::value, T> sub_sat(T x,
1036  T y) __NOEXC {
1037  return __sycl_std::__invoke_u_sub_sat<T>(x, y);
1038 }
1039 
1040 // ugeninteger16bit upsample (ugeninteger8bit hi, ugeninteger8bit lo)
1041 template <typename T>
1042 detail::enable_if_t<detail::is_ugeninteger8bit<T>::value,
1043  detail::make_larger_t<T>>
1044 upsample(T hi, T lo) __NOEXC {
1045  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
1046 }
1047 
1048 // igeninteger16bit upsample (igeninteger8bit hi, ugeninteger8bit lo)
1049 template <typename T, typename T2>
1050 detail::enable_if_t<detail::is_igeninteger8bit<T>::value &&
1051  detail::is_ugeninteger8bit<T2>::value,
1052  detail::make_larger_t<T>>
1053 upsample(T hi, T2 lo) __NOEXC {
1054  detail::check_vector_size<T, T2>();
1055  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
1056 }
1057 
1058 // ugeninteger32bit upsample (ugeninteger16bit hi, ugeninteger16bit lo)
1059 template <typename T>
1060 detail::enable_if_t<detail::is_ugeninteger16bit<T>::value,
1061  detail::make_larger_t<T>>
1062 upsample(T hi, T lo) __NOEXC {
1063  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
1064 }
1065 
1066 // igeninteger32bit upsample (igeninteger16bit hi, ugeninteger16bit lo)
1067 template <typename T, typename T2>
1068 detail::enable_if_t<detail::is_igeninteger16bit<T>::value &&
1069  detail::is_ugeninteger16bit<T2>::value,
1070  detail::make_larger_t<T>>
1071 upsample(T hi, T2 lo) __NOEXC {
1072  detail::check_vector_size<T, T2>();
1073  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
1074 }
1075 
1076 // ugeninteger64bit upsample (ugeninteger32bit hi, ugeninteger32bit lo)
1077 template <typename T>
1078 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value,
1079  detail::make_larger_t<T>>
1080 upsample(T hi, T lo) __NOEXC {
1081  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
1082 }
1083 
1084 // igeninteger64bit upsample (igeninteger32bit hi, ugeninteger32bit lo)
1085 template <typename T, typename T2>
1086 detail::enable_if_t<detail::is_igeninteger32bit<T>::value &&
1087  detail::is_ugeninteger32bit<T2>::value,
1088  detail::make_larger_t<T>>
1089 upsample(T hi, T2 lo) __NOEXC {
1090  detail::check_vector_size<T, T2>();
1091  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
1092 }
1093 
1094 // geninteger popcount (geninteger x)
1095 template <typename T>
1096 detail::enable_if_t<detail::is_geninteger<T>::value, T> popcount(T x) __NOEXC {
1097  return __sycl_std::__invoke_popcount<T>(x);
1098 }
1099 
1100 // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y,
1101 // geninteger32bit z)
1102 template <typename T>
1103 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
1104 mad24(T x, T y, T z) __NOEXC {
1105  return __sycl_std::__invoke_s_mad24<T>(x, y, z);
1106 }
1107 
1108 // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y,
1109 // geninteger32bit z)
1110 template <typename T>
1111 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
1112 mad24(T x, T y, T z) __NOEXC {
1113  return __sycl_std::__invoke_u_mad24<T>(x, y, z);
1114 }
1115 
1116 // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y)
1117 template <typename T>
1118 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
1119 mul24(T x, T y) __NOEXC {
1120  return __sycl_std::__invoke_s_mul24<T>(x, y);
1121 }
1122 
1123 // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y)
1124 template <typename T>
1125 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
1126 mul24(T x, T y) __NOEXC {
1127  return __sycl_std::__invoke_u_mul24<T>(x, y);
1128 }
1129 
1130 /* --------------- 4.13.6 Geometric Functions. ------------------------------*/
1131 // float3 cross (float3 p0, float3 p1)
1132 // float4 cross (float4 p0, float4 p1)
1133 // double3 cross (double3 p0, double3 p1)
1134 // double4 cross (double4 p0, double4 p1)
1135 // half3 cross (half3 p0, half3 p1)
1136 // half4 cross (half4 p0, half4 p1)
1137 template <typename T>
1138 detail::enable_if_t<detail::is_gencross<T>::value, T> cross(T p0,
1139  T p1) __NOEXC {
1140  return __sycl_std::__invoke_cross<T>(p0, p1);
1141 }
1142 
1143 // float dot (float p0, float p1)
1144 // double dot (double p0, double p1)
1145 // half dot (half p0, half p1)
1146 template <typename T>
1147 detail::enable_if_t<detail::is_sgenfloat<T>::value, T> dot(T p0, T p1) __NOEXC {
1148  return p0 * p1;
1149 }
1150 
1151 // float dot (vgengeofloat p0, vgengeofloat p1)
1152 template <typename T>
1153 detail::enable_if_t<detail::is_vgengeofloat<T>::value, float>
1154 dot(T p0, T p1) __NOEXC {
1155  return __sycl_std::__invoke_Dot<float>(p0, p1);
1156 }
1157 
1158 // double dot (vgengeodouble p0, vgengeodouble p1)
1159 template <typename T>
1160 detail::enable_if_t<detail::is_vgengeodouble<T>::value, double>
1161 dot(T p0, T p1) __NOEXC {
1162  return __sycl_std::__invoke_Dot<double>(p0, p1);
1163 }
1164 
1165 // half dot (vgengeohalf p0, vgengeohalf p1)
1166 template <typename T>
1167 detail::enable_if_t<detail::is_vgengeohalf<T>::value, half> dot(T p0,
1168  T p1) __NOEXC {
1169  return __sycl_std::__invoke_Dot<half>(p0, p1);
1170 }
1171 
1172 // float distance (gengeofloat p0, gengeofloat p1)
1173 template <typename T,
1174  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1175 float distance(T p0, T p1) __NOEXC {
1176  return __sycl_std::__invoke_distance<float>(p0, p1);
1177 }
1178 
1179 // double distance (gengeodouble p0, gengeodouble p1)
1180 template <typename T,
1181  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1182 double distance(T p0, T p1) __NOEXC {
1183  return __sycl_std::__invoke_distance<double>(p0, p1);
1184 }
1185 
1186 // half distance (gengeohalf p0, gengeohalf p1)
1187 template <typename T,
1188  typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1189 half distance(T p0, T p1) __NOEXC {
1190  return __sycl_std::__invoke_distance<half>(p0, p1);
1191 }
1192 
1193 // float length (gengeofloat p)
1194 template <typename T,
1195  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1196 float length(T p) __NOEXC {
1197  return __sycl_std::__invoke_length<float>(p);
1198 }
1199 
1200 // double length (gengeodouble p)
1201 template <typename T,
1202  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1203 double length(T p) __NOEXC {
1204  return __sycl_std::__invoke_length<double>(p);
1205 }
1206 
1207 // half length (gengeohalf p)
1208 template <typename T,
1209  typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1210 half length(T p) __NOEXC {
1211  return __sycl_std::__invoke_length<half>(p);
1212 }
1213 
1214 // gengeofloat normalize (gengeofloat p)
1215 template <typename T>
1216 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1217 normalize(T p) __NOEXC {
1218  return __sycl_std::__invoke_normalize<T>(p);
1219 }
1220 
1221 // gengeodouble normalize (gengeodouble p)
1222 template <typename T>
1223 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1224 normalize(T p) __NOEXC {
1225  return __sycl_std::__invoke_normalize<T>(p);
1226 }
1227 
1228 // gengeohalf normalize (gengeohalf p)
1229 template <typename T>
1230 detail::enable_if_t<detail::is_gengeohalf<T>::value, T> normalize(T p) __NOEXC {
1231  return __sycl_std::__invoke_normalize<T>(p);
1232 }
1233 
1234 // float fast_distance (gengeofloat p0, gengeofloat p1)
1235 template <typename T,
1236  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1237 float fast_distance(T p0, T p1) __NOEXC {
1238  return __sycl_std::__invoke_fast_distance<float>(p0, p1);
1239 }
1240 
1241 // double fast_distance (gengeodouble p0, gengeodouble p1)
1242 template <typename T,
1243  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1244 double fast_distance(T p0, T p1) __NOEXC {
1245  return __sycl_std::__invoke_fast_distance<double>(p0, p1);
1246 }
1247 
1248 // float fast_length (gengeofloat p)
1249 template <typename T,
1250  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1251 float fast_length(T p) __NOEXC {
1252  return __sycl_std::__invoke_fast_length<float>(p);
1253 }
1254 
1255 // double fast_length (gengeodouble p)
1256 template <typename T,
1257  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1258 double fast_length(T p) __NOEXC {
1259  return __sycl_std::__invoke_fast_length<double>(p);
1260 }
1261 
1262 // gengeofloat fast_normalize (gengeofloat p)
1263 template <typename T>
1264 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1265 fast_normalize(T p) __NOEXC {
1266  return __sycl_std::__invoke_fast_normalize<T>(p);
1267 }
1268 
1269 // gengeodouble fast_normalize (gengeodouble p)
1270 template <typename T>
1271 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1272 fast_normalize(T p) __NOEXC {
1273  return __sycl_std::__invoke_fast_normalize<T>(p);
1274 }
1275 
1276 /* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/
1277 /* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/
1278 
1279 template <typename T,
1280  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1281 detail::common_rel_ret_t<T> isequal(T x, T y) __NOEXC {
1282  return detail::RelConverter<T>::apply(
1283  __sycl_std::__invoke_FOrdEqual<detail::internal_rel_ret_t<T>>(x, y));
1284 }
1285 
1286 template <typename T,
1287  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1288 detail::common_rel_ret_t<T> isnotequal(T x, T y) __NOEXC {
1289  return detail::RelConverter<T>::apply(
1290  __sycl_std::__invoke_FUnordNotEqual<detail::internal_rel_ret_t<T>>(x, y));
1291 }
1292 
1293 template <typename T,
1294  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1295 detail::common_rel_ret_t<T> isgreater(T x, T y) __NOEXC {
1296  return detail::RelConverter<T>::apply(
1297  __sycl_std::__invoke_FOrdGreaterThan<detail::internal_rel_ret_t<T>>(x,
1298  y));
1299 }
1300 
1301 template <typename T,
1302  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1303 detail::common_rel_ret_t<T> isgreaterequal(T x, T y) __NOEXC {
1304  return detail::RelConverter<T>::apply(
1305  __sycl_std::__invoke_FOrdGreaterThanEqual<detail::internal_rel_ret_t<T>>(
1306  x, y));
1307 }
1308 
1309 template <typename T,
1310  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1311 detail::common_rel_ret_t<T> isless(T x, T y) __NOEXC {
1312  return detail::RelConverter<T>::apply(
1313  __sycl_std::__invoke_FOrdLessThan<detail::internal_rel_ret_t<T>>(x, y));
1314 }
1315 
1316 template <typename T,
1317  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1318 detail::common_rel_ret_t<T> islessequal(T x, T y) __NOEXC {
1319  return detail::RelConverter<T>::apply(
1320  __sycl_std::__invoke_FOrdLessThanEqual<detail::internal_rel_ret_t<T>>(x,
1321  y));
1322 }
1323 
1324 template <typename T,
1325  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1326 detail::common_rel_ret_t<T> islessgreater(T x, T y) __NOEXC {
1327  return detail::RelConverter<T>::apply(
1328  __sycl_std::__invoke_FOrdNotEqual<detail::internal_rel_ret_t<T>>(x, y));
1329 }
1330 
1331 template <typename T,
1332  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1333 detail::common_rel_ret_t<T> isfinite(T x) __NOEXC {
1334  return detail::RelConverter<T>::apply(
1335  __sycl_std::__invoke_IsFinite<detail::internal_rel_ret_t<T>>(x));
1336 }
1337 
1338 template <typename T,
1339  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1340 detail::common_rel_ret_t<T> isinf(T x) __NOEXC {
1341  return detail::RelConverter<T>::apply(
1342  __sycl_std::__invoke_IsInf<detail::internal_rel_ret_t<T>>(x));
1343 }
1344 
1345 template <typename T,
1346  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1347 detail::common_rel_ret_t<T> isnan(T x) __NOEXC {
1348  return detail::RelConverter<T>::apply(
1349  __sycl_std::__invoke_IsNan<detail::internal_rel_ret_t<T>>(x));
1350 }
1351 
1352 template <typename T,
1353  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1354 detail::common_rel_ret_t<T> isnormal(T x) __NOEXC {
1355  return detail::RelConverter<T>::apply(
1356  __sycl_std::__invoke_IsNormal<detail::internal_rel_ret_t<T>>(x));
1357 }
1358 
1359 template <typename T,
1360  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1361 detail::common_rel_ret_t<T> isordered(T x, T y) __NOEXC {
1362  return detail::RelConverter<T>::apply(
1363  __sycl_std::__invoke_Ordered<detail::internal_rel_ret_t<T>>(x, y));
1364 }
1365 
1366 template <typename T,
1367  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1368 detail::common_rel_ret_t<T> isunordered(T x, T y) __NOEXC {
1369  return detail::RelConverter<T>::apply(
1370  __sycl_std::__invoke_Unordered<detail::internal_rel_ret_t<T>>(x, y));
1371 }
1372 
1373 template <typename T,
1374  typename = detail::enable_if_t<detail::is_svgenfloat<T>::value, T>>
1375 detail::common_rel_ret_t<T> signbit(T x) __NOEXC {
1376  return detail::RelConverter<T>::apply(
1377  __sycl_std::__invoke_SignBitSet<detail::internal_rel_ret_t<T>>(x));
1378 }
1379 
1380 namespace detail {
1381 #if defined(SYCL2020_CONFORMANT_APIS) && SYCL_LANGUAGE_VERSION >= 202001
1382 using anyall_ret_t = bool;
1383 #else
1384 using anyall_ret_t = int;
1385 #endif
1386 } // namespace detail
1387 
1388 // int any (sigeninteger x)
1389 template <typename T>
1390 detail::enable_if_t<detail::is_sigeninteger<T>::value, detail::anyall_ret_t>
1391 any(T x) __NOEXC {
1392  return detail::Boolean<1>(int(detail::msbIsSet(x)));
1393 }
1394 
1395 // int any (vigeninteger x)
1396 template <typename T>
1397 detail::enable_if_t<detail::is_vigeninteger<T>::value, detail::anyall_ret_t>
1398 any(T x) __NOEXC {
1399  return detail::rel_sign_bit_test_ret_t<T>(
1400  __sycl_std::__invoke_Any<detail::rel_sign_bit_test_ret_t<T>>(
1401  detail::rel_sign_bit_test_arg_t<T>(x)));
1402 }
1403 
1404 // int all (sigeninteger x)
1405 template <typename T>
1406 detail::enable_if_t<detail::is_sigeninteger<T>::value, detail::anyall_ret_t>
1407 all(T x) __NOEXC {
1408  return detail::Boolean<1>(int(detail::msbIsSet(x)));
1409 }
1410 
1411 // int all (vigeninteger x)
1412 template <typename T>
1413 detail::enable_if_t<detail::is_vigeninteger<T>::value, detail::anyall_ret_t>
1414 all(T x) __NOEXC {
1415  return detail::rel_sign_bit_test_ret_t<T>(
1416  __sycl_std::__invoke_All<detail::rel_sign_bit_test_ret_t<T>>(
1417  detail::rel_sign_bit_test_arg_t<T>(x)));
1418 }
1419 
1420 // gentype bitselect (gentype a, gentype b, gentype c)
1421 template <typename T>
1422 detail::enable_if_t<detail::is_gentype<T>::value, T> bitselect(T a, T b,
1423  T c) __NOEXC {
1424  return __sycl_std::__invoke_bitselect<T>(a, b, c);
1425 }
1426 
1427 // sgentype select (sgentype a, sgentype b, bool c)
1428 template <typename T>
1429 detail::enable_if_t<detail::is_sgentype<T>::value, T> select(T a, T b,
1430  bool c) __NOEXC {
1431  return __sycl_std::__invoke_select<T>(a, b, static_cast<int>(c));
1432 }
1433 
1434 // geninteger select (geninteger a, geninteger b, igeninteger c)
1435 template <typename T, typename T2>
1436 detail::enable_if_t<
1437  detail::is_geninteger<T>::value && detail::is_igeninteger<T2>::value, T>
1438 select(T a, T b, T2 c) __NOEXC {
1439  detail::check_vector_size<T, T2>();
1440  return __sycl_std::__invoke_select<T>(a, b, c);
1441 }
1442 
1443 // geninteger select (geninteger a, geninteger b, ugeninteger c)
1444 template <typename T, typename T2>
1445 detail::enable_if_t<
1446  detail::is_geninteger<T>::value && detail::is_ugeninteger<T2>::value, T>
1447 select(T a, T b, T2 c) __NOEXC {
1448  detail::check_vector_size<T, T2>();
1449  return __sycl_std::__invoke_select<T>(a, b, c);
1450 }
1451 
1452 // svgenfloatf select (svgenfloatf a, svgenfloatf b, genint c)
1453 template <typename T, typename T2>
1454 detail::enable_if_t<
1455  detail::is_svgenfloatf<T>::value && detail::is_genint<T2>::value, T>
1456 select(T a, T b, T2 c) __NOEXC {
1457  detail::check_vector_size<T, T2>();
1458  return __sycl_std::__invoke_select<T>(a, b, c);
1459 }
1460 
1461 // svgenfloatf select (svgenfloatf a, svgenfloatf b, ugenint c)
1462 template <typename T, typename T2>
1463 detail::enable_if_t<
1464  detail::is_svgenfloatf<T>::value && detail::is_ugenint<T2>::value, T>
1465 select(T a, T b, T2 c) __NOEXC {
1466  detail::check_vector_size<T, T2>();
1467  return __sycl_std::__invoke_select<T>(a, b, c);
1468 }
1469 
1470 // svgenfloatd select (svgenfloatd a, svgenfloatd b, igeninteger64 c)
1471 template <typename T, typename T2>
1472 detail::enable_if_t<detail::is_svgenfloatd<T>::value &&
1473  detail::is_igeninteger64bit<T2>::value,
1474  T>
1475 select(T a, T b, T2 c) __NOEXC {
1476  detail::check_vector_size<T, T2>();
1477  return __sycl_std::__invoke_select<T>(a, b, c);
1478 }
1479 
1480 // svgenfloatd select (svgenfloatd a, svgenfloatd b, ugeninteger64 c)
1481 template <typename T, typename T2>
1482 detail::enable_if_t<detail::is_svgenfloatd<T>::value &&
1483  detail::is_ugeninteger64bit<T2>::value,
1484  T>
1485 select(T a, T b, T2 c) __NOEXC {
1486  detail::check_vector_size<T, T2>();
1487  return __sycl_std::__invoke_select<T>(a, b, c);
1488 }
1489 
1490 // svgenfloath select (svgenfloath a, svgenfloath b, igeninteger16 c)
1491 template <typename T, typename T2>
1492 detail::enable_if_t<detail::is_svgenfloath<T>::value &&
1493  detail::is_igeninteger16bit<T2>::value,
1494  T>
1495 select(T a, T b, T2 c) __NOEXC {
1496  detail::check_vector_size<T, T2>();
1497  return __sycl_std::__invoke_select<T>(a, b, c);
1498 }
1499 
1500 // svgenfloath select (svgenfloath a, svgenfloath b, ugeninteger16 c)
1501 template <typename T, typename T2>
1502 detail::enable_if_t<detail::is_svgenfloath<T>::value &&
1503  detail::is_ugeninteger16bit<T2>::value,
1504  T>
1505 select(T a, T b, T2 c) __NOEXC {
1506  detail::check_vector_size<T, T2>();
1507  return __sycl_std::__invoke_select<T>(a, b, c);
1508 }
1509 
1510 namespace native {
1511 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1512 
1513 #define __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(NAME) \
1514  template <size_t N> \
1515  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x) \
1516  __NOEXC { \
1517  marray<float, N> res; \
1518  for (size_t i = 0; i < N / 2; i++) { \
1519  auto partial_res = __sycl_std::__invoke_native_##NAME<vec<float, 2>>( \
1520  detail::to_vec2(x, i * 2)); \
1521  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1522  } \
1523  if (N % 2) { \
1524  res[N - 1] = __sycl_std::__invoke_native_##NAME<float>(x[N - 1]); \
1525  } \
1526  return res; \
1527  }
1528 
1529 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(sin)
1530 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(cos)
1531 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(tan)
1532 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(exp)
1533 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(exp2)
1534 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(exp10)
1535 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(log)
1536 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(log2)
1537 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(log10)
1538 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(sqrt)
1539 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(rsqrt)
1540 __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(recip)
1541 
1542 #undef __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD
1543 
1544 #define __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(NAME) \
1545  template <size_t N> \
1546  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME( \
1547  marray<float, N> x, marray<float, N> y) __NOEXC { \
1548  marray<float, N> res; \
1549  for (size_t i = 0; i < N / 2; i++) { \
1550  auto partial_res = __sycl_std::__invoke_native_##NAME<vec<float, 2>>( \
1551  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
1552  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1553  } \
1554  if (N % 2) { \
1555  res[N - 1] = \
1556  __sycl_std::__invoke_native_##NAME<float>(x[N - 1], y[N - 1]); \
1557  } \
1558  return res; \
1559  }
1560 
1561 __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(divide)
1562 __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(powr)
1563 
1564 #undef __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD
1565 
1566 // svgenfloatf cos (svgenfloatf x)
1567 template <typename T>
1568 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
1569  return __sycl_std::__invoke_native_cos<T>(x);
1570 }
1571 
1572 // svgenfloatf divide (svgenfloatf x, svgenfloatf y)
1573 template <typename T>
1574 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> divide(T x,
1575  T y) __NOEXC {
1576  return __sycl_std::__invoke_native_divide<T>(x, y);
1577 }
1578 
1579 // svgenfloatf exp (svgenfloatf x)
1580 template <typename T>
1581 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {
1582  return __sycl_std::__invoke_native_exp<T>(x);
1583 }
1584 
1585 // svgenfloatf exp2 (svgenfloatf x)
1586 template <typename T>
1587 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp2(T x) __NOEXC {
1588  return __sycl_std::__invoke_native_exp2<T>(x);
1589 }
1590 
1591 // svgenfloatf exp10 (svgenfloatf x)
1592 template <typename T>
1593 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp10(T x) __NOEXC {
1594  return __sycl_std::__invoke_native_exp10<T>(x);
1595 }
1596 
1597 // svgenfloatf log (svgenfloatf x)
1598 template <typename T>
1599 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log(T x) __NOEXC {
1600  return __sycl_std::__invoke_native_log<T>(x);
1601 }
1602 
1603 // svgenfloatf log2 (svgenfloatf x)
1604 template <typename T>
1605 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log2(T x) __NOEXC {
1606  return __sycl_std::__invoke_native_log2<T>(x);
1607 }
1608 
1609 // svgenfloatf log10 (svgenfloatf x)
1610 template <typename T>
1611 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log10(T x) __NOEXC {
1612  return __sycl_std::__invoke_native_log10<T>(x);
1613 }
1614 
1615 // svgenfloatf powr (svgenfloatf x, svgenfloatf y)
1616 template <typename T>
1617 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> powr(T x,
1618  T y) __NOEXC {
1619  return __sycl_std::__invoke_native_powr<T>(x, y);
1620 }
1621 
1622 // svgenfloatf recip (svgenfloatf x)
1623 template <typename T>
1624 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> recip(T x) __NOEXC {
1625  return __sycl_std::__invoke_native_recip<T>(x);
1626 }
1627 
1628 // svgenfloatf rsqrt (svgenfloatf x)
1629 template <typename T>
1630 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> rsqrt(T x) __NOEXC {
1631  return __sycl_std::__invoke_native_rsqrt<T>(x);
1632 }
1633 
1634 // svgenfloatf sin (svgenfloatf x)
1635 template <typename T>
1636 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> sin(T x) __NOEXC {
1637  return __sycl_std::__invoke_native_sin<T>(x);
1638 }
1639 
1640 // svgenfloatf sqrt (svgenfloatf x)
1641 template <typename T>
1642 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> sqrt(T x) __NOEXC {
1643  return __sycl_std::__invoke_native_sqrt<T>(x);
1644 }
1645 
1646 // svgenfloatf tan (svgenfloatf x)
1647 template <typename T>
1648 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
1649  return __sycl_std::__invoke_native_tan<T>(x);
1650 }
1651 
1652 } // namespace native
1653 namespace half_precision {
1654 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1655 #define __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(NAME) \
1656  template <size_t N> \
1657  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x) \
1658  __NOEXC { \
1659  marray<float, N> res; \
1660  for (size_t i = 0; i < N / 2; i++) { \
1661  auto partial_res = __sycl_std::__invoke_half_##NAME<vec<float, 2>>( \
1662  detail::to_vec2(x, i * 2)); \
1663  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1664  } \
1665  if (N % 2) { \
1666  res[N - 1] = __sycl_std::__invoke_half_##NAME<float>(x[N - 1]); \
1667  } \
1668  return res; \
1669  }
1670 
1671 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(sin)
1672 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(cos)
1673 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(tan)
1674 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(exp)
1675 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(exp2)
1676 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(log)
1677 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(log2)
1678 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(log10)
1679 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(sqrt)
1680 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(rsqrt)
1681 __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(recip)
1682 
1683 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD
1684 
1685 #define __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(NAME) \
1686  template <size_t N> \
1687  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME( \
1688  marray<float, N> x, marray<float, N> y) __NOEXC { \
1689  marray<float, N> res; \
1690  for (size_t i = 0; i < N / 2; i++) { \
1691  auto partial_res = __sycl_std::__invoke_half_##NAME<vec<float, 2>>( \
1692  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
1693  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1694  } \
1695  if (N % 2) { \
1696  res[N - 1] = \
1697  __sycl_std::__invoke_half_##NAME<float>(x[N - 1], y[N - 1]); \
1698  } \
1699  return res; \
1700  }
1701 
1702 __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(divide)
1703 __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(powr)
1704 
1705 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD
1706 
1707 // svgenfloatf cos (svgenfloatf x)
1708 template <typename T>
1709 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
1710  return __sycl_std::__invoke_half_cos<T>(x);
1711 }
1712 
1713 // svgenfloatf divide (svgenfloatf x, svgenfloatf y)
1714 template <typename T>
1715 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> divide(T x,
1716  T y) __NOEXC {
1717  return __sycl_std::__invoke_half_divide<T>(x, y);
1718 }
1719 
1720 // svgenfloatf exp (svgenfloatf x)
1721 template <typename T>
1722 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {
1723  return __sycl_std::__invoke_half_exp<T>(x);
1724 }
1725 
1726 // svgenfloatf exp2 (svgenfloatf x)
1727 template <typename T>
1728 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp2(T x) __NOEXC {
1729  return __sycl_std::__invoke_half_exp2<T>(x);
1730 }
1731 
1732 // svgenfloatf exp10 (svgenfloatf x)
1733 template <typename T>
1734 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp10(T x) __NOEXC {
1735  return __sycl_std::__invoke_half_exp10<T>(x);
1736 }
1737 
1738 // svgenfloatf log (svgenfloatf x)
1739 template <typename T>
1740 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log(T x) __NOEXC {
1741  return __sycl_std::__invoke_half_log<T>(x);
1742 }
1743 
1744 // svgenfloatf log2 (svgenfloatf x)
1745 template <typename T>
1746 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log2(T x) __NOEXC {
1747  return __sycl_std::__invoke_half_log2<T>(x);
1748 }
1749 
1750 // svgenfloatf log10 (svgenfloatf x)
1751 template <typename T>
1752 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log10(T x) __NOEXC {
1753  return __sycl_std::__invoke_half_log10<T>(x);
1754 }
1755 
1756 // svgenfloatf powr (svgenfloatf x, svgenfloatf y)
1757 template <typename T>
1758 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> powr(T x,
1759  T y) __NOEXC {
1760  return __sycl_std::__invoke_half_powr<T>(x, y);
1761 }
1762 
1763 // svgenfloatf recip (svgenfloatf x)
1764 template <typename T>
1765 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> recip(T x) __NOEXC {
1766  return __sycl_std::__invoke_half_recip<T>(x);
1767 }
1768 
1769 // svgenfloatf rsqrt (svgenfloatf x)
1770 template <typename T>
1771 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> rsqrt(T x) __NOEXC {
1772  return __sycl_std::__invoke_half_rsqrt<T>(x);
1773 }
1774 
1775 // svgenfloatf sin (svgenfloatf x)
1776 template <typename T>
1777 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> sin(T x) __NOEXC {
1778  return __sycl_std::__invoke_half_sin<T>(x);
1779 }
1780 
1781 // svgenfloatf sqrt (svgenfloatf x)
1782 template <typename T>
1783 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> sqrt(T x) __NOEXC {
1784  return __sycl_std::__invoke_half_sqrt<T>(x);
1785 }
1786 
1787 // svgenfloatf tan (svgenfloatf x)
1788 template <typename T>
1789 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
1790  return __sycl_std::__invoke_half_tan<T>(x);
1791 }
1792 
1793 } // namespace half_precision
1794 
1795 #ifdef __FAST_MATH__
1796 /* ----------------- -ffast-math functions. ---------------------------------*/
1797 
1798 #define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME) \
1799  template <typename T, size_t N> \
1800  inline __SYCL_ALWAYS_INLINE \
1801  std::enable_if_t<std::is_same_v<T, float>, marray<T, N>> \
1802  NAME(marray<T, N> x) __NOEXC { \
1803  return native::NAME(x); \
1804  }
1805 
1806 __SYCL_MATH_FUNCTION_OVERLOAD_FM(sin)
1807 __SYCL_MATH_FUNCTION_OVERLOAD_FM(cos)
1808 __SYCL_MATH_FUNCTION_OVERLOAD_FM(tan)
1809 __SYCL_MATH_FUNCTION_OVERLOAD_FM(exp)
1810 __SYCL_MATH_FUNCTION_OVERLOAD_FM(exp2)
1811 __SYCL_MATH_FUNCTION_OVERLOAD_FM(exp10)
1812 __SYCL_MATH_FUNCTION_OVERLOAD_FM(log)
1813 __SYCL_MATH_FUNCTION_OVERLOAD_FM(log2)
1814 __SYCL_MATH_FUNCTION_OVERLOAD_FM(log10)
1815 __SYCL_MATH_FUNCTION_OVERLOAD_FM(sqrt)
1816 __SYCL_MATH_FUNCTION_OVERLOAD_FM(rsqrt)
1817 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
1818 
1819 template <typename T, size_t N>
1820 inline __SYCL_ALWAYS_INLINE
1821  std::enable_if_t<std::is_same_v<T, float>, marray<T, N>>
1822  powr(marray<T, N> x, marray<T, N> y) __NOEXC {
1823  return native::powr(x, y);
1824 }
1825 
1826 // svgenfloatf cos (svgenfloatf x)
1827 template <typename T>
1828 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
1829  return native::cos(x);
1830 }
1831 
1832 // svgenfloatf exp (svgenfloatf x)
1833 template <typename T>
1834 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {
1835  return native::exp(x);
1836 }
1837 
1838 // svgenfloatf exp2 (svgenfloatf x)
1839 template <typename T>
1840 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp2(T x) __NOEXC {
1841  return native::exp2(x);
1842 }
1843 
1844 // svgenfloatf exp10 (svgenfloatf x)
1845 template <typename T>
1846 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp10(T x) __NOEXC {
1847  return native::exp10(x);
1848 }
1849 
1850 // svgenfloatf log(svgenfloatf x)
1851 template <typename T>
1852 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log(T x) __NOEXC {
1853  return native::log(x);
1854 }
1855 
1856 // svgenfloatf log2 (svgenfloatf x)
1857 template <typename T>
1858 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log2(T x) __NOEXC {
1859  return native::log2(x);
1860 }
1861 
1862 // svgenfloatf log10 (svgenfloatf x)
1863 template <typename T>
1864 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> log10(T x) __NOEXC {
1865  return native::log10(x);
1866 }
1867 
1868 // svgenfloatf powr (svgenfloatf x)
1869 template <typename T>
1870 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> powr(T x,
1871  T y) __NOEXC {
1872  return native::powr(x, y);
1873 }
1874 
1875 // svgenfloatf rsqrt (svgenfloatf x)
1876 template <typename T>
1877 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> rsqrt(T x) __NOEXC {
1878  return native::rsqrt(x);
1879 }
1880 
1881 // svgenfloatf sin (svgenfloatf x)
1882 template <typename T>
1883 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> sin(T x) __NOEXC {
1884  return native::sin(x);
1885 }
1886 
1887 // svgenfloatf sqrt (svgenfloatf x)
1888 template <typename T>
1889 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> sqrt(T x) __NOEXC {
1890  return native::sqrt(x);
1891 }
1892 
1893 // svgenfloatf tan (svgenfloatf x)
1894 template <typename T>
1895 detail::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
1896  return native::tan(x);
1897 }
1898 
1899 #endif // __FAST_MATH__
1900 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1901 } // namespace sycl
1902 
1903 #ifdef __SYCL_DEVICE_ONLY__
1904 extern "C" {
1905 extern SYCL_EXTERNAL int abs(int x);
1906 extern SYCL_EXTERNAL long int labs(long int x);
1907 extern SYCL_EXTERNAL long long int llabs(long long int x);
1908 
1909 extern SYCL_EXTERNAL div_t div(int x, int y);
1910 extern SYCL_EXTERNAL ldiv_t ldiv(long int x, long int y);
1911 extern SYCL_EXTERNAL lldiv_t lldiv(long long int x, long long int y);
1912 extern SYCL_EXTERNAL float scalbnf(float x, int n);
1913 extern SYCL_EXTERNAL double scalbn(double x, int n);
1914 extern SYCL_EXTERNAL float logf(float x);
1915 extern SYCL_EXTERNAL double log(double x);
1916 extern SYCL_EXTERNAL float expf(float x);
1917 extern SYCL_EXTERNAL double exp(double x);
1918 extern SYCL_EXTERNAL float log10f(float x);
1919 extern SYCL_EXTERNAL double log10(double x);
1920 extern SYCL_EXTERNAL float modff(float x, float *intpart);
1921 extern SYCL_EXTERNAL double modf(double x, double *intpart);
1922 extern SYCL_EXTERNAL float exp2f(float x);
1923 extern SYCL_EXTERNAL double exp2(double x);
1924 extern SYCL_EXTERNAL float expm1f(float x);
1925 extern SYCL_EXTERNAL double expm1(double x);
1926 extern SYCL_EXTERNAL int ilogbf(float x);
1927 extern SYCL_EXTERNAL int ilogb(double x);
1928 extern SYCL_EXTERNAL float log1pf(float x);
1929 extern SYCL_EXTERNAL double log1p(double x);
1930 extern SYCL_EXTERNAL float log2f(float x);
1931 extern SYCL_EXTERNAL double log2(double x);
1932 extern SYCL_EXTERNAL float logbf(float x);
1933 extern SYCL_EXTERNAL double logb(double x);
1934 extern SYCL_EXTERNAL float sqrtf(float x);
1935 extern SYCL_EXTERNAL double sqrt(double x);
1936 extern SYCL_EXTERNAL float cbrtf(float x);
1937 extern SYCL_EXTERNAL double cbrt(double x);
1938 extern SYCL_EXTERNAL float erff(float x);
1939 extern SYCL_EXTERNAL double erf(double x);
1940 extern SYCL_EXTERNAL float erfcf(float x);
1941 extern SYCL_EXTERNAL double erfc(double x);
1942 extern SYCL_EXTERNAL float tgammaf(float x);
1943 extern SYCL_EXTERNAL double tgamma(double x);
1944 extern SYCL_EXTERNAL float lgammaf(float x);
1945 extern SYCL_EXTERNAL double lgamma(double x);
1946 extern SYCL_EXTERNAL float fmodf(float x, float y);
1947 extern SYCL_EXTERNAL double fmod(double x, double y);
1948 extern SYCL_EXTERNAL float remainderf(float x, float y);
1949 extern SYCL_EXTERNAL double remainder(double x, double y);
1950 extern SYCL_EXTERNAL float remquof(float x, float y, int *q);
1951 extern SYCL_EXTERNAL double remquo(double x, double y, int *q);
1952 extern SYCL_EXTERNAL float nextafterf(float x, float y);
1953 extern SYCL_EXTERNAL double nextafter(double x, double y);
1954 extern SYCL_EXTERNAL float fdimf(float x, float y);
1955 extern SYCL_EXTERNAL double fdim(double x, double y);
1956 extern SYCL_EXTERNAL float fmaf(float x, float y, float z);
1957 extern SYCL_EXTERNAL double fma(double x, double y, double z);
1958 extern SYCL_EXTERNAL float sinf(float x);
1959 extern SYCL_EXTERNAL double sin(double x);
1960 extern SYCL_EXTERNAL float cosf(float x);
1961 extern SYCL_EXTERNAL double cos(double x);
1962 extern SYCL_EXTERNAL float tanf(float x);
1963 extern SYCL_EXTERNAL double tan(double x);
1964 extern SYCL_EXTERNAL float asinf(float x);
1965 extern SYCL_EXTERNAL double asin(double x);
1966 extern SYCL_EXTERNAL float acosf(float x);
1967 extern SYCL_EXTERNAL double acos(double x);
1968 extern SYCL_EXTERNAL float atanf(float x);
1969 extern SYCL_EXTERNAL double atan(double x);
1970 extern SYCL_EXTERNAL float powf(float x, float y);
1971 extern SYCL_EXTERNAL double pow(double x, double y);
1972 extern SYCL_EXTERNAL float atan2f(float x, float y);
1973 extern SYCL_EXTERNAL double atan2(double x, double y);
1974 
1975 extern SYCL_EXTERNAL float sinhf(float x);
1976 extern SYCL_EXTERNAL double sinh(double x);
1977 extern SYCL_EXTERNAL float coshf(float x);
1978 extern SYCL_EXTERNAL double cosh(double x);
1979 extern SYCL_EXTERNAL float tanhf(float x);
1980 extern SYCL_EXTERNAL double tanh(double x);
1981 extern SYCL_EXTERNAL float asinhf(float x);
1982 extern SYCL_EXTERNAL double asinh(double x);
1983 extern SYCL_EXTERNAL float acoshf(float x);
1984 extern SYCL_EXTERNAL double acosh(double x);
1985 extern SYCL_EXTERNAL float atanhf(float x);
1986 extern SYCL_EXTERNAL double atanh(double x);
1987 extern SYCL_EXTERNAL double frexp(double x, int *exp);
1988 extern SYCL_EXTERNAL double ldexp(double x, int exp);
1989 extern SYCL_EXTERNAL double hypot(double x, double y);
1990 
1991 extern SYCL_EXTERNAL void *memcpy(void *dest, const void *src, size_t n);
1992 extern SYCL_EXTERNAL void *memset(void *dest, int c, size_t n);
1993 extern SYCL_EXTERNAL int memcmp(const void *s1, const void *s2, size_t n);
1994 extern SYCL_EXTERNAL long long int __imf_llmax(long long int x,
1995  long long int y);
1996 extern SYCL_EXTERNAL long long int __imf_llmin(long long int x,
1997  long long int y);
1998 extern SYCL_EXTERNAL unsigned long long int
1999 __imf_ullmax(unsigned long long int x, unsigned long long int y);
2000 extern SYCL_EXTERNAL unsigned long long int
2001 __imf_ullmin(unsigned long long int x, unsigned long long int y);
2002 extern SYCL_EXTERNAL unsigned int __imf_umax(unsigned int x, unsigned int y);
2003 extern SYCL_EXTERNAL unsigned int __imf_umin(unsigned int x, unsigned int y);
2004 extern SYCL_EXTERNAL unsigned int __imf_brev(unsigned int x);
2005 extern SYCL_EXTERNAL unsigned long long int
2006 __imf_brevll(unsigned long long int x);
2007 extern SYCL_EXTERNAL unsigned int
2008 __imf_byte_perm(unsigned int x, unsigned int y, unsigned int s);
2009 extern SYCL_EXTERNAL int __imf_ffs(int x);
2010 extern SYCL_EXTERNAL int __imf_ffsll(long long int x);
2011 extern SYCL_EXTERNAL int __imf_clz(int x);
2012 extern SYCL_EXTERNAL int __imf_clzll(long long int x);
2013 extern SYCL_EXTERNAL int __imf_popc(unsigned int x);
2014 extern SYCL_EXTERNAL int __imf_popcll(unsigned long long int x);
2015 extern SYCL_EXTERNAL unsigned int __imf_sad(int x, int y, unsigned int z);
2016 extern SYCL_EXTERNAL unsigned int __imf_usad(unsigned int x, unsigned int y,
2017  unsigned int z);
2018 extern SYCL_EXTERNAL int __imf_rhadd(int x, int y);
2019 extern SYCL_EXTERNAL unsigned int __imf_urhadd(unsigned int x, unsigned int y);
2020 extern SYCL_EXTERNAL unsigned int __imf_uhadd(unsigned int x, unsigned int y);
2021 extern SYCL_EXTERNAL int __imf_mul24(int x, int y);
2022 extern SYCL_EXTERNAL unsigned int __imf_umul24(unsigned int x, unsigned int y);
2023 extern SYCL_EXTERNAL int __imf_mulhi(int x, int y);
2024 extern SYCL_EXTERNAL unsigned int __imf_umulhi(unsigned int x, unsigned int y);
2025 extern SYCL_EXTERNAL long long int __imf_mul64hi(long long int x,
2026  long long int y);
2027 extern SYCL_EXTERNAL unsigned long long int
2028 __imf_umul64hi(unsigned long long int x, unsigned long long int y);
2029 extern SYCL_EXTERNAL float __imf_saturatef(float x);
2030 extern SYCL_EXTERNAL float __imf_fmaf(float x, float y, float z);
2031 extern SYCL_EXTERNAL float __imf_fabsf(float x);
2032 extern SYCL_EXTERNAL float __imf_floorf(float x);
2033 extern SYCL_EXTERNAL float __imf_ceilf(float x);
2034 extern SYCL_EXTERNAL float __imf_truncf(float x);
2035 extern SYCL_EXTERNAL float __imf_rintf(float x);
2036 extern SYCL_EXTERNAL float __imf_nearbyintf(float x);
2037 extern SYCL_EXTERNAL float __imf_sqrtf(float x);
2038 extern SYCL_EXTERNAL float __imf_rsqrtf(float x);
2039 extern SYCL_EXTERNAL float __imf_invf(float x);
2040 extern SYCL_EXTERNAL float __imf_fmaxf(float x, float y);
2041 extern SYCL_EXTERNAL float __imf_fminf(float x, float y);
2042 extern SYCL_EXTERNAL float __imf_copysignf(float x, float y);
2043 extern SYCL_EXTERNAL int __imf_float2int_rd(float x);
2044 extern SYCL_EXTERNAL int __imf_float2int_rn(float x);
2045 extern SYCL_EXTERNAL int __imf_float2int_ru(float x);
2046 extern SYCL_EXTERNAL int __imf_float2int_rz(float x);
2047 extern SYCL_EXTERNAL unsigned int __imf_float2uint_rd(float x);
2048 extern SYCL_EXTERNAL unsigned int __imf_float2uint_rn(float x);
2049 extern SYCL_EXTERNAL unsigned int __imf_float2uint_ru(float x);
2050 extern SYCL_EXTERNAL unsigned int __imf_float2uint_rz(float x);
2051 extern SYCL_EXTERNAL long long int __imf_float2ll_rd(float x);
2052 extern SYCL_EXTERNAL long long int __imf_float2ll_rn(float x);
2053 extern SYCL_EXTERNAL long long int __imf_float2ll_ru(float x);
2054 extern SYCL_EXTERNAL long long int __imf_float2ll_rz(float x);
2055 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rd(float x);
2056 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rn(float x);
2057 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_ru(float x);
2058 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rz(float x);
2059 extern SYCL_EXTERNAL int __imf_float_as_int(float x);
2060 extern SYCL_EXTERNAL unsigned int __imf_float_as_uint(float x);
2061 extern SYCL_EXTERNAL float __imf_int2float_rd(int x);
2062 extern SYCL_EXTERNAL float __imf_int2float_rn(int x);
2063 extern SYCL_EXTERNAL float __imf_int2float_ru(int x);
2064 extern SYCL_EXTERNAL float __imf_int2float_rz(int x);
2065 extern SYCL_EXTERNAL float __imf_int_as_float(int x);
2066 extern SYCL_EXTERNAL float __imf_ll2float_rd(long long int x);
2067 extern SYCL_EXTERNAL float __imf_ll2float_rn(long long int x);
2068 extern SYCL_EXTERNAL float __imf_ll2float_ru(long long int x);
2069 extern SYCL_EXTERNAL float __imf_ll2float_rz(long long int x);
2070 extern SYCL_EXTERNAL float __imf_uint2float_rd(unsigned int x);
2071 extern SYCL_EXTERNAL float __imf_uint2float_rn(unsigned int x);
2072 extern SYCL_EXTERNAL float __imf_uint2float_ru(unsigned int x);
2073 extern SYCL_EXTERNAL float __imf_uint2float_rz(unsigned int x);
2074 extern SYCL_EXTERNAL float __imf_uint_as_float(unsigned int x);
2075 extern SYCL_EXTERNAL float __imf_ull2float_rd(unsigned long long int x);
2076 extern SYCL_EXTERNAL float __imf_ull2float_rn(unsigned long long int x);
2077 extern SYCL_EXTERNAL float __imf_ull2float_ru(unsigned long long int x);
2078 extern SYCL_EXTERNAL float __imf_ull2float_rz(unsigned long long int x);
2079 extern SYCL_EXTERNAL float __imf_half2float(_Float16 x);
2080 extern SYCL_EXTERNAL _Float16 __imf_float2half_rd(float x);
2081 extern SYCL_EXTERNAL _Float16 __imf_float2half_rn(float x);
2082 extern SYCL_EXTERNAL _Float16 __imf_float2half_ru(float x);
2083 extern SYCL_EXTERNAL _Float16 __imf_float2half_rz(float x);
2084 extern SYCL_EXTERNAL int __imf_half2int_rd(_Float16 x);
2085 extern SYCL_EXTERNAL int __imf_half2int_rn(_Float16 x);
2086 extern SYCL_EXTERNAL int __imf_half2int_ru(_Float16 x);
2087 extern SYCL_EXTERNAL int __imf_half2int_rz(_Float16 x);
2088 extern SYCL_EXTERNAL long long __imf_half2ll_rd(_Float16 x);
2089 extern SYCL_EXTERNAL long long __imf_half2ll_rn(_Float16 x);
2090 extern SYCL_EXTERNAL long long __imf_half2ll_ru(_Float16 x);
2091 extern SYCL_EXTERNAL long long __imf_half2ll_rz(_Float16 x);
2092 extern SYCL_EXTERNAL short __imf_half2short_rd(_Float16 x);
2093 extern SYCL_EXTERNAL short __imf_half2short_rn(_Float16 x);
2094 extern SYCL_EXTERNAL short __imf_half2short_ru(_Float16 x);
2095 extern SYCL_EXTERNAL short __imf_half2short_rz(_Float16 x);
2096 extern SYCL_EXTERNAL unsigned int __imf_half2uint_rd(_Float16 x);
2097 extern SYCL_EXTERNAL unsigned int __imf_half2uint_rn(_Float16 x);
2098 extern SYCL_EXTERNAL unsigned int __imf_half2uint_ru(_Float16 x);
2099 extern SYCL_EXTERNAL unsigned int __imf_half2uint_rz(_Float16 x);
2100 extern SYCL_EXTERNAL unsigned long long __imf_half2ull_rd(_Float16 x);
2101 extern SYCL_EXTERNAL unsigned long long __imf_half2ull_rn(_Float16 x);
2102 extern SYCL_EXTERNAL unsigned long long __imf_half2ull_ru(_Float16 x);
2103 extern SYCL_EXTERNAL unsigned long long __imf_half2ull_rz(_Float16 x);
2104 extern SYCL_EXTERNAL unsigned short __imf_half2ushort_rd(_Float16 x);
2105 extern SYCL_EXTERNAL unsigned short __imf_half2ushort_rn(_Float16 x);
2106 extern SYCL_EXTERNAL unsigned short __imf_half2ushort_ru(_Float16 x);
2107 extern SYCL_EXTERNAL unsigned short __imf_half2ushort_rz(_Float16 x);
2108 extern SYCL_EXTERNAL short __imf_half_as_short(_Float16 x);
2109 extern SYCL_EXTERNAL unsigned short __imf_half_as_ushort(_Float16 x);
2110 extern SYCL_EXTERNAL _Float16 __imf_int2half_rd(int x);
2111 extern SYCL_EXTERNAL _Float16 __imf_int2half_rn(int x);
2112 extern SYCL_EXTERNAL _Float16 __imf_int2half_ru(int x);
2113 extern SYCL_EXTERNAL _Float16 __imf_int2half_rz(int x);
2114 extern SYCL_EXTERNAL _Float16 __imf_ll2half_rd(long long x);
2115 extern SYCL_EXTERNAL _Float16 __imf_ll2half_rn(long long x);
2116 extern SYCL_EXTERNAL _Float16 __imf_ll2half_ru(long long x);
2117 extern SYCL_EXTERNAL _Float16 __imf_ll2half_rz(long long x);
2118 extern SYCL_EXTERNAL _Float16 __imf_short2half_rd(short x);
2119 extern SYCL_EXTERNAL _Float16 __imf_short2half_rn(short x);
2120 extern SYCL_EXTERNAL _Float16 __imf_short2half_ru(short x);
2121 extern SYCL_EXTERNAL _Float16 __imf_short2half_rz(short x);
2122 extern SYCL_EXTERNAL _Float16 __imf_short_as_half(short x);
2123 extern SYCL_EXTERNAL _Float16 __imf_uint2half_rd(unsigned int x);
2124 extern SYCL_EXTERNAL _Float16 __imf_uint2half_rn(unsigned int x);
2125 extern SYCL_EXTERNAL _Float16 __imf_uint2half_ru(unsigned int x);
2126 extern SYCL_EXTERNAL _Float16 __imf_uint2half_rz(unsigned int x);
2127 extern SYCL_EXTERNAL _Float16 __imf_ull2half_rd(unsigned long long x);
2128 extern SYCL_EXTERNAL _Float16 __imf_ull2half_rn(unsigned long long x);
2129 extern SYCL_EXTERNAL _Float16 __imf_ull2half_ru(unsigned long long x);
2130 extern SYCL_EXTERNAL _Float16 __imf_ull2half_rz(unsigned long long x);
2131 extern SYCL_EXTERNAL _Float16 __imf_ushort2half_rd(unsigned short x);
2132 extern SYCL_EXTERNAL _Float16 __imf_ushort2half_rn(unsigned short x);
2133 extern SYCL_EXTERNAL _Float16 __imf_ushort2half_ru(unsigned short x);
2134 extern SYCL_EXTERNAL _Float16 __imf_ushort2half_rz(unsigned short x);
2135 extern SYCL_EXTERNAL _Float16 __imf_ushort_as_half(unsigned short x);
2136 extern SYCL_EXTERNAL _Float16 __imf_double2half(double x);
2137 
2138 extern SYCL_EXTERNAL _Float16 __imf_fmaf16(_Float16 x, _Float16 y, _Float16 z);
2139 extern SYCL_EXTERNAL _Float16 __imf_fabsf16(_Float16 x);
2140 extern SYCL_EXTERNAL _Float16 __imf_floorf16(_Float16 x);
2141 extern SYCL_EXTERNAL _Float16 __imf_ceilf16(_Float16 x);
2142 extern SYCL_EXTERNAL _Float16 __imf_truncf16(_Float16 x);
2143 extern SYCL_EXTERNAL _Float16 __imf_rintf16(_Float16 x);
2144 extern SYCL_EXTERNAL _Float16 __imf_nearbyintf16(_Float16 x);
2145 extern SYCL_EXTERNAL _Float16 __imf_sqrtf16(_Float16 x);
2146 extern SYCL_EXTERNAL _Float16 __imf_rsqrtf16(_Float16 x);
2147 extern SYCL_EXTERNAL _Float16 __imf_invf16(_Float16 x);
2148 extern SYCL_EXTERNAL _Float16 __imf_fmaxf16(_Float16 x, _Float16 y);
2149 extern SYCL_EXTERNAL _Float16 __imf_fminf16(_Float16 x, _Float16 y);
2150 extern SYCL_EXTERNAL _Float16 __imf_copysignf16(_Float16 x, _Float16 y);
2151 extern SYCL_EXTERNAL float __imf_half2float(_Float16 x);
2152 extern SYCL_EXTERNAL float __imf_bfloat162float(uint16_t x);
2153 extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rd(uint16_t x);
2154 extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rn(uint16_t x);
2155 extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_ru(uint16_t x);
2156 extern SYCL_EXTERNAL unsigned int __imf_bfloat162uint_rz(uint16_t x);
2157 extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_rd(uint16_t x);
2158 extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_rn(uint16_t x);
2159 extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_ru(uint16_t x);
2160 extern SYCL_EXTERNAL unsigned short __imf_bfloat162ushort_rz(uint16_t x);
2161 extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_rd(uint16_t x);
2162 extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_rn(uint16_t x);
2163 extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_ru(uint16_t x);
2164 extern SYCL_EXTERNAL unsigned long long __imf_bfloat162ull_rz(uint16_t x);
2165 extern SYCL_EXTERNAL int __imf_bfloat162int_rd(uint16_t x);
2166 extern SYCL_EXTERNAL int __imf_bfloat162int_rn(uint16_t x);
2167 extern SYCL_EXTERNAL int __imf_bfloat162int_ru(uint16_t x);
2168 extern SYCL_EXTERNAL int __imf_bfloat162int_rz(uint16_t x);
2169 extern SYCL_EXTERNAL short __imf_bfloat162short_rd(uint16_t x);
2170 extern SYCL_EXTERNAL short __imf_bfloat162short_rn(uint16_t x);
2171 extern SYCL_EXTERNAL short __imf_bfloat162short_ru(uint16_t x);
2172 extern SYCL_EXTERNAL short __imf_bfloat162short_rz(uint16_t x);
2173 extern SYCL_EXTERNAL long long __imf_bfloat162ll_rd(uint16_t x);
2174 extern SYCL_EXTERNAL long long __imf_bfloat162ll_rn(uint16_t x);
2175 extern SYCL_EXTERNAL long long __imf_bfloat162ll_ru(uint16_t x);
2176 extern SYCL_EXTERNAL long long __imf_bfloat162ll_rz(uint16_t x);
2177 extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16(float x);
2178 extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rd(float x);
2179 extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rn(float x);
2180 extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_ru(float x);
2181 extern SYCL_EXTERNAL uint16_t __imf_float2bfloat16_rz(float x);
2182 extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_rd(unsigned short x);
2183 extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_rn(unsigned short x);
2184 extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_ru(unsigned short x);
2185 extern SYCL_EXTERNAL uint16_t __imf_ushort2bfloat16_rz(unsigned short x);
2186 extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rd(unsigned int x);
2187 extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rn(unsigned int x);
2188 extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_ru(unsigned int x);
2189 extern SYCL_EXTERNAL uint16_t __imf_uint2bfloat16_rz(unsigned int x);
2190 extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_rd(unsigned long long x);
2191 extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_rn(unsigned long long x);
2192 extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_ru(unsigned long long x);
2193 extern SYCL_EXTERNAL uint16_t __imf_ull2bfloat16_rz(unsigned long long x);
2194 extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rd(short x);
2195 extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rn(short x);
2196 extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_ru(short x);
2197 extern SYCL_EXTERNAL uint16_t __imf_short2bfloat16_rz(short x);
2198 extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rd(int x);
2199 extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rn(int x);
2200 extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_ru(int x);
2201 extern SYCL_EXTERNAL uint16_t __imf_int2bfloat16_rz(int x);
2202 extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rd(long long x);
2203 extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rn(long long x);
2204 extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_ru(long long x);
2205 extern SYCL_EXTERNAL uint16_t __imf_ll2bfloat16_rz(long long x);
2206 extern SYCL_EXTERNAL short __imf_bfloat16_as_short(uint16_t x);
2207 extern SYCL_EXTERNAL unsigned short __imf_bfloat16_as_ushort(uint16_t x);
2208 extern SYCL_EXTERNAL uint16_t __imf_short_as_bfloat16(short x);
2209 extern SYCL_EXTERNAL uint16_t __imf_ushort_as_bfloat16(unsigned short x);
2210 extern SYCL_EXTERNAL uint16_t __imf_fmabf16(uint16_t x, uint16_t y, uint16_t z);
2211 extern SYCL_EXTERNAL uint16_t __imf_fmaxbf16(uint16_t x, uint16_t y);
2212 extern SYCL_EXTERNAL uint16_t __imf_fminbf16(uint16_t x, uint16_t y);
2213 extern SYCL_EXTERNAL uint16_t __imf_fabsbf16(uint16_t x);
2214 extern SYCL_EXTERNAL uint16_t __imf_rintbf16(uint16_t x);
2215 extern SYCL_EXTERNAL uint16_t __imf_floorbf16(uint16_t x);
2216 extern SYCL_EXTERNAL uint16_t __imf_ceilbf16(uint16_t x);
2217 extern SYCL_EXTERNAL uint16_t __imf_truncbf16(uint16_t x);
2218 extern SYCL_EXTERNAL uint16_t __imf_copysignbf16(uint16_t x, uint16_t y);
2219 extern SYCL_EXTERNAL uint16_t __imf_sqrtbf16(uint16_t x);
2220 extern SYCL_EXTERNAL uint16_t __imf_rsqrtbf16(uint16_t x);
2221 extern SYCL_EXTERNAL double __imf_fma(double x, double y, double z);
2222 extern SYCL_EXTERNAL double __imf_fabs(double x);
2223 extern SYCL_EXTERNAL double __imf_floor(double x);
2224 extern SYCL_EXTERNAL double __imf_ceil(double x);
2225 extern SYCL_EXTERNAL double __imf_trunc(double x);
2226 extern SYCL_EXTERNAL double __imf_rint(double x);
2227 extern SYCL_EXTERNAL double __imf_nearbyint(double x);
2228 extern SYCL_EXTERNAL double __imf_sqrt(double x);
2229 extern SYCL_EXTERNAL double __imf_rsqrt(double x);
2230 extern SYCL_EXTERNAL double __imf_inv(double x);
2231 extern SYCL_EXTERNAL double __imf_fmax(double x, double y);
2232 extern SYCL_EXTERNAL double __imf_fmin(double x, double y);
2233 extern SYCL_EXTERNAL double __imf_copysign(double x, double y);
2234 extern SYCL_EXTERNAL float __imf_double2float_rd(double x);
2235 extern SYCL_EXTERNAL float __imf_double2float_rn(double x);
2236 extern SYCL_EXTERNAL float __imf_double2float_ru(double x);
2237 extern SYCL_EXTERNAL float __imf_double2float_rz(double x);
2238 extern SYCL_EXTERNAL int __imf_double2hiint(double x);
2239 extern SYCL_EXTERNAL int __imf_double2loint(double x);
2240 extern SYCL_EXTERNAL int __imf_double2int_rd(double x);
2241 extern SYCL_EXTERNAL int __imf_double2int_rn(double x);
2242 extern SYCL_EXTERNAL int __imf_double2int_ru(double x);
2243 extern SYCL_EXTERNAL int __imf_double2int_rz(double x);
2244 extern SYCL_EXTERNAL double __imf_int2double_rn(int x);
2245 extern SYCL_EXTERNAL unsigned int __imf_double2uint_rd(double x);
2246 extern SYCL_EXTERNAL unsigned int __imf_double2uint_rn(double x);
2247 extern SYCL_EXTERNAL unsigned int __imf_double2uint_ru(double x);
2248 extern SYCL_EXTERNAL unsigned int __imf_double2uint_rz(double x);
2249 extern SYCL_EXTERNAL long long int __imf_double2ll_rd(double x);
2250 extern SYCL_EXTERNAL long long int __imf_double2ll_rn(double x);
2251 extern SYCL_EXTERNAL long long int __imf_double2ll_ru(double x);
2252 extern SYCL_EXTERNAL long long int __imf_double2ll_rz(double x);
2253 extern SYCL_EXTERNAL double __imf_ll2double_rd(long long int x);
2254 extern SYCL_EXTERNAL double __imf_ll2double_rn(long long int x);
2255 extern SYCL_EXTERNAL double __imf_ll2double_ru(long long int x);
2256 extern SYCL_EXTERNAL double __imf_ll2double_rz(long long int x);
2257 extern SYCL_EXTERNAL double __imf_ull2double_rd(unsigned long long int x);
2258 extern SYCL_EXTERNAL double __imf_ull2double_rn(unsigned long long int x);
2259 extern SYCL_EXTERNAL double __imf_ull2double_ru(unsigned long long int x);
2260 extern SYCL_EXTERNAL double __imf_ull2double_rz(unsigned long long int x);
2261 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rd(double x);
2262 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rn(double x);
2263 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_ru(double x);
2264 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rz(double x);
2265 extern SYCL_EXTERNAL long long int __imf_double_as_longlong(double x);
2266 extern SYCL_EXTERNAL double __imf_longlong_as_double(long long int x);
2267 extern SYCL_EXTERNAL double __imf_uint2double_rd(unsigned int x);
2268 extern SYCL_EXTERNAL double __imf_uint2double_rn(unsigned int x);
2269 extern SYCL_EXTERNAL double __imf_uint2double_ru(unsigned int x);
2270 extern SYCL_EXTERNAL double __imf_uint2double_rz(unsigned int x);
2271 extern SYCL_EXTERNAL double __imf_hiloint2double(int hi, int lo);
2272 
2273 extern SYCL_EXTERNAL unsigned int __imf_vabs2(unsigned int x);
2274 extern SYCL_EXTERNAL unsigned int __imf_vabs4(unsigned int x);
2275 extern SYCL_EXTERNAL unsigned int __imf_vabsss2(unsigned int x);
2276 extern SYCL_EXTERNAL unsigned int __imf_vabsss4(unsigned int x);
2277 extern SYCL_EXTERNAL unsigned int __imf_vneg2(unsigned int x);
2278 extern SYCL_EXTERNAL unsigned int __imf_vneg4(unsigned int x);
2279 extern SYCL_EXTERNAL unsigned int __imf_vnegss2(unsigned int x);
2280 extern SYCL_EXTERNAL unsigned int __imf_vnegss4(unsigned int x);
2281 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffs2(unsigned int x,
2282  unsigned int y);
2283 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffs4(unsigned int x,
2284  unsigned int y);
2285 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffu2(unsigned int x,
2286  unsigned int y);
2287 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffu4(unsigned int x,
2288  unsigned int y);
2289 extern SYCL_EXTERNAL unsigned int __imf_vadd2(unsigned int x, unsigned int y);
2290 extern SYCL_EXTERNAL unsigned int __imf_vadd4(unsigned int x, unsigned int y);
2291 extern SYCL_EXTERNAL unsigned int __imf_vaddss2(unsigned int x, unsigned int y);
2292 extern SYCL_EXTERNAL unsigned int __imf_vaddss4(unsigned int x, unsigned int y);
2293 extern SYCL_EXTERNAL unsigned int __imf_vaddus2(unsigned int x, unsigned int y);
2294 extern SYCL_EXTERNAL unsigned int __imf_vaddus4(unsigned int x, unsigned int y);
2295 extern SYCL_EXTERNAL unsigned int __imf_vsub2(unsigned int x, unsigned int y);
2296 extern SYCL_EXTERNAL unsigned int __imf_vsub4(unsigned int x, unsigned int y);
2297 extern SYCL_EXTERNAL unsigned int __imf_vsubss2(unsigned int x, unsigned int y);
2298 extern SYCL_EXTERNAL unsigned int __imf_vsubss4(unsigned int x, unsigned int y);
2299 extern SYCL_EXTERNAL unsigned int __imf_vsubus2(unsigned int x, unsigned int y);
2300 extern SYCL_EXTERNAL unsigned int __imf_vsubus4(unsigned int x, unsigned int y);
2301 extern SYCL_EXTERNAL unsigned int __imf_vavgs2(unsigned int x, unsigned int y);
2302 extern SYCL_EXTERNAL unsigned int __imf_vavgs4(unsigned int x, unsigned int y);
2303 extern SYCL_EXTERNAL unsigned int __imf_vavgu2(unsigned int x, unsigned int y);
2304 extern SYCL_EXTERNAL unsigned int __imf_vavgu4(unsigned int x, unsigned int y);
2305 extern SYCL_EXTERNAL unsigned int __imf_vhaddu2(unsigned int x, unsigned int y);
2306 extern SYCL_EXTERNAL unsigned int __imf_vhaddu4(unsigned int x, unsigned int y);
2307 extern SYCL_EXTERNAL unsigned int __imf_vcmpeq2(unsigned int x, unsigned int y);
2308 extern SYCL_EXTERNAL unsigned int __imf_vcmpeq4(unsigned int x, unsigned int y);
2309 extern SYCL_EXTERNAL unsigned int __imf_vcmpne2(unsigned int x, unsigned int y);
2310 extern SYCL_EXTERNAL unsigned int __imf_vcmpne4(unsigned int x, unsigned int y);
2311 extern SYCL_EXTERNAL unsigned int __imf_vcmpges2(unsigned int x,
2312  unsigned int y);
2313 extern SYCL_EXTERNAL unsigned int __imf_vcmpges4(unsigned int x,
2314  unsigned int y);
2315 extern SYCL_EXTERNAL unsigned int __imf_vcmpgeu2(unsigned int x,
2316  unsigned int y);
2317 extern SYCL_EXTERNAL unsigned int __imf_vcmpgeu4(unsigned int x,
2318  unsigned int y);
2319 extern SYCL_EXTERNAL unsigned int __imf_vcmpgts2(unsigned int x,
2320  unsigned int y);
2321 extern SYCL_EXTERNAL unsigned int __imf_vcmpgts4(unsigned int x,
2322  unsigned int y);
2323 extern SYCL_EXTERNAL unsigned int __imf_vcmpgtu2(unsigned int x,
2324  unsigned int y);
2325 extern SYCL_EXTERNAL unsigned int __imf_vcmpgtu4(unsigned int x,
2326  unsigned int y);
2327 extern SYCL_EXTERNAL unsigned int __imf_vcmples2(unsigned int x,
2328  unsigned int y);
2329 extern SYCL_EXTERNAL unsigned int __imf_vcmples4(unsigned int x,
2330  unsigned int y);
2331 extern SYCL_EXTERNAL unsigned int __imf_vcmpleu2(unsigned int x,
2332  unsigned int y);
2333 extern SYCL_EXTERNAL unsigned int __imf_vcmpleu4(unsigned int x,
2334  unsigned int y);
2335 extern SYCL_EXTERNAL unsigned int __imf_vcmplts2(unsigned int x,
2336  unsigned int y);
2337 extern SYCL_EXTERNAL unsigned int __imf_vcmplts4(unsigned int x,
2338  unsigned int y);
2339 extern SYCL_EXTERNAL unsigned int __imf_vcmpltu2(unsigned int x,
2340  unsigned int y);
2341 extern SYCL_EXTERNAL unsigned int __imf_vcmpltu4(unsigned int x,
2342  unsigned int y);
2343 extern SYCL_EXTERNAL unsigned int __imf_vmaxs2(unsigned int x, unsigned int y);
2344 extern SYCL_EXTERNAL unsigned int __imf_vmaxs4(unsigned int x, unsigned int y);
2345 extern SYCL_EXTERNAL unsigned int __imf_vmaxu2(unsigned int x, unsigned int y);
2346 extern SYCL_EXTERNAL unsigned int __imf_vmaxu4(unsigned int x, unsigned int y);
2347 extern SYCL_EXTERNAL unsigned int __imf_vmins2(unsigned int x, unsigned int y);
2348 extern SYCL_EXTERNAL unsigned int __imf_vmins4(unsigned int x, unsigned int y);
2349 extern SYCL_EXTERNAL unsigned int __imf_vminu2(unsigned int x, unsigned int y);
2350 extern SYCL_EXTERNAL unsigned int __imf_vminu4(unsigned int x, unsigned int y);
2351 extern SYCL_EXTERNAL unsigned int __imf_vseteq2(unsigned int x, unsigned int y);
2352 extern SYCL_EXTERNAL unsigned int __imf_vseteq4(unsigned int x, unsigned int y);
2353 extern SYCL_EXTERNAL unsigned int __imf_vsetne2(unsigned int x, unsigned int y);
2354 extern SYCL_EXTERNAL unsigned int __imf_vsetne4(unsigned int x, unsigned int y);
2355 extern SYCL_EXTERNAL unsigned int __imf_vsetges2(unsigned int x,
2356  unsigned int y);
2357 extern SYCL_EXTERNAL unsigned int __imf_vsetges4(unsigned int x,
2358  unsigned int y);
2359 extern SYCL_EXTERNAL unsigned int __imf_vsetgeu2(unsigned int x,
2360  unsigned int y);
2361 extern SYCL_EXTERNAL unsigned int __imf_vsetgeu4(unsigned int x,
2362  unsigned int y);
2363 extern SYCL_EXTERNAL unsigned int __imf_vsetgts2(unsigned int x,
2364  unsigned int y);
2365 extern SYCL_EXTERNAL unsigned int __imf_vsetgts4(unsigned int x,
2366  unsigned int y);
2367 extern SYCL_EXTERNAL unsigned int __imf_vsetgtu2(unsigned int x,
2368  unsigned int y);
2369 extern SYCL_EXTERNAL unsigned int __imf_vsetgtu4(unsigned int x,
2370  unsigned int y);
2371 extern SYCL_EXTERNAL unsigned int __imf_vsetles2(unsigned int x,
2372  unsigned int y);
2373 extern SYCL_EXTERNAL unsigned int __imf_vsetles4(unsigned int x,
2374  unsigned int y);
2375 extern SYCL_EXTERNAL unsigned int __imf_vsetleu2(unsigned int x,
2376  unsigned int y);
2377 extern SYCL_EXTERNAL unsigned int __imf_vsetleu4(unsigned int x,
2378  unsigned int y);
2379 extern SYCL_EXTERNAL unsigned int __imf_vsetlts2(unsigned int x,
2380  unsigned int y);
2381 extern SYCL_EXTERNAL unsigned int __imf_vsetlts4(unsigned int x,
2382  unsigned int y);
2383 extern SYCL_EXTERNAL unsigned int __imf_vsetltu2(unsigned int x,
2384  unsigned int y);
2385 extern SYCL_EXTERNAL unsigned int __imf_vsetltu4(unsigned int x,
2386  unsigned int y);
2387 extern SYCL_EXTERNAL unsigned int __imf_vsads2(unsigned int x, unsigned int y);
2388 extern SYCL_EXTERNAL unsigned int __imf_vsads4(unsigned int x, unsigned int y);
2389 extern SYCL_EXTERNAL unsigned int __imf_vsadu2(unsigned int x, unsigned int y);
2390 extern SYCL_EXTERNAL unsigned int __imf_vsadu4(unsigned int x, unsigned int y);
2391 }
2392 #ifdef __GLIBC__
2393 extern "C" {
2394 extern SYCL_EXTERNAL void __assert_fail(const char *expr, const char *file,
2395  unsigned int line, const char *func);
2396 extern SYCL_EXTERNAL float frexpf(float x, int *exp);
2397 extern SYCL_EXTERNAL float ldexpf(float x, int exp);
2398 extern SYCL_EXTERNAL float hypotf(float x, float y);
2399 
2400 // MS UCRT supports most of the C standard library but <complex.h> is
2401 // an exception.
2402 extern SYCL_EXTERNAL float cimagf(float __complex__ z);
2403 extern SYCL_EXTERNAL double cimag(double __complex__ z);
2404 extern SYCL_EXTERNAL float crealf(float __complex__ z);
2405 extern SYCL_EXTERNAL double creal(double __complex__ z);
2406 extern SYCL_EXTERNAL float cargf(float __complex__ z);
2407 extern SYCL_EXTERNAL double carg(double __complex__ z);
2408 extern SYCL_EXTERNAL float cabsf(float __complex__ z);
2409 extern SYCL_EXTERNAL double cabs(double __complex__ z);
2410 extern SYCL_EXTERNAL float __complex__ cprojf(float __complex__ z);
2411 extern SYCL_EXTERNAL double __complex__ cproj(double __complex__ z);
2412 extern SYCL_EXTERNAL float __complex__ cexpf(float __complex__ z);
2413 extern SYCL_EXTERNAL double __complex__ cexp(double __complex__ z);
2414 extern SYCL_EXTERNAL float __complex__ clogf(float __complex__ z);
2415 extern SYCL_EXTERNAL double __complex__ clog(double __complex__ z);
2416 extern SYCL_EXTERNAL float __complex__ cpowf(float __complex__ z);
2417 extern SYCL_EXTERNAL double __complex__ cpow(double __complex__ z);
2418 extern SYCL_EXTERNAL float __complex__ csqrtf(float __complex__ z);
2419 extern SYCL_EXTERNAL double __complex__ csqrt(double __complex__ z);
2420 extern SYCL_EXTERNAL float __complex__ csinhf(float __complex__ z);
2421 extern SYCL_EXTERNAL double __complex__ csinh(double __complex__ z);
2422 extern SYCL_EXTERNAL float __complex__ ccoshf(float __complex__ z);
2423 extern SYCL_EXTERNAL double __complex__ ccosh(double __complex__ z);
2424 extern SYCL_EXTERNAL float __complex__ ctanhf(float __complex__ z);
2425 extern SYCL_EXTERNAL double __complex__ ctanh(double __complex__ z);
2426 extern SYCL_EXTERNAL float __complex__ csinf(float __complex__ z);
2427 extern SYCL_EXTERNAL double __complex__ csin(double __complex__ z);
2428 extern SYCL_EXTERNAL float __complex__ ccosf(float __complex__ z);
2429 extern SYCL_EXTERNAL double __complex__ ccos(double __complex__ z);
2430 extern SYCL_EXTERNAL float __complex__ ctanf(float __complex__ z);
2431 extern SYCL_EXTERNAL double __complex__ ctan(double __complex__ z);
2432 extern SYCL_EXTERNAL float __complex__ cacosf(float __complex__ z);
2433 extern SYCL_EXTERNAL double __complex__ cacos(double __complex__ z);
2434 extern SYCL_EXTERNAL float __complex__ cacoshf(float __complex__ z);
2435 extern SYCL_EXTERNAL double __complex__ cacosh(double __complex__ z);
2436 extern SYCL_EXTERNAL float __complex__ casinf(float __complex__ z);
2437 extern SYCL_EXTERNAL double __complex__ casin(double __complex__ z);
2438 extern SYCL_EXTERNAL float __complex__ casinhf(float __complex__ z);
2439 extern SYCL_EXTERNAL double __complex__ casinh(double __complex__ z);
2440 extern SYCL_EXTERNAL float __complex__ catanf(float __complex__ z);
2441 extern SYCL_EXTERNAL double __complex__ catan(double __complex__ z);
2442 extern SYCL_EXTERNAL float __complex__ catanhf(float __complex__ z);
2443 extern SYCL_EXTERNAL double __complex__ catanh(double __complex__ z);
2444 extern SYCL_EXTERNAL float __complex__ cpolarf(float rho, float theta);
2445 extern SYCL_EXTERNAL double __complex__ cpolar(double rho, double theta);
2446 extern SYCL_EXTERNAL float __complex__ __mulsc3(float a, float b, float c,
2447  float d);
2448 extern SYCL_EXTERNAL double __complex__ __muldc3(double a, double b, double c,
2449  double d);
2450 extern SYCL_EXTERNAL float __complex__ __divsc3(float a, float b, float c,
2451  float d);
2452 extern SYCL_EXTERNAL double __complex__ __divdc3(float a, float b, float c,
2453  float d);
2454 }
2455 #elif defined(_WIN32)
2456 extern "C" {
2457 // TODO: documented C runtime library APIs must be recognized as
2458 // builtins by FE. This includes _dpcomp, _dsign, _dtest,
2459 // _fdpcomp, _fdsign, _fdtest, _hypotf, _wassert.
2460 // APIs used by STL, such as _Cosh, are undocumented, even though
2461 // they are open-sourced. Recognizing them as builtins is not
2462 // straightforward currently.
2463 extern SYCL_EXTERNAL double _Cosh(double x, double y);
2464 extern SYCL_EXTERNAL int _dpcomp(double x, double y);
2465 extern SYCL_EXTERNAL int _dsign(double x);
2466 extern SYCL_EXTERNAL short _Dtest(double *px);
2467 extern SYCL_EXTERNAL short _dtest(double *px);
2468 extern SYCL_EXTERNAL short _Exp(double *px, double y, short eoff);
2469 extern SYCL_EXTERNAL float _FCosh(float x, float y);
2470 extern SYCL_EXTERNAL int _fdpcomp(float x, float y);
2471 extern SYCL_EXTERNAL int _fdsign(float x);
2472 extern SYCL_EXTERNAL short _FDtest(float *px);
2473 extern SYCL_EXTERNAL short _fdtest(float *px);
2474 extern SYCL_EXTERNAL short _FExp(float *px, float y, short eoff);
2475 extern SYCL_EXTERNAL float _FSinh(float x, float y);
2476 extern SYCL_EXTERNAL double _Sinh(double x, double y);
2477 extern SYCL_EXTERNAL float _hypotf(float x, float y);
2478 extern SYCL_EXTERNAL void _wassert(const wchar_t *wexpr, const wchar_t *wfile,
2479  unsigned line);
2480 }
2481 #endif
2482 #endif // __SYCL_DEVICE_ONLY__
2483 
2484 #undef __NOEXC
__SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD
#define __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(NAME)
__SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL
#define __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME)
Definition: builtins.hpp:129
__SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD
#define __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(NAME)
__NOEXC
#define __NOEXC
Definition: builtins.hpp:18
__SYCL_MATH_FUNCTION_3_OVERLOAD
#define __SYCL_MATH_FUNCTION_3_OVERLOAD(NAME)
Definition: builtins.hpp:173
__SYCL_NATIVE_MATH_FUNCTION_OVERLOAD
#define __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(NAME)
__SYCL_MATH_FUNCTION_OVERLOAD_FM
#define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME)
Definition: builtins.hpp:106
__FAST_MATH_SGENFLOAT
#define __FAST_MATH_SGENFLOAT(T)
Definition: builtins.hpp:43
__SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD
#define __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(NAME)
__FAST_MATH_GENFLOAT
#define __FAST_MATH_GENFLOAT(T)
Definition: builtins.hpp:42
__SYCL_MATH_FUNCTION_2_OVERLOAD
#define __SYCL_MATH_FUNCTION_2_OVERLOAD(NAME)
Definition: builtins.hpp:141
__SYCL_MATH_FUNCTION_OVERLOAD
#define __SYCL_MATH_FUNCTION_OVERLOAD(NAME)
Definition: builtins.hpp:65
sycl::_V1::marray
Provides a cross-platform math array class template that works on SYCL devices as well as in host C++...
Definition: marray.hpp:24
sycl::_V1::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: types.hpp:558
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:39
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:50
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:25
SYCL_EXTERNAL
#define SYCL_EXTERNAL
Definition: defines_elementary.hpp:30
generic_type_traits.hpp
sycl::_V1::ext::intel::esimd::trunc
__ESIMD_API sycl::ext::intel::esimd::simd< RT, SZ > trunc(const sycl::ext::intel::esimd::simd< float, SZ > &src0, Sat sat={})
Round to integral value using the round to zero rounding mode (vector version).
Definition: math.hpp:614
sycl::_V1::ext::intel::esimd::floor
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > floor(const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
"Floor" operation, vector version - alias of rndd.
Definition: math.hpp:581
sycl::_V1::ext::intel::esimd::ceil
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > ceil(const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
"Ceiling" operation, vector version - alias of rndu.
Definition: math.hpp:594
sycl::_V1::ext::intel::esimd::sqrt
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
Definition: math.hpp:389
sycl::_V1::ext::intel::esimd::rsqrt
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
Definition: math.hpp:397
sycl::_V1::ext::intel::esimd::pow
__ESIMD_API simd< T, N > pow(simd< T, N > src0, simd< U, N > src1, Sat sat={})
Power - calculates src0 in power of src1.
Definition: math.hpp:442
sycl::_V1::ext::intel::esimd::log2
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
Definition: math.hpp:381
sycl::_V1::ext::intel::experimental::esimd::div
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T, SZ > > div(sycl::ext::intel::esimd::simd< T, SZ > &remainder, sycl::ext::intel::esimd::simd< T, SZ > src0, U src1)
Integral division with a vector dividend and a scalar divisor.
Definition: math.hpp:562
sycl::_V1::ext::intel::experimental::esimd::asin
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > > asin(sycl::ext::intel::esimd::simd< T, SZ > src0)
Definition: math.hpp:1243
sycl::_V1::ext::intel::experimental::esimd::sincos
__ESIMD_API sycl::ext::intel::esimd::simd< float, SZ > sincos(sycl::ext::intel::esimd::simd< float, SZ > &dstcos, U src0, Sat sat={})
Definition: math.hpp:1146
sycl::_V1::ext::intel::experimental::esimd::acos
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > > acos(sycl::ext::intel::esimd::simd< T, SZ > src0)
Definition: math.hpp:1205
sycl::_V1::ext::intel::experimental::esimd::atan
ESIMD_DETAIL ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< T, SZ > atan(sycl::ext::intel::esimd::simd< T, SZ > src0)
Definition: math.hpp:1161
sycl::_V1::ext::intel::experimental::esimd::fmod
sycl::ext::intel::esimd::simd< float, N > fmod(sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1386
sycl::_V1::ext::intel::experimental::esimd::line
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > > line(sycl::ext::intel::esimd::simd< T, 4 > src0, sycl::ext::intel::esimd::simd< T, SZ > src1, Sat sat={})
Linear equation.
Definition: math.hpp:921
sycl::_V1::ext::intel::experimental::esimd::atan2
sycl::ext::intel::esimd::simd< float, N > atan2(sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1353
sycl::_V1::ext::intel::esimd::abs
ESIMD_DETAIL __ESIMD_API std::enable_if_t< !std::is_same< std::remove_const_t< TRes >, std::remove_const_t< TArg > >::value, simd< TRes, SZ > > abs(simd< TArg, SZ > src0)
Get absolute value (vector version)
Definition: math.hpp:127
__imf_truncf
float __imf_truncf(float)
__imf_saturatef
float __imf_saturatef(float)
__imf_rsqrtf
float __imf_rsqrtf(float)
__imf_rint
double __imf_rint(double)
__imf_ceil
double __imf_ceil(double)
__imf_sqrt
double __imf_sqrt(double)
__imf_sqrtf
float __imf_sqrtf(float)
__imf_trunc
double __imf_trunc(double)
__imf_rintf16
_iml_half_internal __imf_rintf16(_iml_half_internal)
__imf_ceilf
float __imf_ceilf(float)
__imf_rsqrt
double __imf_rsqrt(double)
__imf_copysignf
float __imf_copysignf(float, float)
__imf_floorf16
_iml_half_internal __imf_floorf16(_iml_half_internal)
__imf_floor
double __imf_floor(double)
__imf_sqrtf16
_iml_half_internal __imf_sqrtf16(_iml_half_internal)
__imf_truncf16
_iml_half_internal __imf_truncf16(_iml_half_internal)
__imf_floorf
float __imf_floorf(float)
__imf_copysign
double __imf_copysign(double, double)
__imf_rintf
float __imf_rintf(float)
__imf_copysignf16
_iml_half_internal __imf_copysignf16(_iml_half_internal, _iml_half_internal)
__imf_rsqrtf16
_iml_half_internal __imf_rsqrtf16(_iml_half_internal)
__imf_ceilf16
_iml_half_internal __imf_ceilf16(_iml_half_internal)
__host_std
Definition: builtins.hpp:106
std::clog
__SYCL_EXTERN_STREAM_ATTRS ostream clog
Linked to standard error (buffered)
sycl::_V1::detail::is_geninteger
is_contained< T, gtl::integer_list > is_geninteger
Definition: generic_type_traits.hpp:149
sycl::_V1::detail::to_vec2
vec< T, 2 > to_vec2(marray< T, N > x, size_t start)
Definition: builtins.hpp:25
sycl::_V1::detail::msbIsSet
constexpr bool msbIsSet(const T x)
Definition: generic_type_traits.hpp:572
sycl::_V1::detail::memcpy
void memcpy(void *Dst, const void *Src, std::size_t Size)
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
sycl::_V1::ext::intel::math::copysign
std::enable_if_t< std::is_same_v< Tp, float >, float > copysign(Tp x, Tp y)
Definition: math.hpp:64
sycl::_V1::ext::intel::math::rint
std::enable_if_t< std::is_same_v< Tp, float >, float > rint(Tp x)
Definition: math.hpp:122
sycl::_V1::ext::intel::math::hadd
sycl::half hadd(sycl::half x, sycl::half y)
Definition: imf_half_trivial.hpp:17
sycl::_V1::ext::oneapi::experimental::matrix::use
use
Definition: matrix-unified-utils.hpp:17
sycl::_V1::ext::oneapi::experimental::native::tanh
__SYCL_ALWAYS_INLINE std::enable_if_t< std::is_same_v< T, half >||std::is_same_v< T, float >, sycl::marray< T, N > > tanh(sycl::marray< T, N > x) __NOEXC
Definition: builtins.hpp:114
sycl::_V1::ext::oneapi::experimental::native::exp2
__SYCL_ALWAYS_INLINE sycl::marray< half, N > exp2(sycl::marray< half, N > x) __NOEXC
Definition: builtins.hpp:157
sycl::_V1::ext::oneapi::experimental::isnan
std::enable_if_t< std::is_same< T, bfloat16 >::value, bool > isnan(T x)
Definition: bfloat16_math.hpp:36
sycl::_V1::ext::oneapi::fma
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fma(T x, T y, T z)
Definition: bf16_storage_builtins.hpp:71
sycl::_V1::ext::oneapi::fmax
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmax(T x, T y)
Definition: bf16_storage_builtins.hpp:60
sycl::_V1::ext::oneapi::fmin
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmin(T x, T y)
Definition: bf16_storage_builtins.hpp:49
sycl::_V1::ext::oneapi::fabs
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fabs(T x)
Definition: bf16_storage_builtins.hpp:39
sycl::_V1::instead
std::uint8_t instead
Definition: aliases.hpp:91
sycl::_V1::log
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > log(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:60
sycl::_V1::backend::all
@ all
sycl::_V1::sin
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > sin(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:38
sycl::_V1::cos
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > cos(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:27
sycl::_V1::exp
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > exp(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:49
sycl::_V1::powr
__SYCL_ALWAYS_INLINE std::enable_if_t< __FAST_MATH_SGENFLOAT(T), marray< T, N > > powr(marray< T, N > x, marray< T, N > y) __NOEXC
Definition: builtins.hpp:168
sycl::_V1::half
sycl::detail::half_impl::half half
Definition: aliases.hpp:99
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
popcount
int popcount(const simd_mask< _Tp, _Abi > &) noexcept
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
clamp
simd< _Tp, _Abi > clamp(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &)
__simd_abi
Definition: simd.hpp:718