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>
15 #include <sycl/types.hpp>
16 
17 // TODO Decide whether to mark functions with this attribute.
18 #define __NOEXC /*noexcept*/
19 
20 namespace sycl {
22 #ifdef __SYCL_DEVICE_ONLY__
23 #define __sycl_std
24 #else
25 namespace __sycl_std = __host_std;
26 #endif
27 
28 #ifdef __FAST_MATH__
29 #define __FAST_MATH_GENFLOAT(T) \
30  (detail::is_genfloatd<T>::value || detail::is_genfloath<T>::value)
31 #else
32 #define __FAST_MATH_GENFLOAT(T) (detail::is_genfloat<T>::value)
33 #endif
34 
35 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
36 // genfloat acos (genfloat x)
37 template <typename T>
39  return __sycl_std::__invoke_acos<T>(x);
40 }
41 
42 // genfloat acosh (genfloat x)
43 template <typename T>
45  return __sycl_std::__invoke_acosh<T>(x);
46 }
47 
48 // genfloat acospi (genfloat x)
49 template <typename T>
51  return __sycl_std::__invoke_acospi<T>(x);
52 }
53 
54 // genfloat asin (genfloat x)
55 template <typename T>
57  return __sycl_std::__invoke_asin<T>(x);
58 }
59 
60 // genfloat asinh (genfloat x)
61 template <typename T>
63  return __sycl_std::__invoke_asinh<T>(x);
64 }
65 
66 // genfloat asinpi (genfloat x)
67 template <typename T>
69  return __sycl_std::__invoke_asinpi<T>(x);
70 }
71 
72 // genfloat atan (genfloat y_over_x)
73 template <typename T>
75  return __sycl_std::__invoke_atan<T>(y_over_x);
76 }
77 
78 // genfloat atan2 (genfloat y, genfloat x)
79 template <typename T>
81  return __sycl_std::__invoke_atan2<T>(y, x);
82 }
83 
84 // genfloat atanh (genfloat x)
85 template <typename T>
87  return __sycl_std::__invoke_atanh<T>(x);
88 }
89 
90 // genfloat atanpi (genfloat x)
91 template <typename T>
93  return __sycl_std::__invoke_atanpi<T>(x);
94 }
95 
96 // genfloat atan2pi (genfloat y, genfloat x)
97 template <typename T>
99  T x) __NOEXC {
100  return __sycl_std::__invoke_atan2pi<T>(y, x);
101 }
102 
103 // genfloat cbrt (genfloat x)
104 template <typename T>
106  return __sycl_std::__invoke_cbrt<T>(x);
107 }
108 
109 // genfloat ceil (genfloat x)
110 template <typename T>
112  return __sycl_std::__invoke_ceil<T>(x);
113 }
114 
115 // genfloat copysign (genfloat x, genfloat y)
116 template <typename T>
118  T y) __NOEXC {
119  return __sycl_std::__invoke_copysign<T>(x, y);
120 }
121 
122 // genfloat cos (genfloat x)
123 template <typename T>
125  return __sycl_std::__invoke_cos<T>(x);
126 }
127 
128 // genfloat cosh (genfloat x)
129 template <typename T>
131  return __sycl_std::__invoke_cosh<T>(x);
132 }
133 
134 // genfloat cospi (genfloat x)
135 template <typename T>
137  return __sycl_std::__invoke_cospi<T>(x);
138 }
139 
140 // genfloat erfc (genfloat x)
141 template <typename T>
143  return __sycl_std::__invoke_erfc<T>(x);
144 }
145 
146 // genfloat erf (genfloat x)
147 template <typename T>
149  return __sycl_std::__invoke_erf<T>(x);
150 }
151 
152 // genfloat exp (genfloat x )
153 template <typename T>
155  return __sycl_std::__invoke_exp<T>(x);
156 }
157 
158 // genfloat exp2 (genfloat x)
159 template <typename T>
161  return __sycl_std::__invoke_exp2<T>(x);
162 }
163 
164 // genfloat exp10 (genfloat x)
165 template <typename T>
167  return __sycl_std::__invoke_exp10<T>(x);
168 }
169 
170 // genfloat expm1 (genfloat x)
171 template <typename T>
173  return __sycl_std::__invoke_expm1<T>(x);
174 }
175 
176 // genfloat fabs (genfloat x)
177 template <typename T>
179  return __sycl_std::__invoke_fabs<T>(x);
180 }
181 
182 // genfloat fdim (genfloat x, genfloat y)
183 template <typename T>
185  return __sycl_std::__invoke_fdim<T>(x, y);
186 }
187 
188 // genfloat floor (genfloat x)
189 template <typename T>
191  return __sycl_std::__invoke_floor<T>(x);
192 }
193 
194 // genfloat fma (genfloat a, genfloat b, genfloat c)
195 template <typename T>
197  T c) __NOEXC {
198  return __sycl_std::__invoke_fma<T>(a, b, c);
199 }
200 
201 // genfloat fmax (genfloat x, genfloat y)
202 template <typename T>
204  return __sycl_std::__invoke_fmax<T>(x, y);
205 }
206 
207 // genfloat fmax (genfloat x, sgenfloat y)
208 template <typename T>
209 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
210 fmax(T x, typename T::element_type y) __NOEXC {
211  return __sycl_std::__invoke_fmax<T>(x, T(y));
212 }
213 
214 // genfloat fmin (genfloat x, genfloat y)
215 template <typename T>
217  return __sycl_std::__invoke_fmin<T>(x, y);
218 }
219 
220 // genfloat fmin (genfloat x, sgenfloat y)
221 template <typename T>
222 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
223 fmin(T x, typename T::element_type y) __NOEXC {
224  return __sycl_std::__invoke_fmin<T>(x, T(y));
225 }
226 
227 // genfloat fmod (genfloat x, genfloat y)
228 template <typename T>
230  return __sycl_std::__invoke_fmod<T>(x, y);
231 }
232 
233 // genfloat fract (genfloat x, genfloatptr iptr)
234 template <typename T, typename T2>
237 fract(T x, T2 iptr) __NOEXC {
238  detail::check_vector_size<T, T2>();
239  return __sycl_std::__invoke_fract<T>(x, iptr);
240 }
241 
242 // genfloat frexp (genfloat x, genintptr exp)
243 template <typename T, typename T2>
245  detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
246 frexp(T x, T2 exp) __NOEXC {
247  detail::check_vector_size<T, T2>();
248  return __sycl_std::__invoke_frexp<T>(x, exp);
249 }
250 
251 // genfloat hypot (genfloat x, genfloat y)
252 template <typename T>
254  return __sycl_std::__invoke_hypot<T>(x, y);
255 }
256 
257 // genint ilogb (genfloat x)
258 template <typename T,
259  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
261  return __sycl_std::__invoke_ilogb<detail::change_base_type_t<T, int>>(x);
262 }
263 
264 // float ldexp (float x, int k)
265 // double ldexp (double x, int k)
266 // half ldexp (half x, int k)
267 template <typename T>
269  int k) __NOEXC {
270  return __sycl_std::__invoke_ldexp<T>(x, k);
271 }
272 
273 // vgenfloat ldexp (vgenfloat x, int k)
274 template <typename T>
276  int k) __NOEXC {
277  return __sycl_std::__invoke_ldexp<T>(x, vec<int, T::size()>(k));
278 }
279 
280 // vgenfloat ldexp (vgenfloat x, genint k)
281 template <typename T, typename T2>
283  detail::is_vgenfloat<T>::value && detail::is_intn<T2>::value, T>
284 ldexp(T x, T2 k) __NOEXC {
285  detail::check_vector_size<T, T2>();
286  return __sycl_std::__invoke_ldexp<T>(x, k);
287 }
288 
289 // genfloat lgamma (genfloat x)
290 template <typename T>
292  return __sycl_std::__invoke_lgamma<T>(x);
293 }
294 
295 // genfloat lgamma_r (genfloat x, genintptr signp)
296 template <typename T, typename T2>
298  detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
299 lgamma_r(T x, T2 signp) __NOEXC {
300  detail::check_vector_size<T, T2>();
301  return __sycl_std::__invoke_lgamma_r<T>(x, signp);
302 }
303 
304 // genfloat log (genfloat x)
305 template <typename T>
307  return __sycl_std::__invoke_log<T>(x);
308 }
309 
310 // genfloat log2 (genfloat x)
311 template <typename T>
313  return __sycl_std::__invoke_log2<T>(x);
314 }
315 
316 // genfloat log10 (genfloat x)
317 template <typename T>
319  return __sycl_std::__invoke_log10<T>(x);
320 }
321 
322 // genfloat log1p (genfloat x)
323 template <typename T>
325  return __sycl_std::__invoke_log1p<T>(x);
326 }
327 
328 // genfloat logb (genfloat x)
329 template <typename T>
331  return __sycl_std::__invoke_logb<T>(x);
332 }
333 
334 // genfloat mad (genfloat a, genfloat b, genfloat c)
335 template <typename T>
337  T c) __NOEXC {
338  return __sycl_std::__invoke_mad<T>(a, b, c);
339 }
340 
341 // genfloat maxmag (genfloat x, genfloat y)
342 template <typename T>
344  return __sycl_std::__invoke_maxmag<T>(x, y);
345 }
346 
347 // genfloat minmag (genfloat x, genfloat y)
348 template <typename T>
350  return __sycl_std::__invoke_minmag<T>(x, y);
351 }
352 
353 // genfloat modf (genfloat x, genfloatptr iptr)
354 template <typename T, typename T2>
356  detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
357 modf(T x, T2 iptr) __NOEXC {
358  detail::check_vector_size<T, T2>();
359  return __sycl_std::__invoke_modf<T>(x, iptr);
360 }
361 
362 template <typename T,
363  typename = detail::enable_if_t<detail::is_nan_type<T>::value, T>>
365  return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
367 }
368 
369 // genfloat nextafter (genfloat x, genfloat y)
370 template <typename T>
372  T y) __NOEXC {
373  return __sycl_std::__invoke_nextafter<T>(x, y);
374 }
375 
376 // genfloat pow (genfloat x, genfloat y)
377 template <typename T>
379  return __sycl_std::__invoke_pow<T>(x, y);
380 }
381 
382 // genfloat pown (genfloat x, genint y)
383 template <typename T, typename T2>
385  detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>
386 pown(T x, T2 y) __NOEXC {
387  detail::check_vector_size<T, T2>();
388  return __sycl_std::__invoke_pown<T>(x, y);
389 }
390 
391 // genfloat powr (genfloat x, genfloat y)
392 template <typename T>
394  return __sycl_std::__invoke_powr<T>(x, y);
395 }
396 
397 // genfloat remainder (genfloat x, genfloat y)
398 template <typename T>
400  T y) __NOEXC {
401  return __sycl_std::__invoke_remainder<T>(x, y);
402 }
403 
404 // genfloat remquo (genfloat x, genfloat y, genintptr quo)
405 template <typename T, typename T2>
407  detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
408 remquo(T x, T y, T2 quo) __NOEXC {
409  detail::check_vector_size<T, T2>();
410  return __sycl_std::__invoke_remquo<T>(x, y, quo);
411 }
412 
413 // genfloat rint (genfloat x)
414 template <typename T>
416  return __sycl_std::__invoke_rint<T>(x);
417 }
418 
419 // genfloat rootn (genfloat x, genint y)
420 template <typename T, typename T2>
422  detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>
423 rootn(T x, T2 y) __NOEXC {
424  detail::check_vector_size<T, T2>();
425  return __sycl_std::__invoke_rootn<T>(x, y);
426 }
427 
428 // genfloat round (genfloat x)
429 template <typename T>
431  return __sycl_std::__invoke_round<T>(x);
432 }
433 
434 // genfloat rsqrt (genfloat x)
435 template <typename T>
437  return __sycl_std::__invoke_rsqrt<T>(x);
438 }
439 
440 // genfloat sin (genfloat x)
441 template <typename T>
443  return __sycl_std::__invoke_sin<T>(x);
444 }
445 
446 // genfloat sincos (genfloat x, genfloatptr cosval)
447 template <typename T, typename T2>
449  detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
450 sincos(T x, T2 cosval) __NOEXC {
451  detail::check_vector_size<T, T2>();
452  return __sycl_std::__invoke_sincos<T>(x, cosval);
453 }
454 
455 // genfloat sinh (genfloat x)
456 template <typename T>
458  return __sycl_std::__invoke_sinh<T>(x);
459 }
460 
461 // genfloat sinpi (genfloat x)
462 template <typename T>
464  return __sycl_std::__invoke_sinpi<T>(x);
465 }
466 
467 // genfloat sqrt (genfloat x)
468 template <typename T>
470  return __sycl_std::__invoke_sqrt<T>(x);
471 }
472 
473 // genfloat tan (genfloat x)
474 template <typename T>
476  return __sycl_std::__invoke_tan<T>(x);
477 }
478 
479 // genfloat tanh (genfloat x)
480 template <typename T>
482  return __sycl_std::__invoke_tanh<T>(x);
483 }
484 
485 // genfloat tanpi (genfloat x)
486 template <typename T>
488  return __sycl_std::__invoke_tanpi<T>(x);
489 }
490 
491 // genfloat tgamma (genfloat x)
492 template <typename T>
494  return __sycl_std::__invoke_tgamma<T>(x);
495 }
496 
497 // genfloat trunc (genfloat x)
498 template <typename T>
500  return __sycl_std::__invoke_trunc<T>(x);
501 }
502 
503 /* --------------- 4.13.5 Common functions. ---------------------------------*/
504 // genfloat clamp (genfloat x, genfloat minval, genfloat maxval)
505 template <typename T>
507  T maxval) __NOEXC {
508  return __sycl_std::__invoke_fclamp<T>(x, minval, maxval);
509 }
510 
511 // genfloath clamp (genfloath x, half minval, half maxval)
512 // genfloatf clamp (genfloatf x, float minval, float maxval)
513 // genfloatd clamp (genfloatd x, double minval, double maxval)
514 template <typename T>
515 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
516 clamp(T x, typename T::element_type minval,
517  typename T::element_type maxval) __NOEXC {
518  return __sycl_std::__invoke_fclamp<T>(x, T(minval), T(maxval));
519 }
520 
521 // genfloat degrees (genfloat radians)
522 template <typename T>
523 detail::enable_if_t<detail::is_genfloat<T>::value, T>
525  return __sycl_std::__invoke_degrees<T>(radians);
526 }
527 
528 // genfloat abs (genfloat x)
529 template <typename T>
531  return __sycl_std::__invoke_fabs<T>(x);
532 }
533 
534 // genfloat max (genfloat x, genfloat y)
535 template <typename T>
537  return __sycl_std::__invoke_fmax_common<T>(x, y);
538 }
539 
540 // genfloatf max (genfloatf x, float y)
541 // genfloatd max (genfloatd x, double y)
542 // genfloath max (genfloath x, half y)
543 template <typename T>
544 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(max)(
545  T x, typename T::element_type y) __NOEXC {
546  return __sycl_std::__invoke_fmax_common<T>(x, T(y));
547 }
548 
549 // genfloat min (genfloat x, genfloat y)
550 template <typename T>
551 detail::enable_if_t<detail::is_genfloat<T>::value, T>(min)(T x, T y) __NOEXC {
552  return __sycl_std::__invoke_fmin_common<T>(x, y);
553 }
554 
555 // genfloatf min (genfloatf x, float y)
556 // genfloatd min (genfloatd x, double y)
557 // genfloath min (genfloath x, half y)
558 template <typename T>
559 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(min)(
560  T x, typename T::element_type y) __NOEXC {
561  return __sycl_std::__invoke_fmin_common<T>(x, T(y));
562 }
563 
564 // genfloat mix (genfloat x, genfloat y, genfloat a)
565 template <typename T>
567  T a) __NOEXC {
568  return __sycl_std::__invoke_mix<T>(x, y, a);
569 }
570 
571 // genfloatf mix (genfloatf x, genfloatf y, float a)
572 // genfloatd mix (genfloatd x, genfloatd y, double a)
573 // genfloatd mix (genfloath x, genfloath y, half a)
574 template <typename T>
575 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
576 mix(T x, T y, typename T::element_type a) __NOEXC {
577  return __sycl_std::__invoke_mix<T>(x, y, T(a));
578 }
579 
580 // genfloat radians (genfloat degrees)
581 template <typename T>
582 detail::enable_if_t<detail::is_genfloat<T>::value, T>
584  return __sycl_std::__invoke_radians<T>(degrees);
585 }
586 
587 // genfloat step (genfloat edge, genfloat x)
588 template <typename T>
590  T x) __NOEXC {
591  return __sycl_std::__invoke_step<T>(edge, x);
592 }
593 
594 // genfloatf step (float edge, genfloatf x)
595 // genfloatd step (double edge, genfloatd x)
596 // genfloatd step (half edge, genfloath x)
597 template <typename T>
598 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
599 step(typename T::element_type edge, T x) __NOEXC {
600  return __sycl_std::__invoke_step<T>(T(edge), x);
601 }
602 
603 // genfloat smoothstep (genfloat edge0, genfloat edge1, genfloat x)
604 template <typename T>
605 detail::enable_if_t<detail::is_genfloat<T>::value, T>
606 smoothstep(T edge0, T edge1, T x) __NOEXC {
607  return __sycl_std::__invoke_smoothstep<T>(edge0, edge1, x);
608 }
609 
610 // genfloatf smoothstep (float edge0, float edge1, genfloatf x)
611 // genfloatd smoothstep (double edge0, double edge1, genfloatd x)
612 // genfloath smoothstep (half edge0, half edge1, genfloath x)
613 template <typename T>
614 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
615 smoothstep(typename T::element_type edge0, typename T::element_type edge1,
616  T x) __NOEXC {
617  return __sycl_std::__invoke_smoothstep<T>(T(edge0), T(edge1), x);
618 }
619 
620 // genfloat sign (genfloat x)
621 template <typename T>
623  return __sycl_std::__invoke_sign<T>(x);
624 }
625 
626 /* --------------- 4.13.4 Integer functions. --------------------------------*/
627 // ugeninteger abs (geninteger x)
628 template <typename T>
630  return __sycl_std::__invoke_u_abs<T>(x);
631 }
632 
633 // ugeninteger abs (geninteger x)
634 template <typename T>
635 detail::enable_if_t<detail::is_igeninteger<T>::value,
636  detail::make_unsigned_t<T>>
637 abs(T x) __NOEXC {
638  return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
639 }
640 
641 // ugeninteger abs_diff (geninteger x, geninteger y)
642 template <typename T>
644  T y) __NOEXC {
645  return __sycl_std::__invoke_u_abs_diff<T>(x, y);
646 }
647 
648 // ugeninteger abs_diff (geninteger x, geninteger y)
649 template <typename T>
650 detail::enable_if_t<detail::is_igeninteger<T>::value,
651  detail::make_unsigned_t<T>>
652 abs_diff(T x, T y) __NOEXC {
653  return __sycl_std::__invoke_s_abs_diff<detail::make_unsigned_t<T>>(x, y);
654 }
655 
656 // geninteger add_sat (geninteger x, geninteger y)
657 template <typename T>
659  T y) __NOEXC {
660  return __sycl_std::__invoke_s_add_sat<T>(x, y);
661 }
662 
663 // geninteger add_sat (geninteger x, geninteger y)
664 template <typename T>
666  T y) __NOEXC {
667  return __sycl_std::__invoke_u_add_sat<T>(x, y);
668 }
669 
670 // geninteger hadd (geninteger x, geninteger y)
671 template <typename T>
673  T y) __NOEXC {
674  return __sycl_std::__invoke_s_hadd<T>(x, y);
675 }
676 
677 // geninteger hadd (geninteger x, geninteger y)
678 template <typename T>
680  T y) __NOEXC {
681  return __sycl_std::__invoke_u_hadd<T>(x, y);
682 }
683 
684 // geninteger rhadd (geninteger x, geninteger y)
685 template <typename T>
687  T y) __NOEXC {
688  return __sycl_std::__invoke_s_rhadd<T>(x, y);
689 }
690 
691 // geninteger rhadd (geninteger x, geninteger y)
692 template <typename T>
694  T y) __NOEXC {
695  return __sycl_std::__invoke_u_rhadd<T>(x, y);
696 }
697 
698 // geninteger clamp (geninteger x, geninteger minval, geninteger maxval)
699 template <typename T>
700 detail::enable_if_t<detail::is_igeninteger<T>::value, T>
701 clamp(T x, T minval, T maxval) __NOEXC {
702  return __sycl_std::__invoke_s_clamp<T>(x, minval, maxval);
703 }
704 
705 // geninteger clamp (geninteger x, geninteger minval, geninteger maxval)
706 template <typename T>
707 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>
708 clamp(T x, T minval, T maxval) __NOEXC {
709  return __sycl_std::__invoke_u_clamp<T>(x, minval, maxval);
710 }
711 
712 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
713 template <typename T>
714 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>
715 clamp(T x, typename T::element_type minval,
716  typename T::element_type maxval) __NOEXC {
717  return __sycl_std::__invoke_s_clamp<T>(x, T(minval), T(maxval));
718 }
719 
720 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
721 template <typename T>
722 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>
723 clamp(T x, typename T::element_type minval,
724  typename T::element_type maxval) __NOEXC {
725  return __sycl_std::__invoke_u_clamp<T>(x, T(minval), T(maxval));
726 }
727 
728 // geninteger clz (geninteger x)
729 template <typename T>
731  return __sycl_std::__invoke_clz<T>(x);
732 }
733 
734 // geninteger ctz (geninteger x)
735 template <typename T>
737  return __sycl_std::__invoke_ctz<T>(x);
738 }
739 
740 // geninteger ctz (geninteger x) for calls with deprecated namespace
741 namespace ext {
742 namespace intel {
743 template <typename T>
745  "'sycl::ext::intel::ctz' is deprecated, use 'sycl::ctz' instead")
746 sycl::detail::enable_if_t<sycl::detail::is_geninteger<T>::value, T> ctz(
747  T x) __NOEXC {
748  return sycl::ctz(x);
749 }
750 } // namespace intel
751 } // namespace ext
752 
753 namespace __SYCL2020_DEPRECATED("use 'ext::intel' instead") intel {
754 using namespace ext::intel;
755 }
756 
757 // geninteger mad_hi (geninteger a, geninteger b, geninteger c)
758 template <typename T>
760  T z) __NOEXC {
761  return __sycl_std::__invoke_s_mad_hi<T>(x, y, z);
762 }
763 
764 // geninteger mad_hi (geninteger a, geninteger b, geninteger c)
765 template <typename T>
767  T z) __NOEXC {
768  return __sycl_std::__invoke_u_mad_hi<T>(x, y, z);
769 }
770 
771 // geninteger mad_sat (geninteger a, geninteger b, geninteger c)
772 template <typename T>
774  T c) __NOEXC {
775  return __sycl_std::__invoke_s_mad_sat<T>(a, b, c);
776 }
777 
778 // geninteger mad_sat (geninteger a, geninteger b, geninteger c)
779 template <typename T>
781  T c) __NOEXC {
782  return __sycl_std::__invoke_u_mad_sat<T>(a, b, c);
783 }
784 
785 // igeninteger max (igeninteger x, igeninteger y)
786 template <typename T>
787 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(max)(T x,
788  T y) __NOEXC {
789  return __sycl_std::__invoke_s_max<T>(x, y);
790 }
791 
792 // ugeninteger max (ugeninteger x, ugeninteger y)
793 template <typename T>
794 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>(max)(T x,
795  T y) __NOEXC {
796  return __sycl_std::__invoke_u_max<T>(x, y);
797 }
798 
799 // igeninteger max (vigeninteger x, sigeninteger y)
800 template <typename T>
801 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>(max)(
802  T x, typename T::element_type y) __NOEXC {
803  return __sycl_std::__invoke_s_max<T>(x, T(y));
804 }
805 
806 // vugeninteger max (vugeninteger x, sugeninteger y)
807 template <typename T>
808 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>(max)(
809  T x, typename T::element_type y) __NOEXC {
810  return __sycl_std::__invoke_u_max<T>(x, T(y));
811 }
812 
813 // igeninteger min (igeninteger x, igeninteger y)
814 template <typename T>
815 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(min)(T x,
816  T y) __NOEXC {
817  return __sycl_std::__invoke_s_min<T>(x, y);
818 }
819 
820 // ugeninteger min (ugeninteger x, ugeninteger y)
821 template <typename T>
822 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>(min)(T x,
823  T y) __NOEXC {
824  return __sycl_std::__invoke_u_min<T>(x, y);
825 }
826 
827 // vigeninteger min (vigeninteger x, sigeninteger y)
828 template <typename T>
829 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>(min)(
830  T x, typename T::element_type y) __NOEXC {
831  return __sycl_std::__invoke_s_min<T>(x, T(y));
832 }
833 
834 // vugeninteger min (vugeninteger x, sugeninteger y)
835 template <typename T>
836 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>(min)(
837  T x, typename T::element_type y) __NOEXC {
838  return __sycl_std::__invoke_u_min<T>(x, T(y));
839 }
840 
841 // geninteger mul_hi (geninteger x, geninteger y)
842 template <typename T>
844  T y) __NOEXC {
845  return __sycl_std::__invoke_s_mul_hi<T>(x, y);
846 }
847 
848 // geninteger mul_hi (geninteger x, geninteger y)
849 template <typename T>
851  T y) __NOEXC {
852  return __sycl_std::__invoke_u_mul_hi<T>(x, y);
853 }
854 
855 // geninteger rotate (geninteger v, geninteger i)
856 template <typename T>
858  T i) __NOEXC {
859  return __sycl_std::__invoke_rotate<T>(v, i);
860 }
861 
862 // geninteger sub_sat (geninteger x, geninteger y)
863 template <typename T>
865  T y) __NOEXC {
866  return __sycl_std::__invoke_s_sub_sat<T>(x, y);
867 }
868 
869 // geninteger sub_sat (geninteger x, geninteger y)
870 template <typename T>
872  T y) __NOEXC {
873  return __sycl_std::__invoke_u_sub_sat<T>(x, y);
874 }
875 
876 // ugeninteger16bit upsample (ugeninteger8bit hi, ugeninteger8bit lo)
877 template <typename T>
878 detail::enable_if_t<detail::is_ugeninteger8bit<T>::value,
879  detail::make_larger_t<T>>
880 upsample(T hi, T lo) __NOEXC {
881  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
882 }
883 
884 // igeninteger16bit upsample (igeninteger8bit hi, ugeninteger8bit lo)
885 template <typename T, typename T2>
886 detail::enable_if_t<detail::is_igeninteger8bit<T>::value &&
887  detail::is_ugeninteger8bit<T2>::value,
888  detail::make_larger_t<T>>
889 upsample(T hi, T2 lo) __NOEXC {
890  detail::check_vector_size<T, T2>();
891  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
892 }
893 
894 // ugeninteger32bit upsample (ugeninteger16bit hi, ugeninteger16bit lo)
895 template <typename T>
896 detail::enable_if_t<detail::is_ugeninteger16bit<T>::value,
897  detail::make_larger_t<T>>
898 upsample(T hi, T lo) __NOEXC {
899  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
900 }
901 
902 // igeninteger32bit upsample (igeninteger16bit hi, ugeninteger16bit lo)
903 template <typename T, typename T2>
904 detail::enable_if_t<detail::is_igeninteger16bit<T>::value &&
905  detail::is_ugeninteger16bit<T2>::value,
906  detail::make_larger_t<T>>
907 upsample(T hi, T2 lo) __NOEXC {
908  detail::check_vector_size<T, T2>();
909  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
910 }
911 
912 // ugeninteger64bit upsample (ugeninteger32bit hi, ugeninteger32bit lo)
913 template <typename T>
914 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value,
915  detail::make_larger_t<T>>
916 upsample(T hi, T lo) __NOEXC {
917  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
918 }
919 
920 // igeninteger64bit upsample (igeninteger32bit hi, ugeninteger32bit lo)
921 template <typename T, typename T2>
922 detail::enable_if_t<detail::is_igeninteger32bit<T>::value &&
923  detail::is_ugeninteger32bit<T2>::value,
924  detail::make_larger_t<T>>
925 upsample(T hi, T2 lo) __NOEXC {
926  detail::check_vector_size<T, T2>();
927  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
928 }
929 
930 // geninteger popcount (geninteger x)
931 template <typename T>
933  return __sycl_std::__invoke_popcount<T>(x);
934 }
935 
936 // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y,
937 // geninteger32bit z)
938 template <typename T>
939 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
940 mad24(T x, T y, T z) __NOEXC {
941  return __sycl_std::__invoke_s_mad24<T>(x, y, z);
942 }
943 
944 // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y,
945 // geninteger32bit z)
946 template <typename T>
947 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
948 mad24(T x, T y, T z) __NOEXC {
949  return __sycl_std::__invoke_u_mad24<T>(x, y, z);
950 }
951 
952 // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y)
953 template <typename T>
954 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
955 mul24(T x, T y) __NOEXC {
956  return __sycl_std::__invoke_s_mul24<T>(x, y);
957 }
958 
959 // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y)
960 template <typename T>
961 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
962 mul24(T x, T y) __NOEXC {
963  return __sycl_std::__invoke_u_mul24<T>(x, y);
964 }
965 
966 /* --------------- 4.13.6 Geometric Functions. ------------------------------*/
967 // float3 cross (float3 p0, float3 p1)
968 // float4 cross (float4 p0, float4 p1)
969 // double3 cross (double3 p0, double3 p1)
970 // double4 cross (double4 p0, double4 p1)
971 // half3 cross (half3 p0, half3 p1)
972 // half4 cross (half4 p0, half4 p1)
973 template <typename T>
975  T p1) __NOEXC {
976  return __sycl_std::__invoke_cross<T>(p0, p1);
977 }
978 
979 // float dot (float p0, float p1)
980 // double dot (double p0, double p1)
981 // half dot (half p0, half p1)
982 template <typename T>
984  return p0 * p1;
985 }
986 
987 // float dot (vgengeofloat p0, vgengeofloat p1)
988 template <typename T>
989 detail::enable_if_t<detail::is_vgengeofloat<T>::value, float>
990 dot(T p0, T p1) __NOEXC {
991  return __sycl_std::__invoke_Dot<float>(p0, p1);
992 }
993 
994 // double dot (vgengeodouble p0, vgengeodouble p1)
995 template <typename T>
996 detail::enable_if_t<detail::is_vgengeodouble<T>::value, double>
997 dot(T p0, T p1) __NOEXC {
998  return __sycl_std::__invoke_Dot<double>(p0, p1);
999 }
1000 
1001 // half dot (vgengeohalf p0, vgengeohalf p1)
1002 template <typename T>
1004  T p1) __NOEXC {
1005  return __sycl_std::__invoke_Dot<half>(p0, p1);
1006 }
1007 
1008 // float distance (gengeofloat p0, gengeofloat p1)
1009 template <typename T,
1010  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1011 float distance(T p0, T p1) __NOEXC {
1012  return __sycl_std::__invoke_distance<float>(p0, p1);
1013 }
1014 
1015 // double distance (gengeodouble p0, gengeodouble p1)
1016 template <typename T,
1017  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1018 double distance(T p0, T p1) __NOEXC {
1019  return __sycl_std::__invoke_distance<double>(p0, p1);
1020 }
1021 
1022 // half distance (gengeohalf p0, gengeohalf p1)
1023 template <typename T,
1024  typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1025 half distance(T p0, T p1) __NOEXC {
1026  return __sycl_std::__invoke_distance<half>(p0, p1);
1027 }
1028 
1029 // float length (gengeofloat p)
1030 template <typename T,
1031  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1032 float length(T p) __NOEXC {
1033  return __sycl_std::__invoke_length<float>(p);
1034 }
1035 
1036 // double length (gengeodouble p)
1037 template <typename T,
1038  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1039 double length(T p) __NOEXC {
1040  return __sycl_std::__invoke_length<double>(p);
1041 }
1042 
1043 // half length (gengeohalf p)
1044 template <typename T,
1045  typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1047  return __sycl_std::__invoke_length<half>(p);
1048 }
1049 
1050 // gengeofloat normalize (gengeofloat p)
1051 template <typename T>
1052 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1054  return __sycl_std::__invoke_normalize<T>(p);
1055 }
1056 
1057 // gengeodouble normalize (gengeodouble p)
1058 template <typename T>
1059 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1061  return __sycl_std::__invoke_normalize<T>(p);
1062 }
1063 
1064 // gengeohalf normalize (gengeohalf p)
1065 template <typename T>
1067  return __sycl_std::__invoke_normalize<T>(p);
1068 }
1069 
1070 // float fast_distance (gengeofloat p0, gengeofloat p1)
1071 template <typename T,
1072  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1073 float fast_distance(T p0, T p1) __NOEXC {
1074  return __sycl_std::__invoke_fast_distance<float>(p0, p1);
1075 }
1076 
1077 // double fast_distance (gengeodouble p0, gengeodouble p1)
1078 template <typename T,
1079  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1080 double fast_distance(T p0, T p1) __NOEXC {
1081  return __sycl_std::__invoke_fast_distance<double>(p0, p1);
1082 }
1083 
1084 // float fast_length (gengeofloat p)
1085 template <typename T,
1086  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1087 float fast_length(T p) __NOEXC {
1088  return __sycl_std::__invoke_fast_length<float>(p);
1089 }
1090 
1091 // double fast_length (gengeodouble p)
1092 template <typename T,
1093  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1094 double fast_length(T p) __NOEXC {
1095  return __sycl_std::__invoke_fast_length<double>(p);
1096 }
1097 
1098 // gengeofloat fast_normalize (gengeofloat p)
1099 template <typename T>
1100 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1102  return __sycl_std::__invoke_fast_normalize<T>(p);
1103 }
1104 
1105 // gengeodouble fast_normalize (gengeodouble p)
1106 template <typename T>
1107 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1109  return __sycl_std::__invoke_fast_normalize<T>(p);
1110 }
1111 
1112 /* --------------- 4.13.7 Relational functions. Device version --------------*/
1113 // int isequal (half x, half y)
1114 // shortn isequal (halfn x, halfn y)
1115 // igeninteger32bit isequal (genfloatf x, genfloatf y)
1116 // int isequal (double x,double y);
1117 // longn isequal (doublen x, doublen y)
1118 template <typename T,
1119  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1122  __sycl_std::__invoke_FOrdEqual<detail::rel_ret_t<T>>(x, y));
1123 }
1124 
1125 // int isnotequal (half x, half y)
1126 // shortn isnotequal (halfn x, halfn y)
1127 // igeninteger32bit isnotequal (genfloatf x, genfloatf y)
1128 // int isnotequal (double x, double y)
1129 // longn isnotequal (doublen x, doublen y)
1130 template <typename T,
1131  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1134  __sycl_std::__invoke_FUnordNotEqual<detail::rel_ret_t<T>>(x, y));
1135 }
1136 
1137 // int isgreater (half x, half y)
1138 // shortn isgreater (halfn x, halfn y)
1139 // igeninteger32bit isgreater (genfloatf x, genfloatf y)
1140 // int isgreater (double x, double y)
1141 // longn isgreater (doublen x, doublen y)
1142 template <typename T,
1143  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1146  __sycl_std::__invoke_FOrdGreaterThan<detail::rel_ret_t<T>>(x, y));
1147 }
1148 
1149 // int isgreaterequal (half x, half y)
1150 // shortn isgreaterequal (halfn x, halfn y)
1151 // igeninteger32bit isgreaterequal (genfloatf x, genfloatf y)
1152 // int isgreaterequal (double x, double y)
1153 // longn isgreaterequal (doublen x, doublen y)
1154 template <typename T,
1155  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1158  __sycl_std::__invoke_FOrdGreaterThanEqual<detail::rel_ret_t<T>>(x, y));
1159 }
1160 
1161 // int isless (half x, half y)
1162 // shortn isless (halfn x, halfn y)
1163 // igeninteger32bit isless (genfloatf x, genfloatf y)
1164 // int isless (long x, long y)
1165 // longn isless (doublen x, doublen y)
1166 template <typename T,
1167  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1170  __sycl_std::__invoke_FOrdLessThan<detail::rel_ret_t<T>>(x, y));
1171 }
1172 
1173 // int islessequal (half x, half y)
1174 // shortn islessequal (halfn x, halfn y)
1175 // igeninteger32bit islessequal (genfloatf x, genfloatf y)
1176 // int islessequal (double x, double y)
1177 // longn islessequal (doublen x, doublen y)
1178 template <typename T,
1179  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1182  __sycl_std::__invoke_FOrdLessThanEqual<detail::rel_ret_t<T>>(x, y));
1183 }
1184 
1185 // int islessgreater (half x, half y)
1186 // shortn islessgreater (halfn x, halfn y)
1187 // igeninteger32bit islessgreater (genfloatf x, genfloatf y)
1188 // int islessgreater (double x, double y)
1189 // longn islessgreater (doublen x, doublen y)
1190 template <typename T,
1191  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1194  __sycl_std::__invoke_FOrdNotEqual<detail::rel_ret_t<T>>(x, y));
1195 }
1196 
1197 // int isfinite (half x)
1198 // shortn isfinite (halfn x)
1199 // igeninteger32bit isfinite (genfloatf x)
1200 // int isfinite (double x)
1201 // longn isfinite (doublen x)
1202 template <typename T,
1203  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1206  __sycl_std::__invoke_IsFinite<detail::rel_ret_t<T>>(x));
1207 }
1208 
1209 // int isinf (half x)
1210 // shortn isinf (halfn x)
1211 // igeninteger32bit isinf (genfloatf x)
1212 // int isinf (double x)
1213 // longn isinf (doublen x)
1214 template <typename T,
1215  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1218  __sycl_std::__invoke_IsInf<detail::rel_ret_t<T>>(x));
1219 }
1220 
1221 // int isnan (half x)
1222 // shortn isnan (halfn x)
1223 // igeninteger32bit isnan (genfloatf x)
1224 // int isnan (double x)
1225 // longn isnan (doublen x)
1226 template <typename T,
1227  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1230  __sycl_std::__invoke_IsNan<detail::rel_ret_t<T>>(x));
1231 }
1232 
1233 // int isnormal (half x)
1234 // shortn isnormal (halfn x)
1235 // igeninteger32bit isnormal (genfloatf x)
1236 // int isnormal (double x)
1237 // longn isnormal (doublen x)
1238 template <typename T,
1239  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1242  __sycl_std::__invoke_IsNormal<detail::rel_ret_t<T>>(x));
1243 }
1244 
1245 // int isordered (half x)
1246 // shortn isordered (halfn x, halfn y)
1247 // igeninteger32bit isordered (genfloatf x, genfloatf y)
1248 // int isordered (double x, double y)
1249 // longn isordered (doublen x, doublen y)
1250 template <typename T,
1251  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1254  __sycl_std::__invoke_Ordered<detail::rel_ret_t<T>>(x, y));
1255 }
1256 
1257 // int isunordered (half x, half y)
1258 // shortn isunordered (halfn x, halfn y)
1259 // igeninteger32bit isunordered (genfloatf x, genfloatf y)
1260 // int isunordered (double x, double y)
1261 // longn isunordered (doublen x, doublen y)
1262 template <typename T,
1263  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1266  __sycl_std::__invoke_Unordered<detail::rel_ret_t<T>>(x, y));
1267 }
1268 
1269 // int signbit (half x)
1270 // shortn signbit (halfn x)
1271 // igeninteger32bit signbit (genfloatf x)
1272 // int signbit (double)
1273 // longn signbit (doublen x)
1274 template <typename T,
1275  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1278  __sycl_std::__invoke_SignBitSet<detail::rel_ret_t<T>>(x));
1279 }
1280 
1281 // int any (sigeninteger x)
1282 template <typename T>
1284  return detail::Boolean<1>(int(detail::msbIsSet(x)));
1285 }
1286 
1287 // int any (vigeninteger x)
1288 template <typename T>
1291  __sycl_std::__invoke_Any<detail::rel_sign_bit_test_ret_t<T>>(
1293 }
1294 
1295 // int all (sigeninteger x)
1296 template <typename T>
1298  return detail::Boolean<1>(int(detail::msbIsSet(x)));
1299 }
1300 
1301 // int all (vigeninteger x)
1302 template <typename T>
1305  __sycl_std::__invoke_All<detail::rel_sign_bit_test_ret_t<T>>(
1307 }
1308 
1309 // gentype bitselect (gentype a, gentype b, gentype c)
1310 template <typename T>
1312  T c) __NOEXC {
1313  return __sycl_std::__invoke_bitselect<T>(a, b, c);
1314 }
1315 
1316 // sgentype select (sgentype a, sgentype b, bool c)
1317 template <typename T>
1319  bool c) __NOEXC {
1320  return __sycl_std::__invoke_select<T>(a, b, static_cast<int>(c));
1321 }
1322 
1323 // geninteger select (geninteger a, geninteger b, igeninteger c)
1324 template <typename T, typename T2>
1326  detail::is_geninteger<T>::value && detail::is_igeninteger<T2>::value, T>
1327 select(T a, T b, T2 c) __NOEXC {
1328  detail::check_vector_size<T, T2>();
1329  return __sycl_std::__invoke_select<T>(a, b, c);
1330 }
1331 
1332 // geninteger select (geninteger a, geninteger b, ugeninteger c)
1333 template <typename T, typename T2>
1335  detail::is_geninteger<T>::value && detail::is_ugeninteger<T2>::value, T>
1336 select(T a, T b, T2 c) __NOEXC {
1337  detail::check_vector_size<T, T2>();
1338  return __sycl_std::__invoke_select<T>(a, b, c);
1339 }
1340 
1341 // genfloatf select (genfloatf a, genfloatf b, genint c)
1342 template <typename T, typename T2>
1344  detail::is_genfloatf<T>::value && detail::is_genint<T2>::value, T>
1345 select(T a, T b, T2 c) __NOEXC {
1346  detail::check_vector_size<T, T2>();
1347  return __sycl_std::__invoke_select<T>(a, b, c);
1348 }
1349 
1350 // genfloatf select (genfloatf a, genfloatf b, ugenint c)
1351 template <typename T, typename T2>
1353  detail::is_genfloatf<T>::value && detail::is_ugenint<T2>::value, T>
1354 select(T a, T b, T2 c) __NOEXC {
1355  detail::check_vector_size<T, T2>();
1356  return __sycl_std::__invoke_select<T>(a, b, c);
1357 }
1358 
1359 // genfloatd select (genfloatd a, genfloatd b, igeninteger64 c)
1360 template <typename T, typename T2>
1362  detail::is_genfloatd<T>::value && detail::is_igeninteger64bit<T2>::value, T>
1363 select(T a, T b, T2 c) __NOEXC {
1364  detail::check_vector_size<T, T2>();
1365  return __sycl_std::__invoke_select<T>(a, b, c);
1366 }
1367 
1368 // genfloatd select (genfloatd a, genfloatd b, ugeninteger64 c)
1369 template <typename T, typename T2>
1371  detail::is_genfloatd<T>::value && detail::is_ugeninteger64bit<T2>::value, T>
1372 select(T a, T b, T2 c) __NOEXC {
1373  detail::check_vector_size<T, T2>();
1374  return __sycl_std::__invoke_select<T>(a, b, c);
1375 }
1376 
1377 // genfloath select (genfloath a, genfloath b, igeninteger16 c)
1378 template <typename T, typename T2>
1380  detail::is_genfloath<T>::value && detail::is_igeninteger16bit<T2>::value, T>
1381 select(T a, T b, T2 c) __NOEXC {
1382  detail::check_vector_size<T, T2>();
1383  return __sycl_std::__invoke_select<T>(a, b, c);
1384 }
1385 
1386 // genfloath select (genfloath a, genfloath b, ugeninteger16 c)
1387 template <typename T, typename T2>
1389  detail::is_genfloath<T>::value && detail::is_ugeninteger16bit<T2>::value, T>
1390 select(T a, T b, T2 c) __NOEXC {
1391  detail::check_vector_size<T, T2>();
1392  return __sycl_std::__invoke_select<T>(a, b, c);
1393 }
1394 
1395 namespace native {
1396 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1397 // genfloatf cos (genfloatf x)
1398 template <typename T>
1400  return __sycl_std::__invoke_native_cos<T>(x);
1401 }
1402 
1403 // genfloatf divide (genfloatf x, genfloatf y)
1404 template <typename T>
1406  T y) __NOEXC {
1407  return __sycl_std::__invoke_native_divide<T>(x, y);
1408 }
1409 
1410 // genfloatf exp (genfloatf x)
1411 template <typename T>
1413  return __sycl_std::__invoke_native_exp<T>(x);
1414 }
1415 
1416 // genfloatf exp2 (genfloatf x)
1417 template <typename T>
1419  return __sycl_std::__invoke_native_exp2<T>(x);
1420 }
1421 
1422 // genfloatf exp10 (genfloatf x)
1423 template <typename T>
1425  return __sycl_std::__invoke_native_exp10<T>(x);
1426 }
1427 
1428 // genfloatf log (genfloatf x)
1429 template <typename T>
1431  return __sycl_std::__invoke_native_log<T>(x);
1432 }
1433 
1434 // genfloatf log2 (genfloatf x)
1435 template <typename T>
1437  return __sycl_std::__invoke_native_log2<T>(x);
1438 }
1439 
1440 // genfloatf log10 (genfloatf x)
1441 template <typename T>
1443  return __sycl_std::__invoke_native_log10<T>(x);
1444 }
1445 
1446 // genfloatf powr (genfloatf x, genfloatf y)
1447 template <typename T>
1449  return __sycl_std::__invoke_native_powr<T>(x, y);
1450 }
1451 
1452 // genfloatf recip (genfloatf x)
1453 template <typename T>
1455  return __sycl_std::__invoke_native_recip<T>(x);
1456 }
1457 
1458 // genfloatf rsqrt (genfloatf x)
1459 template <typename T>
1461  return __sycl_std::__invoke_native_rsqrt<T>(x);
1462 }
1463 
1464 // genfloatf sin (genfloatf x)
1465 template <typename T>
1467  return __sycl_std::__invoke_native_sin<T>(x);
1468 }
1469 
1470 // genfloatf sqrt (genfloatf x)
1471 template <typename T>
1473  return __sycl_std::__invoke_native_sqrt<T>(x);
1474 }
1475 
1476 // genfloatf tan (genfloatf x)
1477 template <typename T>
1479  return __sycl_std::__invoke_native_tan<T>(x);
1480 }
1481 
1482 } // namespace native
1483 namespace half_precision {
1484 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1485 // genfloatf cos (genfloatf x)
1486 template <typename T>
1488  return __sycl_std::__invoke_half_cos<T>(x);
1489 }
1490 
1491 // genfloatf divide (genfloatf x, genfloatf y)
1492 template <typename T>
1494  T y) __NOEXC {
1495  return __sycl_std::__invoke_half_divide<T>(x, y);
1496 }
1497 
1498 // genfloatf exp (genfloatf x)
1499 template <typename T>
1501  return __sycl_std::__invoke_half_exp<T>(x);
1502 }
1503 
1504 // genfloatf exp2 (genfloatf x)
1505 template <typename T>
1507  return __sycl_std::__invoke_half_exp2<T>(x);
1508 }
1509 
1510 // genfloatf exp10 (genfloatf x)
1511 template <typename T>
1513  return __sycl_std::__invoke_half_exp10<T>(x);
1514 }
1515 
1516 // genfloatf log (genfloatf x)
1517 template <typename T>
1519  return __sycl_std::__invoke_half_log<T>(x);
1520 }
1521 
1522 // genfloatf log2 (genfloatf x)
1523 template <typename T>
1525  return __sycl_std::__invoke_half_log2<T>(x);
1526 }
1527 
1528 // genfloatf log10 (genfloatf x)
1529 template <typename T>
1531  return __sycl_std::__invoke_half_log10<T>(x);
1532 }
1533 
1534 // genfloatf powr (genfloatf x, genfloatf y)
1535 template <typename T>
1537  return __sycl_std::__invoke_half_powr<T>(x, y);
1538 }
1539 
1540 // genfloatf recip (genfloatf x)
1541 template <typename T>
1543  return __sycl_std::__invoke_half_recip<T>(x);
1544 }
1545 
1546 // genfloatf rsqrt (genfloatf x)
1547 template <typename T>
1549  return __sycl_std::__invoke_half_rsqrt<T>(x);
1550 }
1551 
1552 // genfloatf sin (genfloatf x)
1553 template <typename T>
1555  return __sycl_std::__invoke_half_sin<T>(x);
1556 }
1557 
1558 // genfloatf sqrt (genfloatf x)
1559 template <typename T>
1561  return __sycl_std::__invoke_half_sqrt<T>(x);
1562 }
1563 
1564 // genfloatf tan (genfloatf x)
1565 template <typename T>
1567  return __sycl_std::__invoke_half_tan<T>(x);
1568 }
1569 
1570 } // namespace half_precision
1571 
1572 #ifdef __FAST_MATH__
1573 /* ----------------- -ffast-math functions. ---------------------------------*/
1574 // genfloatf cos (genfloatf x)
1575 template <typename T>
1576 detail::enable_if_t<detail::is_genfloatf<T>::value, T> cos(T x) __NOEXC {
1577  return native::cos(x);
1578 }
1579 
1580 // genfloatf exp (genfloatf x)
1581 template <typename T>
1582 detail::enable_if_t<detail::is_genfloatf<T>::value, T> exp(T x) __NOEXC {
1583  return native::exp(x);
1584 }
1585 
1586 // genfloatf exp2 (genfloatf x)
1587 template <typename T>
1588 detail::enable_if_t<detail::is_genfloatf<T>::value, T> exp2(T x) __NOEXC {
1589  return native::exp2(x);
1590 }
1591 
1592 // genfloatf exp10 (genfloatf x)
1593 template <typename T>
1594 detail::enable_if_t<detail::is_genfloatf<T>::value, T> exp10(T x) __NOEXC {
1595  return native::exp10(x);
1596 }
1597 
1598 // genfloatf log(genfloatf x)
1599 template <typename T>
1600 detail::enable_if_t<detail::is_genfloatf<T>::value, T> log(T x) __NOEXC {
1601  return native::log(x);
1602 }
1603 
1604 // genfloatf log2 (genfloatf x)
1605 template <typename T>
1606 detail::enable_if_t<detail::is_genfloatf<T>::value, T> log2(T x) __NOEXC {
1607  return native::log2(x);
1608 }
1609 
1610 // genfloatf log10 (genfloatf x)
1611 template <typename T>
1612 detail::enable_if_t<detail::is_genfloatf<T>::value, T> log10(T x) __NOEXC {
1613  return native::log10(x);
1614 }
1615 
1616 // genfloatf powr (genfloatf x)
1617 template <typename T>
1618 detail::enable_if_t<detail::is_genfloatf<T>::value, T> powr(T x, T y) __NOEXC {
1619  return native::powr(x, y);
1620 }
1621 
1622 // genfloatf rsqrt (genfloatf x)
1623 template <typename T>
1624 detail::enable_if_t<detail::is_genfloatf<T>::value, T> rsqrt(T x) __NOEXC {
1625  return native::rsqrt(x);
1626 }
1627 
1628 // genfloatf sin (genfloatf x)
1629 template <typename T>
1630 detail::enable_if_t<detail::is_genfloatf<T>::value, T> sin(T x) __NOEXC {
1631  return native::sin(x);
1632 }
1633 
1634 // genfloatf sqrt (genfloatf x)
1635 template <typename T>
1636 detail::enable_if_t<detail::is_genfloatf<T>::value, T> sqrt(T x) __NOEXC {
1637  return native::sqrt(x);
1638 }
1639 
1640 // genfloatf tan (genfloatf x)
1641 template <typename T>
1642 detail::enable_if_t<detail::is_genfloatf<T>::value, T> tan(T x) __NOEXC {
1643  return native::tan(x);
1644 }
1645 
1646 #endif // __FAST_MATH__
1647 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1648 } // namespace sycl
1649 
1650 #ifdef __SYCL_DEVICE_ONLY__
1651 extern "C" {
1652 extern SYCL_EXTERNAL int abs(int x);
1653 extern SYCL_EXTERNAL long int labs(long int x);
1654 extern SYCL_EXTERNAL long long int llabs(long long int x);
1655 
1656 extern SYCL_EXTERNAL div_t div(int x, int y);
1657 extern SYCL_EXTERNAL ldiv_t ldiv(long int x, long int y);
1658 extern SYCL_EXTERNAL lldiv_t lldiv(long long int x, long long int y);
1659 extern SYCL_EXTERNAL float scalbnf(float x, int n);
1660 extern SYCL_EXTERNAL double scalbn(double x, int n);
1661 extern SYCL_EXTERNAL float logf(float x);
1662 extern SYCL_EXTERNAL double log(double x);
1663 extern SYCL_EXTERNAL float expf(float x);
1664 extern SYCL_EXTERNAL double exp(double x);
1665 extern SYCL_EXTERNAL float log10f(float x);
1666 extern SYCL_EXTERNAL double log10(double x);
1667 extern SYCL_EXTERNAL float modff(float x, float *intpart);
1668 extern SYCL_EXTERNAL double modf(double x, double *intpart);
1669 extern SYCL_EXTERNAL float exp2f(float x);
1670 extern SYCL_EXTERNAL double exp2(double x);
1671 extern SYCL_EXTERNAL float expm1f(float x);
1672 extern SYCL_EXTERNAL double expm1(double x);
1673 extern SYCL_EXTERNAL int ilogbf(float x);
1674 extern SYCL_EXTERNAL int ilogb(double x);
1675 extern SYCL_EXTERNAL float log1pf(float x);
1676 extern SYCL_EXTERNAL double log1p(double x);
1677 extern SYCL_EXTERNAL float log2f(float x);
1678 extern SYCL_EXTERNAL double log2(double x);
1679 extern SYCL_EXTERNAL float logbf(float x);
1680 extern SYCL_EXTERNAL double logb(double x);
1681 extern SYCL_EXTERNAL float sqrtf(float x);
1682 extern SYCL_EXTERNAL double sqrt(double x);
1683 extern SYCL_EXTERNAL float cbrtf(float x);
1684 extern SYCL_EXTERNAL double cbrt(double x);
1685 extern SYCL_EXTERNAL float erff(float x);
1686 extern SYCL_EXTERNAL double erf(double x);
1687 extern SYCL_EXTERNAL float erfcf(float x);
1688 extern SYCL_EXTERNAL double erfc(double x);
1689 extern SYCL_EXTERNAL float tgammaf(float x);
1690 extern SYCL_EXTERNAL double tgamma(double x);
1691 extern SYCL_EXTERNAL float lgammaf(float x);
1692 extern SYCL_EXTERNAL double lgamma(double x);
1693 extern SYCL_EXTERNAL float fmodf(float x, float y);
1694 extern SYCL_EXTERNAL double fmod(double x, double y);
1695 extern SYCL_EXTERNAL float remainderf(float x, float y);
1696 extern SYCL_EXTERNAL double remainder(double x, double y);
1697 extern SYCL_EXTERNAL float remquof(float x, float y, int *q);
1698 extern SYCL_EXTERNAL double remquo(double x, double y, int *q);
1699 extern SYCL_EXTERNAL float nextafterf(float x, float y);
1700 extern SYCL_EXTERNAL double nextafter(double x, double y);
1701 extern SYCL_EXTERNAL float fdimf(float x, float y);
1702 extern SYCL_EXTERNAL double fdim(double x, double y);
1703 extern SYCL_EXTERNAL float fmaf(float x, float y, float z);
1704 extern SYCL_EXTERNAL double fma(double x, double y, double z);
1705 extern SYCL_EXTERNAL float sinf(float x);
1706 extern SYCL_EXTERNAL double sin(double x);
1707 extern SYCL_EXTERNAL float cosf(float x);
1708 extern SYCL_EXTERNAL double cos(double x);
1709 extern SYCL_EXTERNAL float tanf(float x);
1710 extern SYCL_EXTERNAL double tan(double x);
1711 extern SYCL_EXTERNAL float asinf(float x);
1712 extern SYCL_EXTERNAL double asin(double x);
1713 extern SYCL_EXTERNAL float acosf(float x);
1714 extern SYCL_EXTERNAL double acos(double x);
1715 extern SYCL_EXTERNAL float atanf(float x);
1716 extern SYCL_EXTERNAL double atan(double x);
1717 extern SYCL_EXTERNAL float powf(float x, float y);
1718 extern SYCL_EXTERNAL double pow(double x, double y);
1719 extern SYCL_EXTERNAL float atan2f(float x, float y);
1720 extern SYCL_EXTERNAL double atan2(double x, double y);
1721 
1722 extern SYCL_EXTERNAL float sinhf(float x);
1723 extern SYCL_EXTERNAL double sinh(double x);
1724 extern SYCL_EXTERNAL float coshf(float x);
1725 extern SYCL_EXTERNAL double cosh(double x);
1726 extern SYCL_EXTERNAL float tanhf(float x);
1727 extern SYCL_EXTERNAL double tanh(double x);
1728 extern SYCL_EXTERNAL float asinhf(float x);
1729 extern SYCL_EXTERNAL double asinh(double x);
1730 extern SYCL_EXTERNAL float acoshf(float x);
1731 extern SYCL_EXTERNAL double acosh(double x);
1732 extern SYCL_EXTERNAL float atanhf(float x);
1733 extern SYCL_EXTERNAL double atanh(double x);
1734 extern SYCL_EXTERNAL double frexp(double x, int *exp);
1735 extern SYCL_EXTERNAL double ldexp(double x, int exp);
1736 extern SYCL_EXTERNAL double hypot(double x, double y);
1737 
1738 extern SYCL_EXTERNAL void *memcpy(void *dest, const void *src, size_t n);
1739 extern SYCL_EXTERNAL void *memset(void *dest, int c, size_t n);
1740 extern SYCL_EXTERNAL int memcmp(const void *s1, const void *s2, size_t n);
1741 extern SYCL_EXTERNAL long long int __imf_llmax(long long int x,
1742  long long int y);
1743 extern SYCL_EXTERNAL long long int __imf_llmin(long long int x,
1744  long long int y);
1745 extern SYCL_EXTERNAL unsigned long long int
1746 __imf_ullmax(unsigned long long int x, unsigned long long int y);
1747 extern SYCL_EXTERNAL unsigned long long int
1748 __imf_ullmin(unsigned long long int x, unsigned long long int y);
1749 extern SYCL_EXTERNAL unsigned int __imf_umax(unsigned int x, unsigned int y);
1750 extern SYCL_EXTERNAL unsigned int __imf_umin(unsigned int x, unsigned int y);
1751 extern SYCL_EXTERNAL unsigned int __imf_brev(unsigned int x);
1752 extern SYCL_EXTERNAL unsigned long long int
1753 __imf_brevll(unsigned long long int x);
1754 extern SYCL_EXTERNAL unsigned int
1755 __imf_byte_perm(unsigned int x, unsigned int y, unsigned int s);
1756 extern SYCL_EXTERNAL int __imf_ffs(int x);
1757 extern SYCL_EXTERNAL int __imf_ffsll(long long int x);
1758 extern SYCL_EXTERNAL int __imf_clz(int x);
1759 extern SYCL_EXTERNAL int __imf_clzll(long long int x);
1760 extern SYCL_EXTERNAL int __imf_popc(unsigned int x);
1761 extern SYCL_EXTERNAL int __imf_popcll(unsigned long long int x);
1762 extern SYCL_EXTERNAL unsigned int __imf_sad(int x, int y, unsigned int z);
1763 extern SYCL_EXTERNAL unsigned int __imf_usad(unsigned int x, unsigned int y,
1764  unsigned int z);
1765 extern SYCL_EXTERNAL int __imf_rhadd(int x, int y);
1766 extern SYCL_EXTERNAL unsigned int __imf_urhadd(unsigned int x, unsigned int y);
1767 extern SYCL_EXTERNAL unsigned int __imf_uhadd(unsigned int x, unsigned int y);
1768 extern SYCL_EXTERNAL int __imf_mul24(int x, int y);
1769 extern SYCL_EXTERNAL unsigned int __imf_umul24(unsigned int x, unsigned int y);
1770 extern SYCL_EXTERNAL int __imf_mulhi(int x, int y);
1771 extern SYCL_EXTERNAL unsigned int __imf_umulhi(unsigned int x, unsigned int y);
1772 extern SYCL_EXTERNAL long long int __imf_mul64hi(long long int x,
1773  long long int y);
1774 extern SYCL_EXTERNAL unsigned long long int
1775 __imf_umul64hi(unsigned long long int x, unsigned long long int y);
1776 extern SYCL_EXTERNAL float __imf_saturatef(float x);
1777 extern SYCL_EXTERNAL float __imf_fmaf(float x, float y, float z);
1778 extern SYCL_EXTERNAL float __imf_fabsf(float x);
1779 extern SYCL_EXTERNAL float __imf_floorf(float x);
1780 extern SYCL_EXTERNAL float __imf_ceilf(float x);
1781 extern SYCL_EXTERNAL float __imf_truncf(float x);
1782 extern SYCL_EXTERNAL float __imf_rintf(float x);
1783 extern SYCL_EXTERNAL float __imf_nearbyintf(float x);
1784 extern SYCL_EXTERNAL float __imf_sqrtf(float x);
1785 extern SYCL_EXTERNAL float __imf_rsqrtf(float x);
1786 extern SYCL_EXTERNAL float __imf_invf(float x);
1787 extern SYCL_EXTERNAL float __imf_fmaxf(float x, float y);
1788 extern SYCL_EXTERNAL float __imf_fminf(float x, float y);
1789 extern SYCL_EXTERNAL float __imf_copysignf(float x, float y);
1790 extern SYCL_EXTERNAL int __imf_float2int_rd(float x);
1791 extern SYCL_EXTERNAL int __imf_float2int_rn(float x);
1792 extern SYCL_EXTERNAL int __imf_float2int_ru(float x);
1793 extern SYCL_EXTERNAL int __imf_float2int_rz(float x);
1794 extern SYCL_EXTERNAL unsigned int __imf_float2uint_rd(float x);
1795 extern SYCL_EXTERNAL unsigned int __imf_float2uint_rn(float x);
1796 extern SYCL_EXTERNAL unsigned int __imf_float2uint_ru(float x);
1797 extern SYCL_EXTERNAL unsigned int __imf_float2uint_rz(float x);
1798 extern SYCL_EXTERNAL long long int __imf_float2ll_rd(float x);
1799 extern SYCL_EXTERNAL long long int __imf_float2ll_rn(float x);
1800 extern SYCL_EXTERNAL long long int __imf_float2ll_ru(float x);
1801 extern SYCL_EXTERNAL long long int __imf_float2ll_rz(float x);
1802 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rd(float x);
1803 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rn(float x);
1804 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_ru(float x);
1805 extern SYCL_EXTERNAL unsigned long long int __imf_float2ull_rz(float x);
1806 extern SYCL_EXTERNAL int __imf_float_as_int(float x);
1807 extern SYCL_EXTERNAL unsigned int __imf_float_as_uint(float x);
1808 extern SYCL_EXTERNAL float __imf_int2float_rd(int x);
1809 extern SYCL_EXTERNAL float __imf_int2float_rn(int x);
1810 extern SYCL_EXTERNAL float __imf_int2float_ru(int x);
1811 extern SYCL_EXTERNAL float __imf_int2float_rz(int x);
1812 extern SYCL_EXTERNAL float __imf_int_as_float(int x);
1813 extern SYCL_EXTERNAL float __imf_ll2float_rd(long long int x);
1814 extern SYCL_EXTERNAL float __imf_ll2float_rn(long long int x);
1815 extern SYCL_EXTERNAL float __imf_ll2float_ru(long long int x);
1816 extern SYCL_EXTERNAL float __imf_ll2float_rz(long long int x);
1817 extern SYCL_EXTERNAL float __imf_uint2float_rd(unsigned int x);
1818 extern SYCL_EXTERNAL float __imf_uint2float_rn(unsigned int x);
1819 extern SYCL_EXTERNAL float __imf_uint2float_ru(unsigned int x);
1820 extern SYCL_EXTERNAL float __imf_uint2float_rz(unsigned int x);
1821 extern SYCL_EXTERNAL float __imf_uint_as_float(unsigned int x);
1822 extern SYCL_EXTERNAL float __imf_ull2float_rd(unsigned long long int x);
1823 extern SYCL_EXTERNAL float __imf_ull2float_rn(unsigned long long int x);
1824 extern SYCL_EXTERNAL float __imf_ull2float_ru(unsigned long long int x);
1825 extern SYCL_EXTERNAL float __imf_ull2float_rz(unsigned long long int x);
1826 
1827 extern SYCL_EXTERNAL _Float16 __imf_fmaf16(_Float16 x, _Float16 y, _Float16 z);
1828 extern SYCL_EXTERNAL _Float16 __imf_fabsf16(_Float16 x);
1829 extern SYCL_EXTERNAL _Float16 __imf_floorf16(_Float16 x);
1830 extern SYCL_EXTERNAL _Float16 __imf_ceilf16(_Float16 x);
1831 extern SYCL_EXTERNAL _Float16 __imf_truncf16(_Float16 x);
1832 extern SYCL_EXTERNAL _Float16 __imf_rintf16(_Float16 x);
1833 extern SYCL_EXTERNAL _Float16 __imf_nearbyintf16(_Float16 x);
1834 extern SYCL_EXTERNAL _Float16 __imf_sqrtf16(_Float16 x);
1835 extern SYCL_EXTERNAL _Float16 __imf_rsqrtf16(_Float16 x);
1836 extern SYCL_EXTERNAL _Float16 __imf_invf16(_Float16 x);
1837 extern SYCL_EXTERNAL _Float16 __imf_fmaxf16(_Float16 x, _Float16 y);
1838 extern SYCL_EXTERNAL _Float16 __imf_fminf16(_Float16 x, _Float16 y);
1839 extern SYCL_EXTERNAL _Float16 __imf_copysignf16(_Float16 x, _Float16 y);
1840 extern SYCL_EXTERNAL float __imf_half2float(_Float16 x);
1841 extern SYCL_EXTERNAL double __imf_fma(double x, double y, double z);
1842 extern SYCL_EXTERNAL double __imf_fabs(double x);
1843 extern SYCL_EXTERNAL double __imf_floor(double x);
1844 extern SYCL_EXTERNAL double __imf_ceil(double x);
1845 extern SYCL_EXTERNAL double __imf_trunc(double x);
1846 extern SYCL_EXTERNAL double __imf_rint(double x);
1847 extern SYCL_EXTERNAL double __imf_nearbyint(double x);
1848 extern SYCL_EXTERNAL double __imf_sqrt(double x);
1849 extern SYCL_EXTERNAL double __imf_rsqrt(double x);
1850 extern SYCL_EXTERNAL double __imf_inv(double x);
1851 extern SYCL_EXTERNAL double __imf_fmax(double x, double y);
1852 extern SYCL_EXTERNAL double __imf_fmin(double x, double y);
1853 extern SYCL_EXTERNAL double __imf_copysign(double x, double y);
1854 extern SYCL_EXTERNAL float __imf_double2float_rd(double x);
1855 extern SYCL_EXTERNAL float __imf_double2float_rn(double x);
1856 extern SYCL_EXTERNAL float __imf_double2float_ru(double x);
1857 extern SYCL_EXTERNAL float __imf_double2float_rz(double x);
1858 extern SYCL_EXTERNAL int __imf_double2hiint(double x);
1859 extern SYCL_EXTERNAL int __imf_double2loint(double x);
1860 extern SYCL_EXTERNAL int __imf_double2int_rd(double x);
1861 extern SYCL_EXTERNAL int __imf_double2int_rn(double x);
1862 extern SYCL_EXTERNAL int __imf_double2int_ru(double x);
1863 extern SYCL_EXTERNAL int __imf_double2int_rz(double x);
1864 extern SYCL_EXTERNAL double __imf_int2double_rn(int x);
1865 extern SYCL_EXTERNAL unsigned int __imf_double2uint_rd(double x);
1866 extern SYCL_EXTERNAL unsigned int __imf_double2uint_rn(double x);
1867 extern SYCL_EXTERNAL unsigned int __imf_double2uint_ru(double x);
1868 extern SYCL_EXTERNAL unsigned int __imf_double2uint_rz(double x);
1869 extern SYCL_EXTERNAL long long int __imf_double2ll_rd(double x);
1870 extern SYCL_EXTERNAL long long int __imf_double2ll_rn(double x);
1871 extern SYCL_EXTERNAL long long int __imf_double2ll_ru(double x);
1872 extern SYCL_EXTERNAL long long int __imf_double2ll_rz(double x);
1873 extern SYCL_EXTERNAL double __imf_ll2double_rd(long long int x);
1874 extern SYCL_EXTERNAL double __imf_ll2double_rn(long long int x);
1875 extern SYCL_EXTERNAL double __imf_ll2double_ru(long long int x);
1876 extern SYCL_EXTERNAL double __imf_ll2double_rz(long long int x);
1877 extern SYCL_EXTERNAL double __imf_ull2double_rd(unsigned long long int x);
1878 extern SYCL_EXTERNAL double __imf_ull2double_rn(unsigned long long int x);
1879 extern SYCL_EXTERNAL double __imf_ull2double_ru(unsigned long long int x);
1880 extern SYCL_EXTERNAL double __imf_ull2double_rz(unsigned long long int x);
1881 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rd(double x);
1882 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rn(double x);
1883 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_ru(double x);
1884 extern SYCL_EXTERNAL unsigned long long int __imf_double2ull_rz(double x);
1885 extern SYCL_EXTERNAL long long int __imf_double_as_longlong(double x);
1886 extern SYCL_EXTERNAL double __imf_longlong_as_double(long long int x);
1887 extern SYCL_EXTERNAL double __imf_uint2double_rd(unsigned int x);
1888 extern SYCL_EXTERNAL double __imf_uint2double_rn(unsigned int x);
1889 extern SYCL_EXTERNAL double __imf_uint2double_ru(unsigned int x);
1890 extern SYCL_EXTERNAL double __imf_uint2double_rz(unsigned int x);
1891 extern SYCL_EXTERNAL double __imf_hiloint2double(int hi, int lo);
1892 
1893 extern SYCL_EXTERNAL unsigned int __imf_vabs2(unsigned int x);
1894 extern SYCL_EXTERNAL unsigned int __imf_vabs4(unsigned int x);
1895 extern SYCL_EXTERNAL unsigned int __imf_vabsss2(unsigned int x);
1896 extern SYCL_EXTERNAL unsigned int __imf_vabsss4(unsigned int x);
1897 extern SYCL_EXTERNAL unsigned int __imf_vneg2(unsigned int x);
1898 extern SYCL_EXTERNAL unsigned int __imf_vneg4(unsigned int x);
1899 extern SYCL_EXTERNAL unsigned int __imf_vnegss2(unsigned int x);
1900 extern SYCL_EXTERNAL unsigned int __imf_vnegss4(unsigned int x);
1901 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffs2(unsigned int x,
1902  unsigned int y);
1903 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffs4(unsigned int x,
1904  unsigned int y);
1905 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffu2(unsigned int x,
1906  unsigned int y);
1907 extern SYCL_EXTERNAL unsigned int __imf_vabsdiffu4(unsigned int x,
1908  unsigned int y);
1909 extern SYCL_EXTERNAL unsigned int __imf_vadd2(unsigned int x, unsigned int y);
1910 extern SYCL_EXTERNAL unsigned int __imf_vadd4(unsigned int x, unsigned int y);
1911 extern SYCL_EXTERNAL unsigned int __imf_vaddss2(unsigned int x, unsigned int y);
1912 extern SYCL_EXTERNAL unsigned int __imf_vaddss4(unsigned int x, unsigned int y);
1913 extern SYCL_EXTERNAL unsigned int __imf_vaddus2(unsigned int x, unsigned int y);
1914 extern SYCL_EXTERNAL unsigned int __imf_vaddus4(unsigned int x, unsigned int y);
1915 extern SYCL_EXTERNAL unsigned int __imf_vsub2(unsigned int x, unsigned int y);
1916 extern SYCL_EXTERNAL unsigned int __imf_vsub4(unsigned int x, unsigned int y);
1917 extern SYCL_EXTERNAL unsigned int __imf_vsubss2(unsigned int x, unsigned int y);
1918 extern SYCL_EXTERNAL unsigned int __imf_vsubss4(unsigned int x, unsigned int y);
1919 extern SYCL_EXTERNAL unsigned int __imf_vsubus2(unsigned int x, unsigned int y);
1920 extern SYCL_EXTERNAL unsigned int __imf_vsubus4(unsigned int x, unsigned int y);
1921 extern SYCL_EXTERNAL unsigned int __imf_vavgs2(unsigned int x, unsigned int y);
1922 extern SYCL_EXTERNAL unsigned int __imf_vavgs4(unsigned int x, unsigned int y);
1923 extern SYCL_EXTERNAL unsigned int __imf_vavgu2(unsigned int x, unsigned int y);
1924 extern SYCL_EXTERNAL unsigned int __imf_vavgu4(unsigned int x, unsigned int y);
1925 extern SYCL_EXTERNAL unsigned int __imf_vhaddu2(unsigned int x, unsigned int y);
1926 extern SYCL_EXTERNAL unsigned int __imf_vhaddu4(unsigned int x, unsigned int y);
1927 extern SYCL_EXTERNAL unsigned int __imf_vcmpeq2(unsigned int x, unsigned int y);
1928 extern SYCL_EXTERNAL unsigned int __imf_vcmpeq4(unsigned int x, unsigned int y);
1929 extern SYCL_EXTERNAL unsigned int __imf_vcmpne2(unsigned int x, unsigned int y);
1930 extern SYCL_EXTERNAL unsigned int __imf_vcmpne4(unsigned int x, unsigned int y);
1931 extern SYCL_EXTERNAL unsigned int __imf_vcmpges2(unsigned int x,
1932  unsigned int y);
1933 extern SYCL_EXTERNAL unsigned int __imf_vcmpges4(unsigned int x,
1934  unsigned int y);
1935 extern SYCL_EXTERNAL unsigned int __imf_vcmpgeu2(unsigned int x,
1936  unsigned int y);
1937 extern SYCL_EXTERNAL unsigned int __imf_vcmpgeu4(unsigned int x,
1938  unsigned int y);
1939 extern SYCL_EXTERNAL unsigned int __imf_vcmpgts2(unsigned int x,
1940  unsigned int y);
1941 extern SYCL_EXTERNAL unsigned int __imf_vcmpgts4(unsigned int x,
1942  unsigned int y);
1943 extern SYCL_EXTERNAL unsigned int __imf_vcmpgtu2(unsigned int x,
1944  unsigned int y);
1945 extern SYCL_EXTERNAL unsigned int __imf_vcmpgtu4(unsigned int x,
1946  unsigned int y);
1947 extern SYCL_EXTERNAL unsigned int __imf_vcmples2(unsigned int x,
1948  unsigned int y);
1949 extern SYCL_EXTERNAL unsigned int __imf_vcmples4(unsigned int x,
1950  unsigned int y);
1951 extern SYCL_EXTERNAL unsigned int __imf_vcmpleu2(unsigned int x,
1952  unsigned int y);
1953 extern SYCL_EXTERNAL unsigned int __imf_vcmpleu4(unsigned int x,
1954  unsigned int y);
1955 extern SYCL_EXTERNAL unsigned int __imf_vcmplts2(unsigned int x,
1956  unsigned int y);
1957 extern SYCL_EXTERNAL unsigned int __imf_vcmplts4(unsigned int x,
1958  unsigned int y);
1959 extern SYCL_EXTERNAL unsigned int __imf_vcmpltu2(unsigned int x,
1960  unsigned int y);
1961 extern SYCL_EXTERNAL unsigned int __imf_vcmpltu4(unsigned int x,
1962  unsigned int y);
1963 extern SYCL_EXTERNAL unsigned int __imf_vmaxs2(unsigned int x, unsigned int y);
1964 extern SYCL_EXTERNAL unsigned int __imf_vmaxs4(unsigned int x, unsigned int y);
1965 extern SYCL_EXTERNAL unsigned int __imf_vmaxu2(unsigned int x, unsigned int y);
1966 extern SYCL_EXTERNAL unsigned int __imf_vmaxu4(unsigned int x, unsigned int y);
1967 extern SYCL_EXTERNAL unsigned int __imf_vmins2(unsigned int x, unsigned int y);
1968 extern SYCL_EXTERNAL unsigned int __imf_vmins4(unsigned int x, unsigned int y);
1969 extern SYCL_EXTERNAL unsigned int __imf_vminu2(unsigned int x, unsigned int y);
1970 extern SYCL_EXTERNAL unsigned int __imf_vminu4(unsigned int x, unsigned int y);
1971 extern SYCL_EXTERNAL unsigned int __imf_vseteq2(unsigned int x, unsigned int y);
1972 extern SYCL_EXTERNAL unsigned int __imf_vseteq4(unsigned int x, unsigned int y);
1973 extern SYCL_EXTERNAL unsigned int __imf_vsetne2(unsigned int x, unsigned int y);
1974 extern SYCL_EXTERNAL unsigned int __imf_vsetne4(unsigned int x, unsigned int y);
1975 extern SYCL_EXTERNAL unsigned int __imf_vsetges2(unsigned int x,
1976  unsigned int y);
1977 extern SYCL_EXTERNAL unsigned int __imf_vsetges4(unsigned int x,
1978  unsigned int y);
1979 extern SYCL_EXTERNAL unsigned int __imf_vsetgeu2(unsigned int x,
1980  unsigned int y);
1981 extern SYCL_EXTERNAL unsigned int __imf_vsetgeu4(unsigned int x,
1982  unsigned int y);
1983 extern SYCL_EXTERNAL unsigned int __imf_vsetgts2(unsigned int x,
1984  unsigned int y);
1985 extern SYCL_EXTERNAL unsigned int __imf_vsetgts4(unsigned int x,
1986  unsigned int y);
1987 extern SYCL_EXTERNAL unsigned int __imf_vsetgtu2(unsigned int x,
1988  unsigned int y);
1989 extern SYCL_EXTERNAL unsigned int __imf_vsetgtu4(unsigned int x,
1990  unsigned int y);
1991 extern SYCL_EXTERNAL unsigned int __imf_vsetles2(unsigned int x,
1992  unsigned int y);
1993 extern SYCL_EXTERNAL unsigned int __imf_vsetles4(unsigned int x,
1994  unsigned int y);
1995 extern SYCL_EXTERNAL unsigned int __imf_vsetleu2(unsigned int x,
1996  unsigned int y);
1997 extern SYCL_EXTERNAL unsigned int __imf_vsetleu4(unsigned int x,
1998  unsigned int y);
1999 extern SYCL_EXTERNAL unsigned int __imf_vsetlts2(unsigned int x,
2000  unsigned int y);
2001 extern SYCL_EXTERNAL unsigned int __imf_vsetlts4(unsigned int x,
2002  unsigned int y);
2003 extern SYCL_EXTERNAL unsigned int __imf_vsetltu2(unsigned int x,
2004  unsigned int y);
2005 extern SYCL_EXTERNAL unsigned int __imf_vsetltu4(unsigned int x,
2006  unsigned int y);
2007 extern SYCL_EXTERNAL unsigned int __imf_vsads2(unsigned int x, unsigned int y);
2008 extern SYCL_EXTERNAL unsigned int __imf_vsads4(unsigned int x, unsigned int y);
2009 extern SYCL_EXTERNAL unsigned int __imf_vsadu2(unsigned int x, unsigned int y);
2010 extern SYCL_EXTERNAL unsigned int __imf_vsadu4(unsigned int x, unsigned int y);
2011 }
2012 #ifdef __GLIBC__
2013 extern "C" {
2014 extern SYCL_EXTERNAL void __assert_fail(const char *expr, const char *file,
2015  unsigned int line, const char *func);
2016 extern SYCL_EXTERNAL float frexpf(float x, int *exp);
2017 extern SYCL_EXTERNAL float ldexpf(float x, int exp);
2018 extern SYCL_EXTERNAL float hypotf(float x, float y);
2019 
2020 // MS UCRT supports most of the C standard library but <complex.h> is
2021 // an exception.
2022 extern SYCL_EXTERNAL float cimagf(float __complex__ z);
2023 extern SYCL_EXTERNAL double cimag(double __complex__ z);
2024 extern SYCL_EXTERNAL float crealf(float __complex__ z);
2025 extern SYCL_EXTERNAL double creal(double __complex__ z);
2026 extern SYCL_EXTERNAL float cargf(float __complex__ z);
2027 extern SYCL_EXTERNAL double carg(double __complex__ z);
2028 extern SYCL_EXTERNAL float cabsf(float __complex__ z);
2029 extern SYCL_EXTERNAL double cabs(double __complex__ z);
2030 extern SYCL_EXTERNAL float __complex__ cprojf(float __complex__ z);
2031 extern SYCL_EXTERNAL double __complex__ cproj(double __complex__ z);
2032 extern SYCL_EXTERNAL float __complex__ cexpf(float __complex__ z);
2033 extern SYCL_EXTERNAL double __complex__ cexp(double __complex__ z);
2034 extern SYCL_EXTERNAL float __complex__ clogf(float __complex__ z);
2035 extern SYCL_EXTERNAL double __complex__ clog(double __complex__ z);
2036 extern SYCL_EXTERNAL float __complex__ cpowf(float __complex__ z);
2037 extern SYCL_EXTERNAL double __complex__ cpow(double __complex__ z);
2038 extern SYCL_EXTERNAL float __complex__ csqrtf(float __complex__ z);
2039 extern SYCL_EXTERNAL double __complex__ csqrt(double __complex__ z);
2040 extern SYCL_EXTERNAL float __complex__ csinhf(float __complex__ z);
2041 extern SYCL_EXTERNAL double __complex__ csinh(double __complex__ z);
2042 extern SYCL_EXTERNAL float __complex__ ccoshf(float __complex__ z);
2043 extern SYCL_EXTERNAL double __complex__ ccosh(double __complex__ z);
2044 extern SYCL_EXTERNAL float __complex__ ctanhf(float __complex__ z);
2045 extern SYCL_EXTERNAL double __complex__ ctanh(double __complex__ z);
2046 extern SYCL_EXTERNAL float __complex__ csinf(float __complex__ z);
2047 extern SYCL_EXTERNAL double __complex__ csin(double __complex__ z);
2048 extern SYCL_EXTERNAL float __complex__ ccosf(float __complex__ z);
2049 extern SYCL_EXTERNAL double __complex__ ccos(double __complex__ z);
2050 extern SYCL_EXTERNAL float __complex__ ctanf(float __complex__ z);
2051 extern SYCL_EXTERNAL double __complex__ ctan(double __complex__ z);
2052 extern SYCL_EXTERNAL float __complex__ cacosf(float __complex__ z);
2053 extern SYCL_EXTERNAL double __complex__ cacos(double __complex__ z);
2054 extern SYCL_EXTERNAL float __complex__ cacoshf(float __complex__ z);
2055 extern SYCL_EXTERNAL double __complex__ cacosh(double __complex__ z);
2056 extern SYCL_EXTERNAL float __complex__ casinf(float __complex__ z);
2057 extern SYCL_EXTERNAL double __complex__ casin(double __complex__ z);
2058 extern SYCL_EXTERNAL float __complex__ casinhf(float __complex__ z);
2059 extern SYCL_EXTERNAL double __complex__ casinh(double __complex__ z);
2060 extern SYCL_EXTERNAL float __complex__ catanf(float __complex__ z);
2061 extern SYCL_EXTERNAL double __complex__ catan(double __complex__ z);
2062 extern SYCL_EXTERNAL float __complex__ catanhf(float __complex__ z);
2063 extern SYCL_EXTERNAL double __complex__ catanh(double __complex__ z);
2064 extern SYCL_EXTERNAL float __complex__ cpolarf(float rho, float theta);
2065 extern SYCL_EXTERNAL double __complex__ cpolar(double rho, double theta);
2066 extern SYCL_EXTERNAL float __complex__ __mulsc3(float a, float b, float c,
2067  float d);
2068 extern SYCL_EXTERNAL double __complex__ __muldc3(double a, double b, double c,
2069  double d);
2070 extern SYCL_EXTERNAL float __complex__ __divsc3(float a, float b, float c,
2071  float d);
2072 extern SYCL_EXTERNAL double __complex__ __divdc3(float a, float b, float c,
2073  float d);
2074 }
2075 #elif defined(_WIN32)
2076 extern "C" {
2077 // TODO: documented C runtime library APIs must be recognized as
2078 // builtins by FE. This includes _dpcomp, _dsign, _dtest,
2079 // _fdpcomp, _fdsign, _fdtest, _hypotf, _wassert.
2080 // APIs used by STL, such as _Cosh, are undocumented, even though
2081 // they are open-sourced. Recognizing them as builtins is not
2082 // straightforward currently.
2083 extern SYCL_EXTERNAL double _Cosh(double x, double y);
2084 extern SYCL_EXTERNAL int _dpcomp(double x, double y);
2085 extern SYCL_EXTERNAL int _dsign(double x);
2086 extern SYCL_EXTERNAL short _Dtest(double *px);
2087 extern SYCL_EXTERNAL short _dtest(double *px);
2088 extern SYCL_EXTERNAL short _Exp(double *px, double y, short eoff);
2089 extern SYCL_EXTERNAL float _FCosh(float x, float y);
2090 extern SYCL_EXTERNAL int _fdpcomp(float x, float y);
2091 extern SYCL_EXTERNAL int _fdsign(float x);
2092 extern SYCL_EXTERNAL short _FDtest(float *px);
2093 extern SYCL_EXTERNAL short _fdtest(float *px);
2094 extern SYCL_EXTERNAL short _FExp(float *px, float y, short eoff);
2095 extern SYCL_EXTERNAL float _FSinh(float x, float y);
2096 extern SYCL_EXTERNAL double _Sinh(double x, double y);
2097 extern SYCL_EXTERNAL float _hypotf(float x, float y);
2098 extern SYCL_EXTERNAL void _wassert(const wchar_t *wexpr, const wchar_t *wfile,
2099  unsigned line);
2100 }
2101 #endif
2102 #endif // __SYCL_DEVICE_ONLY__
2103 
2104 #undef __NOEXC
sycl::_V1::length
half length(T p) __NOEXC
Definition: builtins.hpp:1046
sycl::_V1::isnotequal
detail::common_rel_ret_t< T > isnotequal(T x, T y) __NOEXC
Definition: builtins.hpp:1132
sycl::_V1::half_precision::rsqrt
detail::enable_if_t< detail::is_genfloatf< T >::value, T > rsqrt(T x) __NOEXC
Definition: builtins.hpp:1548
sycl::_V1::nextafter
detail::enable_if_t< detail::is_genfloat< T >::value, T > nextafter(T x, T y) __NOEXC
Definition: builtins.hpp:371
sycl::_V1::clz
detail::enable_if_t< detail::is_geninteger< T >::value, T > clz(T x) __NOEXC
Definition: builtins.hpp:730
sycl::_V1::atan2
detail::enable_if_t< detail::is_genfloat< T >::value, T > atan2(T y, T x) __NOEXC
Definition: builtins.hpp:80
sycl::_V1::dot
detail::enable_if_t< detail::is_vgengeohalf< T >::value, half > dot(T p0, T p1) __NOEXC
Definition: builtins.hpp:1003
sycl::_V1::half_precision::recip
detail::enable_if_t< detail::is_genfloatf< T >::value, T > recip(T x) __NOEXC
Definition: builtins.hpp:1542
__imf_copysign
double __imf_copysign(double, double)
sycl::_V1::all
detail::enable_if_t< detail::is_sigeninteger< T >::value, int > all(T x) __NOEXC
Definition: builtins.hpp:1297
__NOEXC
#define __NOEXC
Definition: builtins.hpp:18
sycl::_V1::cospi
detail::enable_if_t< detail::is_genfloat< T >::value, T > cospi(T x) __NOEXC
Definition: builtins.hpp:136
T
sycl::_V1::mul24
detail::enable_if_t< detail::is_ugeninteger32bit< T >::value, T > mul24(T x, T y) __NOEXC
Definition: builtins.hpp:962
sycl::_V1::fma
detail::enable_if_t< detail::is_genfloat< T >::value, T > fma(T a, T b, T c) __NOEXC
Definition: builtins.hpp:196
sycl::_V1::half_precision::tan
detail::enable_if_t< detail::is_genfloatf< T >::value, T > tan(T x) __NOEXC
Definition: builtins.hpp:1566
sycl::_V1::hypot
detail::enable_if_t< detail::is_genfloat< T >::value, T > hypot(T x, T y) __NOEXC
Definition: builtins.hpp:253
sycl::_V1::ctz
detail::enable_if_t< detail::is_geninteger< T >::value, T > ctz(T x) __NOEXC
Definition: builtins.hpp:736
sycl::_V1::modf
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genfloatptr< T2 >::value, T > modf(T x, T2 iptr) __NOEXC
Definition: builtins.hpp:357
sycl::_V1::minmag
detail::enable_if_t< detail::is_genfloat< T >::value, T > minmag(T x, T y) __NOEXC
Definition: builtins.hpp:349
sycl::_V1::any
detail::enable_if_t< detail::is_vigeninteger< T >::value, int > any(T x) __NOEXC
Definition: builtins.hpp:1289
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:13
sycl::_V1::detail::msbIsSet
constexpr bool msbIsSet(const T x)
Definition: generic_type_traits.hpp:517
__FAST_MATH_GENFLOAT
#define __FAST_MATH_GENFLOAT(T)
Definition: builtins.hpp:32
types.hpp
sycl::_V1::erf
detail::enable_if_t< detail::is_genfloat< T >::value, T > erf(T x) __NOEXC
Definition: builtins.hpp:148
sycl::_V1::abs
detail::enable_if_t< detail::is_igeninteger< T >::value, detail::make_unsigned_t< T > > abs(T x) __NOEXC
Definition: builtins.hpp:637
sycl::_V1::degrees
detail::enable_if_t< detail::is_genfloat< T >::value, T > degrees(T radians) __NOEXC
Definition: builtins.hpp:524
sycl::_V1::detail::rel_ret_t
typename RelationalReturnType< T >::type rel_ret_t
Definition: generic_type_traits.hpp:551
sycl::_V1::erfc
detail::enable_if_t< detail::is_genfloat< T >::value, T > erfc(T x) __NOEXC
Definition: builtins.hpp:142
sycl::_V1::mad24
detail::enable_if_t< detail::is_ugeninteger32bit< T >::value, T > mad24(T x, T y, T z) __NOEXC
Definition: builtins.hpp:948
sycl::_V1::half_precision::cos
detail::enable_if_t< detail::is_genfloatf< T >::value, T > cos(T x) __NOEXC
Definition: builtins.hpp:1487
sycl::_V1::detail::is_geninteger
is_contained< T, gtl::integer_list > is_geninteger
Definition: generic_type_traits.hpp:147
sycl::_V1::nan
detail::nan_return_t< T > nan(T nancode) __NOEXC
Definition: builtins.hpp:364
__host_std
Definition: builtins.hpp:106
sycl::_V1::half_precision::exp10
detail::enable_if_t< detail::is_genfloatf< T >::value, T > exp10(T x) __NOEXC
Definition: builtins.hpp:1512
sycl::_V1::atanh
detail::enable_if_t< detail::is_genfloat< T >::value, T > atanh(T x) __NOEXC
Definition: builtins.hpp:86
sycl::_V1::sincos
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genfloatptr< T2 >::value, T > sincos(T x, T2 cosval) __NOEXC
Definition: builtins.hpp:450
sycl::_V1::hadd
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > hadd(T x, T y) __NOEXC
Definition: builtins.hpp:679
sycl::_V1::acosh
detail::enable_if_t< detail::is_genfloat< T >::value, T > acosh(T x) __NOEXC
Definition: builtins.hpp:44
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:13
sycl::_V1::isless
detail::common_rel_ret_t< T > isless(T x, T y) __NOEXC
Definition: builtins.hpp:1168
sycl::_V1::isfinite
detail::common_rel_ret_t< T > isfinite(T x) __NOEXC
Definition: builtins.hpp:1204
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::_V1::rootn
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genint< T2 >::value, T > rootn(T x, T2 y) __NOEXC
Definition: builtins.hpp:423
sycl::_V1::acospi
detail::enable_if_t< detail::is_genfloat< T >::value, T > acospi(T x) __NOEXC
Definition: builtins.hpp:50
sycl::_V1::detail::memcpy
void memcpy(void *Dst, const void *Src, std::size_t Size)
sycl::_V1::detail::is_genfloatptr
bool_constant< is_pointer< T >::value &&is_genfloat< remove_pointer_t< T > >::value &&is_address_space_compliant< T, gvl::nonconst_address_space_list >::value > is_genfloatptr
Definition: generic_type_traits.hpp:225
sycl::_V1::ext::oneapi::experimental::native::tanh
__SYCL_ALWAYS_INLINE sycl::detail::enable_if_t< sycl::detail::is_genfloatf< T >::value||sycl::detail::is_genfloath< T >::value, T > tanh(T x) __NOEXC
Definition: builtins.hpp:101
sycl::_V1::half_precision::divide
detail::enable_if_t< detail::is_genfloatf< T >::value, T > divide(T x, T y) __NOEXC
Definition: builtins.hpp:1493
sycl::_V1::ext::oneapi::experimental::fabs
sycl::marray< bfloat16, N > fabs(sycl::marray< bfloat16, N > x)
Definition: builtins.hpp:139
sycl::_V1::sinh
detail::enable_if_t< detail::is_genfloat< T >::value, T > sinh(T x) __NOEXC
Definition: builtins.hpp:457
std::clog
__SYCL_EXTERN_STREAM_ATTRS ostream clog
Linked to standard error (buffered)
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:45
sycl::_V1::round
detail::enable_if_t< detail::is_genfloat< T >::value, T > round(T x) __NOEXC
Definition: builtins.hpp:430
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
sycl::_V1::maxmag
detail::enable_if_t< detail::is_genfloat< T >::value, T > maxmag(T x, T y) __NOEXC
Definition: builtins.hpp:343
sycl::_V1::half_precision::log2
detail::enable_if_t< detail::is_genfloatf< T >::value, T > log2(T x) __NOEXC
Definition: builtins.hpp:1524
sycl::_V1::detail::rel_sign_bit_test_arg_t
typename RelationalTestForSignBitType< T >::argument_type rel_sign_bit_test_arg_t
Definition: generic_type_traits.hpp:570
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::mix
detail::enable_if_t< detail::is_vgenfloat< T >::value, T > mix(T x, T y, typename T::element_type a) __NOEXC
Definition: builtins.hpp:576
sycl::_V1::radians
detail::enable_if_t< detail::is_genfloat< T >::value, T > radians(T degrees) __NOEXC
Definition: builtins.hpp:583
sycl::_V1::sign
detail::enable_if_t< detail::is_genfloat< T >::value, T > sign(T x) __NOEXC
Definition: builtins.hpp:622
sycl::_V1::fast_normalize
detail::enable_if_t< detail::is_gengeodouble< T >::value, T > fast_normalize(T p) __NOEXC
Definition: builtins.hpp:1108
sycl::_V1::isequal
detail::common_rel_ret_t< T > isequal(T x, T y) __NOEXC
Definition: builtins.hpp:1120
sycl::_V1::distance
half distance(T p0, T p1) __NOEXC
Definition: builtins.hpp:1025
sycl::_V1::ldexp
detail::enable_if_t< detail::is_vgenfloat< T >::value &&detail::is_intn< T2 >::value, T > ldexp(T x, T2 k) __NOEXC
Definition: builtins.hpp:284
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::detail::change_base_type_t
typename change_base_type< T, B >::type change_base_type_t
Definition: type_traits.hpp:114
sycl::_V1::half_precision::sin
detail::enable_if_t< detail::is_genfloatf< T >::value, T > sin(T x) __NOEXC
Definition: builtins.hpp:1554
__imf_saturatef
float __imf_saturatef(float)
sycl::_V1::ext::oneapi::experimental::fma
sycl::marray< bfloat16, N > fma(sycl::marray< bfloat16, N > x, sycl::marray< bfloat16, N > y, sycl::marray< bfloat16, N > z)
Definition: builtins.hpp:248
sycl::_V1::abs_diff
detail::enable_if_t< detail::is_igeninteger< T >::value, detail::make_unsigned_t< T > > abs_diff(T x, T y) __NOEXC
Definition: builtins.hpp:652
generic_type_traits.hpp
sycl::_V1::tanpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > tanpi(T x) __NOEXC
Definition: builtins.hpp:487
sycl::_V1::mad_sat
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > mad_sat(T a, T b, T c) __NOEXC
Definition: builtins.hpp:780
sycl::_V1::ilogb
detail::change_base_type_t< T, int > ilogb(T x) __NOEXC
Definition: builtins.hpp:260
sycl::_V1::remainder
detail::enable_if_t< detail::is_genfloat< T >::value, T > remainder(T x, T y) __NOEXC
Definition: builtins.hpp:399
sycl::_V1::cbrt
detail::enable_if_t< detail::is_genfloat< T >::value, T > cbrt(T x) __NOEXC
Definition: builtins.hpp:105
sycl::_V1::isgreater
detail::common_rel_ret_t< T > isgreater(T x, T y) __NOEXC
Definition: builtins.hpp:1144
sycl::_V1::remquo
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genintptr< T2 >::value, T > remquo(T x, T y, T2 quo) __NOEXC
Definition: builtins.hpp:408
sycl::_V1::isordered
detail::common_rel_ret_t< T > isordered(T x, T y) __NOEXC
Definition: builtins.hpp:1252
sycl::_V1::rint
detail::enable_if_t< detail::is_genfloat< T >::value, T > rint(T x) __NOEXC
Definition: builtins.hpp:415
SYCL_EXTERNAL
#define SYCL_EXTERNAL
Definition: defines_elementary.hpp:32
sycl::_V1::smoothstep
detail::enable_if_t< detail::is_vgenfloat< T >::value, T > smoothstep(typename T::element_type edge0, typename T::element_type edge1, T x) __NOEXC
Definition: builtins.hpp:615
sycl::_V1::instead
std::uint8_t instead
Definition: aliases.hpp:69
common.hpp
sycl::_V1::detail::convert_data_type_impl
Definition: generic_type_traits.hpp:276
sycl::_V1::fdim
detail::enable_if_t< detail::is_genfloat< T >::value, T > fdim(T x, T y) __NOEXC
Definition: builtins.hpp:184
sycl::_V1::asin
detail::enable_if_t< detail::is_genfloat< T >::value, T > asin(T x) __NOEXC
Definition: builtins.hpp:56
sycl::_V1::fast_length
double fast_length(T p) __NOEXC
Definition: builtins.hpp:1094
sycl::_V1::asinh
detail::enable_if_t< detail::is_genfloat< T >::value, T > asinh(T x) __NOEXC
Definition: builtins.hpp:62
sycl::_V1::fast_distance
double fast_distance(T p0, T p1) __NOEXC
Definition: builtins.hpp:1080
sycl::_V1::isnormal
detail::common_rel_ret_t< T > isnormal(T x) __NOEXC
Definition: builtins.hpp:1240
sycl::_V1::sinpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > sinpi(T x) __NOEXC
Definition: builtins.hpp:463
sycl::_V1::sub_sat
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > sub_sat(T x, T y) __NOEXC
Definition: builtins.hpp:871
sycl::_V1::half
sycl::detail::half_impl::half half
Definition: aliases.hpp:77
sycl::_V1::half_precision::powr
detail::enable_if_t< detail::is_genfloatf< T >::value, T > powr(T x, T y) __NOEXC
Definition: builtins.hpp:1536
sycl::_V1::trunc
detail::enable_if_t< detail::is_genfloat< T >::value, T > trunc(T x) __NOEXC
Definition: builtins.hpp:499
sycl::_V1::detail::rel_sign_bit_test_ret_t
typename RelationalTestForSignBitType< T >::return_type rel_sign_bit_test_ret_t
Definition: generic_type_traits.hpp:566
sycl::_V1::expm1
detail::enable_if_t< detail::is_genfloat< T >::value, T > expm1(T x) __NOEXC
Definition: builtins.hpp:172
sycl::_V1::acos
detail::enable_if_t< detail::is_genfloat< T >::value, T > acos(T x) __NOEXC
Definition: builtins.hpp:38
sycl::_V1::pow
detail::enable_if_t< detail::is_genfloat< T >::value, T > pow(T x, T y) __NOEXC
Definition: builtins.hpp:378
sycl::_V1::detail::Boolean< 1 >
Definition: boolean.hpp:103
clamp
simd< _Tp, _Abi > clamp(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &)
__imf_copysignf
float __imf_copysignf(float, float)
sycl::_V1::atanpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > atanpi(T x) __NOEXC
Definition: builtins.hpp:92
sycl::_V1::detail::nan_return_t
typename nan_types< T, T >::ret_type nan_return_t
Definition: generic_type_traits.hpp:260
sycl::_V1::asinpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > asinpi(T x) __NOEXC
Definition: builtins.hpp:68
sycl::_V1::detail::common_rel_ret_t
conditional_t< is_vgentype< T >::value, make_singed_integer_t< T >, int > common_rel_ret_t
Definition: generic_type_traits.hpp:523
__simd_abi
Definition: simd.hpp:718
sycl::_V1::islessgreater
detail::common_rel_ret_t< T > islessgreater(T x, T y) __NOEXC
Definition: builtins.hpp:1192
boolean.hpp
sycl::_V1::rotate
detail::enable_if_t< detail::is_geninteger< T >::value, T > rotate(T v, T i) __NOEXC
Definition: builtins.hpp:857
sycl::_V1::mad_hi
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > mad_hi(T x, T y, T z) __NOEXC
Definition: builtins.hpp:766
sycl::_V1::detail::RelConverter
Definition: generic_type_traits.hpp:572
sycl::_V1::ext::oneapi::experimental::native::exp2
__SYCL_ALWAYS_INLINE sycl::detail::enable_if_t< sycl::detail::is_genfloath< T >::value, T > exp2(T x) __NOEXC
Definition: builtins.hpp:115
builtins.hpp
sycl::_V1::lgamma
detail::enable_if_t< detail::is_genfloat< T >::value, T > lgamma(T x) __NOEXC
Definition: builtins.hpp:291
sycl::_V1::cosh
detail::enable_if_t< detail::is_genfloat< T >::value, T > cosh(T x) __NOEXC
Definition: builtins.hpp:130
sycl::_V1::floor
detail::enable_if_t< detail::is_genfloat< T >::value, T > floor(T x) __NOEXC
Definition: builtins.hpp:190
sycl::_V1::isunordered
detail::common_rel_ret_t< T > isunordered(T x, T y) __NOEXC
Definition: builtins.hpp:1264
sycl::_V1::logb
detail::enable_if_t< detail::is_genfloat< T >::value, T > logb(T x) __NOEXC
Definition: builtins.hpp:330
sycl::_V1::fmod
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmod(T x, T y) __NOEXC
Definition: builtins.hpp:229
sycl::_V1::half_precision::exp
detail::enable_if_t< detail::is_genfloatf< T >::value, T > exp(T x) __NOEXC
Definition: builtins.hpp:1500
sycl::_V1::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:19
sycl::_V1::ceil
detail::enable_if_t< detail::is_genfloat< T >::value, T > ceil(T x) __NOEXC
Definition: builtins.hpp:111
sycl::_V1::isnan
detail::common_rel_ret_t< T > isnan(T x) __NOEXC
Definition: builtins.hpp:1228
sycl::_V1::detail::is_contained
Definition: type_list.hpp:54
sycl::_V1::islessequal
detail::common_rel_ret_t< T > islessequal(T x, T y) __NOEXC
Definition: builtins.hpp:1180
sycl::_V1::step
detail::enable_if_t< detail::is_vgenfloat< T >::value, T > step(typename T::element_type edge, T x) __NOEXC
Definition: builtins.hpp:599
sycl::_V1::copysign
detail::enable_if_t< detail::is_genfloat< T >::value, T > copysign(T x, T y) __NOEXC
Definition: builtins.hpp:117
sycl::_V1::mad
detail::enable_if_t< detail::is_genfloat< T >::value, T > mad(T a, T b, T c) __NOEXC
Definition: builtins.hpp:336
sycl::_V1::upsample
detail::enable_if_t< detail::is_igeninteger32bit< T >::value &&detail::is_ugeninteger32bit< T2 >::value, detail::make_larger_t< T > > upsample(T hi, T2 lo) __NOEXC
Definition: builtins.hpp:925
sycl::_V1::atan
detail::enable_if_t< detail::is_genfloat< T >::value, T > atan(T y_over_x) __NOEXC
Definition: builtins.hpp:74
sycl::_V1::isgreaterequal
detail::common_rel_ret_t< T > isgreaterequal(T x, T y) __NOEXC
Definition: builtins.hpp:1156
sycl::_V1::ext::oneapi::experimental::matrix::use
use
Definition: matrix-jit-use.hpp:44
sycl::_V1::fract
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genfloatptr< T2 >::value, T > fract(T x, T2 iptr) __NOEXC
Definition: builtins.hpp:237
sycl::_V1::ext::oneapi::experimental::fmin
sycl::marray< bfloat16, N > fmin(sycl::marray< bfloat16, N > x, sycl::marray< bfloat16, N > y)
Definition: builtins.hpp:172
sycl::_V1::signbit
detail::common_rel_ret_t< T > signbit(T x) __NOEXC
Definition: builtins.hpp:1276
sycl::_V1::tgamma
detail::enable_if_t< detail::is_genfloat< T >::value, T > tgamma(T x) __NOEXC
Definition: builtins.hpp:493
sycl::_V1::log1p
detail::enable_if_t< detail::is_genfloat< T >::value, T > log1p(T x) __NOEXC
Definition: builtins.hpp:324
sycl::_V1::normalize
detail::enable_if_t< detail::is_gengeohalf< T >::value, T > normalize(T p) __NOEXC
Definition: builtins.hpp:1066
sycl::_V1::add_sat
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > add_sat(T x, T y) __NOEXC
Definition: builtins.hpp:665
sycl::_V1::rhadd
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > rhadd(T x, T y) __NOEXC
Definition: builtins.hpp:693
sycl::_V1::mul_hi
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > mul_hi(T x, T y) __NOEXC
Definition: builtins.hpp:850
sycl::_V1::pown
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genint< T2 >::value, T > pown(T x, T2 y) __NOEXC
Definition: builtins.hpp:386
sycl::_V1::cross
detail::enable_if_t< detail::is_gencross< T >::value, T > cross(T p0, T p1) __NOEXC
Definition: builtins.hpp:974
sycl::_V1::lgamma_r
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genintptr< T2 >::value, T > lgamma_r(T x, T2 signp) __NOEXC
Definition: builtins.hpp:299
sycl::_V1::isinf
detail::common_rel_ret_t< T > isinf(T x) __NOEXC
Definition: builtins.hpp:1216
sycl::_V1::frexp
detail::enable_if_t< detail::is_genfloat< T >::value &&detail::is_genintptr< T2 >::value, T > frexp(T x, T2 exp) __NOEXC
Definition: builtins.hpp:246
sycl::_V1::half_precision::sqrt
detail::enable_if_t< detail::is_genfloatf< T >::value, T > sqrt(T x) __NOEXC
Definition: builtins.hpp:1560
sycl::_V1::half_precision::log
detail::enable_if_t< detail::is_genfloatf< T >::value, T > log(T x) __NOEXC
Definition: builtins.hpp:1518
sycl::_V1::bitselect
detail::enable_if_t< detail::is_gentype< T >::value, T > bitselect(T a, T b, T c) __NOEXC
Definition: builtins.hpp:1311
sycl::_V1::atan2pi
detail::enable_if_t< detail::is_genfloat< T >::value, T > atan2pi(T y, T x) __NOEXC
Definition: builtins.hpp:98
__imf_copysignf16
_iml_half_internal __imf_copysignf16(_iml_half_internal, _iml_half_internal)
sycl::_V1::ext::oneapi::experimental::fmax
sycl::marray< bfloat16, N > fmax(sycl::marray< bfloat16, N > x, sycl::marray< bfloat16, N > y)
Definition: builtins.hpp:210
sycl::_V1::popcount
detail::enable_if_t< detail::is_geninteger< T >::value, T > popcount(T x) __NOEXC
Definition: builtins.hpp:932
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:54
sycl::_V1::half_precision::log10
detail::enable_if_t< detail::is_genfloatf< T >::value, T > log10(T x) __NOEXC
Definition: builtins.hpp:1530
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::_V1::select
detail::enable_if_t< detail::is_genfloath< T >::value &&detail::is_ugeninteger16bit< T2 >::value, T > select(T a, T b, T2 c) __NOEXC
Definition: builtins.hpp:1390