clang  14.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__) && !defined(__OPENMP_AMDGCN__)
13 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
14 #endif
15 
16 #if !defined(__HIPCC_RTC__)
17 #if defined(__cplusplus)
18 #include <algorithm>
19 #endif
20 #include <limits.h>
21 #include <stdint.h>
22 #ifdef __OPENMP_AMDGCN__
23 #include <omp.h>
24 #endif
25 #endif // !defined(__HIPCC_RTC__)
26 
27 #pragma push_macro("__DEVICE__")
28 
29 #ifdef __OPENMP_AMDGCN__
30 #define __DEVICE__ static inline __attribute__((always_inline, nothrow))
31 #else
32 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
33 #endif
34 
35 // A few functions return bool type starting only in C++11.
36 #pragma push_macro("__RETURN_TYPE")
37 #ifdef __OPENMP_AMDGCN__
38 #define __RETURN_TYPE int
39 #else
40 #if defined(__cplusplus)
41 #define __RETURN_TYPE bool
42 #else
43 #define __RETURN_TYPE int
44 #endif
45 #endif // __OPENMP_AMDGCN__
46 
47 #if defined (__cplusplus) && __cplusplus < 201103L
48 // emulate static_assert on type sizes
49 template<bool>
50 struct __compare_result{};
51 template<>
52 struct __compare_result<true> {
53  static const __device__ bool valid;
54 };
55 
57 void __suppress_unused_warning(bool b){};
58 template <unsigned int S, unsigned int T>
59 __DEVICE__ void __static_assert_equal_size() {
60  __suppress_unused_warning(__compare_result<S == T>::valid);
61 }
62 
63 #define __static_assert_type_size_equal(A, B) \
64  __static_assert_equal_size<A,B>()
65 
66 #else
67 #define __static_assert_type_size_equal(A,B) \
68  static_assert((A) == (B), "")
69 
70 #endif
71 
73 uint64_t __make_mantissa_base8(const char *__tagp) {
74  uint64_t __r = 0;
75  while (__tagp) {
76  char __tmp = *__tagp;
77 
78  if (__tmp >= '0' && __tmp <= '7')
79  __r = (__r * 8u) + __tmp - '0';
80  else
81  return 0;
82 
83  ++__tagp;
84  }
85 
86  return __r;
87 }
88 
90 uint64_t __make_mantissa_base10(const char *__tagp) {
91  uint64_t __r = 0;
92  while (__tagp) {
93  char __tmp = *__tagp;
94 
95  if (__tmp >= '0' && __tmp <= '9')
96  __r = (__r * 10u) + __tmp - '0';
97  else
98  return 0;
99 
100  ++__tagp;
101  }
102 
103  return __r;
104 }
105 
107 uint64_t __make_mantissa_base16(const char *__tagp) {
108  uint64_t __r = 0;
109  while (__tagp) {
110  char __tmp = *__tagp;
111 
112  if (__tmp >= '0' && __tmp <= '9')
113  __r = (__r * 16u) + __tmp - '0';
114  else if (__tmp >= 'a' && __tmp <= 'f')
115  __r = (__r * 16u) + __tmp - 'a' + 10;
116  else if (__tmp >= 'A' && __tmp <= 'F')
117  __r = (__r * 16u) + __tmp - 'A' + 10;
118  else
119  return 0;
120 
121  ++__tagp;
122  }
123 
124  return __r;
125 }
126 
128 uint64_t __make_mantissa(const char *__tagp) {
129  if (!__tagp)
130  return 0u;
131 
132  if (*__tagp == '0') {
133  ++__tagp;
134 
135  if (*__tagp == 'x' || *__tagp == 'X')
136  return __make_mantissa_base16(__tagp);
137  else
138  return __make_mantissa_base8(__tagp);
139  }
140 
141  return __make_mantissa_base10(__tagp);
142 }
143 
144 // BEGIN FLOAT
145 #if defined(__cplusplus)
147 int abs(int __x) {
148  int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1);
149  return (__x ^ __sgn) - __sgn;
150 }
152 long labs(long __x) {
153  long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1);
154  return (__x ^ __sgn) - __sgn;
155 }
157 long long llabs(long long __x) {
158  long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
159  return (__x ^ __sgn) - __sgn;
160 }
161 #endif
162 
164 float acosf(float __x) { return __ocml_acos_f32(__x); }
165 
167 float acoshf(float __x) { return __ocml_acosh_f32(__x); }
168 
170 float asinf(float __x) { return __ocml_asin_f32(__x); }
171 
173 float asinhf(float __x) { return __ocml_asinh_f32(__x); }
174 
176 float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); }
177 
179 float atanf(float __x) { return __ocml_atan_f32(__x); }
180 
182 float atanhf(float __x) { return __ocml_atanh_f32(__x); }
183 
185 float cbrtf(float __x) { return __ocml_cbrt_f32(__x); }
186 
188 float ceilf(float __x) { return __ocml_ceil_f32(__x); }
189 
191 float copysignf(float __x, float __y) { return __ocml_copysign_f32(__x, __y); }
192 
194 float cosf(float __x) { return __ocml_cos_f32(__x); }
195 
197 float coshf(float __x) { return __ocml_cosh_f32(__x); }
198 
200 float cospif(float __x) { return __ocml_cospi_f32(__x); }
201 
203 float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); }
204 
206 float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); }
207 
209 float erfcf(float __x) { return __ocml_erfc_f32(__x); }
210 
212 float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); }
213 
215 float erfcxf(float __x) { return __ocml_erfcx_f32(__x); }
216 
218 float erff(float __x) { return __ocml_erf_f32(__x); }
219 
221 float erfinvf(float __x) { return __ocml_erfinv_f32(__x); }
222 
224 float exp10f(float __x) { return __ocml_exp10_f32(__x); }
225 
227 float exp2f(float __x) { return __ocml_exp2_f32(__x); }
228 
230 float expf(float __x) { return __ocml_exp_f32(__x); }
231 
233 float expm1f(float __x) { return __ocml_expm1_f32(__x); }
234 
236 float fabsf(float __x) { return __ocml_fabs_f32(__x); }
237 
239 float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); }
240 
242 float fdividef(float __x, float __y) { return __x / __y; }
243 
245 float floorf(float __x) { return __ocml_floor_f32(__x); }
246 
248 float fmaf(float __x, float __y, float __z) {
249  return __ocml_fma_f32(__x, __y, __z);
250 }
251 
253 float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); }
254 
256 float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); }
257 
259 float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
260 
262 float frexpf(float __x, int *__nptr) {
263  int __tmp;
264 #ifdef __OPENMP_AMDGCN__
265 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
266 #endif
267  float __r =
268  __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp);
269  *__nptr = __tmp;
270 
271  return __r;
272 }
273 
275 float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); }
276 
278 int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }
279 
281 __RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); }
282 
284 __RETURN_TYPE __isinff(float __x) { return __ocml_isinf_f32(__x); }
285 
287 __RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); }
288 
290 float j0f(float __x) { return __ocml_j0_f32(__x); }
291 
293 float j1f(float __x) { return __ocml_j1_f32(__x); }
294 
296 float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication
297  // and the Miller & Brown algorithm
298  // for linear recurrences to get O(log n) steps, but it's unclear if
299  // it'd be beneficial in this case.
300  if (__n == 0)
301  return j0f(__x);
302  if (__n == 1)
303  return j1f(__x);
304 
305  float __x0 = j0f(__x);
306  float __x1 = j1f(__x);
307  for (int __i = 1; __i < __n; ++__i) {
308  float __x2 = (2 * __i) / __x * __x1 - __x0;
309  __x0 = __x1;
310  __x1 = __x2;
311  }
312 
313  return __x1;
314 }
315 
317 float ldexpf(float __x, int __e) { return __ocml_ldexp_f32(__x, __e); }
318 
320 float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
321 
323 long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
324 
326 long long int llroundf(float __x) { return __ocml_round_f32(__x); }
327 
329 float log10f(float __x) { return __ocml_log10_f32(__x); }
330 
332 float log1pf(float __x) { return __ocml_log1p_f32(__x); }
333 
335 float log2f(float __x) { return __ocml_log2_f32(__x); }
336 
338 float logbf(float __x) { return __ocml_logb_f32(__x); }
339 
341 float logf(float __x) { return __ocml_log_f32(__x); }
342 
344 long int lrintf(float __x) { return __ocml_rint_f32(__x); }
345 
347 long int lroundf(float __x) { return __ocml_round_f32(__x); }
348 
350 float modff(float __x, float *__iptr) {
351  float __tmp;
352 #ifdef __OPENMP_AMDGCN__
353 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
354 #endif
355  float __r =
356  __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
357  *__iptr = __tmp;
358  return __r;
359 }
360 
362 float nanf(const char *__tagp) {
363  union {
364  float val;
365  struct ieee_float {
366  unsigned int mantissa : 22;
367  unsigned int quiet : 1;
368  unsigned int exponent : 8;
369  unsigned int sign : 1;
370  } bits;
371  } __tmp;
372  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
373 
374  __tmp.bits.sign = 0u;
375  __tmp.bits.exponent = ~0u;
376  __tmp.bits.quiet = 1u;
377  __tmp.bits.mantissa = __make_mantissa(__tagp);
378 
379  return __tmp.val;
380 }
381 
383 float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); }
384 
386 float nextafterf(float __x, float __y) {
387  return __ocml_nextafter_f32(__x, __y);
388 }
389 
391 float norm3df(float __x, float __y, float __z) {
392  return __ocml_len3_f32(__x, __y, __z);
393 }
394 
396 float norm4df(float __x, float __y, float __z, float __w) {
397  return __ocml_len4_f32(__x, __y, __z, __w);
398 }
399 
401 float normcdff(float __x) { return __ocml_ncdf_f32(__x); }
402 
404 float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); }
405 
407 float normf(int __dim,
408  const float *__a) { // TODO: placeholder until OCML adds support.
409  float __r = 0;
410  while (__dim--) {
411  __r += __a[0] * __a[0];
412  ++__a;
413  }
414 
415  return __ocml_sqrt_f32(__r);
416 }
417 
419 float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
420 
422 float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); }
423 
425 float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }
426 
428 float remainderf(float __x, float __y) {
429  return __ocml_remainder_f32(__x, __y);
430 }
431 
433 float remquof(float __x, float __y, int *__quo) {
434  int __tmp;
435 #ifdef __OPENMP_AMDGCN__
436 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
437 #endif
438  float __r = __ocml_remquo_f32(
439  __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
440  *__quo = __tmp;
441 
442  return __r;
443 }
444 
446 float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }
447 
449 float rintf(float __x) { return __ocml_rint_f32(__x); }
450 
452 float rnorm3df(float __x, float __y, float __z) {
453  return __ocml_rlen3_f32(__x, __y, __z);
454 }
455 
457 float rnorm4df(float __x, float __y, float __z, float __w) {
458  return __ocml_rlen4_f32(__x, __y, __z, __w);
459 }
460 
462 float rnormf(int __dim,
463  const float *__a) { // TODO: placeholder until OCML adds support.
464  float __r = 0;
465  while (__dim--) {
466  __r += __a[0] * __a[0];
467  ++__a;
468  }
469 
470  return __ocml_rsqrt_f32(__r);
471 }
472 
474 float roundf(float __x) { return __ocml_round_f32(__x); }
475 
477 float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
478 
480 float scalblnf(float __x, long int __n) {
481  return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n)
482  : __ocml_scalb_f32(__x, __n);
483 }
484 
486 float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
487 
489 __RETURN_TYPE __signbitf(float __x) { return __ocml_signbit_f32(__x); }
490 
492 void sincosf(float __x, float *__sinptr, float *__cosptr) {
493  float __tmp;
494 #ifdef __OPENMP_AMDGCN__
495 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
496 #endif
497  *__sinptr =
498  __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
499  *__cosptr = __tmp;
500 }
501 
503 void sincospif(float __x, float *__sinptr, float *__cosptr) {
504  float __tmp;
505 #ifdef __OPENMP_AMDGCN__
506 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
507 #endif
508  *__sinptr = __ocml_sincospi_f32(
509  __x, (__attribute__((address_space(5))) float *)&__tmp);
510  *__cosptr = __tmp;
511 }
512 
514 float sinf(float __x) { return __ocml_sin_f32(__x); }
515 
517 float sinhf(float __x) { return __ocml_sinh_f32(__x); }
518 
520 float sinpif(float __x) { return __ocml_sinpi_f32(__x); }
521 
523 float sqrtf(float __x) { return __ocml_sqrt_f32(__x); }
524 
526 float tanf(float __x) { return __ocml_tan_f32(__x); }
527 
529 float tanhf(float __x) { return __ocml_tanh_f32(__x); }
530 
532 float tgammaf(float __x) { return __ocml_tgamma_f32(__x); }
533 
535 float truncf(float __x) { return __ocml_trunc_f32(__x); }
536 
538 float y0f(float __x) { return __ocml_y0_f32(__x); }
539 
541 float y1f(float __x) { return __ocml_y1_f32(__x); }
542 
544 float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication
545  // and the Miller & Brown algorithm
546  // for linear recurrences to get O(log n) steps, but it's unclear if
547  // it'd be beneficial in this case. Placeholder until OCML adds
548  // support.
549  if (__n == 0)
550  return y0f(__x);
551  if (__n == 1)
552  return y1f(__x);
553 
554  float __x0 = y0f(__x);
555  float __x1 = y1f(__x);
556  for (int __i = 1; __i < __n; ++__i) {
557  float __x2 = (2 * __i) / __x * __x1 - __x0;
558  __x0 = __x1;
559  __x1 = __x2;
560  }
561 
562  return __x1;
563 }
564 
565 // BEGIN INTRINSICS
566 
568 float __cosf(float __x) { return __ocml_native_cos_f32(__x); }
569 
571 float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }
572 
574 float __expf(float __x) { return __ocml_native_exp_f32(__x); }
575 
576 #if defined OCML_BASIC_ROUNDED_OPERATIONS
578 float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); }
580 float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); }
582 float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); }
584 float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); }
585 #else
587 float __fadd_rn(float __x, float __y) { return __x + __y; }
588 #endif
589 
590 #if defined OCML_BASIC_ROUNDED_OPERATIONS
592 float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); }
594 float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); }
596 float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); }
598 float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); }
599 #else
601 float __fdiv_rn(float __x, float __y) { return __x / __y; }
602 #endif
603 
605 float __fdividef(float __x, float __y) { return __x / __y; }
606 
607 #if defined OCML_BASIC_ROUNDED_OPERATIONS
609 float __fmaf_rd(float __x, float __y, float __z) {
610  return __ocml_fma_rtn_f32(__x, __y, __z);
611 }
613 float __fmaf_rn(float __x, float __y, float __z) {
614  return __ocml_fma_rte_f32(__x, __y, __z);
615 }
617 float __fmaf_ru(float __x, float __y, float __z) {
618  return __ocml_fma_rtp_f32(__x, __y, __z);
619 }
621 float __fmaf_rz(float __x, float __y, float __z) {
622  return __ocml_fma_rtz_f32(__x, __y, __z);
623 }
624 #else
626 float __fmaf_rn(float __x, float __y, float __z) {
627  return __ocml_fma_f32(__x, __y, __z);
628 }
629 #endif
630 
631 #if defined OCML_BASIC_ROUNDED_OPERATIONS
633 float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }
635 float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }
637 float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); }
639 float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); }
640 #else
642 float __fmul_rn(float __x, float __y) { return __x * __y; }
643 #endif
644 
645 #if defined OCML_BASIC_ROUNDED_OPERATIONS
647 float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); }
649 float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); }
651 float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); }
653 float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }
654 #else
656 float __frcp_rn(float __x) { return 1.0f / __x; }
657 #endif
658 
660 float __frsqrt_rn(float __x) { return __llvm_amdgcn_rsq_f32(__x); }
661 
662 #if defined OCML_BASIC_ROUNDED_OPERATIONS
664 float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
666 float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }
668 float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }
670 float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }
671 #else
673 float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }
674 #endif
675 
676 #if defined OCML_BASIC_ROUNDED_OPERATIONS
678 float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); }
680 float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); }
682 float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); }
684 float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); }
685 #else
687 float __fsub_rn(float __x, float __y) { return __x - __y; }
688 #endif
689 
691 float __log10f(float __x) { return __ocml_native_log10_f32(__x); }
692 
694 float __log2f(float __x) { return __ocml_native_log2_f32(__x); }
695 
697 float __logf(float __x) { return __ocml_native_log_f32(__x); }
698 
700 float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
701 
703 float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }
704 
706 void __sincosf(float __x, float *__sinptr, float *__cosptr) {
707  *__sinptr = __ocml_native_sin_f32(__x);
708  *__cosptr = __ocml_native_cos_f32(__x);
709 }
710 
712 float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
713 
715 float __tanf(float __x) { return __ocml_tan_f32(__x); }
716 // END INTRINSICS
717 // END FLOAT
718 
719 // BEGIN DOUBLE
721 double acos(double __x) { return __ocml_acos_f64(__x); }
722 
724 double acosh(double __x) { return __ocml_acosh_f64(__x); }
725 
727 double asin(double __x) { return __ocml_asin_f64(__x); }
728 
730 double asinh(double __x) { return __ocml_asinh_f64(__x); }
731 
733 double atan(double __x) { return __ocml_atan_f64(__x); }
734 
736 double atan2(double __x, double __y) { return __ocml_atan2_f64(__x, __y); }
737 
739 double atanh(double __x) { return __ocml_atanh_f64(__x); }
740 
742 double cbrt(double __x) { return __ocml_cbrt_f64(__x); }
743 
745 double ceil(double __x) { return __ocml_ceil_f64(__x); }
746 
748 double copysign(double __x, double __y) {
749  return __ocml_copysign_f64(__x, __y);
750 }
751 
753 double cos(double __x) { return __ocml_cos_f64(__x); }
754 
756 double cosh(double __x) { return __ocml_cosh_f64(__x); }
757 
759 double cospi(double __x) { return __ocml_cospi_f64(__x); }
760 
762 double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); }
763 
765 double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); }
766 
768 double erf(double __x) { return __ocml_erf_f64(__x); }
769 
771 double erfc(double __x) { return __ocml_erfc_f64(__x); }
772 
774 double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); }
775 
777 double erfcx(double __x) { return __ocml_erfcx_f64(__x); }
778 
780 double erfinv(double __x) { return __ocml_erfinv_f64(__x); }
781 
783 double exp(double __x) { return __ocml_exp_f64(__x); }
784 
786 double exp10(double __x) { return __ocml_exp10_f64(__x); }
787 
789 double exp2(double __x) { return __ocml_exp2_f64(__x); }
790 
792 double expm1(double __x) { return __ocml_expm1_f64(__x); }
793 
795 double fabs(double __x) { return __ocml_fabs_f64(__x); }
796 
798 double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }
799 
801 double floor(double __x) { return __ocml_floor_f64(__x); }
802 
804 double fma(double __x, double __y, double __z) {
805  return __ocml_fma_f64(__x, __y, __z);
806 }
807 
809 double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); }
810 
812 double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); }
813 
815 double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
816 
818 double frexp(double __x, int *__nptr) {
819  int __tmp;
820 #ifdef __OPENMP_AMDGCN__
821 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
822 #endif
823  double __r =
824  __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp);
825  *__nptr = __tmp;
826  return __r;
827 }
828 
830 double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); }
831 
833 int ilogb(double __x) { return __ocml_ilogb_f64(__x); }
834 
836 __RETURN_TYPE __finite(double __x) { return __ocml_isfinite_f64(__x); }
837 
839 __RETURN_TYPE __isinf(double __x) { return __ocml_isinf_f64(__x); }
840 
842 __RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); }
843 
845 double j0(double __x) { return __ocml_j0_f64(__x); }
846 
848 double j1(double __x) { return __ocml_j1_f64(__x); }
849 
851 double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication
852  // and the Miller & Brown algorithm
853  // for linear recurrences to get O(log n) steps, but it's unclear if
854  // it'd be beneficial in this case. Placeholder until OCML adds
855  // support.
856  if (__n == 0)
857  return j0(__x);
858  if (__n == 1)
859  return j1(__x);
860 
861  double __x0 = j0(__x);
862  double __x1 = j1(__x);
863  for (int __i = 1; __i < __n; ++__i) {
864  double __x2 = (2 * __i) / __x * __x1 - __x0;
865  __x0 = __x1;
866  __x1 = __x2;
867  }
868  return __x1;
869 }
870 
872 double ldexp(double __x, int __e) { return __ocml_ldexp_f64(__x, __e); }
873 
875 double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
876 
878 long long int llrint(double __x) { return __ocml_rint_f64(__x); }
879 
881 long long int llround(double __x) { return __ocml_round_f64(__x); }
882 
884 double log(double __x) { return __ocml_log_f64(__x); }
885 
887 double log10(double __x) { return __ocml_log10_f64(__x); }
888 
890 double log1p(double __x) { return __ocml_log1p_f64(__x); }
891 
893 double log2(double __x) { return __ocml_log2_f64(__x); }
894 
896 double logb(double __x) { return __ocml_logb_f64(__x); }
897 
899 long int lrint(double __x) { return __ocml_rint_f64(__x); }
900 
902 long int lround(double __x) { return __ocml_round_f64(__x); }
903 
905 double modf(double __x, double *__iptr) {
906  double __tmp;
907 #ifdef __OPENMP_AMDGCN__
908 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
909 #endif
910  double __r =
911  __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
912  *__iptr = __tmp;
913 
914  return __r;
915 }
916 
918 double nan(const char *__tagp) {
919 #if !_WIN32
920  union {
921  double val;
922  struct ieee_double {
923  uint64_t mantissa : 51;
924  uint32_t quiet : 1;
925  uint32_t exponent : 11;
926  uint32_t sign : 1;
927  } bits;
928  } __tmp;
929  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
930 
931  __tmp.bits.sign = 0u;
932  __tmp.bits.exponent = ~0u;
933  __tmp.bits.quiet = 1u;
934  __tmp.bits.mantissa = __make_mantissa(__tagp);
935 
936  return __tmp.val;
937 #else
938  __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double));
939  uint64_t __val = __make_mantissa(__tagp);
940  __val |= 0xFFF << 51;
941  return *reinterpret_cast<double *>(&__val);
942 #endif
943 }
944 
946 double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); }
947 
949 double nextafter(double __x, double __y) {
950  return __ocml_nextafter_f64(__x, __y);
951 }
952 
954 double norm(int __dim,
955  const double *__a) { // TODO: placeholder until OCML adds support.
956  double __r = 0;
957  while (__dim--) {
958  __r += __a[0] * __a[0];
959  ++__a;
960  }
961 
962  return __ocml_sqrt_f64(__r);
963 }
964 
966 double norm3d(double __x, double __y, double __z) {
967  return __ocml_len3_f64(__x, __y, __z);
968 }
969 
971 double norm4d(double __x, double __y, double __z, double __w) {
972  return __ocml_len4_f64(__x, __y, __z, __w);
973 }
974 
976 double normcdf(double __x) { return __ocml_ncdf_f64(__x); }
977 
979 double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); }
980 
982 double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); }
983 
985 double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); }
986 
988 double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }
989 
991 double remainder(double __x, double __y) {
992  return __ocml_remainder_f64(__x, __y);
993 }
994 
996 double remquo(double __x, double __y, int *__quo) {
997  int __tmp;
998 #ifdef __OPENMP_AMDGCN__
999 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
1000 #endif
1001  double __r = __ocml_remquo_f64(
1002  __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
1003  *__quo = __tmp;
1004 
1005  return __r;
1006 }
1007 
1008 __DEVICE__
1009 double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }
1010 
1011 __DEVICE__
1012 double rint(double __x) { return __ocml_rint_f64(__x); }
1013 
1014 __DEVICE__
1015 double rnorm(int __dim,
1016  const double *__a) { // TODO: placeholder until OCML adds support.
1017  double __r = 0;
1018  while (__dim--) {
1019  __r += __a[0] * __a[0];
1020  ++__a;
1021  }
1022 
1023  return __ocml_rsqrt_f64(__r);
1024 }
1025 
1026 __DEVICE__
1027 double rnorm3d(double __x, double __y, double __z) {
1028  return __ocml_rlen3_f64(__x, __y, __z);
1029 }
1030 
1031 __DEVICE__
1032 double rnorm4d(double __x, double __y, double __z, double __w) {
1033  return __ocml_rlen4_f64(__x, __y, __z, __w);
1034 }
1035 
1036 __DEVICE__
1037 double round(double __x) { return __ocml_round_f64(__x); }
1038 
1039 __DEVICE__
1040 double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }
1041 
1042 __DEVICE__
1043 double scalbln(double __x, long int __n) {
1044  return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n)
1045  : __ocml_scalb_f64(__x, __n);
1046 }
1047 __DEVICE__
1048 double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); }
1049 
1050 __DEVICE__
1051 __RETURN_TYPE __signbit(double __x) { return __ocml_signbit_f64(__x); }
1052 
1053 __DEVICE__
1054 double sin(double __x) { return __ocml_sin_f64(__x); }
1055 
1056 __DEVICE__
1057 void sincos(double __x, double *__sinptr, double *__cosptr) {
1058  double __tmp;
1059 #ifdef __OPENMP_AMDGCN__
1060 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
1061 #endif
1062  *__sinptr = __ocml_sincos_f64(
1063  __x, (__attribute__((address_space(5))) double *)&__tmp);
1064  *__cosptr = __tmp;
1065 }
1066 
1067 __DEVICE__
1068 void sincospi(double __x, double *__sinptr, double *__cosptr) {
1069  double __tmp;
1070 #ifdef __OPENMP_AMDGCN__
1071 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
1072 #endif
1073  *__sinptr = __ocml_sincospi_f64(
1074  __x, (__attribute__((address_space(5))) double *)&__tmp);
1075  *__cosptr = __tmp;
1076 }
1077 
1078 __DEVICE__
1079 double sinh(double __x) { return __ocml_sinh_f64(__x); }
1080 
1081 __DEVICE__
1082 double sinpi(double __x) { return __ocml_sinpi_f64(__x); }
1083 
1084 __DEVICE__
1085 double sqrt(double __x) { return __ocml_sqrt_f64(__x); }
1086 
1087 __DEVICE__
1088 double tan(double __x) { return __ocml_tan_f64(__x); }
1089 
1090 __DEVICE__
1091 double tanh(double __x) { return __ocml_tanh_f64(__x); }
1092 
1093 __DEVICE__
1094 double tgamma(double __x) { return __ocml_tgamma_f64(__x); }
1095 
1096 __DEVICE__
1097 double trunc(double __x) { return __ocml_trunc_f64(__x); }
1098 
1099 __DEVICE__
1100 double y0(double __x) { return __ocml_y0_f64(__x); }
1101 
1102 __DEVICE__
1103 double y1(double __x) { return __ocml_y1_f64(__x); }
1104 
1105 __DEVICE__
1106 double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication
1107  // and the Miller & Brown algorithm
1108  // for linear recurrences to get O(log n) steps, but it's unclear if
1109  // it'd be beneficial in this case. Placeholder until OCML adds
1110  // support.
1111  if (__n == 0)
1112  return y0(__x);
1113  if (__n == 1)
1114  return y1(__x);
1115 
1116  double __x0 = y0(__x);
1117  double __x1 = y1(__x);
1118  for (int __i = 1; __i < __n; ++__i) {
1119  double __x2 = (2 * __i) / __x * __x1 - __x0;
1120  __x0 = __x1;
1121  __x1 = __x2;
1122  }
1123 
1124  return __x1;
1125 }
1126 
1127 // BEGIN INTRINSICS
1128 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1129 __DEVICE__
1130 double __dadd_rd(double __x, double __y) {
1131  return __ocml_add_rtn_f64(__x, __y);
1132 }
1133 __DEVICE__
1134 double __dadd_rn(double __x, double __y) {
1135  return __ocml_add_rte_f64(__x, __y);
1136 }
1137 __DEVICE__
1138 double __dadd_ru(double __x, double __y) {
1139  return __ocml_add_rtp_f64(__x, __y);
1140 }
1141 __DEVICE__
1142 double __dadd_rz(double __x, double __y) {
1143  return __ocml_add_rtz_f64(__x, __y);
1144 }
1145 #else
1146 __DEVICE__
1147 double __dadd_rn(double __x, double __y) { return __x + __y; }
1148 #endif
1149 
1150 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1151 __DEVICE__
1152 double __ddiv_rd(double __x, double __y) {
1153  return __ocml_div_rtn_f64(__x, __y);
1154 }
1155 __DEVICE__
1156 double __ddiv_rn(double __x, double __y) {
1157  return __ocml_div_rte_f64(__x, __y);
1158 }
1159 __DEVICE__
1160 double __ddiv_ru(double __x, double __y) {
1161  return __ocml_div_rtp_f64(__x, __y);
1162 }
1163 __DEVICE__
1164 double __ddiv_rz(double __x, double __y) {
1165  return __ocml_div_rtz_f64(__x, __y);
1166 }
1167 #else
1168 __DEVICE__
1169 double __ddiv_rn(double __x, double __y) { return __x / __y; }
1170 #endif
1171 
1172 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1173 __DEVICE__
1174 double __dmul_rd(double __x, double __y) {
1175  return __ocml_mul_rtn_f64(__x, __y);
1176 }
1177 __DEVICE__
1178 double __dmul_rn(double __x, double __y) {
1179  return __ocml_mul_rte_f64(__x, __y);
1180 }
1181 __DEVICE__
1182 double __dmul_ru(double __x, double __y) {
1183  return __ocml_mul_rtp_f64(__x, __y);
1184 }
1185 __DEVICE__
1186 double __dmul_rz(double __x, double __y) {
1187  return __ocml_mul_rtz_f64(__x, __y);
1188 }
1189 #else
1190 __DEVICE__
1191 double __dmul_rn(double __x, double __y) { return __x * __y; }
1192 #endif
1193 
1194 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1195 __DEVICE__
1196 double __drcp_rd(double __x) { return __ocml_div_rtn_f64(1.0, __x); }
1197 __DEVICE__
1198 double __drcp_rn(double __x) { return __ocml_div_rte_f64(1.0, __x); }
1199 __DEVICE__
1200 double __drcp_ru(double __x) { return __ocml_div_rtp_f64(1.0, __x); }
1201 __DEVICE__
1202 double __drcp_rz(double __x) { return __ocml_div_rtz_f64(1.0, __x); }
1203 #else
1204 __DEVICE__
1205 double __drcp_rn(double __x) { return 1.0 / __x; }
1206 #endif
1207 
1208 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1209 __DEVICE__
1210 double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); }
1211 __DEVICE__
1212 double __dsqrt_rn(double __x) { return __ocml_sqrt_rte_f64(__x); }
1213 __DEVICE__
1214 double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); }
1215 __DEVICE__
1216 double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); }
1217 #else
1218 __DEVICE__
1219 double __dsqrt_rn(double __x) { return __ocml_sqrt_f64(__x); }
1220 #endif
1221 
1222 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1223 __DEVICE__
1224 double __dsub_rd(double __x, double __y) {
1225  return __ocml_sub_rtn_f64(__x, __y);
1226 }
1227 __DEVICE__
1228 double __dsub_rn(double __x, double __y) {
1229  return __ocml_sub_rte_f64(__x, __y);
1230 }
1231 __DEVICE__
1232 double __dsub_ru(double __x, double __y) {
1233  return __ocml_sub_rtp_f64(__x, __y);
1234 }
1235 __DEVICE__
1236 double __dsub_rz(double __x, double __y) {
1237  return __ocml_sub_rtz_f64(__x, __y);
1238 }
1239 #else
1240 __DEVICE__
1241 double __dsub_rn(double __x, double __y) { return __x - __y; }
1242 #endif
1243 
1244 #if defined OCML_BASIC_ROUNDED_OPERATIONS
1245 __DEVICE__
1246 double __fma_rd(double __x, double __y, double __z) {
1247  return __ocml_fma_rtn_f64(__x, __y, __z);
1248 }
1249 __DEVICE__
1250 double __fma_rn(double __x, double __y, double __z) {
1251  return __ocml_fma_rte_f64(__x, __y, __z);
1252 }
1253 __DEVICE__
1254 double __fma_ru(double __x, double __y, double __z) {
1255  return __ocml_fma_rtp_f64(__x, __y, __z);
1256 }
1257 __DEVICE__
1258 double __fma_rz(double __x, double __y, double __z) {
1259  return __ocml_fma_rtz_f64(__x, __y, __z);
1260 }
1261 #else
1262 __DEVICE__
1263 double __fma_rn(double __x, double __y, double __z) {
1264  return __ocml_fma_f64(__x, __y, __z);
1265 }
1266 #endif
1267 // END INTRINSICS
1268 // END DOUBLE
1269 
1270 // C only macros
1271 #if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L
1272 #define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x)
1273 #define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x)
1274 #define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x)
1275 #define signbit(__x) \
1276  _Generic((__x), float : __signbitf, double : __signbit)(__x)
1277 #endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L
1278 
1279 #if defined(__cplusplus)
1280 template <class T> __DEVICE__ T min(T __arg1, T __arg2) {
1281  return (__arg1 < __arg2) ? __arg1 : __arg2;
1282 }
1283 
1284 template <class T> __DEVICE__ T max(T __arg1, T __arg2) {
1285  return (__arg1 > __arg2) ? __arg1 : __arg2;
1286 }
1287 
1288 __DEVICE__ int min(int __arg1, int __arg2) {
1289  return (__arg1 < __arg2) ? __arg1 : __arg2;
1290 }
1291 __DEVICE__ int max(int __arg1, int __arg2) {
1292  return (__arg1 > __arg2) ? __arg1 : __arg2;
1293 }
1294 
1295 __DEVICE__
1296 float max(float __x, float __y) { return fmaxf(__x, __y); }
1297 
1298 __DEVICE__
1299 double max(double __x, double __y) { return fmax(__x, __y); }
1300 
1301 __DEVICE__
1302 float min(float __x, float __y) { return fminf(__x, __y); }
1303 
1304 __DEVICE__
1305 double min(double __x, double __y) { return fmin(__x, __y); }
1306 
1307 #if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
1308 __host__ inline static int min(int __arg1, int __arg2) {
1309  return std::min(__arg1, __arg2);
1310 }
1311 
1312 __host__ inline static int max(int __arg1, int __arg2) {
1313  return std::max(__arg1, __arg2);
1314 }
1315 #endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
1316 #endif
1317 
1318 #pragma pop_macro("__DEVICE__")
1319 #pragma pop_macro("__RETURN_TYPE")
1320 
1321 #endif // __CLANG_HIP_MATH_H__
__fsub_ru
__DEVICE__ float __fsub_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:330
y1f
__DEVICE__ float y1f(float __x)
Definition: __clang_hip_math.h:541
lgammaf
__DEVICE__ float lgammaf(float __x)
Definition: __clang_hip_math.h:320
__ocml_y0_f32
__device__ float __ocml_y0_f32(float)
cbrt
__DEVICE__ double cbrt(double __x)
Definition: __clang_hip_math.h:742
ynf
__DEVICE__ float ynf(int __n, float __x)
Definition: __clang_hip_math.h:544
exp2f
__DEVICE__ float exp2f(float __x)
Definition: __clang_hip_math.h:227
asinhf
__DEVICE__ float asinhf(float __x)
Definition: __clang_hip_math.h:173
__dmul_rd
__DEVICE__ double __dmul_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:85
__fmul_rn
__DEVICE__ float __fmul_rn(float __x, float __y)
Definition: __clang_hip_math.h:642
__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:786
max
__DEVICE__ int max(int __a, int __b)
Definition: __clang_cuda_math.h:196
__fdiv_rz
__DEVICE__ float __fdiv_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:218
atan2f
__DEVICE__ float atan2f(float __x, float __y)
Definition: __clang_hip_math.h:176
__ddiv_rz
__DEVICE__ double __ddiv_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:82
fmaxf
__DEVICE__ float fmaxf(float __x, float __y)
Definition: __clang_hip_math.h:253
__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:477
expm1f
__DEVICE__ float expm1f(float __x)
Definition: __clang_hip_math.h:233
erfcf
__DEVICE__ float erfcf(float __x)
Definition: __clang_hip_math.h:209
powf
__DEVICE__ float powf(float __x, float __y)
Definition: __clang_hip_math.h:419
__frcp_ru
__DEVICE__ float __frcp_ru(float __a)
Definition: __clang_cuda_device_functions.h:317
acosh
__DEVICE__ double acosh(double __x)
Definition: __clang_hip_math.h:724
round
__DEVICE__ double round(double __x)
Definition: __clang_hip_math.h:1037
j0
__DEVICE__ double j0(double __x)
Definition: __clang_hip_math.h:845
__dmul_rz
__DEVICE__ double __dmul_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:94
__drcp_rd
__DEVICE__ double __drcp_rd(double __a)
Definition: __clang_cuda_device_functions.h:154
__fdividef
__DEVICE__ float __fdividef(float __x, float __y)
Definition: __clang_hip_math.h:605
__finitef
__DEVICE__ __RETURN_TYPE __finitef(float __x)
Definition: __clang_hip_math.h:281
normcdf
__DEVICE__ double normcdf(double __x)
Definition: __clang_hip_math.h:976
rhypotf
__DEVICE__ float rhypotf(float __x, float __y)
Definition: __clang_hip_math.h:446
__finite
__DEVICE__ __RETURN_TYPE __finite(double __x)
Definition: __clang_hip_math.h:836
__dadd_rd
__DEVICE__ double __dadd_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:61
__fadd_rz
__DEVICE__ float __fadd_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:206
tanh
__DEVICE__ double tanh(double __x)
Definition: __clang_hip_math.h:1091
__fsub_rd
__DEVICE__ float __fsub_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:324
__dmul_rn
__DEVICE__ double __dmul_rn(double __x, double __y)
Definition: __clang_hip_math.h:1191
__ocml_y1_f32
__device__ float __ocml_y1_f32(float)
log10f
__DEVICE__ float log10f(float __x)
Definition: __clang_hip_math.h:329
__dsub_rz
__DEVICE__ double __dsub_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:171
ilogb
__DEVICE__ int ilogb(double __x)
Definition: __clang_hip_math.h:833
INT_MAX
#define INT_MAX
Definition: limits.h:46
asin
__DEVICE__ double asin(double __x)
Definition: __clang_hip_math.h:727
__log2f
__DEVICE__ float __log2f(float __x)
Definition: __clang_hip_math.h:694
__fdiv_rd
__DEVICE__ float __fdiv_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:209
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:309
__ocml_i1_f32
__device__ float __ocml_i1_f32(float)
sincospi
__DEVICE__ void sincospi(double __x, double *__sinptr, double *__cosptr)
Definition: __clang_hip_math.h:1068
fmin
__DEVICE__ double fmin(double __x, double __y)
Definition: __clang_hip_math.h:812
__fsqrt_ru
__DEVICE__ float __fsqrt_ru(float __a)
Definition: __clang_cuda_device_functions.h:322
erfinv
__DEVICE__ double erfinv(double __x)
Definition: __clang_hip_math.h:780
int
__device__ int
Definition: __clang_hip_libdevice_declares.h:63
floor
__DEVICE__ double floor(double __x)
Definition: __clang_hip_math.h:801
truncf
__DEVICE__ float truncf(float __x)
Definition: __clang_hip_math.h:535
__fmul_rz
__DEVICE__ float __fmul_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:312
acoshf
__DEVICE__ float acoshf(float __x)
Definition: __clang_hip_math.h:167
nextafter
__DEVICE__ double nextafter(double __x, double __y)
Definition: __clang_hip_math.h:949
copysignf
__DEVICE__ float copysignf(float __x, float __y)
Definition: __clang_hip_math.h:191
normf
__DEVICE__ float normf(int __dim, const float *__a)
Definition: __clang_hip_math.h:407
modf
__DEVICE__ double modf(double __x, double *__iptr)
Definition: __clang_hip_math.h:905
__dsqrt_rz
__DEVICE__ double __dsqrt_rz(double __a)
Definition: __clang_cuda_device_functions.h:161
__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:297
__DEVICE__
#define __DEVICE__
Definition: __clang_hip_math.h:32
__ocml_j0_f64
__device__ double __ocml_j0_f64(double)
tan
__DEVICE__ double tan(double __x)
Definition: __clang_hip_math.h:1088
exp
__DEVICE__ double exp(double __x)
Definition: __clang_hip_math.h:783
__logf
__DEVICE__ float __logf(float __x)
Definition: __clang_hip_math.h:697
sinpi
__DEVICE__ double sinpi(double __x)
Definition: __clang_hip_math.h:1082
__a
static __inline__ void int __a
Definition: emmintrin.h:4189
copysign
__DEVICE__ double copysign(double __x, double __y)
Definition: __clang_hip_math.h:748
__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:745
sincospif
__DEVICE__ void sincospif(float __x, float *__sinptr, float *__cosptr)
Definition: __clang_hip_math.h:503
__dmul_ru
__DEVICE__ double __dmul_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:91
sinf
__DEVICE__ float sinf(float __x)
Definition: __clang_hip_math.h:514
__exp10f
__DEVICE__ float __exp10f(float __x)
Definition: __clang_hip_math.h:571
tgammaf
__DEVICE__ float tgammaf(float __x)
Definition: __clang_hip_math.h:532
__RETURN_TYPE
#define __RETURN_TYPE
Definition: __clang_hip_math.h:43
remquo
__DEVICE__ double remquo(double __x, double __y, int *__quo)
Definition: __clang_hip_math.h:996
norm3d
__DEVICE__ double norm3d(double __x, double __y, double __z)
Definition: __clang_hip_math.h:966
floorf
__DEVICE__ float floorf(float __x)
Definition: __clang_hip_math.h:245
coshf
__DEVICE__ float coshf(float __x)
Definition: __clang_hip_math.h:197
__fma_rd
__DEVICE__ double __fma_rd(double __a, double __b, double __c)
Definition: __clang_cuda_device_functions.h:267
__dsub_rn
__DEVICE__ double __dsub_rn(double __x, double __y)
Definition: __clang_hip_math.h:1241
log10
__DEVICE__ double log10(double __x)
Definition: __clang_hip_math.h:887
b
__device__ __2f16 b
Definition: __clang_hip_libdevice_declares.h:314
rnorm4df
__DEVICE__ float rnorm4df(float __x, float __y, float __z, float __w)
Definition: __clang_hip_math.h:457
llrint
__DEVICE__ long long int llrint(double __x)
Definition: __clang_hip_math.h:878
fmod
__DEVICE__ double fmod(double __x, double __y)
Definition: __clang_hip_math.h:815
norm
__DEVICE__ double norm(int __dim, const double *__a)
Definition: __clang_hip_math.h:954
erfcxf
__DEVICE__ float erfcxf(float __x)
Definition: __clang_hip_math.h:215
__fma_rz
__DEVICE__ double __fma_rz(double __a, double __b, double __c)
Definition: __clang_cuda_device_functions.h:276
asinf
__DEVICE__ float asinf(float __x)
Definition: __clang_hip_math.h:170
sinh
__DEVICE__ double sinh(double __x)
Definition: __clang_hip_math.h:1079
hypotf
__DEVICE__ float hypotf(float __x, float __y)
Definition: __clang_hip_math.h:275
rcbrtf
__DEVICE__ float rcbrtf(float __x)
Definition: __clang_hip_math.h:425
yn
__DEVICE__ double yn(int __n, double __x)
Definition: __clang_hip_math.h:1106
ldexpf
__DEVICE__ float ldexpf(float __x, int __e)
Definition: __clang_hip_math.h:317
frexpf
__DEVICE__ float frexpf(float __x, int *__nptr)
Definition: __clang_hip_math.h:262
erfcinvf
__DEVICE__ float erfcinvf(float __x)
Definition: __clang_hip_math.h:212
cbrtf
__DEVICE__ float cbrtf(float __x)
Definition: __clang_hip_math.h:185
log2
__DEVICE__ double log2(double __x)
Definition: __clang_hip_math.h:893
min
__DEVICE__ int min(int __a, int __b)
Definition: __clang_cuda_math.h:197
__drcp_rn
__DEVICE__ double __drcp_rn(double __x)
Definition: __clang_hip_math.h:1205
cospi
__DEVICE__ double cospi(double __x)
Definition: __clang_hip_math.h:759
__fsqrt_rn
__DEVICE__ float __fsqrt_rn(float __x)
Definition: __clang_hip_math.h:673
norm4df
__DEVICE__ float norm4df(float __x, float __y, float __z, float __w)
Definition: __clang_hip_math.h:396
__isnanf
__DEVICE__ __RETURN_TYPE __isnanf(float __x)
Definition: __clang_hip_math.h:287
cosf
__DEVICE__ float cosf(float __x)
Definition: __clang_hip_math.h:194
asinh
__DEVICE__ double asinh(double __x)
Definition: __clang_hip_math.h:730
__ocml_lgamma_f64
__device__ double __ocml_lgamma_f64(double)
ceilf
__DEVICE__ float ceilf(float __x)
Definition: __clang_hip_math.h:188
lroundf
__DEVICE__ long int lroundf(float __x)
Definition: __clang_hip_math.h:347
tanf
__DEVICE__ float tanf(float __x)
Definition: __clang_hip_math.h:526
__ocml_sin_f32
__device__ float __ocml_sin_f32(float)
logbf
__DEVICE__ float logbf(float __x)
Definition: __clang_hip_math.h:338
llrintf
__DEVICE__ long long int llrintf(float __x)
Definition: __clang_hip_math.h:323
exp2
__DEVICE__ double exp2(double __x)
Definition: __clang_hip_math.h:789
remquof
__DEVICE__ float remquof(float __x, float __y, int *__quo)
Definition: __clang_hip_math.h:433
nearbyint
__DEVICE__ double nearbyint(double __x)
Definition: __clang_hip_math.h:946
nan
__DEVICE__ double nan(const char *__tagp)
Definition: __clang_hip_math.h:918
frexp
__DEVICE__ double frexp(double __x, int *__nptr)
Definition: __clang_hip_math.h:818
rhypot
__DEVICE__ double rhypot(double __x, double __y)
Definition: __clang_hip_math.h:1009
remainder
__DEVICE__ double remainder(double __x, double __y)
Definition: __clang_hip_math.h:991
__ocml_cospi_f32
__device__ float __ocml_cospi_f32(float)
nearbyintf
__DEVICE__ float nearbyintf(float __x)
Definition: __clang_hip_math.h:383
limits.h
j1
__DEVICE__ double j1(double __x)
Definition: __clang_hip_math.h:848
__ocml_j0_f32
__device__ float __ocml_j0_f32(float)
fdividef
__DEVICE__ float fdividef(float __x, float __y)
Definition: __clang_hip_math.h:242
__signbit
__DEVICE__ __RETURN_TYPE __signbit(double __x)
Definition: __clang_hip_math.h:1051
erfcx
__DEVICE__ double erfcx(double __x)
Definition: __clang_hip_math.h:777
nextafterf
__DEVICE__ float nextafterf(float __x, float __y)
Definition: __clang_hip_math.h:386
__dadd_ru
__DEVICE__ double __dadd_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:67
erfinvf
__DEVICE__ float erfinvf(float __x)
Definition: __clang_hip_math.h:221
__ocml_sinpi_f32
__device__ float __ocml_sinpi_f32(float)
__expf
__DEVICE__ float __expf(float __x)
Definition: __clang_hip_math.h:574
__fma_rn
__DEVICE__ double __fma_rn(double __x, double __y, double __z)
Definition: __clang_hip_math.h:1263
__dadd_rz
__DEVICE__ double __dadd_rz(double __a, double __b)
Definition: __clang_cuda_device_functions.h:70
rcbrt
__DEVICE__ double rcbrt(double __x)
Definition: __clang_hip_math.h:988
sincosf
__DEVICE__ void sincosf(float __x, float *__sinptr, float *__cosptr)
Definition: __clang_hip_math.h:492
log
__DEVICE__ double log(double __x)
Definition: __clang_hip_math.h:884
erfc
__DEVICE__ double erfc(double __x)
Definition: __clang_hip_math.h:771
__ocml_native_sin_f32
__device__ float __ocml_native_sin_f32(float)
normcdff
__DEVICE__ float normcdff(float __x)
Definition: __clang_hip_math.h:401
labs
__DEVICE__ long labs(long __a)
Definition: __clang_cuda_math.h:152
lrintf
__DEVICE__ long int lrintf(float __x)
Definition: __clang_hip_math.h:344
__frsqrt_rn
__DEVICE__ float __frsqrt_rn(float __x)
Definition: __clang_hip_math.h:660
__ocml_i0_f32
__device__ float __ocml_i0_f32(float)
rnormf
__DEVICE__ float rnormf(int __dim, const float *__a)
Definition: __clang_hip_math.h:462
__drcp_rz
__DEVICE__ double __drcp_rz(double __a)
Definition: __clang_cuda_device_functions.h:157
__frcp_rz
__DEVICE__ float __frcp_rz(float __a)
Definition: __clang_cuda_device_functions.h:318
j1f
__DEVICE__ float j1f(float __x)
Definition: __clang_hip_math.h:293
jnf
__DEVICE__ float jnf(int __n, float __x)
Definition: __clang_hip_math.h:296
log1p
__DEVICE__ double log1p(double __x)
Definition: __clang_hip_math.h:890
__cosf
__DEVICE__ float __cosf(float __x)
Definition: __clang_hip_math.h:568
__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:721
fmaf
__DEVICE__ float fmaf(float __x, float __y, float __z)
Definition: __clang_hip_math.h:248
rintf
__DEVICE__ float rintf(float __x)
Definition: __clang_hip_math.h:449
__tanf
__DEVICE__ float __tanf(float __x)
Definition: __clang_hip_math.h:715
roundf
__DEVICE__ float roundf(float __x)
Definition: __clang_hip_math.h:474
rint
__DEVICE__ double rint(double __x)
Definition: __clang_hip_math.h:1012
atanhf
__DEVICE__ float atanhf(float __x)
Definition: __clang_hip_math.h:182
lrint
__DEVICE__ long int lrint(double __x)
Definition: __clang_hip_math.h:899
__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:736
__sincosf
__DEVICE__ void __sincosf(float __x, float *__sinptr, float *__cosptr)
Definition: __clang_hip_math.h:706
sqrt
__DEVICE__ double sqrt(double __x)
Definition: __clang_hip_math.h:1085
sin
__DEVICE__ double sin(double __x)
Definition: __clang_hip_math.h:1054
__ocml_i0_f64
__device__ double __ocml_i0_f64(double)
ilogbf
__DEVICE__ int ilogbf(float __x)
Definition: __clang_hip_math.h:278
fabs
__DEVICE__ double fabs(double __x)
Definition: __clang_hip_math.h:795
__ddiv_ru
__DEVICE__ double __ddiv_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:79
cosh
__DEVICE__ double cosh(double __x)
Definition: __clang_hip_math.h:756
fdimf
__DEVICE__ float fdimf(float __x, float __y)
Definition: __clang_hip_math.h:239
__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:656
__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:1015
exp10f
__DEVICE__ float exp10f(float __x)
Definition: __clang_hip_math.h:224
erfcinv
__DEVICE__ double erfcinv(double __x)
Definition: __clang_hip_math.h:774
jn
__DEVICE__ double jn(int __n, double __x)
Definition: __clang_hip_math.h:851
cos
__DEVICE__ double cos(double __x)
Definition: __clang_hip_math.h:753
__dsub_ru
__DEVICE__ double __dsub_ru(double __a, double __b)
Definition: __clang_cuda_device_functions.h:168
__make_mantissa_base16
__DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp)
Definition: __clang_hip_math.h:107
fma
__DEVICE__ double fma(double __x, double __y, double __z)
Definition: __clang_hip_math.h:804
__fsqrt_rz
__DEVICE__ float __fsqrt_rz(float __a)
Definition: __clang_cuda_device_functions.h:323
tgamma
__DEVICE__ double tgamma(double __x)
Definition: __clang_hip_math.h:1094
ldexp
__DEVICE__ double ldexp(double __x, int __e)
Definition: __clang_hip_math.h:872
lround
__DEVICE__ long int lround(double __x)
Definition: __clang_hip_math.h:902
stdint.h
__powf
__DEVICE__ float __powf(float __x, float __y)
Definition: __clang_hip_math.h:700
__isnan
__DEVICE__ __RETURN_TYPE __isnan(double __x)
Definition: __clang_hip_math.h:842
sincos
__DEVICE__ void sincos(double __x, double *__sinptr, double *__cosptr)
Definition: __clang_hip_math.h:1057
cyl_bessel_i0f
__DEVICE__ float cyl_bessel_i0f(float __x)
Definition: __clang_hip_math.h:203
sqrtf
__DEVICE__ float sqrtf(float __x)
Definition: __clang_hip_math.h:523
__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:156
__ocml_j1_f32
__device__ float __ocml_j1_f32(float)
__dsqrt_ru
__DEVICE__ double __dsqrt_ru(double __a)
Definition: __clang_cuda_device_functions.h:160
scalblnf
__DEVICE__ float scalblnf(float __x, long int __n)
Definition: __clang_hip_math.h:480
nanf
__DEVICE__ float nanf(const char *__tagp)
Definition: __clang_hip_math.h:362
__make_mantissa_base8
__DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp)
Definition: __clang_hip_math.h:73
erff
__DEVICE__ float erff(float __x)
Definition: __clang_hip_math.h:218
rsqrt
__DEVICE__ double rsqrt(double __x)
Definition: __clang_hip_math.h:1040
cospif
__DEVICE__ float cospif(float __x)
Definition: __clang_hip_math.h:200
cyl_bessel_i1
__DEVICE__ double cyl_bessel_i1(double __x)
Definition: __clang_hip_math.h:765
scalbnf
__DEVICE__ float scalbnf(float __x, int __n)
Definition: __clang_hip_math.h:486
__make_mantissa
__DEVICE__ uint64_t __make_mantissa(const char *__tagp)
Definition: __clang_hip_math.h:128
__fsub_rn
__DEVICE__ float __fsub_rn(float __x, float __y)
Definition: __clang_hip_math.h:687
__ocml_y1_f64
__device__ double __ocml_y1_f64(double)
log2f
__DEVICE__ float log2f(float __x)
Definition: __clang_hip_math.h:335
__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:601
rnorm3df
__DEVICE__ float rnorm3df(float __x, float __y, float __z)
Definition: __clang_hip_math.h:452
pow
__DEVICE__ double pow(double __x, double __y)
Definition: __clang_hip_math.h:982
cyl_bessel_i0
__DEVICE__ double cyl_bessel_i0(double __x)
Definition: __clang_hip_math.h:762
__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:315
erf
__DEVICE__ double erf(double __x)
Definition: __clang_hip_math.h:768
__ocml_tan_f32
__device__ float __ocml_tan_f32(float)
logf
__DEVICE__ float logf(float __x)
Definition: __clang_hip_math.h:341
atanf
__DEVICE__ float atanf(float __x)
Definition: __clang_hip_math.h:179
y0
__DEVICE__ double y0(double __x)
Definition: __clang_hip_math.h:1100
y0f
__DEVICE__ float y0f(float __x)
Definition: __clang_hip_math.h:538
y1
__DEVICE__ double y1(double __x)
Definition: __clang_hip_math.h:1103
__fadd_ru
__DEVICE__ float __fadd_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:203
__ocml_j1_f64
__device__ double __ocml_j1_f64(double)
__isinf
__DEVICE__ __RETURN_TYPE __isinf(double __x)
Definition: __clang_hip_math.h:839
__dsqrt_rn
__DEVICE__ double __dsqrt_rn(double __x)
Definition: __clang_hip_math.h:1219
atanh
__DEVICE__ double atanh(double __x)
Definition: __clang_hip_math.h:739
trunc
__DEVICE__ double trunc(double __x)
Definition: __clang_hip_math.h:1097
__fmaf_rz
__DEVICE__ float __fmaf_rz(float __a, float __b, float __c)
Definition: __clang_cuda_device_functions.h:300
__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:985
hypot
__DEVICE__ double hypot(double __x, double __y)
Definition: __clang_hip_math.h:830
norm3df
__DEVICE__ float norm3df(float __x, float __y, float __z)
Definition: __clang_hip_math.h:391
__signbitf
__DEVICE__ __RETURN_TYPE __signbitf(float __x)
Definition: __clang_hip_math.h:489
fminf
__DEVICE__ float fminf(float __x, float __y)
Definition: __clang_hip_math.h:256
__ocml_cos_f64
__device__ double __ocml_cos_f64(double)
atan
__DEVICE__ double atan(double __x)
Definition: __clang_hip_math.h:733
normcdfinvf
__DEVICE__ float normcdfinvf(float __x)
Definition: __clang_hip_math.h:404
remainderf
__DEVICE__ float remainderf(float __x, float __y)
Definition: __clang_hip_math.h:428
__ddiv_rn
__DEVICE__ double __ddiv_rn(double __x, double __y)
Definition: __clang_hip_math.h:1169
__make_mantissa_base10
__DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp)
Definition: __clang_hip_math.h:90
scalbln
__DEVICE__ double scalbln(double __x, long int __n)
Definition: __clang_hip_math.h:1043
norm4d
__DEVICE__ double norm4d(double __x, double __y, double __z, double __w)
Definition: __clang_hip_math.h:971
__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:311
cyl_bessel_i1f
__DEVICE__ float cyl_bessel_i1f(float __x)
Definition: __clang_hip_math.h:206
powif
__DEVICE__ float powif(float __x, int __y)
Definition: __clang_hip_math.h:422
lgamma
__DEVICE__ double lgamma(double __x)
Definition: __clang_hip_math.h:875
logb
__DEVICE__ double logb(double __x)
Definition: __clang_hip_math.h:896
__fsqrt_rd
__DEVICE__ float __fsqrt_rd(float __a)
Definition: __clang_cuda_device_functions.h:320
__fadd_rd
__DEVICE__ float __fadd_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:197
__fdiv_ru
__DEVICE__ float __fdiv_ru(float __a, float __b)
Definition: __clang_cuda_device_functions.h:215
sinpif
__DEVICE__ float sinpif(float __x)
Definition: __clang_hip_math.h:520
expm1
__DEVICE__ double expm1(double __x)
Definition: __clang_hip_math.h:792
__dadd_rn
__DEVICE__ double __dadd_rn(double __x, double __y)
Definition: __clang_hip_math.h:1147
__fadd_rn
__DEVICE__ float __fadd_rn(float __x, float __y)
Definition: __clang_hip_math.h:587
__ocml_sin_f64
__device__ double __ocml_sin_f64(double)
llround
__DEVICE__ long long int llround(double __x)
Definition: __clang_hip_math.h:881
__dsub_rd
__DEVICE__ double __dsub_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:162
__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:529
__fmul_rd
__DEVICE__ float __fmul_rd(float __a, float __b)
Definition: __clang_cuda_device_functions.h:303
scalbn
__DEVICE__ double scalbn(double __x, int __n)
Definition: __clang_hip_math.h:1048
__log10f
__DEVICE__ float __log10f(float __x)
Definition: __clang_hip_math.h:691
j0f
__DEVICE__ float j0f(float __x)
Definition: __clang_hip_math.h:290
__fsub_rz
__DEVICE__ float __fsub_rz(float __a, float __b)
Definition: __clang_cuda_device_functions.h:333
rnorm3d
__DEVICE__ double rnorm3d(double __x, double __y, double __z)
Definition: __clang_hip_math.h:1027
__fma_ru
__DEVICE__ double __fma_ru(double __a, double __b, double __c)
Definition: __clang_cuda_device_functions.h:273
__fmaf_rd
__DEVICE__ float __fmaf_rd(float __a, float __b, float __c)
Definition: __clang_cuda_device_functions.h:291
fmax
__DEVICE__ double fmax(double __x, double __y)
Definition: __clang_hip_math.h:809
__ocml_tgamma_f64
__device__ double __ocml_tgamma_f64(double)
__saturatef
__DEVICE__ float __saturatef(float __x)
Definition: __clang_hip_math.h:703
llroundf
__DEVICE__ long long int llroundf(float __x)
Definition: __clang_hip_math.h:326
expf
__DEVICE__ float expf(float __x)
Definition: __clang_hip_math.h:230
acosf
__DEVICE__ float acosf(float __x)
Definition: __clang_hip_math.h:164
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:67
__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:350
normcdfinv
__DEVICE__ double normcdfinv(double __x)
Definition: __clang_hip_math.h:979
rnorm4d
__DEVICE__ double rnorm4d(double __x, double __y, double __z, double __w)
Definition: __clang_hip_math.h:1032
__ocml_tgamma_f32
__device__ float __ocml_tgamma_f32(float)
fmodf
__DEVICE__ float fmodf(float __x, float __y)
Definition: __clang_hip_math.h:259
__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:158
fdim
__DEVICE__ double fdim(double __x, double __y)
Definition: __clang_hip_math.h:798
fabsf
__DEVICE__ float fabsf(float __x)
Definition: __clang_hip_math.h:236
__sinf
__DEVICE__ float __sinf(float __x)
Definition: __clang_hip_math.h:712
__fmaf_rn
__DEVICE__ float __fmaf_rn(float __x, float __y, float __z)
Definition: __clang_hip_math.h:626
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:332
__ddiv_rd
__DEVICE__ double __ddiv_rd(double __a, double __b)
Definition: __clang_cuda_device_functions.h:73
__isinff
__DEVICE__ __RETURN_TYPE __isinff(float __x)
Definition: __clang_hip_math.h:284
sinhf
__DEVICE__ float sinhf(float __x)
Definition: __clang_hip_math.h:517