DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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 
15 #include <CL/sycl/types.hpp>
16 
17 // TODO Decide whether to mark functions with this attribute.
18 #define __NOEXC /*noexcept*/
19 
21 namespace sycl {
22 #ifdef __SYCL_DEVICE_ONLY__
23 #define __sycl_std
24 #else
25 namespace __sycl_std = __host_std;
26 #endif
27 
28 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
29 // genfloat acos (genfloat x)
30 template <typename T>
32  return __sycl_std::__invoke_acos<T>(x);
33 }
34 
35 // genfloat acosh (genfloat x)
36 template <typename T>
38  return __sycl_std::__invoke_acosh<T>(x);
39 }
40 
41 // genfloat acospi (genfloat x)
42 template <typename T>
44  return __sycl_std::__invoke_acospi<T>(x);
45 }
46 
47 // genfloat asin (genfloat x)
48 template <typename T>
50  return __sycl_std::__invoke_asin<T>(x);
51 }
52 
53 // genfloat asinh (genfloat x)
54 template <typename T>
56  return __sycl_std::__invoke_asinh<T>(x);
57 }
58 
59 // genfloat asinpi (genfloat x)
60 template <typename T>
62  return __sycl_std::__invoke_asinpi<T>(x);
63 }
64 
65 // genfloat atan (genfloat y_over_x)
66 template <typename T>
68  return __sycl_std::__invoke_atan<T>(y_over_x);
69 }
70 
71 // genfloat atan2 (genfloat y, genfloat x)
72 template <typename T>
74  return __sycl_std::__invoke_atan2<T>(y, x);
75 }
76 
77 // genfloat atanh (genfloat x)
78 template <typename T>
80  return __sycl_std::__invoke_atanh<T>(x);
81 }
82 
83 // genfloat atanpi (genfloat x)
84 template <typename T>
86  return __sycl_std::__invoke_atanpi<T>(x);
87 }
88 
89 // genfloat atan2pi (genfloat y, genfloat x)
90 template <typename T>
92  T x) __NOEXC {
93  return __sycl_std::__invoke_atan2pi<T>(y, x);
94 }
95 
96 // genfloat cbrt (genfloat x)
97 template <typename T>
99  return __sycl_std::__invoke_cbrt<T>(x);
100 }
101 
102 // genfloat ceil (genfloat x)
103 template <typename T>
105  return __sycl_std::__invoke_ceil<T>(x);
106 }
107 
108 // genfloat copysign (genfloat x, genfloat y)
109 template <typename T>
111  T y) __NOEXC {
112  return __sycl_std::__invoke_copysign<T>(x, y);
113 }
114 
115 // genfloat cos (genfloat x)
116 template <typename T>
118  return __sycl_std::__invoke_cos<T>(x);
119 }
120 
121 // genfloat cosh (genfloat x)
122 template <typename T>
124  return __sycl_std::__invoke_cosh<T>(x);
125 }
126 
127 // genfloat cospi (genfloat x)
128 template <typename T>
130  return __sycl_std::__invoke_cospi<T>(x);
131 }
132 
133 // genfloat erfc (genfloat x)
134 template <typename T>
136  return __sycl_std::__invoke_erfc<T>(x);
137 }
138 
139 // genfloat erf (genfloat x)
140 template <typename T>
142  return __sycl_std::__invoke_erf<T>(x);
143 }
144 
145 // genfloat exp (genfloat x )
146 template <typename T>
148  return __sycl_std::__invoke_exp<T>(x);
149 }
150 
151 // genfloat exp2 (genfloat x)
152 template <typename T>
154  return __sycl_std::__invoke_exp2<T>(x);
155 }
156 
157 // genfloat exp10 (genfloat x)
158 template <typename T>
160  return __sycl_std::__invoke_exp10<T>(x);
161 }
162 
163 // genfloat expm1 (genfloat x)
164 template <typename T>
166  return __sycl_std::__invoke_expm1<T>(x);
167 }
168 
169 // genfloat fabs (genfloat x)
170 template <typename T>
172  return __sycl_std::__invoke_fabs<T>(x);
173 }
174 
175 // genfloat fdim (genfloat x, genfloat y)
176 template <typename T>
178  return __sycl_std::__invoke_fdim<T>(x, y);
179 }
180 
181 // genfloat floor (genfloat x)
182 template <typename T>
184  return __sycl_std::__invoke_floor<T>(x);
185 }
186 
187 // genfloat fma (genfloat a, genfloat b, genfloat c)
188 template <typename T>
190  T c) __NOEXC {
191  return __sycl_std::__invoke_fma<T>(a, b, c);
192 }
193 
194 // genfloat fmax (genfloat x, genfloat y)
195 template <typename T>
197  return __sycl_std::__invoke_fmax<T>(x, y);
198 }
199 
200 // genfloat fmax (genfloat x, sgenfloat y)
201 template <typename T>
202 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
203 fmax(T x, typename T::element_type y) __NOEXC {
204  return __sycl_std::__invoke_fmax<T>(x, T(y));
205 }
206 
207 // genfloat fmin (genfloat x, genfloat y)
208 template <typename T>
210  return __sycl_std::__invoke_fmin<T>(x, y);
211 }
212 
213 // genfloat fmin (genfloat x, sgenfloat y)
214 template <typename T>
215 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
216 fmin(T x, typename T::element_type y) __NOEXC {
217  return __sycl_std::__invoke_fmin<T>(x, T(y));
218 }
219 
220 // genfloat fmod (genfloat x, genfloat y)
221 template <typename T>
223  return __sycl_std::__invoke_fmod<T>(x, y);
224 }
225 
226 // genfloat fract (genfloat x, genfloatptr iptr)
227 template <typename T, typename T2>
229  detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
230 fract(T x, T2 iptr) __NOEXC {
231  detail::check_vector_size<T, T2>();
232  return __sycl_std::__invoke_fract<T>(x, iptr);
233 }
234 
235 // genfloat frexp (genfloat x, genintptr exp)
236 template <typename T, typename T2>
238  detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
239 frexp(T x, T2 exp) __NOEXC {
240  detail::check_vector_size<T, T2>();
241  return __sycl_std::__invoke_frexp<T>(x, exp);
242 }
243 
244 // genfloat hypot (genfloat x, genfloat y)
245 template <typename T>
247  return __sycl_std::__invoke_hypot<T>(x, y);
248 }
249 
250 // genint ilogb (genfloat x)
251 template <typename T,
252  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
254  return __sycl_std::__invoke_ilogb<detail::change_base_type_t<T, int>>(x);
255 }
256 
257 // float ldexp (float x, int k)
258 // double ldexp (double x, int k)
259 // half ldexp (half x, int k)
260 template <typename T>
262  int k) __NOEXC {
263  return __sycl_std::__invoke_ldexp<T>(x, k);
264 }
265 
266 // vgenfloat ldexp (vgenfloat x, int k)
267 template <typename T>
269  int k) __NOEXC {
270  return __sycl_std::__invoke_ldexp<T>(x, vec<int, T::size()>(k));
271 }
272 
273 // vgenfloat ldexp (vgenfloat x, genint k)
274 template <typename T, typename T2>
276  detail::is_vgenfloat<T>::value && detail::is_intn<T2>::value, T>
277 ldexp(T x, T2 k) __NOEXC {
278  detail::check_vector_size<T, T2>();
279  return __sycl_std::__invoke_ldexp<T>(x, k);
280 }
281 
282 // genfloat lgamma (genfloat x)
283 template <typename T>
285  return __sycl_std::__invoke_lgamma<T>(x);
286 }
287 
288 // genfloat lgamma_r (genfloat x, genintptr signp)
289 template <typename T, typename T2>
291  detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
292 lgamma_r(T x, T2 signp) __NOEXC {
293  detail::check_vector_size<T, T2>();
294  return __sycl_std::__invoke_lgamma_r<T>(x, signp);
295 }
296 
297 // genfloat log (genfloat x)
298 template <typename T>
300  return __sycl_std::__invoke_log<T>(x);
301 }
302 
303 // genfloat log2 (genfloat x)
304 template <typename T>
306  return __sycl_std::__invoke_log2<T>(x);
307 }
308 
309 // genfloat log10 (genfloat x)
310 template <typename T>
312  return __sycl_std::__invoke_log10<T>(x);
313 }
314 
315 // genfloat log1p (genfloat x)
316 template <typename T>
318  return __sycl_std::__invoke_log1p<T>(x);
319 }
320 
321 // genfloat logb (genfloat x)
322 template <typename T>
324  return __sycl_std::__invoke_logb<T>(x);
325 }
326 
327 // genfloat mad (genfloat a, genfloat b, genfloat c)
328 template <typename T>
330  T c) __NOEXC {
331  return __sycl_std::__invoke_mad<T>(a, b, c);
332 }
333 
334 // genfloat maxmag (genfloat x, genfloat y)
335 template <typename T>
337  return __sycl_std::__invoke_maxmag<T>(x, y);
338 }
339 
340 // genfloat minmag (genfloat x, genfloat y)
341 template <typename T>
343  return __sycl_std::__invoke_minmag<T>(x, y);
344 }
345 
346 // genfloat modf (genfloat x, genfloatptr iptr)
347 template <typename T, typename T2>
349  detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
350 modf(T x, T2 iptr) __NOEXC {
351  detail::check_vector_size<T, T2>();
352  return __sycl_std::__invoke_modf<T>(x, iptr);
353 }
354 
355 template <typename T,
356  typename = detail::enable_if_t<detail::is_nan_type<T>::value, T>>
358  return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
360 }
361 
362 // genfloat nextafter (genfloat x, genfloat y)
363 template <typename T>
365  T y) __NOEXC {
366  return __sycl_std::__invoke_nextafter<T>(x, y);
367 }
368 
369 // genfloat pow (genfloat x, genfloat y)
370 template <typename T>
372  return __sycl_std::__invoke_pow<T>(x, y);
373 }
374 
375 // genfloat pown (genfloat x, genint y)
376 template <typename T, typename T2>
378  detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>
379 pown(T x, T2 y) __NOEXC {
380  detail::check_vector_size<T, T2>();
381  return __sycl_std::__invoke_pown<T>(x, y);
382 }
383 
384 // genfloat powr (genfloat x, genfloat y)
385 template <typename T>
387  return __sycl_std::__invoke_powr<T>(x, y);
388 }
389 
390 // genfloat remainder (genfloat x, genfloat y)
391 template <typename T>
393  T y) __NOEXC {
394  return __sycl_std::__invoke_remainder<T>(x, y);
395 }
396 
397 // genfloat remquo (genfloat x, genfloat y, genintptr quo)
398 template <typename T, typename T2>
400  detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
401 remquo(T x, T y, T2 quo) __NOEXC {
402  detail::check_vector_size<T, T2>();
403  return __sycl_std::__invoke_remquo<T>(x, y, quo);
404 }
405 
406 // genfloat rint (genfloat x)
407 template <typename T>
409  return __sycl_std::__invoke_rint<T>(x);
410 }
411 
412 // genfloat rootn (genfloat x, genint y)
413 template <typename T, typename T2>
415  detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>
416 rootn(T x, T2 y) __NOEXC {
417  detail::check_vector_size<T, T2>();
418  return __sycl_std::__invoke_rootn<T>(x, y);
419 }
420 
421 // genfloat round (genfloat x)
422 template <typename T>
424  return __sycl_std::__invoke_round<T>(x);
425 }
426 
427 // genfloat rsqrt (genfloat x)
428 template <typename T>
430  return __sycl_std::__invoke_rsqrt<T>(x);
431 }
432 
433 // genfloat sin (genfloat x)
434 template <typename T>
436  return __sycl_std::__invoke_sin<T>(x);
437 }
438 
439 // genfloat sincos (genfloat x, genfloatptr cosval)
440 template <typename T, typename T2>
442  detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
443 sincos(T x, T2 cosval) __NOEXC {
444  detail::check_vector_size<T, T2>();
445  return __sycl_std::__invoke_sincos<T>(x, cosval);
446 }
447 
448 // genfloat sinh (genfloat x)
449 template <typename T>
451  return __sycl_std::__invoke_sinh<T>(x);
452 }
453 
454 // genfloat sinpi (genfloat x)
455 template <typename T>
457  return __sycl_std::__invoke_sinpi<T>(x);
458 }
459 
460 // genfloat sqrt (genfloat x)
461 template <typename T>
463  return __sycl_std::__invoke_sqrt<T>(x);
464 }
465 
466 // genfloat tan (genfloat x)
467 template <typename T>
469  return __sycl_std::__invoke_tan<T>(x);
470 }
471 
472 // genfloat tanh (genfloat x)
473 template <typename T>
475  return __sycl_std::__invoke_tanh<T>(x);
476 }
477 
478 // genfloat tanpi (genfloat x)
479 template <typename T>
481  return __sycl_std::__invoke_tanpi<T>(x);
482 }
483 
484 // genfloat tgamma (genfloat x)
485 template <typename T>
487  return __sycl_std::__invoke_tgamma<T>(x);
488 }
489 
490 // genfloat trunc (genfloat x)
491 template <typename T>
493  return __sycl_std::__invoke_trunc<T>(x);
494 }
495 
496 /* --------------- 4.13.5 Common functions. ---------------------------------*/
497 // genfloat clamp (genfloat x, genfloat minval, genfloat maxval)
498 template <typename T>
500  T maxval) __NOEXC {
501  return __sycl_std::__invoke_fclamp<T>(x, minval, maxval);
502 }
503 
504 // genfloath clamp (genfloath x, half minval, half maxval)
505 // genfloatf clamp (genfloatf x, float minval, float maxval)
506 // genfloatd clamp (genfloatd x, double minval, double maxval)
507 template <typename T>
508 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
509 clamp(T x, typename T::element_type minval,
510  typename T::element_type maxval) __NOEXC {
511  return __sycl_std::__invoke_fclamp<T>(x, T(minval), T(maxval));
512 }
513 
514 // genfloat degrees (genfloat radians)
515 template <typename T>
516 detail::enable_if_t<detail::is_genfloat<T>::value, T>
518  return __sycl_std::__invoke_degrees<T>(radians);
519 }
520 
521 // genfloat abs (genfloat x)
522 template <typename T>
524  return __sycl_std::__invoke_fabs<T>(x);
525 }
526 
527 // genfloat max (genfloat x, genfloat y)
528 template <typename T>
530  return __sycl_std::__invoke_fmax_common<T>(x, y);
531 }
532 
533 // genfloatf max (genfloatf x, float y)
534 // genfloatd max (genfloatd x, double y)
535 // genfloath max (genfloath x, half y)
536 template <typename T>
537 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(max)(
538  T x, typename T::element_type y) __NOEXC {
539  return __sycl_std::__invoke_fmax_common<T>(x, T(y));
540 }
541 
542 // genfloat min (genfloat x, genfloat y)
543 template <typename T>
544 detail::enable_if_t<detail::is_genfloat<T>::value, T>(min)(T x, T y) __NOEXC {
545  return __sycl_std::__invoke_fmin_common<T>(x, y);
546 }
547 
548 // genfloatf min (genfloatf x, float y)
549 // genfloatd min (genfloatd x, double y)
550 // genfloath min (genfloath x, half y)
551 template <typename T>
552 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(min)(
553  T x, typename T::element_type y) __NOEXC {
554  return __sycl_std::__invoke_fmin_common<T>(x, T(y));
555 }
556 
557 // genfloat mix (genfloat x, genfloat y, genfloat a)
558 template <typename T>
560  T a) __NOEXC {
561  return __sycl_std::__invoke_mix<T>(x, y, a);
562 }
563 
564 // genfloatf mix (genfloatf x, genfloatf y, float a)
565 // genfloatd mix (genfloatd x, genfloatd y, double a)
566 // genfloatd mix (genfloath x, genfloath y, half a)
567 template <typename T>
568 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
569 mix(T x, T y, typename T::element_type a) __NOEXC {
570  return __sycl_std::__invoke_mix<T>(x, y, T(a));
571 }
572 
573 // genfloat radians (genfloat degrees)
574 template <typename T>
575 detail::enable_if_t<detail::is_genfloat<T>::value, T>
577  return __sycl_std::__invoke_radians<T>(degrees);
578 }
579 
580 // genfloat step (genfloat edge, genfloat x)
581 template <typename T>
583  T x) __NOEXC {
584  return __sycl_std::__invoke_step<T>(edge, x);
585 }
586 
587 // genfloatf step (float edge, genfloatf x)
588 // genfloatd step (double edge, genfloatd x)
589 // genfloatd step (half edge, genfloath x)
590 template <typename T>
591 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
592 step(typename T::element_type edge, T x) __NOEXC {
593  return __sycl_std::__invoke_step<T>(T(edge), x);
594 }
595 
596 // genfloat smoothstep (genfloat edge0, genfloat edge1, genfloat x)
597 template <typename T>
598 detail::enable_if_t<detail::is_genfloat<T>::value, T>
599 smoothstep(T edge0, T edge1, T x) __NOEXC {
600  return __sycl_std::__invoke_smoothstep<T>(edge0, edge1, x);
601 }
602 
603 // genfloatf smoothstep (float edge0, float edge1, genfloatf x)
604 // genfloatd smoothstep (double edge0, double edge1, genfloatd x)
605 // genfloath smoothstep (half edge0, half edge1, genfloath x)
606 template <typename T>
607 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
608 smoothstep(typename T::element_type edge0, typename T::element_type edge1,
609  T x) __NOEXC {
610  return __sycl_std::__invoke_smoothstep<T>(T(edge0), T(edge1), x);
611 }
612 
613 // genfloat sign (genfloat x)
614 template <typename T>
616  return __sycl_std::__invoke_sign<T>(x);
617 }
618 
619 /* --------------- 4.13.4 Integer functions. --------------------------------*/
620 // ugeninteger abs (geninteger x)
621 template <typename T>
623  return __sycl_std::__invoke_u_abs<T>(x);
624 }
625 
626 // ugeninteger abs (geninteger x)
627 template <typename T>
628 detail::enable_if_t<detail::is_igeninteger<T>::value,
629  detail::make_unsigned_t<T>>
630 abs(T x) __NOEXC {
631  return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
632 }
633 
634 // ugeninteger abs_diff (geninteger x, geninteger y)
635 template <typename T>
637  T y) __NOEXC {
638  return __sycl_std::__invoke_u_abs_diff<T>(x, y);
639 }
640 
641 // ugeninteger abs_diff (geninteger x, geninteger y)
642 template <typename T>
643 detail::enable_if_t<detail::is_igeninteger<T>::value,
644  detail::make_unsigned_t<T>>
645 abs_diff(T x, T y) __NOEXC {
646  return __sycl_std::__invoke_s_abs_diff<detail::make_unsigned_t<T>>(x, y);
647 }
648 
649 // geninteger add_sat (geninteger x, geninteger y)
650 template <typename T>
652  T y) __NOEXC {
653  return __sycl_std::__invoke_s_add_sat<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_u_add_sat<T>(x, y);
661 }
662 
663 // geninteger hadd (geninteger x, geninteger y)
664 template <typename T>
666  T y) __NOEXC {
667  return __sycl_std::__invoke_s_hadd<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_u_hadd<T>(x, y);
675 }
676 
677 // geninteger rhadd (geninteger x, geninteger y)
678 template <typename T>
680  T y) __NOEXC {
681  return __sycl_std::__invoke_s_rhadd<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_u_rhadd<T>(x, y);
689 }
690 
691 // geninteger clamp (geninteger x, geninteger minval, geninteger maxval)
692 template <typename T>
693 detail::enable_if_t<detail::is_igeninteger<T>::value, T>
694 clamp(T x, T minval, T maxval) __NOEXC {
695  return __sycl_std::__invoke_s_clamp<T>(x, minval, maxval);
696 }
697 
698 // geninteger clamp (geninteger x, geninteger minval, geninteger maxval)
699 template <typename T>
700 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>
701 clamp(T x, T minval, T maxval) __NOEXC {
702  return __sycl_std::__invoke_u_clamp<T>(x, minval, maxval);
703 }
704 
705 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
706 template <typename T>
707 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>
708 clamp(T x, typename T::element_type minval,
709  typename T::element_type maxval) __NOEXC {
710  return __sycl_std::__invoke_s_clamp<T>(x, T(minval), T(maxval));
711 }
712 
713 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
714 template <typename T>
715 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>
716 clamp(T x, typename T::element_type minval,
717  typename T::element_type maxval) __NOEXC {
718  return __sycl_std::__invoke_u_clamp<T>(x, T(minval), T(maxval));
719 }
720 
721 // geninteger clz (geninteger x)
722 template <typename T>
724  return __sycl_std::__invoke_clz<T>(x);
725 }
726 
727 namespace ext {
728 namespace intel {
729 // geninteger ctz (geninteger x)
730 template <typename T>
732 ctz(T x) __NOEXC {
733  return __sycl_std::__invoke_ctz<T>(x);
734 }
735 } // namespace intel
736 } // namespace ext
737 
738 namespace __SYCL2020_DEPRECATED("use 'ext::intel' instead") intel {
739  using namespace ext::intel;
740 }
741 
742 // geninteger mad_hi (geninteger a, geninteger b, geninteger c)
743 template <typename T>
745  T z) __NOEXC {
746  return __sycl_std::__invoke_s_mad_hi<T>(x, y, z);
747 }
748 
749 // geninteger mad_hi (geninteger a, geninteger b, geninteger c)
750 template <typename T>
752  T z) __NOEXC {
753  return __sycl_std::__invoke_u_mad_hi<T>(x, y, z);
754 }
755 
756 // geninteger mad_sat (geninteger a, geninteger b, geninteger c)
757 template <typename T>
759  T c) __NOEXC {
760  return __sycl_std::__invoke_s_mad_sat<T>(a, b, c);
761 }
762 
763 // geninteger mad_sat (geninteger a, geninteger b, geninteger c)
764 template <typename T>
766  T c) __NOEXC {
767  return __sycl_std::__invoke_u_mad_sat<T>(a, b, c);
768 }
769 
770 // igeninteger max (igeninteger x, igeninteger y)
771 template <typename T>
772 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(max)(T x,
773  T y) __NOEXC {
774  return __sycl_std::__invoke_s_max<T>(x, y);
775 }
776 
777 // ugeninteger max (ugeninteger x, ugeninteger y)
778 template <typename T>
779 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>(max)(T x,
780  T y) __NOEXC {
781  return __sycl_std::__invoke_u_max<T>(x, y);
782 }
783 
784 // igeninteger max (vigeninteger x, sigeninteger y)
785 template <typename T>
786 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>(max)(
787  T x, typename T::element_type y) __NOEXC {
788  return __sycl_std::__invoke_s_max<T>(x, T(y));
789 }
790 
791 // vugeninteger max (vugeninteger x, sugeninteger y)
792 template <typename T>
793 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>(max)(
794  T x, typename T::element_type y) __NOEXC {
795  return __sycl_std::__invoke_u_max<T>(x, T(y));
796 }
797 
798 // igeninteger min (igeninteger x, igeninteger y)
799 template <typename T>
800 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(min)(T x,
801  T y) __NOEXC {
802  return __sycl_std::__invoke_s_min<T>(x, y);
803 }
804 
805 // ugeninteger min (ugeninteger x, ugeninteger y)
806 template <typename T>
807 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>(min)(T x,
808  T y) __NOEXC {
809  return __sycl_std::__invoke_u_min<T>(x, y);
810 }
811 
812 // vigeninteger min (vigeninteger x, sigeninteger y)
813 template <typename T>
814 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>(min)(
815  T x, typename T::element_type y) __NOEXC {
816  return __sycl_std::__invoke_s_min<T>(x, T(y));
817 }
818 
819 // vugeninteger min (vugeninteger x, sugeninteger y)
820 template <typename T>
821 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>(min)(
822  T x, typename T::element_type y) __NOEXC {
823  return __sycl_std::__invoke_u_min<T>(x, T(y));
824 }
825 
826 // geninteger mul_hi (geninteger x, geninteger y)
827 template <typename T>
829  T y) __NOEXC {
830  return __sycl_std::__invoke_s_mul_hi<T>(x, y);
831 }
832 
833 // geninteger mul_hi (geninteger x, geninteger y)
834 template <typename T>
836  T y) __NOEXC {
837  return __sycl_std::__invoke_u_mul_hi<T>(x, y);
838 }
839 
840 // geninteger rotate (geninteger v, geninteger i)
841 template <typename T>
843  T i) __NOEXC {
844  return __sycl_std::__invoke_rotate<T>(v, i);
845 }
846 
847 // geninteger sub_sat (geninteger x, geninteger y)
848 template <typename T>
850  T y) __NOEXC {
851  return __sycl_std::__invoke_s_sub_sat<T>(x, y);
852 }
853 
854 // geninteger sub_sat (geninteger x, geninteger y)
855 template <typename T>
857  T y) __NOEXC {
858  return __sycl_std::__invoke_u_sub_sat<T>(x, y);
859 }
860 
861 // ugeninteger16bit upsample (ugeninteger8bit hi, ugeninteger8bit lo)
862 template <typename T>
863 detail::enable_if_t<detail::is_ugeninteger8bit<T>::value,
864  detail::make_larger_t<T>>
865 upsample(T hi, T lo) __NOEXC {
866  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
867 }
868 
869 // igeninteger16bit upsample (igeninteger8bit hi, ugeninteger8bit lo)
870 template <typename T, typename T2>
871 detail::enable_if_t<detail::is_igeninteger8bit<T>::value &&
872  detail::is_ugeninteger8bit<T2>::value,
873  detail::make_larger_t<T>>
874 upsample(T hi, T2 lo) __NOEXC {
875  detail::check_vector_size<T, T2>();
876  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
877 }
878 
879 // ugeninteger32bit upsample (ugeninteger16bit hi, ugeninteger16bit lo)
880 template <typename T>
881 detail::enable_if_t<detail::is_ugeninteger16bit<T>::value,
882  detail::make_larger_t<T>>
883 upsample(T hi, T lo) __NOEXC {
884  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
885 }
886 
887 // igeninteger32bit upsample (igeninteger16bit hi, ugeninteger16bit lo)
888 template <typename T, typename T2>
889 detail::enable_if_t<detail::is_igeninteger16bit<T>::value &&
890  detail::is_ugeninteger16bit<T2>::value,
891  detail::make_larger_t<T>>
892 upsample(T hi, T2 lo) __NOEXC {
893  detail::check_vector_size<T, T2>();
894  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
895 }
896 
897 // ugeninteger64bit upsample (ugeninteger32bit hi, ugeninteger32bit lo)
898 template <typename T>
899 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value,
900  detail::make_larger_t<T>>
901 upsample(T hi, T lo) __NOEXC {
902  return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
903 }
904 
905 // igeninteger64bit upsample (igeninteger32bit hi, ugeninteger32bit lo)
906 template <typename T, typename T2>
907 detail::enable_if_t<detail::is_igeninteger32bit<T>::value &&
908  detail::is_ugeninteger32bit<T2>::value,
909  detail::make_larger_t<T>>
910 upsample(T hi, T2 lo) __NOEXC {
911  detail::check_vector_size<T, T2>();
912  return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
913 }
914 
915 // geninteger popcount (geninteger x)
916 template <typename T>
918  return __sycl_std::__invoke_popcount<T>(x);
919 }
920 
921 // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y,
922 // geninteger32bit z)
923 template <typename T>
924 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
925 mad24(T x, T y, T z) __NOEXC {
926  return __sycl_std::__invoke_s_mad24<T>(x, y, z);
927 }
928 
929 // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y,
930 // geninteger32bit z)
931 template <typename T>
932 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
933 mad24(T x, T y, T z) __NOEXC {
934  return __sycl_std::__invoke_u_mad24<T>(x, y, z);
935 }
936 
937 // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y)
938 template <typename T>
939 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
940 mul24(T x, T y) __NOEXC {
941  return __sycl_std::__invoke_s_mul24<T>(x, y);
942 }
943 
944 // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y)
945 template <typename T>
946 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
947 mul24(T x, T y) __NOEXC {
948  return __sycl_std::__invoke_u_mul24<T>(x, y);
949 }
950 
951 /* --------------- 4.13.6 Geometric Functions. ------------------------------*/
952 // float3 cross (float3 p0, float3 p1)
953 // float4 cross (float4 p0, float4 p1)
954 // double3 cross (double3 p0, double3 p1)
955 // double4 cross (double4 p0, double4 p1)
956 // half3 cross (half3 p0, half3 p1)
957 // half4 cross (half4 p0, half4 p1)
958 template <typename T>
960  T p1) __NOEXC {
961  return __sycl_std::__invoke_cross<T>(p0, p1);
962 }
963 
964 // float dot (float p0, float p1)
965 // double dot (double p0, double p1)
966 // half dot (half p0, half p1)
967 template <typename T>
969  return p0 * p1;
970 }
971 
972 // float dot (vgengeofloat p0, vgengeofloat p1)
973 template <typename T>
974 detail::enable_if_t<detail::is_vgengeofloat<T>::value, float>
975 dot(T p0, T p1) __NOEXC {
976  return __sycl_std::__invoke_Dot<float>(p0, p1);
977 }
978 
979 // double dot (vgengeodouble p0, vgengeodouble p1)
980 template <typename T>
981 detail::enable_if_t<detail::is_vgengeodouble<T>::value, double>
982 dot(T p0, T p1) __NOEXC {
983  return __sycl_std::__invoke_Dot<double>(p0, p1);
984 }
985 
986 // half dot (vgengeohalf p0, vgengeohalf p1)
987 template <typename T>
989  T p1) __NOEXC {
990  return __sycl_std::__invoke_Dot<half>(p0, p1);
991 }
992 
993 // float distance (gengeofloat p0, gengeofloat p1)
994 template <typename T,
995  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
996 float distance(T p0, T p1) __NOEXC {
997  return __sycl_std::__invoke_distance<float>(p0, p1);
998 }
999 
1000 // double distance (gengeodouble p0, gengeodouble p1)
1001 template <typename T,
1002  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1003 double distance(T p0, T p1) __NOEXC {
1004  return __sycl_std::__invoke_distance<double>(p0, p1);
1005 }
1006 
1007 // half distance (gengeohalf p0, gengeohalf p1)
1008 template <typename T,
1009  typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1010 half distance(T p0, T p1) __NOEXC {
1011  return __sycl_std::__invoke_distance<half>(p0, p1);
1012 }
1013 
1014 // float length (gengeofloat p)
1015 template <typename T,
1016  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1017 float length(T p) __NOEXC {
1018  return __sycl_std::__invoke_length<float>(p);
1019 }
1020 
1021 // double length (gengeodouble p)
1022 template <typename T,
1023  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1024 double length(T p) __NOEXC {
1025  return __sycl_std::__invoke_length<double>(p);
1026 }
1027 
1028 // half length (gengeohalf p)
1029 template <typename T,
1030  typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1032  return __sycl_std::__invoke_length<half>(p);
1033 }
1034 
1035 // gengeofloat normalize (gengeofloat p)
1036 template <typename T>
1037 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1039  return __sycl_std::__invoke_normalize<T>(p);
1040 }
1041 
1042 // gengeodouble normalize (gengeodouble p)
1043 template <typename T>
1044 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1046  return __sycl_std::__invoke_normalize<T>(p);
1047 }
1048 
1049 // gengeohalf normalize (gengeohalf p)
1050 template <typename T>
1052  return __sycl_std::__invoke_normalize<T>(p);
1053 }
1054 
1055 // float fast_distance (gengeofloat p0, gengeofloat p1)
1056 template <typename T,
1057  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1058 float fast_distance(T p0, T p1) __NOEXC {
1059  return __sycl_std::__invoke_fast_distance<float>(p0, p1);
1060 }
1061 
1062 // double fast_distance (gengeodouble p0, gengeodouble p1)
1063 template <typename T,
1064  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1065 double fast_distance(T p0, T p1) __NOEXC {
1066  return __sycl_std::__invoke_fast_distance<double>(p0, p1);
1067 }
1068 
1069 // float fast_length (gengeofloat p)
1070 template <typename T,
1071  typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1072 float fast_length(T p) __NOEXC {
1073  return __sycl_std::__invoke_fast_length<float>(p);
1074 }
1075 
1076 // double fast_length (gengeodouble p)
1077 template <typename T,
1078  typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1079 double fast_length(T p) __NOEXC {
1080  return __sycl_std::__invoke_fast_length<double>(p);
1081 }
1082 
1083 // gengeofloat fast_normalize (gengeofloat p)
1084 template <typename T>
1085 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1087  return __sycl_std::__invoke_fast_normalize<T>(p);
1088 }
1089 
1090 // gengeodouble fast_normalize (gengeodouble p)
1091 template <typename T>
1092 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1094  return __sycl_std::__invoke_fast_normalize<T>(p);
1095 }
1096 
1097 /* --------------- 4.13.7 Relational functions. Device version --------------*/
1098 // int isequal (half x, half y)
1099 // shortn isequal (halfn x, halfn y)
1100 // igeninteger32bit isequal (genfloatf x, genfloatf y)
1101 // int isequal (double x,double y);
1102 // longn isequal (doublen x, doublen y)
1103 template <typename T,
1104  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1107  __sycl_std::__invoke_FOrdEqual<detail::rel_ret_t<T>>(x, y));
1108 }
1109 
1110 // int isnotequal (half x, half y)
1111 // shortn isnotequal (halfn x, halfn y)
1112 // igeninteger32bit isnotequal (genfloatf x, genfloatf y)
1113 // int isnotequal (double x, double y)
1114 // longn isnotequal (doublen x, doublen y)
1115 template <typename T,
1116  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1119  __sycl_std::__invoke_FUnordNotEqual<detail::rel_ret_t<T>>(x, y));
1120 }
1121 
1122 // int isgreater (half x, half y)
1123 // shortn isgreater (halfn x, halfn y)
1124 // igeninteger32bit isgreater (genfloatf x, genfloatf y)
1125 // int isgreater (double x, double y)
1126 // longn isgreater (doublen x, doublen y)
1127 template <typename T,
1128  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1131  __sycl_std::__invoke_FOrdGreaterThan<detail::rel_ret_t<T>>(x, y));
1132 }
1133 
1134 // int isgreaterequal (half x, half y)
1135 // shortn isgreaterequal (halfn x, halfn y)
1136 // igeninteger32bit isgreaterequal (genfloatf x, genfloatf y)
1137 // int isgreaterequal (double x, double y)
1138 // longn isgreaterequal (doublen x, doublen y)
1139 template <typename T,
1140  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1143  __sycl_std::__invoke_FOrdGreaterThanEqual<detail::rel_ret_t<T>>(x, y));
1144 }
1145 
1146 // int isless (half x, half y)
1147 // shortn isless (halfn x, halfn y)
1148 // igeninteger32bit isless (genfloatf x, genfloatf y)
1149 // int isless (long x, long y)
1150 // longn isless (doublen x, doublen y)
1151 template <typename T,
1152  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1155  __sycl_std::__invoke_FOrdLessThan<detail::rel_ret_t<T>>(x, y));
1156 }
1157 
1158 // int islessequal (half x, half y)
1159 // shortn islessequal (halfn x, halfn y)
1160 // igeninteger32bit islessequal (genfloatf x, genfloatf y)
1161 // int islessequal (double x, double y)
1162 // longn islessequal (doublen x, doublen y)
1163 template <typename T,
1164  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1167  __sycl_std::__invoke_FOrdLessThanEqual<detail::rel_ret_t<T>>(x, y));
1168 }
1169 
1170 // int islessgreater (half x, half y)
1171 // shortn islessgreater (halfn x, halfn y)
1172 // igeninteger32bit islessgreater (genfloatf x, genfloatf y)
1173 // int islessgreater (double x, double y)
1174 // longn islessgreater (doublen x, doublen y)
1175 template <typename T,
1176  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1179  __sycl_std::__invoke_LessOrGreater<detail::rel_ret_t<T>>(x, y));
1180 }
1181 
1182 // int isfinite (half x)
1183 // shortn isfinite (halfn x)
1184 // igeninteger32bit isfinite (genfloatf x)
1185 // int isfinite (double x)
1186 // longn isfinite (doublen x)
1187 template <typename T,
1188  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1191  __sycl_std::__invoke_IsFinite<detail::rel_ret_t<T>>(x));
1192 }
1193 
1194 // int isinf (half x)
1195 // shortn isinf (halfn x)
1196 // igeninteger32bit isinf (genfloatf x)
1197 // int isinf (double x)
1198 // longn isinf (doublen x)
1199 template <typename T,
1200  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1203  __sycl_std::__invoke_IsInf<detail::rel_ret_t<T>>(x));
1204 }
1205 
1206 // int isnan (half x)
1207 // shortn isnan (halfn x)
1208 // igeninteger32bit isnan (genfloatf x)
1209 // int isnan (double x)
1210 // longn isnan (doublen x)
1211 template <typename T,
1212  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1215  __sycl_std::__invoke_IsNan<detail::rel_ret_t<T>>(x));
1216 }
1217 
1218 // int isnormal (half x)
1219 // shortn isnormal (halfn x)
1220 // igeninteger32bit isnormal (genfloatf x)
1221 // int isnormal (double x)
1222 // longn isnormal (doublen x)
1223 template <typename T,
1224  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1227  __sycl_std::__invoke_IsNormal<detail::rel_ret_t<T>>(x));
1228 }
1229 
1230 // int isordered (half x)
1231 // shortn isordered (halfn x, halfn y)
1232 // igeninteger32bit isordered (genfloatf x, genfloatf y)
1233 // int isordered (double x, double y)
1234 // longn isordered (doublen x, doublen y)
1235 template <typename T,
1236  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1239  __sycl_std::__invoke_Ordered<detail::rel_ret_t<T>>(x, y));
1240 }
1241 
1242 // int isunordered (half x, half y)
1243 // shortn isunordered (halfn x, halfn y)
1244 // igeninteger32bit isunordered (genfloatf x, genfloatf y)
1245 // int isunordered (double x, double y)
1246 // longn isunordered (doublen x, doublen y)
1247 template <typename T,
1248  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1251  __sycl_std::__invoke_Unordered<detail::rel_ret_t<T>>(x, y));
1252 }
1253 
1254 // int signbit (half x)
1255 // shortn signbit (halfn x)
1256 // igeninteger32bit signbit (genfloatf x)
1257 // int signbit (double)
1258 // longn signbit (doublen x)
1259 template <typename T,
1260  typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1263  __sycl_std::__invoke_SignBitSet<detail::rel_ret_t<T>>(x));
1264 }
1265 
1266 // int any (sigeninteger x)
1267 template <typename T>
1269  return detail::Boolean<1>(int(detail::msbIsSet(x)));
1270 }
1271 
1272 // int any (vigeninteger x)
1273 template <typename T>
1276  __sycl_std::__invoke_Any<detail::rel_sign_bit_test_ret_t<T>>(
1278 }
1279 
1280 // int all (sigeninteger x)
1281 template <typename T>
1283  return detail::Boolean<1>(int(detail::msbIsSet(x)));
1284 }
1285 
1286 // int all (vigeninteger x)
1287 template <typename T>
1290  __sycl_std::__invoke_All<detail::rel_sign_bit_test_ret_t<T>>(
1292 }
1293 
1294 // gentype bitselect (gentype a, gentype b, gentype c)
1295 template <typename T>
1297  T c) __NOEXC {
1298  return __sycl_std::__invoke_bitselect<T>(a, b, c);
1299 }
1300 
1301 // geninteger select (geninteger a, geninteger b, igeninteger c)
1302 template <typename T, typename T2>
1304  detail::is_geninteger<T>::value && detail::is_igeninteger<T2>::value, T>
1305 select(T a, T b, T2 c) __NOEXC {
1306  detail::check_vector_size<T, T2>();
1307  return __sycl_std::__invoke_select<T>(a, b, c);
1308 }
1309 
1310 // geninteger select (geninteger a, geninteger b, ugeninteger c)
1311 template <typename T, typename T2>
1313  detail::is_geninteger<T>::value && detail::is_ugeninteger<T2>::value, T>
1314 select(T a, T b, T2 c) __NOEXC {
1315  detail::check_vector_size<T, T2>();
1316  return __sycl_std::__invoke_select<T>(a, b, c);
1317 }
1318 
1319 // genfloatf select (genfloatf a, genfloatf b, genint c)
1320 template <typename T, typename T2>
1322  detail::is_genfloatf<T>::value && detail::is_genint<T2>::value, T>
1323 select(T a, T b, T2 c) __NOEXC {
1324  detail::check_vector_size<T, T2>();
1325  return __sycl_std::__invoke_select<T>(a, b, c);
1326 }
1327 
1328 // genfloatf select (genfloatf a, genfloatf b, ugenint c)
1329 template <typename T, typename T2>
1331  detail::is_genfloatf<T>::value && detail::is_ugenint<T2>::value, T>
1332 select(T a, T b, T2 c) __NOEXC {
1333  detail::check_vector_size<T, T2>();
1334  return __sycl_std::__invoke_select<T>(a, b, c);
1335 }
1336 
1337 // genfloatd select (genfloatd a, genfloatd b, igeninteger64 c)
1338 template <typename T, typename T2>
1340  detail::is_genfloatd<T>::value && detail::is_igeninteger64bit<T2>::value, T>
1341 select(T a, T b, T2 c) __NOEXC {
1342  detail::check_vector_size<T, T2>();
1343  return __sycl_std::__invoke_select<T>(a, b, c);
1344 }
1345 
1346 // genfloatd select (genfloatd a, genfloatd b, ugeninteger64 c)
1347 template <typename T, typename T2>
1349  detail::is_genfloatd<T>::value && detail::is_ugeninteger64bit<T2>::value, T>
1350 select(T a, T b, T2 c) __NOEXC {
1351  detail::check_vector_size<T, T2>();
1352  return __sycl_std::__invoke_select<T>(a, b, c);
1353 }
1354 
1355 // genfloath select (genfloath a, genfloath b, igeninteger16 c)
1356 template <typename T, typename T2>
1358  detail::is_genfloath<T>::value && detail::is_igeninteger16bit<T2>::value, T>
1359 select(T a, T b, T2 c) __NOEXC {
1360  detail::check_vector_size<T, T2>();
1361  return __sycl_std::__invoke_select<T>(a, b, c);
1362 }
1363 
1364 // genfloath select (genfloath a, genfloath b, ugeninteger16 c)
1365 template <typename T, typename T2>
1367  detail::is_genfloath<T>::value && detail::is_ugeninteger16bit<T2>::value, T>
1368 select(T a, T b, T2 c) __NOEXC {
1369  detail::check_vector_size<T, T2>();
1370  return __sycl_std::__invoke_select<T>(a, b, c);
1371 }
1372 
1373 namespace native {
1374 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1375 // genfloatf cos (genfloatf x)
1376 template <typename T>
1378  return __sycl_std::__invoke_native_cos<T>(x);
1379 }
1380 
1381 // genfloatf divide (genfloatf x, genfloatf y)
1382 template <typename T>
1384  T y) __NOEXC {
1385  return __sycl_std::__invoke_native_divide<T>(x, y);
1386 }
1387 
1388 // genfloatf exp (genfloatf x)
1389 template <typename T>
1391  return __sycl_std::__invoke_native_exp<T>(x);
1392 }
1393 
1394 // genfloatf exp2 (genfloatf x)
1395 template <typename T>
1397  return __sycl_std::__invoke_native_exp2<T>(x);
1398 }
1399 
1400 // genfloatf exp10 (genfloatf x)
1401 template <typename T>
1403  return __sycl_std::__invoke_native_exp10<T>(x);
1404 }
1405 
1406 // genfloatf log (genfloatf x)
1407 template <typename T>
1409  return __sycl_std::__invoke_native_log<T>(x);
1410 }
1411 
1412 // genfloatf log2 (genfloatf x)
1413 template <typename T>
1415  return __sycl_std::__invoke_native_log2<T>(x);
1416 }
1417 
1418 // genfloatf log10 (genfloatf x)
1419 template <typename T>
1421  return __sycl_std::__invoke_native_log10<T>(x);
1422 }
1423 
1424 // genfloatf powr (genfloatf x, genfloatf y)
1425 template <typename T>
1427  return __sycl_std::__invoke_native_powr<T>(x, y);
1428 }
1429 
1430 // genfloatf recip (genfloatf x)
1431 template <typename T>
1433  return __sycl_std::__invoke_native_recip<T>(x);
1434 }
1435 
1436 // genfloatf rsqrt (genfloatf x)
1437 template <typename T>
1439  return __sycl_std::__invoke_native_rsqrt<T>(x);
1440 }
1441 
1442 // genfloatf sin (genfloatf x)
1443 template <typename T>
1445  return __sycl_std::__invoke_native_sin<T>(x);
1446 }
1447 
1448 // genfloatf sqrt (genfloatf x)
1449 template <typename T>
1451  return __sycl_std::__invoke_native_sqrt<T>(x);
1452 }
1453 
1454 // genfloatf tan (genfloatf x)
1455 template <typename T>
1457  return __sycl_std::__invoke_native_tan<T>(x);
1458 }
1459 
1460 } // namespace native
1461 namespace half_precision {
1462 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1463 // genfloatf cos (genfloatf x)
1464 template <typename T>
1466  return __sycl_std::__invoke_half_cos<T>(x);
1467 }
1468 
1469 // genfloatf divide (genfloatf x, genfloatf y)
1470 template <typename T>
1472  T y) __NOEXC {
1473  return __sycl_std::__invoke_half_divide<T>(x, y);
1474 }
1475 
1476 // genfloatf exp (genfloatf x)
1477 template <typename T>
1479  return __sycl_std::__invoke_half_exp<T>(x);
1480 }
1481 
1482 // genfloatf exp2 (genfloatf x)
1483 template <typename T>
1485  return __sycl_std::__invoke_half_exp2<T>(x);
1486 }
1487 
1488 // genfloatf exp10 (genfloatf x)
1489 template <typename T>
1491  return __sycl_std::__invoke_half_exp10<T>(x);
1492 }
1493 
1494 // genfloatf log (genfloatf x)
1495 template <typename T>
1497  return __sycl_std::__invoke_half_log<T>(x);
1498 }
1499 
1500 // genfloatf log2 (genfloatf x)
1501 template <typename T>
1503  return __sycl_std::__invoke_half_log2<T>(x);
1504 }
1505 
1506 // genfloatf log10 (genfloatf x)
1507 template <typename T>
1509  return __sycl_std::__invoke_half_log10<T>(x);
1510 }
1511 
1512 // genfloatf powr (genfloatf x, genfloatf y)
1513 template <typename T>
1515  return __sycl_std::__invoke_half_powr<T>(x, y);
1516 }
1517 
1518 // genfloatf recip (genfloatf x)
1519 template <typename T>
1521  return __sycl_std::__invoke_half_recip<T>(x);
1522 }
1523 
1524 // genfloatf rsqrt (genfloatf x)
1525 template <typename T>
1527  return __sycl_std::__invoke_half_rsqrt<T>(x);
1528 }
1529 
1530 // genfloatf sin (genfloatf x)
1531 template <typename T>
1533  return __sycl_std::__invoke_half_sin<T>(x);
1534 }
1535 
1536 // genfloatf sqrt (genfloatf x)
1537 template <typename T>
1539  return __sycl_std::__invoke_half_sqrt<T>(x);
1540 }
1541 
1542 // genfloatf tan (genfloatf x)
1543 template <typename T>
1545  return __sycl_std::__invoke_half_tan<T>(x);
1546 }
1547 
1548 } // namespace half_precision
1549 } // namespace sycl
1550 } // __SYCL_INLINE_NAMESPACE(cl)
1551 
1552 #ifdef __SYCL_DEVICE_ONLY__
1553 extern "C" {
1554 extern SYCL_EXTERNAL int abs(int x);
1555 extern SYCL_EXTERNAL long int labs(long int x);
1556 extern SYCL_EXTERNAL long long int llabs(long long int x);
1557 
1558 extern SYCL_EXTERNAL div_t div(int x, int y);
1559 extern SYCL_EXTERNAL ldiv_t ldiv(long int x, long int y);
1560 extern SYCL_EXTERNAL lldiv_t lldiv(long long int x, long long int y);
1561 extern SYCL_EXTERNAL float scalbnf(float x, int n);
1562 extern SYCL_EXTERNAL double scalbn(double x, int n);
1563 extern SYCL_EXTERNAL float logf(float x);
1564 extern SYCL_EXTERNAL double log(double x);
1565 extern SYCL_EXTERNAL float expf(float x);
1566 extern SYCL_EXTERNAL double exp(double x);
1567 extern SYCL_EXTERNAL float log10f(float x);
1568 extern SYCL_EXTERNAL double log10(double x);
1569 extern SYCL_EXTERNAL float modff(float x, float *intpart);
1570 extern SYCL_EXTERNAL double modf(double x, double *intpart);
1571 extern SYCL_EXTERNAL float exp2f(float x);
1572 extern SYCL_EXTERNAL double exp2(double x);
1573 extern SYCL_EXTERNAL float expm1f(float x);
1574 extern SYCL_EXTERNAL double expm1(double x);
1575 extern SYCL_EXTERNAL int ilogbf(float x);
1576 extern SYCL_EXTERNAL int ilogb(double x);
1577 extern SYCL_EXTERNAL float log1pf(float x);
1578 extern SYCL_EXTERNAL double log1p(double x);
1579 extern SYCL_EXTERNAL float log2f(float x);
1580 extern SYCL_EXTERNAL double log2(double x);
1581 extern SYCL_EXTERNAL float logbf(float x);
1582 extern SYCL_EXTERNAL double logb(double x);
1583 extern SYCL_EXTERNAL float sqrtf(float x);
1584 extern SYCL_EXTERNAL double sqrt(double x);
1585 extern SYCL_EXTERNAL float cbrtf(float x);
1586 extern SYCL_EXTERNAL double cbrt(double x);
1587 extern SYCL_EXTERNAL float erff(float x);
1588 extern SYCL_EXTERNAL double erf(double x);
1589 extern SYCL_EXTERNAL float erfcf(float x);
1590 extern SYCL_EXTERNAL double erfc(double x);
1591 extern SYCL_EXTERNAL float tgammaf(float x);
1592 extern SYCL_EXTERNAL double tgamma(double x);
1593 extern SYCL_EXTERNAL float lgammaf(float x);
1594 extern SYCL_EXTERNAL double lgamma(double x);
1595 extern SYCL_EXTERNAL float fmodf(float x, float y);
1596 extern SYCL_EXTERNAL double fmod(double x, double y);
1597 extern SYCL_EXTERNAL float remainderf(float x, float y);
1598 extern SYCL_EXTERNAL double remainder(double x, double y);
1599 extern SYCL_EXTERNAL float remquof(float x, float y, int *q);
1600 extern SYCL_EXTERNAL double remquo(double x, double y, int *q);
1601 extern SYCL_EXTERNAL float nextafterf(float x, float y);
1602 extern SYCL_EXTERNAL double nextafter(double x, double y);
1603 extern SYCL_EXTERNAL float fdimf(float x, float y);
1604 extern SYCL_EXTERNAL double fdim(double x, double y);
1605 extern SYCL_EXTERNAL float fmaf(float x, float y, float z);
1606 extern SYCL_EXTERNAL double fma(double x, double y, double z);
1607 extern SYCL_EXTERNAL float sinf(float x);
1608 extern SYCL_EXTERNAL double sin(double x);
1609 extern SYCL_EXTERNAL float cosf(float x);
1610 extern SYCL_EXTERNAL double cos(double x);
1611 extern SYCL_EXTERNAL float tanf(float x);
1612 extern SYCL_EXTERNAL double tan(double x);
1613 extern SYCL_EXTERNAL float asinf(float x);
1614 extern SYCL_EXTERNAL double asin(double x);
1615 extern SYCL_EXTERNAL float acosf(float x);
1616 extern SYCL_EXTERNAL double acos(double x);
1617 extern SYCL_EXTERNAL float atanf(float x);
1618 extern SYCL_EXTERNAL double atan(double x);
1619 extern SYCL_EXTERNAL float powf(float x, float y);
1620 extern SYCL_EXTERNAL double pow(double x, double y);
1621 extern SYCL_EXTERNAL float atan2f(float x, float y);
1622 extern SYCL_EXTERNAL double atan2(double x, double y);
1623 
1624 extern SYCL_EXTERNAL float sinhf(float x);
1625 extern SYCL_EXTERNAL double sinh(double x);
1626 extern SYCL_EXTERNAL float coshf(float x);
1627 extern SYCL_EXTERNAL double cosh(double x);
1628 extern SYCL_EXTERNAL float tanhf(float x);
1629 extern SYCL_EXTERNAL double tanh(double x);
1630 extern SYCL_EXTERNAL float asinhf(float x);
1631 extern SYCL_EXTERNAL double asinh(double x);
1632 extern SYCL_EXTERNAL float acoshf(float x);
1633 extern SYCL_EXTERNAL double acosh(double x);
1634 extern SYCL_EXTERNAL float atanhf(float x);
1635 extern SYCL_EXTERNAL double atanh(double x);
1636 extern SYCL_EXTERNAL double frexp(double x, int *exp);
1637 extern SYCL_EXTERNAL double ldexp(double x, int exp);
1638 extern SYCL_EXTERNAL double hypot(double x, double y);
1639 
1640 extern SYCL_EXTERNAL void *memcpy(void *dest, const void *src, size_t n);
1641 extern SYCL_EXTERNAL void *memset(void *dest, int c, size_t n);
1642 extern SYCL_EXTERNAL int memcmp(const void *s1, const void *s2, size_t n);
1643 }
1644 #ifdef __GLIBC__
1645 extern "C" {
1646 extern SYCL_EXTERNAL void __assert_fail(const char *expr, const char *file,
1647  unsigned int line, const char *func);
1648 extern SYCL_EXTERNAL float frexpf(float x, int *exp);
1649 extern SYCL_EXTERNAL float ldexpf(float x, int exp);
1650 extern SYCL_EXTERNAL float hypotf(float x, float y);
1651 
1652 // MS UCRT supports most of the C standard library but <complex.h> is
1653 // an exception.
1654 extern SYCL_EXTERNAL float cimagf(float __complex__ z);
1655 extern SYCL_EXTERNAL double cimag(double __complex__ z);
1656 extern SYCL_EXTERNAL float crealf(float __complex__ z);
1657 extern SYCL_EXTERNAL double creal(double __complex__ z);
1658 extern SYCL_EXTERNAL float cargf(float __complex__ z);
1659 extern SYCL_EXTERNAL double carg(double __complex__ z);
1660 extern SYCL_EXTERNAL float cabsf(float __complex__ z);
1661 extern SYCL_EXTERNAL double cabs(double __complex__ z);
1662 extern SYCL_EXTERNAL float __complex__ cprojf(float __complex__ z);
1663 extern SYCL_EXTERNAL double __complex__ cproj(double __complex__ z);
1664 extern SYCL_EXTERNAL float __complex__ cexpf(float __complex__ z);
1665 extern SYCL_EXTERNAL double __complex__ cexp(double __complex__ z);
1666 extern SYCL_EXTERNAL float __complex__ clogf(float __complex__ z);
1667 extern SYCL_EXTERNAL double __complex__ clog(double __complex__ z);
1668 extern SYCL_EXTERNAL float __complex__ cpowf(float __complex__ z);
1669 extern SYCL_EXTERNAL double __complex__ cpow(double __complex__ z);
1670 extern SYCL_EXTERNAL float __complex__ csqrtf(float __complex__ z);
1671 extern SYCL_EXTERNAL double __complex__ csqrt(double __complex__ z);
1672 extern SYCL_EXTERNAL float __complex__ csinhf(float __complex__ z);
1673 extern SYCL_EXTERNAL double __complex__ csinh(double __complex__ z);
1674 extern SYCL_EXTERNAL float __complex__ ccoshf(float __complex__ z);
1675 extern SYCL_EXTERNAL double __complex__ ccosh(double __complex__ z);
1676 extern SYCL_EXTERNAL float __complex__ ctanhf(float __complex__ z);
1677 extern SYCL_EXTERNAL double __complex__ ctanh(double __complex__ z);
1678 extern SYCL_EXTERNAL float __complex__ csinf(float __complex__ z);
1679 extern SYCL_EXTERNAL double __complex__ csin(double __complex__ z);
1680 extern SYCL_EXTERNAL float __complex__ ccosf(float __complex__ z);
1681 extern SYCL_EXTERNAL double __complex__ ccos(double __complex__ z);
1682 extern SYCL_EXTERNAL float __complex__ ctanf(float __complex__ z);
1683 extern SYCL_EXTERNAL double __complex__ ctan(double __complex__ z);
1684 extern SYCL_EXTERNAL float __complex__ cacosf(float __complex__ z);
1685 extern SYCL_EXTERNAL double __complex__ cacos(double __complex__ z);
1686 extern SYCL_EXTERNAL float __complex__ cacoshf(float __complex__ z);
1687 extern SYCL_EXTERNAL double __complex__ cacosh(double __complex__ z);
1688 extern SYCL_EXTERNAL float __complex__ casinf(float __complex__ z);
1689 extern SYCL_EXTERNAL double __complex__ casin(double __complex__ z);
1690 extern SYCL_EXTERNAL float __complex__ casinhf(float __complex__ z);
1691 extern SYCL_EXTERNAL double __complex__ casinh(double __complex__ z);
1692 extern SYCL_EXTERNAL float __complex__ catanf(float __complex__ z);
1693 extern SYCL_EXTERNAL double __complex__ catan(double __complex__ z);
1694 extern SYCL_EXTERNAL float __complex__ catanhf(float __complex__ z);
1695 extern SYCL_EXTERNAL double __complex__ catanh(double __complex__ z);
1696 extern SYCL_EXTERNAL float __complex__ cpolarf(float rho, float theta);
1697 extern SYCL_EXTERNAL double __complex__ cpolar(double rho, double theta);
1698 extern SYCL_EXTERNAL float __complex__ __mulsc3(float a, float b, float c,
1699  float d);
1700 extern SYCL_EXTERNAL double __complex__ __muldc3(double a, double b, double c,
1701  double d);
1702 extern SYCL_EXTERNAL float __complex__ __divsc3(float a, float b, float c,
1703  float d);
1704 extern SYCL_EXTERNAL double __complex__ __divdc3(float a, float b, float c,
1705  float d);
1706 }
1707 #elif defined(_WIN32)
1708 extern "C" {
1709 // TODO: documented C runtime library APIs must be recognized as
1710 // builtins by FE. This includes _dpcomp, _dsign, _dtest,
1711 // _fdpcomp, _fdsign, _fdtest, _hypotf, _wassert.
1712 // APIs used by STL, such as _Cosh, are undocumented, even though
1713 // they are open-sourced. Recognizing them as builtins is not
1714 // straightforward currently.
1715 extern SYCL_EXTERNAL double _Cosh(double x, double y);
1716 extern SYCL_EXTERNAL int _dpcomp(double x, double y);
1717 extern SYCL_EXTERNAL int _dsign(double x);
1718 extern SYCL_EXTERNAL short _Dtest(double *px);
1719 extern SYCL_EXTERNAL short _dtest(double *px);
1720 extern SYCL_EXTERNAL short _Exp(double *px, double y, short eoff);
1721 extern SYCL_EXTERNAL float _FCosh(float x, float y);
1722 extern SYCL_EXTERNAL int _fdpcomp(float x, float y);
1723 extern SYCL_EXTERNAL int _fdsign(float x);
1724 extern SYCL_EXTERNAL short _FDtest(float *px);
1725 extern SYCL_EXTERNAL short _fdtest(float *px);
1726 extern SYCL_EXTERNAL short _FExp(float *px, float y, short eoff);
1727 extern SYCL_EXTERNAL float _FSinh(float x, float y);
1728 extern SYCL_EXTERNAL double _Sinh(double x, double y);
1729 extern SYCL_EXTERNAL float _hypotf(float x, float y);
1730 extern SYCL_EXTERNAL void _wassert(const wchar_t *wexpr, const wchar_t *wfile,
1731  unsigned line);
1732 }
1733 #endif
1734 #endif // __SYCL_DEVICE_ONLY__
1735 
1736 #undef __NOEXC
cl::sycl::maxmag
detail::enable_if_t< detail::is_genfloat< T >::value, T > maxmag(T x, T y) __NOEXC
Definition: builtins.hpp:336
cl::sycl::mul24
detail::enable_if_t< detail::is_ugeninteger32bit< T >::value, T > mul24(T x, T y) __NOEXC
Definition: builtins.hpp:947
cl::sycl::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:379
cl::sycl::fast_normalize
detail::enable_if_t< detail::is_gengeodouble< T >::value, T > fast_normalize(T p) __NOEXC
Definition: builtins.hpp:1093
cl::sycl::tanh
detail::enable_if_t< detail::is_genfloat< T >::value, T > tanh(T x) __NOEXC
Definition: builtins.hpp:474
cl::sycl::mul_hi
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > mul_hi(T x, T y) __NOEXC
Definition: builtins.hpp:835
cl::sycl::popcount
detail::enable_if_t< detail::is_geninteger< T >::value, T > popcount(T x) __NOEXC
Definition: builtins.hpp:917
cl::sycl::half_precision::log
detail::enable_if_t< detail::is_genfloatf< T >::value, T > log(T x) __NOEXC
Definition: builtins.hpp:1496
cl::sycl::half_precision::rsqrt
detail::enable_if_t< detail::is_genfloatf< T >::value, T > rsqrt(T x) __NOEXC
Definition: builtins.hpp:1526
cl::sycl::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:350
cl::sycl::isnotequal
detail::common_rel_ret_t< T > isnotequal(T x, T y) __NOEXC
Definition: builtins.hpp:1117
T
cl::sycl::cbrt
detail::enable_if_t< detail::is_genfloat< T >::value, T > cbrt(T x) __NOEXC
Definition: builtins.hpp:98
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:56
cl::sycl::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:645
cl::sycl::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:230
cl::sycl::fmax
detail::enable_if_t< detail::is_vgenfloat< T >::value, T > fmax(T x, typename T::element_type y) __NOEXC
Definition: builtins.hpp:203
cl::sycl::expm1
detail::enable_if_t< detail::is_genfloat< T >::value, T > expm1(T x) __NOEXC
Definition: builtins.hpp:165
cl::sycl::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:1368
cl::sycl::nextafter
detail::enable_if_t< detail::is_genfloat< T >::value, T > nextafter(T x, T y) __NOEXC
Definition: builtins.hpp:364
cl::sycl::cospi
detail::enable_if_t< detail::is_genfloat< T >::value, T > cospi(T x) __NOEXC
Definition: builtins.hpp:129
SYCL_EXTERNAL
#define SYCL_EXTERNAL
Definition: defines_elementary.hpp:34
cl::sycl::any
detail::enable_if_t< detail::is_vigeninteger< T >::value, int > any(T x) __NOEXC
Definition: builtins.hpp:1274
cl::sycl::detail::nan_return_t
typename nan_types< T, T >::ret_type nan_return_t
Definition: generic_type_traits.hpp:260
cl::sycl::tanpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > tanpi(T x) __NOEXC
Definition: builtins.hpp:480
cl::sycl::detail::msbIsSet
constexpr bool msbIsSet(const T x)
Definition: generic_type_traits.hpp:500
cl::sycl::atan2
detail::enable_if_t< detail::is_genfloat< T >::value, T > atan2(T y, T x) __NOEXC
Definition: builtins.hpp:73
cl::sycl::trunc
detail::enable_if_t< detail::is_genfloat< T >::value, T > trunc(T x) __NOEXC
Definition: builtins.hpp:492
cl::sycl::lgamma
detail::enable_if_t< detail::is_genfloat< T >::value, T > lgamma(T x) __NOEXC
Definition: builtins.hpp:284
cl::sycl::dot
detail::enable_if_t< detail::is_vgengeohalf< T >::value, half > dot(T p0, T p1) __NOEXC
Definition: builtins.hpp:988
cl::sycl::half_precision::exp2
detail::enable_if_t< detail::is_genfloatf< T >::value, T > exp2(T x) __NOEXC
Definition: builtins.hpp:1484
cl::sycl::half_precision::sqrt
detail::enable_if_t< detail::is_genfloatf< T >::value, T > sqrt(T x) __NOEXC
Definition: builtins.hpp:1538
cl::sycl::asinh
detail::enable_if_t< detail::is_genfloat< T >::value, T > asinh(T x) __NOEXC
Definition: builtins.hpp:55
cl::sycl::detail::rel_sign_bit_test_ret_t
typename RelationalTestForSignBitType< T >::return_type rel_sign_bit_test_ret_t
Definition: generic_type_traits.hpp:549
cl::sycl::mad
detail::enable_if_t< detail::is_genfloat< T >::value, T > mad(T a, T b, T c) __NOEXC
Definition: builtins.hpp:329
builtins.hpp
cl::sycl::normalize
detail::enable_if_t< detail::is_gengeohalf< T >::value, T > normalize(T p) __NOEXC
Definition: builtins.hpp:1051
cl::sycl::atan2pi
detail::enable_if_t< detail::is_genfloat< T >::value, T > atan2pi(T y, T x) __NOEXC
Definition: builtins.hpp:91
cl::sycl::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:401
cl::sycl::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:292
cl::sycl::cross
detail::enable_if_t< detail::is_gencross< T >::value, T > cross(T p0, T p1) __NOEXC
Definition: builtins.hpp:959
cl::sycl::atan
detail::enable_if_t< detail::is_genfloat< T >::value, T > atan(T y_over_x) __NOEXC
Definition: builtins.hpp:67
cl::sycl::fast_distance
double fast_distance(T p0, T p1) __NOEXC
Definition: builtins.hpp:1065
cl::sycl::isordered
detail::common_rel_ret_t< T > isordered(T x, T y) __NOEXC
Definition: builtins.hpp:1237
cl::sycl::half_precision::log10
detail::enable_if_t< detail::is_genfloatf< T >::value, T > log10(T x) __NOEXC
Definition: builtins.hpp:1508
cl::sycl::minmag
detail::enable_if_t< detail::is_genfloat< T >::value, T > minmag(T x, T y) __NOEXC
Definition: builtins.hpp:342
cl::sycl::detail::convert_data_type_impl
Definition: generic_type_traits.hpp:276
cl::sycl::floor
detail::enable_if_t< detail::is_genfloat< T >::value, T > floor(T x) __NOEXC
Definition: builtins.hpp:183
cl::sycl::logb
detail::enable_if_t< detail::is_genfloat< T >::value, T > logb(T x) __NOEXC
Definition: builtins.hpp:323
cl::sycl::step
detail::enable_if_t< detail::is_vgenfloat< T >::value, T > step(typename T::element_type edge, T x) __NOEXC
Definition: builtins.hpp:592
cl::sycl::clamp
detail::enable_if_t< detail::is_vugeninteger< T >::value, T > clamp(T x, typename T::element_type minval, typename T::element_type maxval) __NOEXC
Definition: builtins.hpp:716
cl::sycl::log1p
detail::enable_if_t< detail::is_genfloat< T >::value, T > log1p(T x) __NOEXC
Definition: builtins.hpp:317
cl::sycl::isnan
detail::common_rel_ret_t< T > isnan(T x) __NOEXC
Definition: builtins.hpp:1213
cl::sycl::fdim
detail::enable_if_t< detail::is_genfloat< T >::value, T > fdim(T x, T y) __NOEXC
Definition: builtins.hpp:177
cl::sycl::add_sat
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > add_sat(T x, T y) __NOEXC
Definition: builtins.hpp:658
cl::sycl::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:416
cl::sycl::acos
detail::enable_if_t< detail::is_genfloat< T >::value, T > acos(T x) __NOEXC
Definition: builtins.hpp:31
cl::sycl::detail::memcpy
void memcpy(void *Dst, const void *Src, std::size_t Size)
cl::sycl::ext::intel::experimental::esimd::atomic_op::min
@ min
cl::sycl::half_precision::log2
detail::enable_if_t< detail::is_genfloatf< T >::value, T > log2(T x) __NOEXC
Definition: builtins.hpp:1502
cl::sycl::fast_length
double fast_length(T p) __NOEXC
Definition: builtins.hpp:1079
cl::sycl::isequal
detail::common_rel_ret_t< T > isequal(T x, T y) __NOEXC
Definition: builtins.hpp:1105
cl::sycl::detail::Boolean< 1 >
Definition: boolean.hpp:103
cl::sycl::half_precision::tan
detail::enable_if_t< detail::is_genfloatf< T >::value, T > tan(T x) __NOEXC
Definition: builtins.hpp:1544
cl::sycl::detail::half_impl::half
Definition: half_type.hpp:319
cl::sycl::fabs
detail::enable_if_t< detail::is_genfloat< T >::value, T > fabs(T x) __NOEXC
Definition: builtins.hpp:171
cl::sycl::half_precision::recip
detail::enable_if_t< detail::is_genfloatf< T >::value, T > recip(T x) __NOEXC
Definition: builtins.hpp:1520
cl::sycl::ext::intel::experimental::esimd::atomic_op::max
@ max
cl::sycl::radians
detail::enable_if_t< detail::is_genfloat< T >::value, T > radians(T degrees) __NOEXC
Definition: builtins.hpp:576
cl::sycl::fma
detail::enable_if_t< detail::is_genfloat< T >::value, T > fma(T a, T b, T c) __NOEXC
Definition: builtins.hpp:189
generic_type_traits.hpp
cl::sycl::erf
detail::enable_if_t< detail::is_genfloat< T >::value, T > erf(T x) __NOEXC
Definition: builtins.hpp:141
cl::sycl::islessgreater
detail::common_rel_ret_t< T > islessgreater(T x, T y) __NOEXC
Definition: builtins.hpp:1177
cl::sycl::distance
half distance(T p0, T p1) __NOEXC
Definition: builtins.hpp:1010
cl::sycl::isgreaterequal
detail::common_rel_ret_t< T > isgreaterequal(T x, T y) __NOEXC
Definition: builtins.hpp:1141
cl::sycl::acospi
detail::enable_if_t< detail::is_genfloat< T >::value, T > acospi(T x) __NOEXC
Definition: builtins.hpp:43
cl::sycl::length
half length(T p) __NOEXC
Definition: builtins.hpp:1031
cl::sycl::hypot
detail::enable_if_t< detail::is_genfloat< T >::value, T > hypot(T x, T y) __NOEXC
Definition: builtins.hpp:246
cl::sycl::isunordered
detail::common_rel_ret_t< T > isunordered(T x, T y) __NOEXC
Definition: builtins.hpp:1249
__NOEXC
#define __NOEXC
Definition: builtins.hpp:18
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::signbit
detail::common_rel_ret_t< T > signbit(T x) __NOEXC
Definition: builtins.hpp:1261
cl::sycl::sign
detail::enable_if_t< detail::is_genfloat< T >::value, T > sign(T x) __NOEXC
Definition: builtins.hpp:615
cl::sycl::half_precision::exp
detail::enable_if_t< detail::is_genfloatf< T >::value, T > exp(T x) __NOEXC
Definition: builtins.hpp:1478
cl::sycl::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:506
cl::sycl::hadd
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > hadd(T x, T y) __NOEXC
Definition: builtins.hpp:672
cl::sycl::isless
detail::common_rel_ret_t< T > isless(T x, T y) __NOEXC
Definition: builtins.hpp:1153
cl::sycl::half_precision::divide
detail::enable_if_t< detail::is_genfloatf< T >::value, T > divide(T x, T y) __NOEXC
Definition: builtins.hpp:1471
cl::sycl::degrees
detail::enable_if_t< detail::is_genfloat< T >::value, T > degrees(T radians) __NOEXC
Definition: builtins.hpp:517
cl::sycl::isinf
detail::common_rel_ret_t< T > isinf(T x) __NOEXC
Definition: builtins.hpp:1201
cl::sycl::cosh
detail::enable_if_t< detail::is_genfloat< T >::value, T > cosh(T x) __NOEXC
Definition: builtins.hpp:123
cl::sycl::mad24
detail::enable_if_t< detail::is_ugeninteger32bit< T >::value, T > mad24(T x, T y, T z) __NOEXC
Definition: builtins.hpp:933
cl::sycl::isgreater
detail::common_rel_ret_t< T > isgreater(T x, T y) __NOEXC
Definition: builtins.hpp:1129
cl::sycl::image_channel_order::a
@ a
cl::sycl::fmod
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmod(T x, T y) __NOEXC
Definition: builtins.hpp:222
cl::sycl::detail::change_base_type_t
typename change_base_type< T, B >::type change_base_type_t
Definition: type_traits.hpp:110
boolean.hpp
cl::sycl::rhadd
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > rhadd(T x, T y) __NOEXC
Definition: builtins.hpp:686
cl::sycl::ceil
detail::enable_if_t< detail::is_genfloat< T >::value, T > ceil(T x) __NOEXC
Definition: builtins.hpp:104
cl::sycl::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:277
cl::sycl::asinpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > asinpi(T x) __NOEXC
Definition: builtins.hpp:61
cl::sycl::detail::RelConverter
Definition: generic_type_traits.hpp:555
cl::sycl::pow
detail::enable_if_t< detail::is_genfloat< T >::value, T > pow(T x, T y) __NOEXC
Definition: builtins.hpp:371
cl::sycl::sub_sat
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > sub_sat(T x, T y) __NOEXC
Definition: builtins.hpp:856
cl::sycl::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:18
cl::sycl::ext::intel::ctz
sycl::detail::enable_if_t< sycl::detail::is_geninteger< T >::value, T > ctz(T x) __NOEXC
Definition: builtins.hpp:732
cl::sycl::fmin
detail::enable_if_t< detail::is_vgenfloat< T >::value, T > fmin(T x, typename T::element_type y) __NOEXC
Definition: builtins.hpp:216
cl::sycl::rint
detail::enable_if_t< detail::is_genfloat< T >::value, T > rint(T x) __NOEXC
Definition: builtins.hpp:408
cl::sycl::sinh
detail::enable_if_t< detail::is_genfloat< T >::value, T > sinh(T x) __NOEXC
Definition: builtins.hpp:450
cl::sycl::detail::rel_ret_t
typename RelationalReturnType< T >::type rel_ret_t
Definition: generic_type_traits.hpp:534
cl::sycl::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:569
cl::sycl::copysign
detail::enable_if_t< detail::is_genfloat< T >::value, T > copysign(T x, T y) __NOEXC
Definition: builtins.hpp:110
cl::sycl::abs
detail::enable_if_t< detail::is_igeninteger< T >::value, detail::make_unsigned_t< T > > abs(T x) __NOEXC
Definition: builtins.hpp:630
cl::sycl::sinpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > sinpi(T x) __NOEXC
Definition: builtins.hpp:456
cl::sycl::acosh
detail::enable_if_t< detail::is_genfloat< T >::value, T > acosh(T x) __NOEXC
Definition: builtins.hpp:37
cl::sycl::atanpi
detail::enable_if_t< detail::is_genfloat< T >::value, T > atanpi(T x) __NOEXC
Definition: builtins.hpp:85
cl::sycl::mad_hi
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > mad_hi(T x, T y, T z) __NOEXC
Definition: builtins.hpp:751
cl::sycl::tgamma
detail::enable_if_t< detail::is_genfloat< T >::value, T > tgamma(T x) __NOEXC
Definition: builtins.hpp:486
cl::sycl::atanh
detail::enable_if_t< detail::is_genfloat< T >::value, T > atanh(T x) __NOEXC
Definition: builtins.hpp:79
cl::sycl::ilogb
detail::change_base_type_t< T, int > ilogb(T x) __NOEXC
Definition: builtins.hpp:253
cl::sycl::half_precision::sin
detail::enable_if_t< detail::is_genfloatf< T >::value, T > sin(T x) __NOEXC
Definition: builtins.hpp:1532
cl::sycl::half_precision::cos
detail::enable_if_t< detail::is_genfloatf< T >::value, T > cos(T x) __NOEXC
Definition: builtins.hpp:1465
common.hpp
cl::sycl::round
detail::enable_if_t< detail::is_genfloat< T >::value, T > round(T x) __NOEXC
Definition: builtins.hpp:423
cl::sycl::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:239
cl::sycl::isnormal
detail::common_rel_ret_t< T > isnormal(T x) __NOEXC
Definition: builtins.hpp:1225
cl::sycl::islessequal
detail::common_rel_ret_t< T > islessequal(T x, T y) __NOEXC
Definition: builtins.hpp:1165
cl::sycl::instead
std::uint8_t instead
Definition: aliases.hpp:68
cl::sycl::asin
detail::enable_if_t< detail::is_genfloat< T >::value, T > asin(T x) __NOEXC
Definition: builtins.hpp:49
cl::sycl::erfc
detail::enable_if_t< detail::is_genfloat< T >::value, T > erfc(T x) __NOEXC
Definition: builtins.hpp:135
types.hpp
cl::sycl::half_precision::exp10
detail::enable_if_t< detail::is_genfloatf< T >::value, T > exp10(T x) __NOEXC
Definition: builtins.hpp:1490
cl::sycl::rotate
detail::enable_if_t< detail::is_geninteger< T >::value, T > rotate(T v, T i) __NOEXC
Definition: builtins.hpp:842
cl::sycl::all
detail::enable_if_t< detail::is_vigeninteger< T >::value, int > all(T x) __NOEXC
Definition: builtins.hpp:1288
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::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:910
cl::sycl::mad_sat
detail::enable_if_t< detail::is_ugeninteger< T >::value, T > mad_sat(T a, T b, T c) __NOEXC
Definition: builtins.hpp:765
cl::sycl::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:608
cl::sycl::isfinite
detail::common_rel_ret_t< T > isfinite(T x) __NOEXC
Definition: builtins.hpp:1189
cl::sycl::half_precision::powr
detail::enable_if_t< detail::is_genfloatf< T >::value, T > powr(T x, T y) __NOEXC
Definition: builtins.hpp:1514
cl::sycl::remainder
detail::enable_if_t< detail::is_genfloat< T >::value, T > remainder(T x, T y) __NOEXC
Definition: builtins.hpp:392
cl::sycl::clz
detail::enable_if_t< detail::is_geninteger< T >::value, T > clz(T x) __NOEXC
Definition: builtins.hpp:723
cl::sycl::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:443
cl::sycl::bitselect
detail::enable_if_t< detail::is_gentype< T >::value, T > bitselect(T a, T b, T c) __NOEXC
Definition: builtins.hpp:1296
cl::sycl::nan
detail::nan_return_t< T > nan(T nancode) __NOEXC
Definition: builtins.hpp:357
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12
cl::sycl::detail::rel_sign_bit_test_arg_t
typename RelationalTestForSignBitType< T >::argument_type rel_sign_bit_test_arg_t
Definition: generic_type_traits.hpp:553