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