clang  13.0.0git
__clang_hip_math.h
Go to the documentation of this file.
1 /*===---- __clang_hip_math.h - Device-side HIP math support ----------------===
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 #ifndef __CLANG_HIP_MATH_H__
10 #define __CLANG_HIP_MATH_H__
11 
12 #if !defined(__HIP__)
13 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
14 #endif
15 
16 #if defined(__cplusplus)
17 #include <algorithm>
18 #endif
19 #include <limits.h>
20 #include <stdint.h>
21 
22 #pragma push_macro("__DEVICE__")
23 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
24 
25 // A few functions return bool type starting only in C++11.
26 #pragma push_macro("__RETURN_TYPE")
27 #if defined(__cplusplus)
28 #define __RETURN_TYPE bool
29 #else
30 #define __RETURN_TYPE int
31 #endif
32 
33 #if defined (__cplusplus) && __cplusplus < 201103L
34 // emulate static_assert on type sizes
35 template<bool>
36 struct __compare_result{};
37 template<>
38 struct __compare_result<true> {
39  static const bool valid;
40 };
41 
43 void __suppress_unused_warning(bool b){};
44 template <unsigned int S, unsigned int T>
45 __DEVICE__ void __static_assert_equal_size() {
46  __suppress_unused_warning(__compare_result<S == T>::valid);
47 }
48 
49 #define __static_assert_type_size_equal(A, B) \
50  __static_assert_equal_size<A,B>()
51 
52 #else
53 #define __static_assert_type_size_equal(A,B) \
54  static_assert((A) == (B), "")
55 
56 #endif
57 
59 uint64_t __make_mantissa_base8(const char *__tagp) {
60  uint64_t __r = 0;
61  while (__tagp) {
62  char __tmp = *__tagp;
63 
64  if (__tmp >= '0' && __tmp <= '7')
65  __r = (__r * 8u) + __tmp - '0';
66  else
67  return 0;
68 
69  ++__tagp;
70  }
71 
72  return __r;
73 }
74 
76 uint64_t __make_mantissa_base10(const char *__tagp) {
77  uint64_t __r = 0;
78  while (__tagp) {
79  char __tmp = *__tagp;
80 
81  if (__tmp >= '0' && __tmp <= '9')
82  __r = (__r * 10u) + __tmp - '0';
83  else
84  return 0;
85 
86  ++__tagp;
87  }
88 
89  return __r;
90 }
91 
93 uint64_t __make_mantissa_base16(const char *__tagp) {
94  uint64_t __r = 0;
95  while (__tagp) {
96  char __tmp = *__tagp;
97 
98  if (__tmp >= '0' && __tmp <= '9')
99  __r = (__r * 16u) + __tmp - '0';
100  else if (__tmp >= 'a' && __tmp <= 'f')
101  __r = (__r * 16u) + __tmp - 'a' + 10;
102  else if (__tmp >= 'A' && __tmp <= 'F')
103  __r = (__r * 16u) + __tmp - 'A' + 10;
104  else
105  return 0;
106 
107  ++__tagp;
108  }
109 
110  return __r;
111 }
112 
114 uint64_t __make_mantissa(const char *__tagp) {
115  if (!__tagp)
116  return 0u;
117 
118  if (*__tagp == '0') {
119  ++__tagp;
120 
121  if (*__tagp == 'x' || *__tagp == 'X')
122  return __make_mantissa_base16(__tagp);
123  else
124  return __make_mantissa_base8(__tagp);
125  }
126 
127  return __make_mantissa_base10(__tagp);
128 }
129 
130 // BEGIN FLOAT
131 #if defined(__cplusplus)
133 int abs(int __x) {
134  int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1);
135  return (__x ^ __sgn) - __sgn;
136 }
138 long labs(long __x) {
139  long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1);
140  return (__x ^ __sgn) - __sgn;
141 }
143 long long llabs(long long __x) {
144  long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
145  return (__x ^ __sgn) - __sgn;
146 }
147 #endif
148 
150 float acosf(float __x) { return __ocml_acos_f32(__x); }
151 
153 float acoshf(float __x) { return __ocml_acosh_f32(__x); }
154 
156 float asinf(float __x) { return __ocml_asin_f32(__x); }
157 
159 float asinhf(float __x) { return __ocml_asinh_f32(__x); }
160 
162 float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); }
163 
165 float atanf(float __x) { return __ocml_atan_f32(__x); }
166 
168 float atanhf(float __x) { return __ocml_atanh_f32(__x); }
169 
171 float cbrtf(float __x) { return __ocml_cbrt_f32(__x); }
172 
174 float ceilf(float __x) { return __ocml_ceil_f32(__x); }
175 
177 float copysignf(float __x, float __y) { return __ocml_copysign_f32(__x, __y); }
178 
180 float cosf(float __x) { return __ocml_cos_f32(__x); }
181 
183 float coshf(float __x) { return __ocml_cosh_f32(__x); }
184 
186 float cospif(float __x) { return __ocml_cospi_f32(__x); }
187 
189 float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); }
190 
192 float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); }
193 
195 float erfcf(float __x) { return __ocml_erfc_f32(__x); }
196 
198 float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); }
199 
201 float erfcxf(float __x) { return __ocml_erfcx_f32(__x); }
202 
204 float erff(float __x) { return __ocml_erf_f32(__x); }
205 
207 float erfinvf(float __x) { return __ocml_erfinv_f32(__x); }
208 
210 float exp10f(float __x) { return __ocml_exp10_f32(__x); }
211 
213 float exp2f(float __x) { return __ocml_exp2_f32(__x); }
214 
216 float expf(float __x) { return __ocml_exp_f32(__x); }
217 
219 float expm1f(float __x) { return __ocml_expm1_f32(__x); }
220 
222 float fabsf(float __x) { return __ocml_fabs_f32(__x); }
223 
225 float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); }
226 
228 float fdividef(float __x, float __y) { return __x / __y; }
229 
231 float floorf(float __x) { return __ocml_floor_f32(__x); }
232 
234 float fmaf(float __x, float __y, float __z) {
235  return __ocml_fma_f32(__x, __y, __z);
236 }
237 
239 float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); }
240 
242 float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); }
243 
245 float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
246 
248 float frexpf(float __x, int *__nptr) {
249  int __tmp;
250  float __r =
251  __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp);
252  *__nptr = __tmp;
253 
254  return __r;
255 }
256 
258 float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); }
259 
261 int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }
262 
264 __RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); }
265 
267 __RETURN_TYPE __isinff(float __x) { return __ocml_isinf_f32(__x); }
268 
270 __RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); }
271 
273 float j0f(float __x) { return __ocml_j0_f32(__x); }
274 
276 float j1f(float __x) { return __ocml_j1_f32(__x); }
277 
279 float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication
280  // and the Miller & Brown algorithm
281  // for linear recurrences to get O(log n) steps, but it's unclear if
282  // it'd be beneficial in this case.
283  if (__n == 0)
284  return j0f(__x);
285  if (__n == 1)
286  return j1f(__x);
287 
288  float __x0 = j0f(__x);
289  float __x1 = j1f(__x);
290  for (int __i = 1; __i < __n; ++__i) {
291  float __x2 = (2 * __i) / __x * __x1 - __x0;
292  __x0 = __x1;
293  __x1 = __x2;
294  }
295 
296  return __x1;
297 }
298 
300 float ldexpf(float __x, int __e) { return __ocml_ldexp_f32(__x, __e); }
301 
303 float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
304 
306 long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
307 
309 long long int llroundf(float __x) { return __ocml_round_f32(__x); }
310 
312 float log10f(float __x) { return __ocml_log10_f32(__x); }
313 
315 float log1pf(float __x) { return __ocml_log1p_f32(__x); }
316 
318 float log2f(float __x) { return __ocml_log2_f32(__x); }
319 
321 float logbf(float __x) { return __ocml_logb_f32(__x); }
322 
324 float logf(float __x) { return __ocml_log_f32(__x); }
325 
327 long int lrintf(float __x) { return __ocml_rint_f32(__x); }
328 
330 long int lroundf(float __x) { return __ocml_round_f32(__x); }
331 
333 float modff(float __x, float *__iptr) {
334  float __tmp;
335  float __r =
336  __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
337  *__iptr = __tmp;
338  return __r;
339 }
340 
342 float nanf(const char *__tagp) {
343  union {
344  float val;
345  struct ieee_float {
346  unsigned int mantissa : 22;
347  unsigned int quiet : 1;
348  unsigned int exponent : 8;
349  unsigned int sign : 1;
350  } bits;
351  } __tmp;
352  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
353 
354  __tmp.bits.sign = 0u;
355  __tmp.bits.exponent = ~0u;
356  __tmp.bits.quiet = 1u;
357  __tmp.bits.mantissa = __make_mantissa(__tagp);
358 
359  return __tmp.val;
360 }
361 
363 float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); }
364 
366 float nextafterf(float __x, float __y) {
367  return __ocml_nextafter_f32(__x, __y);
368 }
369 
371 float norm3df(float __x, float __y, float __z) {
372  return __ocml_len3_f32(__x, __y, __z);
373 }
374 
376 float norm4df(float __x, float __y, float __z, float __w) {
377  return __ocml_len4_f32(__x, __y, __z, __w);
378 }
379 
381 float normcdff(float __x) { return __ocml_ncdf_f32(__x); }
382 
384 float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); }
385 
387 float normf(int __dim,
388  const float *__a) { // TODO: placeholder until OCML adds support.
389  float __r = 0;
390  while (__dim--) {
391  __r += __a[0] * __a[0];
392  ++__a;
393  }
394 
395  return __ocml_sqrt_f32(__r);
396 }
397 
399 float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
400 
402 float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); }
403 
405 float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }
406 
408 float remainderf(float __x, float __y) {
409  return __ocml_remainder_f32(__x, __y);
410 }
411 
413 float remquof(float __x, float __y, int *__quo) {
414  int __tmp;
415  float __r = __ocml_remquo_f32(
416  __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
417  *__quo = __tmp;
418 
419  return __r;
420 }
421 
423 float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }
424 
426 float rintf(float __x) { return __ocml_rint_f32(__x); }
427 
429 float rnorm3df(float __x, float __y, float __z) {
430  return __ocml_rlen3_f32(__x, __y, __z);
431 }
432 
434 float rnorm4df(float __x, float __y, float __z, float __w) {
435  return __ocml_rlen4_f32(__x, __y, __z, __w);
436 }
437 
439 float rnormf(int __dim,
440  const float *__a) { // TODO: placeholder until OCML adds support.
441  float __r = 0;
442  while (__dim--) {
443  __r += __a[0] * __a[0];
444  ++__a;
445  }
446 
447  return __ocml_rsqrt_f32(__r);
448 }
449 
451 float roundf(float __x) { return __ocml_round_f32(__x); }
452 
454 float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
455 
457 float scalblnf(float __x, long int __n) {
458  return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n)
459  : __ocml_scalb_f32(__x, __n);
460 }
461 
463 float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
464 
466 __RETURN_TYPE __signbitf(float __x) { return __ocml_signbit_f32(__x); }
467 
469 void sincosf(float __x, float *__sinptr, float *__cosptr) {
470  float __tmp;
471  *__sinptr =
472  __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
473  *__cosptr = __tmp;
474 }
475 
477 void sincospif(float __x, float *__sinptr, float *__cosptr) {
478  float __tmp;
479  *__sinptr = __ocml_sincospi_f32(
480  __x, (__attribute__((address_space(5))) float *)&__tmp);
481  *__cosptr = __tmp;
482 }
483 
485 float sinf(float __x) { return __ocml_sin_f32(__x); }
486 
488 float sinhf(float __x) { return __ocml_sinh_f32(__x); }
489 
491 float sinpif(float __x) { return __ocml_sinpi_f32(__x); }
492 
494 float sqrtf(float __x) { return __ocml_sqrt_f32(__x); }
495 
497 float tanf(float __x) { return __ocml_tan_f32(__x); }
498 
500 float tanhf(float __x) { return __ocml_tanh_f32(__x); }
501 
503 float tgammaf(float __x) { return __ocml_tgamma_f32(__x); }
504 
506 float truncf(float __x) { return __ocml_trunc_f32(__x); }
507 
509 float y0f(float __x) { return __ocml_y0_f32(__x); }
510 
512 float y1f(float __x) { return __ocml_y1_f32(__x); }
513 
515 float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication
516  // and the Miller & Brown algorithm
517  // for linear recurrences to get O(log n) steps, but it's unclear if
518  // it'd be beneficial in this case. Placeholder until OCML adds
519  // support.
520  if (__n == 0)
521  return y0f(__x);
522  if (__n == 1)
523  return y1f(__x);
524 
525  float __x0 = y0f(__x);
526  float __x1 = y1f(__x);
527  for (int __i = 1; __i < __n; ++__i) {
528  float __x2 = (2 * __i) / __x * __x1 - __x0;
529  __x0 = __x1;
530  __x1 = __x2;
531  }
532 
533  return __x1;
534 }
535 
536 // BEGIN INTRINSICS
537 
539 float __cosf(float __x) { return __ocml_native_cos_f32(__x); }
540 
542 float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }
543 
545 float __expf(float __x) { return __ocml_native_exp_f32(__x); }
546 
547 #if defined OCML_BASIC_ROUNDED_OPERATIONS
549 float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); }
551 float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); }
553 float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); }
555 float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); }
556 #else
558 float __fadd_rn(float __x, float __y) { return __x + __y; }
559 #endif
560 
561 #if defined OCML_BASIC_ROUNDED_OPERATIONS
563 float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); }
565 float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); }
567 float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); }
569 float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); }
570 #else
572 float __fdiv_rn(float __x, float __y) { return __x / __y; }
573 #endif
574 
576 float __fdividef(float __x, float __y) { return __x / __y; }
577 
578 #if defined OCML_BASIC_ROUNDED_OPERATIONS
580 float __fmaf_rd(float __x, float __y, float __z) {
581  return __ocml_fma_rtn_f32(__x, __y, __z);
582 }
584 float __fmaf_rn(float __x, float __y, float __z) {
585  return __ocml_fma_rte_f32(__x, __y, __z);
586 }
588 float __fmaf_ru(float __x, float __y, float __z) {
589  return __ocml_fma_rtp_f32(__x, __y, __z);
590 }
592 float __fmaf_rz(float __x, float __y, float __z) {
593  return __ocml_fma_rtz_f32(__x, __y, __z);
594 }
595 #else
597 float __fmaf_rn(float __x, float __y, float __z) {
598  return __ocml_fma_f32(__x, __y, __z);
599 }
600 #endif
601 
602 #if defined OCML_BASIC_ROUNDED_OPERATIONS
604 float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }
606 float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }
608 float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); }
610 float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); }
611 #else
613 float __fmul_rn(float __x, float __y) { return __x * __y; }
614 #endif
615 
616 #if defined OCML_BASIC_ROUNDED_OPERATIONS
618 float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); }
620 float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); }
622 float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); }
624 float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }
625 #else
627 float __frcp_rn(float __x) { return 1.0f / __x; }
628 #endif
629 
631 float __frsqrt_rn(float __x) { return __llvm_amdgcn_rsq_f32(__x); }
632 
633 #if defined OCML_BASIC_ROUNDED_OPERATIONS
635 float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
637 float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }
639 float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }
641 float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }
642 #else
644 float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }
645 #endif
646 
647 #if defined OCML_BASIC_ROUNDED_OPERATIONS
649 float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); }
651 float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); }
653 float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); }
655 float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); }
656 #else
658 float __fsub_rn(float __x, float __y) { return __x - __y; }
659 #endif
660 
662 float __log10f(float __x) { return __ocml_native_log10_f32(__x); }
663 
665 float __log2f(float __x) { return __ocml_native_log2_f32(__x); }
666 
668 float __logf(float __x) { return __ocml_native_log_f32(__x); }
669 
671 float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
672 
674 float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }
675 
677 void __sincosf(float __x, float *__sinptr, float *__cosptr) {
678  *__sinptr = __ocml_native_sin_f32(__x);
679  *__cosptr = __ocml_native_cos_f32(__x);
680 }
681 
683 float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
684 
686 float __tanf(float __x) { return __ocml_tan_f32(__x); }
687 // END INTRINSICS
688 // END FLOAT
689 
690 // BEGIN DOUBLE
692 double acos(double __x) { return __ocml_acos_f64(__x); }
693 
695 double acosh(double __x) { return __ocml_acosh_f64(__x); }
696 
698 double asin(double __x) { return __ocml_asin_f64(__x); }
699 
701 double asinh(double __x) { return __ocml_asinh_f64(__x); }
702 
704 double atan(double __x) { return __ocml_atan_f64(__x); }
705 
707 double atan2(double __x, double __y) { return __ocml_atan2_f64(__x, __y); }
708 
710 double atanh(double __x) { return __ocml_atanh_f64(__x); }
711 
713 double cbrt(double __x) { return __ocml_cbrt_f64(__x); }
714 
716 double ceil(double __x) { return __ocml_ceil_f64(__x); }
717 
719 double copysign(double __x, double __y) {
720  return __ocml_copysign_f64(__x, __y);
721 }
722 
724 double cos(double __x) { return __ocml_cos_f64(__x); }
725 
727 double cosh(double __x) { return __ocml_cosh_f64(__x); }
728 
730 double cospi(double __x) { return __ocml_cospi_f64(__x); }
731 
733 double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); }
734 
736 double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); }
737 
739 double erf(double __x) { return __ocml_erf_f64(__x); }
740 
742 double erfc(double __x) { return __ocml_erfc_f64(__x); }
743 
745 double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); }
746 
748 double erfcx(double __x) { return __ocml_erfcx_f64(__x); }
749 
751 double erfinv(double __x) { return __ocml_erfinv_f64(__x); }
752 
754 double exp(double __x) { return __ocml_exp_f64(__x); }
755 
757 double exp10(double __x) { return __ocml_exp10_f64(__x); }
758 
760 double exp2(double __x) { return __ocml_exp2_f64(__x); }
761 
763 double expm1(double __x) { return __ocml_expm1_f64(__x); }
764 
766 double fabs(double __x) { return __ocml_fabs_f64(__x); }
767 
769 double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }
770 
772 double floor(double __x) { return __ocml_floor_f64(__x); }
773 
775 double fma(double __x, double __y, double __z) {
776  return __ocml_fma_f64(__x, __y, __z);
777 }
778 
780 double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); }
781 
783 double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); }
784 
786 double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
787 
789 double frexp(double __x, int *__nptr) {
790  int __tmp;
791  double __r =
792  __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp);
793  *__nptr = __tmp;
794  return __r;
795 }
796 
798 double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); }
799 
801 int ilogb(double __x) { return __ocml_ilogb_f64(__x); }
802 
804 __RETURN_TYPE __finite(double __x) { return __ocml_isfinite_f64(__x); }
805 
807 __RETURN_TYPE __isinf(double __x) { return __ocml_isinf_f64(__x); }
808 
810 __RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); }
811 
813 double j0(double __x) { return __ocml_j0_f64(__x); }
814 
816 double j1(double __x) { return __ocml_j1_f64(__x); }
817 
819 double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication
820  // and the Miller & Brown algorithm
821  // for linear recurrences to get O(log n) steps, but it's unclear if
822  // it'd be beneficial in this case. Placeholder until OCML adds
823  // support.
824  if (__n == 0)
825  return j0(__x);
826  if (__n == 1)
827  return j1(__x);
828 
829  double __x0 = j0(__x);
830  double __x1 = j1(__x);
831  for (int __i = 1; __i < __n; ++__i) {
832  double __x2 = (2 * __i) / __x * __x1 - __x0;
833  __x0 = __x1;
834  __x1 = __x2;
835  }
836  return __x1;
837 }
838 
840 double ldexp(double __x, int __e) { return __ocml_ldexp_f64(__x, __e); }
841 
843 double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
844 
846 long long int llrint(double __x) { return __ocml_rint_f64(__x); }
847 
849 long long int llround(double __x) { return __ocml_round_f64(__x); }
850 
852 double log(double __x) { return __ocml_log_f64(__x); }
853 
855 double log10(double __x) { return __ocml_log10_f64(__x); }
856 
858 double log1p(double __x) { return __ocml_log1p_f64(__x); }
859 
861 double log2(double __x) { return __ocml_log2_f64(__x); }
862 
864 double logb(double __x) { return __ocml_logb_f64(__x); }
865 
867 long int lrint(double __x) { return __ocml_rint_f64(__x); }
868 
870 long int lround(double __x) { return __ocml_round_f64(__x); }
871 
873 double modf(double __x, double *__iptr) {
874  double __tmp;
875  double __r =
876  __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
877  *__iptr = __tmp;
878 
879  return __r;
880 }
881 
883 double nan(const char *__tagp) {
884 #if !_WIN32
885  union {
886  double val;
887  struct ieee_double {
888  uint64_t mantissa : 51;
889  uint32_t quiet : 1;
890  uint32_t exponent : 11;
891  uint32_t sign : 1;
892  } bits;
893  } __tmp;
894  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
895 
896  __tmp.bits.sign = 0u;
897  __tmp.bits.exponent = ~0u;
898  __tmp.bits.quiet = 1u;
899  __tmp.bits.mantissa = __make_mantissa(__tagp);
900 
901  return __tmp.val;
902 #else
903  __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double));
904  uint64_t __val = __make_mantissa(__tagp);
905  __val |= 0xFFF << 51;
906  return *reinterpret_cast<double *>(&__val);
907 #endif
908 }
909 
911 double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); }
912 
914 double nextafter(double __x, double __y) {
915  return __ocml_nextafter_f64(__x, __y);
916 }
917 
919 double norm(int __dim,
920  const double *__a) { // TODO: placeholder until OCML adds support.
921  double __r = 0;
922  while (__dim--) {
923  __r += __a[0] * __a[0];
924  ++__a;
925  }
926 
927  return __ocml_sqrt_f64(__r);
928 }
929 
931 double norm3d(double __x, double __y, double __z) {
932  return __ocml_len3_f64(__x, __y, __z);
933 }
934 
936 double norm4d(double __x, double __y, double __z, double __w) {
937  return __ocml_len4_f64(__x, __y, __z, __w);
938 }
939 
941 double normcdf(double __x) { return __ocml_ncdf_f64(__x); }
942 
944 double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); }
945 
947 double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); }
948 
950 double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); }
951 
953 double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }
954 
956 double remainder(double __x, double __y) {
957  return __ocml_remainder_f64(__x, __y);
958 }
959 
961 double remquo(double __x, double __y, int *__quo) {
962  int __tmp;
963  double __r = __ocml_remquo_f64(
964  __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
965  *__quo = __tmp;
966 
967  return __r;
968 }
969 
971 double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }
972 
974 double rint(double __x) { return __ocml_rint_f64(__x); }
975 
977 double rnorm(int __dim,
978  const double *__a) { // TODO: placeholder until OCML adds support.
979  double __r = 0;
980  while (__dim--) {
981  __r += __a[0] * __a[0];
982  ++__a;
983  }
984 
985  return __ocml_rsqrt_f64(__r);
986 }
987 
989 double rnorm3d(double __x, double __y, double __z) {
990  return __ocml_rlen3_f64(__x, __y, __z);
991 }
992 
994 double rnorm4d(double __x, double __y, double __z, double __w) {
995  return __ocml_rlen4_f64(__x, __y, __z, __w);
996 }
997 
999 double round(double __x) { return __ocml_round_f64(__x); }
1000 
1001 __DEVICE__
1002 double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }
1003 
1004 __DEVICE__
1005 double scalbln(double __x, long int __n) {
1006  return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n)
1007  : __ocml_scalb_f64(__x, __n);
1008 }
1009 __DEVICE__
1010 double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); }
1011 
1012 __DEVICE__
1013 __RETURN_TYPE __signbit(double __x) { return __ocml_signbit_f64(__x); }
1014 
1015 __DEVICE__
1016 double sin(double __x) { return __ocml_sin_f64(__x); }
1017 
1018 __DEVICE__
1019 void sincos(double __x, double *__sinptr, double *__cosptr) {
1020  double __tmp;
1021  *__sinptr = __ocml_sincos_f64(
1022  __x, (__attribute__((address_space(5))) double *)&__tmp);
1023  *__cosptr = __tmp;
1024 }
1025 
1026 __DEVICE__
1027 void sincospi(double __x, double *__sinptr, double *__cosptr) {
1028  double __tmp;
1029  *__sinptr = __ocml_sincospi_f64(
1030  __x, (__attribute__((address_space(5))) double *)&__tmp);
1031  *__cosptr = __tmp;
1032 }
1033 
1034 __DEVICE__
1035 double sinh(double __x) { return __ocml_sinh_f64(__x); }
1036 
1037 __DEVICE__
1038 double sinpi(double __x) { return __ocml_sinpi_f64(__x); }
1039 
1040 __DEVICE__
1041 double sqrt(double __x) { return __ocml_sqrt_f64(__x); }
1042 
1043 __DEVICE__
1044 double tan(double __x) { return __ocml_tan_f64(__x); }
1045 
1046 __DEVICE__
1047 double tanh(double __x) { return __ocml_tanh_f64(__x); }
1048 
1049 __DEVICE__
1050 double tgamma(double __x) { return __ocml_tgamma_f64(__x); }
1051 
1052 __DEVICE__
1053 double trunc(double __x) { return __ocml_trunc_f64(__x); }
1054 
1055 __DEVICE__
1056 double y0(double __x) { return __ocml_y0_f64(__x); }
1057 
1058 __DEVICE__
1059 double y1(double __x) { return __ocml_y1_f64(__x); }
1060 
1061 __DEVICE__
1062 double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication
1063  // and the Miller & Brown algorithm
1064  // for linear recurrences to get O(log n) steps, but it's unclear if
1065  // it'd be beneficial in this case. Placeholder until OCML adds
1066  // support.
1067  if (__n == 0)
1068  return y0(__x);
1069  if (__n == 1)
1070  return y1(__x);
1071 
1072  double __x0 = y0(__x);
1073  double __x1 = y1(__x);
1074  for (int __i = 1; __i < __n; ++__i) {
1075  double __x2 = (2 * __i) / __x * __x1 - __x0;
1076  __x0 = __x1;
1077  __x1 = __x2;
1078  }
1079 
1080  return __x1;
1081 }
1082 
1083 // BEGIN INTRINSICS
1084 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1085 __DEVICE__
1086 double __dadd_rd(double __x, double __y) {
1087  return __ocml_add_rtn_f64(__x, __y);
1088 }
1089 __DEVICE__
1090 double __dadd_rn(double __x, double __y) {
1091  return __ocml_add_rte_f64(__x, __y);
1092 }
1093 __DEVICE__
1094 double __dadd_ru(double __x, double __y) {
1095  return __ocml_add_rtp_f64(__x, __y);
1096 }
1097 __DEVICE__
1098 double __dadd_rz(double __x, double __y) {
1099  return __ocml_add_rtz_f64(__x, __y);
1100 }
1101 #else
1102 __DEVICE__
1103 double __dadd_rn(double __x, double __y) { return __x + __y; }
1104 #endif
1105 
1106 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1107 __DEVICE__
1108 double __ddiv_rd(double __x, double __y) {
1109  return __ocml_div_rtn_f64(__x, __y);
1110 }
1111 __DEVICE__
1112 double __ddiv_rn(double __x, double __y) {
1113  return __ocml_div_rte_f64(__x, __y);
1114 }
1115 __DEVICE__
1116 double __ddiv_ru(double __x, double __y) {
1117  return __ocml_div_rtp_f64(__x, __y);
1118 }
1119 __DEVICE__
1120 double __ddiv_rz(double __x, double __y) {
1121  return __ocml_div_rtz_f64(__x, __y);
1122 }
1123 #else
1124 __DEVICE__
1125 double __ddiv_rn(double __x, double __y) { return __x / __y; }
1126 #endif
1127 
1128 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1129 __DEVICE__
1130 double __dmul_rd(double __x, double __y) {
1131  return __ocml_mul_rtn_f64(__x, __y);
1132 }
1133 __DEVICE__
1134 double __dmul_rn(double __x, double __y) {
1135  return __ocml_mul_rte_f64(__x, __y);
1136 }
1137 __DEVICE__
1138 double __dmul_ru(double __x, double __y) {
1139  return __ocml_mul_rtp_f64(__x, __y);
1140 }
1141 __DEVICE__
1142 double __dmul_rz(double __x, double __y) {
1143  return __ocml_mul_rtz_f64(__x, __y);
1144 }
1145 #else
1146 __DEVICE__
1147 double __dmul_rn(double __x, double __y) { return __x * __y; }
1148 #endif
1149 
1150 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1151 __DEVICE__
1152 double __drcp_rd(double __x) { return __ocml_div_rtn_f64(1.0, __x); }
1153 __DEVICE__
1154 double __drcp_rn(double __x) { return __ocml_div_rte_f64(1.0, __x); }
1155 __DEVICE__
1156 double __drcp_ru(double __x) { return __ocml_div_rtp_f64(1.0, __x); }
1157 __DEVICE__
1158 double __drcp_rz(double __x) { return __ocml_div_rtz_f64(1.0, __x); }
1159 #else
1160 __DEVICE__
1161 double __drcp_rn(double __x) { return 1.0 / __x; }
1162 #endif
1163 
1164 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1165 __DEVICE__
1166 double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); }
1167 __DEVICE__
1168 double __dsqrt_rn(double __x) { return __ocml_sqrt_rte_f64(__x); }
1169 __DEVICE__
1170 double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); }
1171 __DEVICE__
1172 double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); }
1173 #else
1174 __DEVICE__
1175 double __dsqrt_rn(double __x) { return __ocml_sqrt_f64(__x); }
1176 #endif
1177 
1178 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1179 __DEVICE__
1180 double __dsub_rd(double __x, double __y) {
1181  return __ocml_sub_rtn_f64(__x, __y);
1182 }
1183 __DEVICE__
1184 double __dsub_rn(double __x, double __y) {
1185  return __ocml_sub_rte_f64(__x, __y);
1186 }
1187 __DEVICE__
1188 double __dsub_ru(double __x, double __y) {
1189  return __ocml_sub_rtp_f64(__x, __y);
1190 }
1191 __DEVICE__
1192 double __dsub_rz(double __x, double __y) {
1193  return __ocml_sub_rtz_f64(__x, __y);
1194 }
1195 #else
1196 __DEVICE__
1197 double __dsub_rn(double __x, double __y) { return __x - __y; }
1198 #endif
1199 
1200 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1201 __DEVICE__
1202 double __fma_rd(double __x, double __y, double __z) {
1203  return __ocml_fma_rtn_f64(__x, __y, __z);
1204 }
1205 __DEVICE__
1206 double __fma_rn(double __x, double __y, double __z) {
1207  return __ocml_fma_rte_f64(__x, __y, __z);
1208 }
1209 __DEVICE__
1210 double __fma_ru(double __x, double __y, double __z) {
1211  return __ocml_fma_rtp_f64(__x, __y, __z);
1212 }
1213 __DEVICE__
1214 double __fma_rz(double __x, double __y, double __z) {
1215  return __ocml_fma_rtz_f64(__x, __y, __z);
1216 }
1217 #else
1218 __DEVICE__
1219 double __fma_rn(double __x, double __y, double __z) {
1220  return __ocml_fma_f64(__x, __y, __z);
1221 }
1222 #endif
1223 // END INTRINSICS
1224 // END DOUBLE
1225 
1226 // C only macros
1227 #if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L
1228 #define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x)
1229 #define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x)
1230 #define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x)
1231 #define signbit(__x) \
1232  _Generic((__x), float : __signbitf, double : __signbit)(__x)
1233 #endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L
1234 
1235 #if defined(__cplusplus)
1236 template <class T> __DEVICE__ T min(T __arg1, T __arg2) {
1237  return (__arg1 < __arg2) ? __arg1 : __arg2;
1238 }
1239 
1240 template <class T> __DEVICE__ T max(T __arg1, T __arg2) {
1241  return (__arg1 > __arg2) ? __arg1 : __arg2;
1242 }
1243 
1244 __DEVICE__ int min(int __arg1, int __arg2) {
1245  return (__arg1 < __arg2) ? __arg1 : __arg2;
1246 }
1247 __DEVICE__ int max(int __arg1, int __arg2) {
1248  return (__arg1 > __arg2) ? __arg1 : __arg2;
1249 }
1250 
1251 __DEVICE__
1252 float max(float __x, float __y) { return fmaxf(__x, __y); }
1253 
1254 __DEVICE__
1255 double max(double __x, double __y) { return fmax(__x, __y); }
1256 
1257 __DEVICE__
1258 float min(float __x, float __y) { return fminf(__x, __y); }
1259 
1260 __DEVICE__
1261 double min(double __x, double __y) { return fmin(__x, __y); }
1262 
1263 __host__ inline static int min(int __arg1, int __arg2) {
1264  return std::min(__arg1, __arg2);
1265 }
1266 
1267 __host__ inline static int max(int __arg1, int __arg2) {
1268  return std::max(__arg1, __arg2);
1269 }
1270 #endif
1271 
1272 #pragma pop_macro("__DEVICE__")
1273 #pragma pop_macro("__RETURN_TYPE")
1274 
1275 #endif // __CLANG_HIP_MATH_H__
__fsub_ru
__DEVICE__ float __fsub_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:328
y1f
__DEVICE__ float y1f(float __x)
Definition: __clang_hip_math.h:512
lgammaf
__DEVICE__ float lgammaf(float __x)
Definition: __clang_hip_math.h:303
__ocml_y0_f32
__device__ float __ocml_y0_f32(float)
cbrt
__DEVICE__ double cbrt(double __x)
Definition: __clang_hip_math.h:713
ynf
__DEVICE__ float ynf(int __n, float __x)
Definition: __clang_hip_math.h:515
exp2f
__DEVICE__ float exp2f(float __x)
Definition: __clang_hip_math.h:213
asinhf
__DEVICE__ float asinhf(float __x)
Definition: __clang_hip_math.h:159
__dmul_rd
__DEVICE__ double __dmul_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:83
__fmul_rn
__DEVICE__ float __fmul_rn(float __x, float __y)
Definition: __clang_hip_math.h:613
__ocml_modf_f64
__device__ double __ocml_modf_f64(double, __attribute__((address_space(5))) double *)
exp10
__DEVICE__ double exp10(double __x)
Definition: __clang_hip_math.h:757
max
__DEVICE__ int max(int __a, int __b)
Definition: __clang_cuda_math.h:194
__fdiv_rz
__DEVICE__ float __fdiv_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:216
atan2f
__DEVICE__ float atan2f(float __x, float __y)
Definition: __clang_hip_math.h:162
__ddiv_rz
__DEVICE__ double __ddiv_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:80
fmaxf
__DEVICE__ float fmaxf(float __x, float __y)
Definition: __clang_hip_math.h:239
__ocml_y0_f64
__device__ double __ocml_y0_f64(double)
__x
static __inline unsigned char unsigned int __x
Definition: adxintrin.h:22
__ocml_sincospi_f64
__device__ double __ocml_sincospi_f64(double, __attribute__((address_space(5))) double *)
rsqrtf
__DEVICE__ float rsqrtf(float __x)
Definition: __clang_hip_math.h:454
expm1f
__DEVICE__ float expm1f(float __x)
Definition: __clang_hip_math.h:219
erfcf
__DEVICE__ float erfcf(float __x)
Definition: __clang_hip_math.h:195
powf
__DEVICE__ float powf(float __x, float __y)
Definition: __clang_hip_math.h:399
__frcp_ru
__DEVICE__ float __frcp_ru(float __a)
Definition: __clang_cuda_device_functions.h:315
acosh
__DEVICE__ double acosh(double __x)
Definition: __clang_hip_math.h:695
round
__DEVICE__ double round(double __x)
Definition: __clang_hip_math.h:999
j0
__DEVICE__ double j0(double __x)
Definition: __clang_hip_math.h:813
__dmul_rz
__DEVICE__ double __dmul_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:92
__drcp_rd
__DEVICE__ double __drcp_rd(double __a)
Definition: __clang_cuda_device_functions.h:152
__fdividef
__DEVICE__ float __fdividef(float __x, float __y)
Definition: __clang_hip_math.h:576
__finitef
__DEVICE__ __RETURN_TYPE __finitef(float __x)
Definition: __clang_hip_math.h:264
normcdf
__DEVICE__ double normcdf(double __x)
Definition: __clang_hip_math.h:941
rhypotf
__DEVICE__ float rhypotf(float __x, float __y)
Definition: __clang_hip_math.h:423
__finite
__DEVICE__ __RETURN_TYPE __finite(double __x)
Definition: __clang_hip_math.h:804
__dadd_rd
__DEVICE__ double __dadd_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:59
__fadd_rz
__DEVICE__ float __fadd_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:204
tanh
__DEVICE__ double tanh(double __x)
Definition: __clang_hip_math.h:1047
__fsub_rd
__DEVICE__ float __fsub_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:322
__dmul_rn
__DEVICE__ double __dmul_rn(double __x, double __y)
Definition: __clang_hip_math.h:1147
__ocml_y1_f32
__device__ float __ocml_y1_f32(float)
log10f
__DEVICE__ float log10f(float __x)
Definition: __clang_hip_math.h:312
__dsub_rz
__DEVICE__ double __dsub_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:169
ilogb
__DEVICE__ int ilogb(double __x)
Definition: __clang_hip_math.h:801
INT_MAX
#define INT_MAX
Definition: limits.h:46
asin
__DEVICE__ double asin(double __x)
Definition: __clang_hip_math.h:698
__log2f
__DEVICE__ float __log2f(float __x)
Definition: __clang_hip_math.h:665
__fdiv_rd
__DEVICE__ float __fdiv_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:207
CHAR_BIT
#define CHAR_BIT
Definition: limits.h:63
__fmul_ru
__DEVICE__ float __fmul_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:307
__ocml_i1_f32
__device__ float __ocml_i1_f32(float)
sincospi
__DEVICE__ void sincospi(double __x, double *__sinptr, double *__cosptr)
Definition: __clang_hip_math.h:1027
fmin
__DEVICE__ double fmin(double __x, double __y)
Definition: __clang_hip_math.h:783
__fsqrt_ru
__DEVICE__ float __fsqrt_ru(float __a)
Definition: __clang_cuda_device_functions.h:320
erfinv
__DEVICE__ double erfinv(double __x)
Definition: __clang_hip_math.h:751
int
__device__ int
Definition: __clang_hip_libdevice_declares.h:63
floor
__DEVICE__ double floor(double __x)
Definition: __clang_hip_math.h:772
truncf
__DEVICE__ float truncf(float __x)
Definition: __clang_hip_math.h:506
__fmul_rz
__DEVICE__ float __fmul_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:310
acoshf
__DEVICE__ float acoshf(float __x)
Definition: __clang_hip_math.h:153
nextafter
__DEVICE__ double nextafter(double __x, double __y)
Definition: __clang_hip_math.h:914
copysignf
__DEVICE__ float copysignf(float __x, float __y)
Definition: __clang_hip_math.h:177
normf
__DEVICE__ float normf(int __dim, const float *__a)
Definition: __clang_hip_math.h:387
modf
__DEVICE__ double modf(double __x, double *__iptr)
Definition: __clang_hip_math.h:873
__dsqrt_rz
__DEVICE__ double __dsqrt_rz(double __a)
Definition: __clang_cuda_device_functions.h:159
__ocml_i1_f64
__device__ double __ocml_i1_f64(double)
__fmaf_ru
__DEVICE__ float __fmaf_ru(float __a, float __b, float __c)
Definition: __clang_cuda_device_functions.h:295
__DEVICE__
#define __DEVICE__
Definition: __clang_hip_math.h:23
__ocml_j0_f64
__device__ double __ocml_j0_f64(double)
tan
__DEVICE__ double tan(double __x)
Definition: __clang_hip_math.h:1044
exp
__DEVICE__ double exp(double __x)
Definition: __clang_hip_math.h:754
__logf
__DEVICE__ float __logf(float __x)
Definition: __clang_hip_math.h:668
sinpi
__DEVICE__ double sinpi(double __x)
Definition: __clang_hip_math.h:1038
__a
static __inline__ void int __a
Definition: emmintrin.h:4185
copysign
__DEVICE__ double copysign(double __x, double __y)
Definition: __clang_hip_math.h:719
__ocml_tan_f64
__device__ double __ocml_tan_f64(double)
__ocml_sinpi_f64
__device__ double __ocml_sinpi_f64(double)
ceil
__DEVICE__ double ceil(double __x)
Definition: __clang_hip_math.h:716
sincospif
__DEVICE__ void sincospif(float __x, float *__sinptr, float *__cosptr)
Definition: __clang_hip_math.h:477
__dmul_ru
__DEVICE__ double __dmul_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:89
sinf
__DEVICE__ float sinf(float __x)
Definition: __clang_hip_math.h:485
__exp10f
__DEVICE__ float __exp10f(float __x)
Definition: __clang_hip_math.h:542
tgammaf
__DEVICE__ float tgammaf(float __x)
Definition: __clang_hip_math.h:503
__RETURN_TYPE
#define __RETURN_TYPE
Definition: __clang_hip_math.h:30
remquo
__DEVICE__ double remquo(double __x, double __y, int *__quo)
Definition: __clang_hip_math.h:961
norm3d
__DEVICE__ double norm3d(double __x, double __y, double __z)
Definition: __clang_hip_math.h:931
floorf
__DEVICE__ float floorf(float __x)
Definition: __clang_hip_math.h:231
coshf
__DEVICE__ float coshf(float __x)
Definition: __clang_hip_math.h:183
__fma_rd
__DEVICE__ double __fma_rd(double __a, double __b, double __c)
Definition: __clang_cuda_device_functions.h:265
__dsub_rn
__DEVICE__ double __dsub_rn(double __x, double __y)
Definition: __clang_hip_math.h:1197
log10
__DEVICE__ double log10(double __x)
Definition: __clang_hip_math.h:855
b
__device__ __2f16 b
Definition: __clang_hip_libdevice_declares.h:302
rnorm4df
__DEVICE__ float rnorm4df(float __x, float __y, float __z, float __w)
Definition: __clang_hip_math.h:434
llrint
__DEVICE__ long long int llrint(double __x)
Definition: __clang_hip_math.h:846
fmod
__DEVICE__ double fmod(double __x, double __y)
Definition: __clang_hip_math.h:786
norm
__DEVICE__ double norm(int __dim, const double *__a)
Definition: __clang_hip_math.h:919
erfcxf
__DEVICE__ float erfcxf(float __x)
Definition: __clang_hip_math.h:201
__fma_rz
__DEVICE__ double __fma_rz(double __a, double __b, double __c)
Definition: __clang_cuda_device_functions.h:274
asinf
__DEVICE__ float asinf(float __x)
Definition: __clang_hip_math.h:156
sinh
__DEVICE__ double sinh(double __x)
Definition: __clang_hip_math.h:1035
hypotf
__DEVICE__ float hypotf(float __x, float __y)
Definition: __clang_hip_math.h:258
rcbrtf
__DEVICE__ float rcbrtf(float __x)
Definition: __clang_hip_math.h:405
yn
__DEVICE__ double yn(int __n, double __x)
Definition: __clang_hip_math.h:1062
ldexpf
__DEVICE__ float ldexpf(float __x, int __e)
Definition: __clang_hip_math.h:300
frexpf
__DEVICE__ float frexpf(float __x, int *__nptr)
Definition: __clang_hip_math.h:248
erfcinvf
__DEVICE__ float erfcinvf(float __x)
Definition: __clang_hip_math.h:198
cbrtf
__DEVICE__ float cbrtf(float __x)
Definition: __clang_hip_math.h:171
log2
__DEVICE__ double log2(double __x)
Definition: __clang_hip_math.h:861
min
__DEVICE__ int min(int __a, int __b)
Definition: __clang_cuda_math.h:195
__drcp_rn
__DEVICE__ double __drcp_rn(double __x)
Definition: __clang_hip_math.h:1161
cospi
__DEVICE__ double cospi(double __x)
Definition: __clang_hip_math.h:730
__fsqrt_rn
__DEVICE__ float __fsqrt_rn(float __x)
Definition: __clang_hip_math.h:644
norm4df
__DEVICE__ float norm4df(float __x, float __y, float __z, float __w)
Definition: __clang_hip_math.h:376
__isnanf
__DEVICE__ __RETURN_TYPE __isnanf(float __x)
Definition: __clang_hip_math.h:270
cosf
__DEVICE__ float cosf(float __x)
Definition: __clang_hip_math.h:180
asinh
__DEVICE__ double asinh(double __x)
Definition: __clang_hip_math.h:701
__ocml_lgamma_f64
__device__ double __ocml_lgamma_f64(double)
ceilf
__DEVICE__ float ceilf(float __x)
Definition: __clang_hip_math.h:174
lroundf
__DEVICE__ long int lroundf(float __x)
Definition: __clang_hip_math.h:330
tanf
__DEVICE__ float tanf(float __x)
Definition: __clang_hip_math.h:497
__ocml_sin_f32
__device__ float __ocml_sin_f32(float)
logbf
__DEVICE__ float logbf(float __x)
Definition: __clang_hip_math.h:321
llrintf
__DEVICE__ long long int llrintf(float __x)
Definition: __clang_hip_math.h:306
exp2
__DEVICE__ double exp2(double __x)
Definition: __clang_hip_math.h:760
remquof
__DEVICE__ float remquof(float __x, float __y, int *__quo)
Definition: __clang_hip_math.h:413
nearbyint
__DEVICE__ double nearbyint(double __x)
Definition: __clang_hip_math.h:911
nan
__DEVICE__ double nan(const char *__tagp)
Definition: __clang_hip_math.h:883
frexp
__DEVICE__ double frexp(double __x, int *__nptr)
Definition: __clang_hip_math.h:789
rhypot
__DEVICE__ double rhypot(double __x, double __y)
Definition: __clang_hip_math.h:971
remainder
__DEVICE__ double remainder(double __x, double __y)
Definition: __clang_hip_math.h:956
__ocml_cospi_f32
__device__ float __ocml_cospi_f32(float)
nearbyintf
__DEVICE__ float nearbyintf(float __x)
Definition: __clang_hip_math.h:363
limits.h
j1
__DEVICE__ double j1(double __x)
Definition: __clang_hip_math.h:816
__ocml_j0_f32
__device__ float __ocml_j0_f32(float)
fdividef
__DEVICE__ float fdividef(float __x, float __y)
Definition: __clang_hip_math.h:228
__signbit
__DEVICE__ __RETURN_TYPE __signbit(double __x)
Definition: __clang_hip_math.h:1013
erfcx
__DEVICE__ double erfcx(double __x)
Definition: __clang_hip_math.h:748
nextafterf
__DEVICE__ float nextafterf(float __x, float __y)
Definition: __clang_hip_math.h:366
__dadd_ru
__DEVICE__ double __dadd_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:65
erfinvf
__DEVICE__ float erfinvf(float __x)
Definition: __clang_hip_math.h:207
__ocml_sinpi_f32
__device__ float __ocml_sinpi_f32(float)
__expf
__DEVICE__ float __expf(float __x)
Definition: __clang_hip_math.h:545
__fma_rn
__DEVICE__ double __fma_rn(double __x, double __y, double __z)
Definition: __clang_hip_math.h:1219
__dadd_rz
__DEVICE__ double __dadd_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:68
rcbrt
__DEVICE__ double rcbrt(double __x)
Definition: __clang_hip_math.h:953
sincosf
__DEVICE__ void sincosf(float __x, float *__sinptr, float *__cosptr)
Definition: __clang_hip_math.h:469
log
__DEVICE__ double log(double __x)
Definition: __clang_hip_math.h:852
erfc
__DEVICE__ double erfc(double __x)
Definition: __clang_hip_math.h:742
__ocml_native_sin_f32
__device__ float __ocml_native_sin_f32(float)
normcdff
__DEVICE__ float normcdff(float __x)
Definition: __clang_hip_math.h:381
labs
__DEVICE__ long labs(long __a)
Definition: __clang_cuda_math.h:152
lrintf
__DEVICE__ long int lrintf(float __x)
Definition: __clang_hip_math.h:327
__frsqrt_rn
__DEVICE__ float __frsqrt_rn(float __x)
Definition: __clang_hip_math.h:631
__ocml_i0_f32
__device__ float __ocml_i0_f32(float)
rnormf
__DEVICE__ float rnormf(int __dim, const float *__a)
Definition: __clang_hip_math.h:439
__drcp_rz
__DEVICE__ double __drcp_rz(double __a)
Definition: __clang_cuda_device_functions.h:155
__frcp_rz
__DEVICE__ float __frcp_rz(float __a)
Definition: __clang_cuda_device_functions.h:316
j1f
__DEVICE__ float j1f(float __x)
Definition: __clang_hip_math.h:276
jnf
__DEVICE__ float jnf(int __n, float __x)
Definition: __clang_hip_math.h:279
log1p
__DEVICE__ double log1p(double __x)
Definition: __clang_hip_math.h:858
__cosf
__DEVICE__ float __cosf(float __x)
Definition: __clang_hip_math.h:539
__ocml_modf_f32
__device__ float __ocml_modf_f32(float, __attribute__((address_space(5))) float *)
acos
__DEVICE__ double acos(double __x)
Definition: __clang_hip_math.h:692
fmaf
__DEVICE__ float fmaf(float __x, float __y, float __z)
Definition: __clang_hip_math.h:234
rintf
__DEVICE__ float rintf(float __x)
Definition: __clang_hip_math.h:426
__tanf
__DEVICE__ float __tanf(float __x)
Definition: __clang_hip_math.h:686
roundf
__DEVICE__ float roundf(float __x)
Definition: __clang_hip_math.h:451
rint
__DEVICE__ double rint(double __x)
Definition: __clang_hip_math.h:974
atanhf
__DEVICE__ float atanhf(float __x)
Definition: __clang_hip_math.h:168
lrint
__DEVICE__ long int lrint(double __x)
Definition: __clang_hip_math.h:867
__y
static __inline unsigned char unsigned int unsigned int __y
Definition: adxintrin.h:22
atan2
__DEVICE__ double atan2(double __x, double __y)
Definition: __clang_hip_math.h:707
__sincosf
__DEVICE__ void __sincosf(float __x, float *__sinptr, float *__cosptr)
Definition: __clang_hip_math.h:677
sqrt
__DEVICE__ double sqrt(double __x)
Definition: __clang_hip_math.h:1041
sin
__DEVICE__ double sin(double __x)
Definition: __clang_hip_math.h:1016
__ocml_i0_f64
__device__ double __ocml_i0_f64(double)
ilogbf
__DEVICE__ int ilogbf(float __x)
Definition: __clang_hip_math.h:261
fabs
__DEVICE__ double fabs(double __x)
Definition: __clang_hip_math.h:766
__ddiv_ru
__DEVICE__ double __ddiv_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:77
cosh
__DEVICE__ double cosh(double __x)
Definition: __clang_hip_math.h:727
fdimf
__DEVICE__ float fdimf(float __x, float __y)
Definition: __clang_hip_math.h:225
__ocml_remquo_f32
__device__ float __ocml_remquo_f32(float, float, __attribute__((address_space(5))) int *)
__frcp_rn
__DEVICE__ float __frcp_rn(float __x)
Definition: __clang_hip_math.h:627
__ocml_cos_f32
__device__ float __ocml_cos_f32(float)
sign
float __ovld __cnfn sign(float x)
Returns 1.0 if x > 0, -0.0 if x = -0.0, +0.0 if x = +0.0, or -1.0 if x < 0.
rnorm
__DEVICE__ double rnorm(int __dim, const double *__a)
Definition: __clang_hip_math.h:977
exp10f
__DEVICE__ float exp10f(float __x)
Definition: __clang_hip_math.h:210
erfcinv
__DEVICE__ double erfcinv(double __x)
Definition: __clang_hip_math.h:745
jn
__DEVICE__ double jn(int __n, double __x)
Definition: __clang_hip_math.h:819
cos
__DEVICE__ double cos(double __x)
Definition: __clang_hip_math.h:724
__dsub_ru
__DEVICE__ double __dsub_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:166
__make_mantissa_base16
__DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp)
Definition: __clang_hip_math.h:93
fma
__DEVICE__ double fma(double __x, double __y, double __z)
Definition: __clang_hip_math.h:775
__fsqrt_rz
__DEVICE__ float __fsqrt_rz(float __a)
Definition: __clang_cuda_device_functions.h:321
tgamma
__DEVICE__ double tgamma(double __x)
Definition: __clang_hip_math.h:1050
ldexp
__DEVICE__ double ldexp(double __x, int __e)
Definition: __clang_hip_math.h:840
lround
__DEVICE__ long int lround(double __x)
Definition: __clang_hip_math.h:870
stdint.h
__powf
__DEVICE__ float __powf(float __x, float __y)
Definition: __clang_hip_math.h:671
__isnan
__DEVICE__ __RETURN_TYPE __isnan(double __x)
Definition: __clang_hip_math.h:810
sincos
__DEVICE__ void sincos(double __x, double *__sinptr, double *__cosptr)
Definition: __clang_hip_math.h:1019
cyl_bessel_i0f
__DEVICE__ float cyl_bessel_i0f(float __x)
Definition: __clang_hip_math.h:189
sqrtf
__DEVICE__ float sqrtf(float __x)
Definition: __clang_hip_math.h:494
__ocml_frexp_f64
__device__ double __ocml_frexp_f64(double, __attribute__((address_space(5))) int *)
__drcp_ru
__DEVICE__ double __drcp_ru(double __a)
Definition: __clang_cuda_device_functions.h:154
__ocml_j1_f32
__device__ float __ocml_j1_f32(float)
__dsqrt_ru
__DEVICE__ double __dsqrt_ru(double __a)
Definition: __clang_cuda_device_functions.h:158
scalblnf
__DEVICE__ float scalblnf(float __x, long int __n)
Definition: __clang_hip_math.h:457
nanf
__DEVICE__ float nanf(const char *__tagp)
Definition: __clang_hip_math.h:342
__make_mantissa_base8
__DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp)
Definition: __clang_hip_math.h:59
erff
__DEVICE__ float erff(float __x)
Definition: __clang_hip_math.h:204
rsqrt
__DEVICE__ double rsqrt(double __x)
Definition: __clang_hip_math.h:1002
cospif
__DEVICE__ float cospif(float __x)
Definition: __clang_hip_math.h:186
cyl_bessel_i1
__DEVICE__ double cyl_bessel_i1(double __x)
Definition: __clang_hip_math.h:736
scalbnf
__DEVICE__ float scalbnf(float __x, int __n)
Definition: __clang_hip_math.h:463
__make_mantissa
__DEVICE__ uint64_t __make_mantissa(const char *__tagp)
Definition: __clang_hip_math.h:114
__fsub_rn
__DEVICE__ float __fsub_rn(float __x, float __y)
Definition: __clang_hip_math.h:658
__ocml_y1_f64
__device__ double __ocml_y1_f64(double)
log2f
__DEVICE__ float log2f(float __x)
Definition: __clang_hip_math.h:318
__ocml_frexp_f32
__device__ float __ocml_frexp_f32(float, __attribute__((address_space(5))) int *)
__fdiv_rn
__DEVICE__ float __fdiv_rn(float __x, float __y)
Definition: __clang_hip_math.h:572
rnorm3df
__DEVICE__ float rnorm3df(float __x, float __y, float __z)
Definition: __clang_hip_math.h:429
pow
__DEVICE__ double pow(double __x, double __y)
Definition: __clang_hip_math.h:947
cyl_bessel_i0
__DEVICE__ double cyl_bessel_i0(double __x)
Definition: __clang_hip_math.h:733
__ocml_native_cos_f32
__device__ float __ocml_native_cos_f32(float)
__frcp_rd
__DEVICE__ float __frcp_rd(float __a)
Definition: __clang_cuda_device_functions.h:313
erf
__DEVICE__ double erf(double __x)
Definition: __clang_hip_math.h:739
__ocml_tan_f32
__device__ float __ocml_tan_f32(float)
logf
__DEVICE__ float logf(float __x)
Definition: __clang_hip_math.h:324
atanf
__DEVICE__ float atanf(float __x)
Definition: __clang_hip_math.h:165
y0
__DEVICE__ double y0(double __x)
Definition: __clang_hip_math.h:1056
y0f
__DEVICE__ float y0f(float __x)
Definition: __clang_hip_math.h:509
y1
__DEVICE__ double y1(double __x)
Definition: __clang_hip_math.h:1059
__fadd_ru
__DEVICE__ float __fadd_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:201
__ocml_j1_f64
__device__ double __ocml_j1_f64(double)
__isinf
__DEVICE__ __RETURN_TYPE __isinf(double __x)
Definition: __clang_hip_math.h:807
__dsqrt_rn
__DEVICE__ double __dsqrt_rn(double __x)
Definition: __clang_hip_math.h:1175
atanh
__DEVICE__ double atanh(double __x)
Definition: __clang_hip_math.h:710
trunc
__DEVICE__ double trunc(double __x)
Definition: __clang_hip_math.h:1053
__fmaf_rz
__DEVICE__ float __fmaf_rz(float __a, float __b, float __c)
Definition: __clang_cuda_device_functions.h:298
__ocml_sincos_f32
__device__ float __ocml_sincos_f32(float, __attribute__((address_space(5))) float *)
llabs
__DEVICE__ long long llabs(long long __a)
Definition: __clang_cuda_math.h:158
powi
__DEVICE__ double powi(double __x, int __y)
Definition: __clang_hip_math.h:950
hypot
__DEVICE__ double hypot(double __x, double __y)
Definition: __clang_hip_math.h:798
norm3df
__DEVICE__ float norm3df(float __x, float __y, float __z)
Definition: __clang_hip_math.h:371
__signbitf
__DEVICE__ __RETURN_TYPE __signbitf(float __x)
Definition: __clang_hip_math.h:466
fminf
__DEVICE__ float fminf(float __x, float __y)
Definition: __clang_hip_math.h:242
__ocml_cos_f64
__device__ double __ocml_cos_f64(double)
atan
__DEVICE__ double atan(double __x)
Definition: __clang_hip_math.h:704
normcdfinvf
__DEVICE__ float normcdfinvf(float __x)
Definition: __clang_hip_math.h:384
remainderf
__DEVICE__ float remainderf(float __x, float __y)
Definition: __clang_hip_math.h:408
__ddiv_rn
__DEVICE__ double __ddiv_rn(double __x, double __y)
Definition: __clang_hip_math.h:1125
__make_mantissa_base10
__DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp)
Definition: __clang_hip_math.h:76
scalbln
__DEVICE__ double scalbln(double __x, long int __n)
Definition: __clang_hip_math.h:1005
norm4d
__DEVICE__ double norm4d(double __x, double __y, double __z, double __w)
Definition: __clang_hip_math.h:936
__ocml_cospi_f64
__device__ double __ocml_cospi_f64(double)
__attribute__
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
Definition: __clang_hip_libdevice_declares.h:299
cyl_bessel_i1f
__DEVICE__ float cyl_bessel_i1f(float __x)
Definition: __clang_hip_math.h:192
powif
__DEVICE__ float powif(float __x, int __y)
Definition: __clang_hip_math.h:402
lgamma
__DEVICE__ double lgamma(double __x)
Definition: __clang_hip_math.h:843
logb
__DEVICE__ double logb(double __x)
Definition: __clang_hip_math.h:864
__fsqrt_rd
__DEVICE__ float __fsqrt_rd(float __a)
Definition: __clang_cuda_device_functions.h:318
__fadd_rd
__DEVICE__ float __fadd_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:195
__fdiv_ru
__DEVICE__ float __fdiv_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:213
sinpif
__DEVICE__ float sinpif(float __x)
Definition: __clang_hip_math.h:491
expm1
__DEVICE__ double expm1(double __x)
Definition: __clang_hip_math.h:763
__dadd_rn
__DEVICE__ double __dadd_rn(double __x, double __y)
Definition: __clang_hip_math.h:1103
__fadd_rn
__DEVICE__ float __fadd_rn(float __x, float __y)
Definition: __clang_hip_math.h:558
__ocml_sin_f64
__device__ double __ocml_sin_f64(double)
llround
__DEVICE__ long long int llround(double __x)
Definition: __clang_hip_math.h:849
__dsub_rd
__DEVICE__ double __dsub_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:160
__ocml_sincos_f64
__device__ double __ocml_sincos_f64(double, __attribute__((address_space(5))) double *)
tanhf
__DEVICE__ float tanhf(float __x)
Definition: __clang_hip_math.h:500
__fmul_rd
__DEVICE__ float __fmul_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:301
scalbn
__DEVICE__ double scalbn(double __x, int __n)
Definition: __clang_hip_math.h:1010
__log10f
__DEVICE__ float __log10f(float __x)
Definition: __clang_hip_math.h:662
j0f
__DEVICE__ float j0f(float __x)
Definition: __clang_hip_math.h:273
__fsub_rz
__DEVICE__ float __fsub_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:331
rnorm3d
__DEVICE__ double rnorm3d(double __x, double __y, double __z)
Definition: __clang_hip_math.h:989
__fma_ru
__DEVICE__ double __fma_ru(double __a, double __b, double __c)
Definition: __clang_cuda_device_functions.h:271
__fmaf_rd
__DEVICE__ float __fmaf_rd(float __a, float __b, float __c)
Definition: __clang_cuda_device_functions.h:289
fmax
__DEVICE__ double fmax(double __x, double __y)
Definition: __clang_hip_math.h:780
__ocml_tgamma_f64
__device__ double __ocml_tgamma_f64(double)
__saturatef
__DEVICE__ float __saturatef(float __x)
Definition: __clang_hip_math.h:674
llroundf
__DEVICE__ long long int llroundf(float __x)
Definition: __clang_hip_math.h:309
expf
__DEVICE__ float expf(float __x)
Definition: __clang_hip_math.h:216
acosf
__DEVICE__ float acosf(float __x)
Definition: __clang_hip_math.h:150
true
#define true
Definition: stdbool.h:16
__static_assert_type_size_equal
#define __static_assert_type_size_equal(A, B)
Definition: __clang_hip_math.h:53
__ocml_remquo_f64
__device__ double __ocml_remquo_f64(double, double, __attribute__((address_space(5))) int *)
modff
__DEVICE__ float modff(float __x, float *__iptr)
Definition: __clang_hip_math.h:333
normcdfinv
__DEVICE__ double normcdfinv(double __x)
Definition: __clang_hip_math.h:944
rnorm4d
__DEVICE__ double rnorm4d(double __x, double __y, double __z, double __w)
Definition: __clang_hip_math.h:994
__ocml_tgamma_f32
__device__ float __ocml_tgamma_f32(float)
fmodf
__DEVICE__ float fmodf(float __x, float __y)
Definition: __clang_hip_math.h:245
__ocml_lgamma_f32
__device__ float __ocml_lgamma_f32(float)
__ocml_sincospi_f32
__device__ float __ocml_sincospi_f32(float, __attribute__((address_space(5))) float *)
__dsqrt_rd
__DEVICE__ double __dsqrt_rd(double __a)
Definition: __clang_cuda_device_functions.h:156
fdim
__DEVICE__ double fdim(double __x, double __y)
Definition: __clang_hip_math.h:769
fabsf
__DEVICE__ float fabsf(float __x)
Definition: __clang_hip_math.h:222
__sinf
__DEVICE__ float __sinf(float __x)
Definition: __clang_hip_math.h:683
__fmaf_rn
__DEVICE__ float __fmaf_rn(float __x, float __y, float __z)
Definition: __clang_hip_math.h:597
abs
__DEVICE__ long long abs(long long __n)
Definition: __clang_cuda_cmath.h:41
log1pf
__DEVICE__ float log1pf(float __x)
Definition: __clang_hip_math.h:315
__ddiv_rd
__DEVICE__ double __ddiv_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:71
__isinff
__DEVICE__ __RETURN_TYPE __isinff(float __x)
Definition: __clang_hip_math.h:267
sinhf
__DEVICE__ float sinhf(float __x)
Definition: __clang_hip_math.h:488