ETISS 0.8.0
Extendable Translating Instruction Set Simulator (version 0.8.0)
Loading...
Searching...
No Matches
__clang_hip_math.h
Go to the documentation of this file.
1/*===---- __clang_hip_math.h - HIP math decls -------------------------------===
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
10#ifndef __CLANG_HIP_MATH_H__
11#define __CLANG_HIP_MATH_H__
12
13#include <algorithm>
14#include <limits.h>
15#include <limits>
16#include <stdint.h>
17
18#pragma push_macro("__DEVICE__")
19#pragma push_macro("__RETURN_TYPE")
20
21// to be consistent with __clang_cuda_math_forward_declares
22#define __DEVICE__ static __device__
23#define __RETURN_TYPE bool
24
26inline uint64_t __make_mantissa_base8(const char *__tagp) {
27 uint64_t __r = 0;
28 while (__tagp) {
29 char __tmp = *__tagp;
30
31 if (__tmp >= '0' && __tmp <= '7')
32 __r = (__r * 8u) + __tmp - '0';
33 else
34 return 0;
35
36 ++__tagp;
37 }
38
39 return __r;
40}
41
43inline uint64_t __make_mantissa_base10(const char *__tagp) {
44 uint64_t __r = 0;
45 while (__tagp) {
46 char __tmp = *__tagp;
47
48 if (__tmp >= '0' && __tmp <= '9')
49 __r = (__r * 10u) + __tmp - '0';
50 else
51 return 0;
52
53 ++__tagp;
54 }
55
56 return __r;
57}
58
60inline uint64_t __make_mantissa_base16(const char *__tagp) {
61 uint64_t __r = 0;
62 while (__tagp) {
63 char __tmp = *__tagp;
64
65 if (__tmp >= '0' && __tmp <= '9')
66 __r = (__r * 16u) + __tmp - '0';
67 else if (__tmp >= 'a' && __tmp <= 'f')
68 __r = (__r * 16u) + __tmp - 'a' + 10;
69 else if (__tmp >= 'A' && __tmp <= 'F')
70 __r = (__r * 16u) + __tmp - 'A' + 10;
71 else
72 return 0;
73
74 ++__tagp;
75 }
76
77 return __r;
78}
79
81inline uint64_t __make_mantissa(const char *__tagp) {
82 if (!__tagp)
83 return 0u;
84
85 if (*__tagp == '0') {
86 ++__tagp;
87
88 if (*__tagp == 'x' || *__tagp == 'X')
89 return __make_mantissa_base16(__tagp);
90 else
91 return __make_mantissa_base8(__tagp);
92 }
93
94 return __make_mantissa_base10(__tagp);
95}
96
97// BEGIN FLOAT
99inline float abs(float __x) { return __ocml_fabs_f32(__x); }
101inline float acosf(float __x) { return __ocml_acos_f32(__x); }
103inline float acoshf(float __x) { return __ocml_acosh_f32(__x); }
105inline float asinf(float __x) { return __ocml_asin_f32(__x); }
107inline float asinhf(float __x) { return __ocml_asinh_f32(__x); }
109inline float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); }
111inline float atanf(float __x) { return __ocml_atan_f32(__x); }
113inline float atanhf(float __x) { return __ocml_atanh_f32(__x); }
115inline float cbrtf(float __x) { return __ocml_cbrt_f32(__x); }
117inline float ceilf(float __x) { return __ocml_ceil_f32(__x); }
119inline float copysignf(float __x, float __y) {
120 return __ocml_copysign_f32(__x, __y);
121}
123inline float cosf(float __x) { return __ocml_cos_f32(__x); }
125inline float coshf(float __x) { return __ocml_cosh_f32(__x); }
127inline float cospif(float __x) { return __ocml_cospi_f32(__x); }
129inline float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); }
131inline float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); }
133inline float erfcf(float __x) { return __ocml_erfc_f32(__x); }
135inline float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); }
137inline float erfcxf(float __x) { return __ocml_erfcx_f32(__x); }
139inline float erff(float __x) { return __ocml_erf_f32(__x); }
141inline float erfinvf(float __x) { return __ocml_erfinv_f32(__x); }
143inline float exp10f(float __x) { return __ocml_exp10_f32(__x); }
145inline float exp2f(float __x) { return __ocml_exp2_f32(__x); }
147inline float expf(float __x) { return __ocml_exp_f32(__x); }
149inline float expm1f(float __x) { return __ocml_expm1_f32(__x); }
151inline float fabsf(float __x) { return __ocml_fabs_f32(__x); }
153inline float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); }
155inline float fdividef(float __x, float __y) { return __x / __y; }
157inline float floorf(float __x) { return __ocml_floor_f32(__x); }
159inline float fmaf(float __x, float __y, float __z) {
160 return __ocml_fma_f32(__x, __y, __z);
161}
163inline float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); }
165inline float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); }
167inline float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
169inline float frexpf(float __x, int *__nptr) {
170 int __tmp;
171 float __r =
172 __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp);
173 *__nptr = __tmp;
174
175 return __r;
176}
178inline float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); }
180inline int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }
182inline __RETURN_TYPE isfinite(float __x) { return __ocml_isfinite_f32(__x); }
184inline __RETURN_TYPE isinf(float __x) { return __ocml_isinf_f32(__x); }
186inline __RETURN_TYPE isnan(float __x) { return __ocml_isnan_f32(__x); }
188inline float j0f(float __x) { return __ocml_j0_f32(__x); }
190inline float j1f(float __x) { return __ocml_j1_f32(__x); }
192inline float jnf(int __n,
193 float __x) { // TODO: we could use Ahmes multiplication
194 // and the Miller & Brown algorithm
195 // for linear recurrences to get O(log n) steps, but it's unclear if
196 // it'd be beneficial in this case.
197 if (__n == 0)
198 return j0f(__x);
199 if (__n == 1)
200 return j1f(__x);
201
202 float __x0 = j0f(__x);
203 float __x1 = j1f(__x);
204 for (int __i = 1; __i < __n; ++__i) {
205 float __x2 = (2 * __i) / __x * __x1 - __x0;
206 __x0 = __x1;
207 __x1 = __x2;
208 }
209
210 return __x1;
211}
213inline float ldexpf(float __x, int __e) { return __ocml_ldexp_f32(__x, __e); }
215inline float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
217inline long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
219inline long long int llroundf(float __x) { return __ocml_round_f32(__x); }
221inline float log10f(float __x) { return __ocml_log10_f32(__x); }
223inline float log1pf(float __x) { return __ocml_log1p_f32(__x); }
225inline float log2f(float __x) { return __ocml_log2_f32(__x); }
227inline float logbf(float __x) { return __ocml_logb_f32(__x); }
229inline float logf(float __x) { return __ocml_log_f32(__x); }
231inline long int lrintf(float __x) { return __ocml_rint_f32(__x); }
233inline long int lroundf(float __x) { return __ocml_round_f32(__x); }
235inline float modff(float __x, float *__iptr) {
236 float __tmp;
237 float __r =
238 __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
239 *__iptr = __tmp;
240
241 return __r;
242}
244inline float nanf(const char *__tagp) {
245 union {
246 float val;
247 struct ieee_float {
248 uint32_t mantissa : 22;
249 uint32_t quiet : 1;
250 uint32_t exponent : 8;
251 uint32_t sign : 1;
252 } bits;
253
254 static_assert(sizeof(float) == sizeof(ieee_float), "");
255 } __tmp;
256
257 __tmp.bits.sign = 0u;
258 __tmp.bits.exponent = ~0u;
259 __tmp.bits.quiet = 1u;
260 __tmp.bits.mantissa = __make_mantissa(__tagp);
261
262 return __tmp.val;
263}
265inline float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); }
267inline float nextafterf(float __x, float __y) {
268 return __ocml_nextafter_f32(__x, __y);
269}
271inline float norm3df(float __x, float __y, float __z) {
272 return __ocml_len3_f32(__x, __y, __z);
273}
275inline float norm4df(float __x, float __y, float __z, float __w) {
276 return __ocml_len4_f32(__x, __y, __z, __w);
277}
279inline float normcdff(float __x) { return __ocml_ncdf_f32(__x); }
281inline float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); }
283inline float
284normf(int __dim,
285 const float *__a) { // TODO: placeholder until OCML adds support.
286 float __r = 0;
287 while (__dim--) {
288 __r += __a[0] * __a[0];
289 ++__a;
290 }
291
292 return __ocml_sqrt_f32(__r);
293}
295inline float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
297inline float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }
299inline float remainderf(float __x, float __y) {
300 return __ocml_remainder_f32(__x, __y);
301}
303inline float remquof(float __x, float __y, int *__quo) {
304 int __tmp;
305 float __r = __ocml_remquo_f32(
306 __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
307 *__quo = __tmp;
308
309 return __r;
310}
312inline float rhypotf(float __x, float __y) {
313 return __ocml_rhypot_f32(__x, __y);
314}
316inline float rintf(float __x) { return __ocml_rint_f32(__x); }
318inline float rnorm3df(float __x, float __y, float __z) {
319 return __ocml_rlen3_f32(__x, __y, __z);
320}
321
323inline float rnorm4df(float __x, float __y, float __z, float __w) {
324 return __ocml_rlen4_f32(__x, __y, __z, __w);
325}
327inline float
328rnormf(int __dim,
329 const float *__a) { // TODO: placeholder until OCML adds support.
330 float __r = 0;
331 while (__dim--) {
332 __r += __a[0] * __a[0];
333 ++__a;
334 }
335
336 return __ocml_rsqrt_f32(__r);
337}
339inline float roundf(float __x) { return __ocml_round_f32(__x); }
341inline float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
343inline float scalblnf(float __x, long int __n) {
344 return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n)
345 : __ocml_scalb_f32(__x, __n);
346}
348inline float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
350inline __RETURN_TYPE signbit(float __x) { return __ocml_signbit_f32(__x); }
352inline void sincosf(float __x, float *__sinptr, float *__cosptr) {
353 float __tmp;
354
355 *__sinptr =
356 __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
357 *__cosptr = __tmp;
358}
360inline void sincospif(float __x, float *__sinptr, float *__cosptr) {
361 float __tmp;
362
363 *__sinptr = __ocml_sincospi_f32(
364 __x, (__attribute__((address_space(5))) float *)&__tmp);
365 *__cosptr = __tmp;
366}
368inline float sinf(float __x) { return __ocml_sin_f32(__x); }
370inline float sinhf(float __x) { return __ocml_sinh_f32(__x); }
372inline float sinpif(float __x) { return __ocml_sinpi_f32(__x); }
374inline float sqrtf(float __x) { return __ocml_sqrt_f32(__x); }
376inline float tanf(float __x) { return __ocml_tan_f32(__x); }
378inline float tanhf(float __x) { return __ocml_tanh_f32(__x); }
380inline float tgammaf(float __x) { return __ocml_tgamma_f32(__x); }
382inline float truncf(float __x) { return __ocml_trunc_f32(__x); }
384inline float y0f(float __x) { return __ocml_y0_f32(__x); }
386inline float y1f(float __x) { return __ocml_y1_f32(__x); }
388inline float ynf(int __n,
389 float __x) { // TODO: we could use Ahmes multiplication
390 // and the Miller & Brown algorithm
391 // for linear recurrences to get O(log n) steps, but it's unclear if
392 // it'd be beneficial in this case. Placeholder until OCML adds
393 // support.
394 if (__n == 0)
395 return y0f(__x);
396 if (__n == 1)
397 return y1f(__x);
398
399 float __x0 = y0f(__x);
400 float __x1 = y1f(__x);
401 for (int __i = 1; __i < __n; ++__i) {
402 float __x2 = (2 * __i) / __x * __x1 - __x0;
403 __x0 = __x1;
404 __x1 = __x2;
405 }
406
407 return __x1;
408}
409
410// BEGIN INTRINSICS
412inline float __cosf(float __x) { return __ocml_native_cos_f32(__x); }
414inline float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }
416inline float __expf(float __x) { return __ocml_native_exp_f32(__x); }
417#if defined OCML_BASIC_ROUNDED_OPERATIONS
419inline float __fadd_rd(float __x, float __y) {
420 return __ocml_add_rtn_f32(__x, __y);
421}
422#endif
424inline float __fadd_rn(float __x, float __y) { return __x + __y; }
425#if defined OCML_BASIC_ROUNDED_OPERATIONS
427inline float __fadd_ru(float __x, float __y) {
428 return __ocml_add_rtp_f32(__x, __y);
429}
431inline float __fadd_rz(float __x, float __y) {
432 return __ocml_add_rtz_f32(__x, __y);
433}
435inline float __fdiv_rd(float __x, float __y) {
436 return __ocml_div_rtn_f32(__x, __y);
437}
438#endif
440inline float __fdiv_rn(float __x, float __y) { return __x / __y; }
441#if defined OCML_BASIC_ROUNDED_OPERATIONS
443inline float __fdiv_ru(float __x, float __y) {
444 return __ocml_div_rtp_f32(__x, __y);
445}
447inline float __fdiv_rz(float __x, float __y) {
448 return __ocml_div_rtz_f32(__x, __y);
449}
450#endif
452inline float __fdividef(float __x, float __y) { return __x / __y; }
453#if defined OCML_BASIC_ROUNDED_OPERATIONS
455inline float __fmaf_rd(float __x, float __y, float __z) {
456 return __ocml_fma_rtn_f32(__x, __y, __z);
457}
458#endif
460inline float __fmaf_rn(float __x, float __y, float __z) {
461 return __ocml_fma_f32(__x, __y, __z);
462}
463#if defined OCML_BASIC_ROUNDED_OPERATIONS
465inline float __fmaf_ru(float __x, float __y, float __z) {
466 return __ocml_fma_rtp_f32(__x, __y, __z);
467}
469inline float __fmaf_rz(float __x, float __y, float __z) {
470 return __ocml_fma_rtz_f32(__x, __y, __z);
471}
473inline float __fmul_rd(float __x, float __y) {
474 return __ocml_mul_rtn_f32(__x, __y);
475}
476#endif
478inline float __fmul_rn(float __x, float __y) { return __x * __y; }
479#if defined OCML_BASIC_ROUNDED_OPERATIONS
481inline float __fmul_ru(float __x, float __y) {
482 return __ocml_mul_rtp_f32(__x, __y);
483}
485inline float __fmul_rz(float __x, float __y) {
486 return __ocml_mul_rtz_f32(__x, __y);
487}
489inline float __frcp_rd(float __x) { return __llvm_amdgcn_rcp_f32(__x); }
490#endif
492inline float __frcp_rn(float __x) { return __llvm_amdgcn_rcp_f32(__x); }
493#if defined OCML_BASIC_ROUNDED_OPERATIONS
495inline float __frcp_ru(float __x) { return __llvm_amdgcn_rcp_f32(__x); }
497inline float __frcp_rz(float __x) { return __llvm_amdgcn_rcp_f32(__x); }
498#endif
500inline float __frsqrt_rn(float __x) { return __llvm_amdgcn_rsq_f32(__x); }
501#if defined OCML_BASIC_ROUNDED_OPERATIONS
503inline float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
504#endif
506inline float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }
507#if defined OCML_BASIC_ROUNDED_OPERATIONS
509inline float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }
511inline float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }
513inline float __fsub_rd(float __x, float __y) {
514 return __ocml_sub_rtn_f32(__x, __y);
515}
516#endif
518inline float __fsub_rn(float __x, float __y) { return __x - __y; }
519#if defined OCML_BASIC_ROUNDED_OPERATIONS
521inline float __fsub_ru(float __x, float __y) {
522 return __ocml_sub_rtp_f32(__x, __y);
523}
525inline float __fsub_rz(float __x, float __y) {
526 return __ocml_sub_rtz_f32(__x, __y);
527}
528#endif
530inline float __log10f(float __x) { return __ocml_native_log10_f32(__x); }
532inline float __log2f(float __x) { return __ocml_native_log2_f32(__x); }
534inline float __logf(float __x) { return __ocml_native_log_f32(__x); }
536inline float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
538inline float __saturatef(float __x) {
539 return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x);
540}
542inline void __sincosf(float __x, float *__sinptr, float *__cosptr) {
543 *__sinptr = __ocml_native_sin_f32(__x);
544 *__cosptr = __ocml_native_cos_f32(__x);
545}
547inline float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
549inline float __tanf(float __x) { return __ocml_tan_f32(__x); }
550// END INTRINSICS
551// END FLOAT
552
553// BEGIN DOUBLE
555inline double abs(double __x) { return __ocml_fabs_f64(__x); }
557inline double acos(double __x) { return __ocml_acos_f64(__x); }
559inline double acosh(double __x) { return __ocml_acosh_f64(__x); }
561inline double asin(double __x) { return __ocml_asin_f64(__x); }
563inline double asinh(double __x) { return __ocml_asinh_f64(__x); }
565inline double atan(double __x) { return __ocml_atan_f64(__x); }
567inline double atan2(double __x, double __y) {
568 return __ocml_atan2_f64(__x, __y);
569}
571inline double atanh(double __x) { return __ocml_atanh_f64(__x); }
573inline double cbrt(double __x) { return __ocml_cbrt_f64(__x); }
575inline double ceil(double __x) { return __ocml_ceil_f64(__x); }
577inline double copysign(double __x, double __y) {
578 return __ocml_copysign_f64(__x, __y);
579}
581inline double cos(double __x) { return __ocml_cos_f64(__x); }
583inline double cosh(double __x) { return __ocml_cosh_f64(__x); }
585inline double cospi(double __x) { return __ocml_cospi_f64(__x); }
587inline double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); }
589inline double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); }
591inline double erf(double __x) { return __ocml_erf_f64(__x); }
593inline double erfc(double __x) { return __ocml_erfc_f64(__x); }
595inline double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); }
597inline double erfcx(double __x) { return __ocml_erfcx_f64(__x); }
599inline double erfinv(double __x) { return __ocml_erfinv_f64(__x); }
601inline double exp(double __x) { return __ocml_exp_f64(__x); }
603inline double exp10(double __x) { return __ocml_exp10_f64(__x); }
605inline double exp2(double __x) { return __ocml_exp2_f64(__x); }
607inline double expm1(double __x) { return __ocml_expm1_f64(__x); }
609inline double fabs(double __x) { return __ocml_fabs_f64(__x); }
611inline double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }
613inline double floor(double __x) { return __ocml_floor_f64(__x); }
615inline double fma(double __x, double __y, double __z) {
616 return __ocml_fma_f64(__x, __y, __z);
617}
619inline double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); }
621inline double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); }
623inline double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
625inline double frexp(double __x, int *__nptr) {
626 int __tmp;
627 double __r =
628 __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp);
629 *__nptr = __tmp;
630
631 return __r;
632}
634inline double hypot(double __x, double __y) {
635 return __ocml_hypot_f64(__x, __y);
636}
638inline int ilogb(double __x) { return __ocml_ilogb_f64(__x); }
640inline __RETURN_TYPE isfinite(double __x) { return __ocml_isfinite_f64(__x); }
642inline __RETURN_TYPE isinf(double __x) { return __ocml_isinf_f64(__x); }
644inline __RETURN_TYPE isnan(double __x) { return __ocml_isnan_f64(__x); }
646inline double j0(double __x) { return __ocml_j0_f64(__x); }
648inline double j1(double __x) { return __ocml_j1_f64(__x); }
650inline double jn(int __n,
651 double __x) { // TODO: we could use Ahmes multiplication
652 // and the Miller & Brown algorithm
653 // for linear recurrences to get O(log n) steps, but it's unclear if
654 // it'd be beneficial in this case. Placeholder until OCML adds
655 // support.
656 if (__n == 0)
657 return j0f(__x);
658 if (__n == 1)
659 return j1f(__x);
660
661 double __x0 = j0f(__x);
662 double __x1 = j1f(__x);
663 for (int __i = 1; __i < __n; ++__i) {
664 double __x2 = (2 * __i) / __x * __x1 - __x0;
665 __x0 = __x1;
666 __x1 = __x2;
667 }
668
669 return __x1;
670}
672inline double ldexp(double __x, int __e) { return __ocml_ldexp_f64(__x, __e); }
674inline double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
676inline long long int llrint(double __x) { return __ocml_rint_f64(__x); }
678inline long long int llround(double __x) { return __ocml_round_f64(__x); }
680inline double log(double __x) { return __ocml_log_f64(__x); }
682inline double log10(double __x) { return __ocml_log10_f64(__x); }
684inline double log1p(double __x) { return __ocml_log1p_f64(__x); }
686inline double log2(double __x) { return __ocml_log2_f64(__x); }
688inline double logb(double __x) { return __ocml_logb_f64(__x); }
690inline long int lrint(double __x) { return __ocml_rint_f64(__x); }
692inline long int lround(double __x) { return __ocml_round_f64(__x); }
694inline double modf(double __x, double *__iptr) {
695 double __tmp;
696 double __r =
697 __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
698 *__iptr = __tmp;
699
700 return __r;
701}
703inline double nan(const char *__tagp) {
704#if !_WIN32
705 union {
706 double val;
707 struct ieee_double {
708 uint64_t mantissa : 51;
709 uint32_t quiet : 1;
710 uint32_t exponent : 11;
711 uint32_t sign : 1;
712 } bits;
713 static_assert(sizeof(double) == sizeof(ieee_double), "");
714 } __tmp;
715
716 __tmp.bits.sign = 0u;
717 __tmp.bits.exponent = ~0u;
718 __tmp.bits.quiet = 1u;
719 __tmp.bits.mantissa = __make_mantissa(__tagp);
720
721 return __tmp.val;
722#else
723 static_assert(sizeof(uint64_t) == sizeof(double));
724 uint64_t val = __make_mantissa(__tagp);
725 val |= 0xFFF << 51;
726 return *reinterpret_cast<double *>(&val);
727#endif
728}
730inline double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); }
732inline double nextafter(double __x, double __y) {
733 return __ocml_nextafter_f64(__x, __y);
734}
736inline double
737norm(int __dim,
738 const double *__a) { // TODO: placeholder until OCML adds support.
739 double __r = 0;
740 while (__dim--) {
741 __r += __a[0] * __a[0];
742 ++__a;
743 }
744
745 return __ocml_sqrt_f64(__r);
746}
748inline double norm3d(double __x, double __y, double __z) {
749 return __ocml_len3_f64(__x, __y, __z);
750}
752inline double norm4d(double __x, double __y, double __z, double __w) {
753 return __ocml_len4_f64(__x, __y, __z, __w);
754}
756inline double normcdf(double __x) { return __ocml_ncdf_f64(__x); }
758inline double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); }
760inline double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); }
762inline double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }
764inline double remainder(double __x, double __y) {
765 return __ocml_remainder_f64(__x, __y);
766}
768inline double remquo(double __x, double __y, int *__quo) {
769 int __tmp;
770 double __r = __ocml_remquo_f64(
771 __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
772 *__quo = __tmp;
773
774 return __r;
775}
777inline double rhypot(double __x, double __y) {
778 return __ocml_rhypot_f64(__x, __y);
779}
781inline double rint(double __x) { return __ocml_rint_f64(__x); }
783inline double
784rnorm(int __dim,
785 const double *__a) { // TODO: placeholder until OCML adds support.
786 double __r = 0;
787 while (__dim--) {
788 __r += __a[0] * __a[0];
789 ++__a;
790 }
791
792 return __ocml_rsqrt_f64(__r);
793}
795inline double rnorm3d(double __x, double __y, double __z) {
796 return __ocml_rlen3_f64(__x, __y, __z);
797}
799inline double rnorm4d(double __x, double __y, double __z, double __w) {
800 return __ocml_rlen4_f64(__x, __y, __z, __w);
801}
803inline double round(double __x) { return __ocml_round_f64(__x); }
805inline double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }
807inline double scalbln(double __x, long int __n) {
808 return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n)
809 : __ocml_scalb_f64(__x, __n);
810}
812inline double scalbn(double __x, int __n) {
813 return __ocml_scalbn_f64(__x, __n);
814}
816inline __RETURN_TYPE signbit(double __x) { return __ocml_signbit_f64(__x); }
818inline double sin(double __x) { return __ocml_sin_f64(__x); }
820inline void sincos(double __x, double *__sinptr, double *__cosptr) {
821 double __tmp;
822 *__sinptr = __ocml_sincos_f64(
823 __x, (__attribute__((address_space(5))) double *)&__tmp);
824 *__cosptr = __tmp;
825}
827inline void sincospi(double __x, double *__sinptr, double *__cosptr) {
828 double __tmp;
829 *__sinptr = __ocml_sincospi_f64(
830 __x, (__attribute__((address_space(5))) double *)&__tmp);
831 *__cosptr = __tmp;
832}
834inline double sinh(double __x) { return __ocml_sinh_f64(__x); }
836inline double sinpi(double __x) { return __ocml_sinpi_f64(__x); }
838inline double sqrt(double __x) { return __ocml_sqrt_f64(__x); }
840inline double tan(double __x) { return __ocml_tan_f64(__x); }
842inline double tanh(double __x) { return __ocml_tanh_f64(__x); }
844inline double tgamma(double __x) { return __ocml_tgamma_f64(__x); }
846inline double trunc(double __x) { return __ocml_trunc_f64(__x); }
848inline double y0(double __x) { return __ocml_y0_f64(__x); }
850inline double y1(double __x) { return __ocml_y1_f64(__x); }
852inline double yn(int __n,
853 double __x) { // TODO: we could use Ahmes multiplication
854 // and the Miller & Brown algorithm
855 // for linear recurrences to get O(log n) steps, but it's unclear if
856 // it'd be beneficial in this case. Placeholder until OCML adds
857 // support.
858 if (__n == 0)
859 return j0f(__x);
860 if (__n == 1)
861 return j1f(__x);
862
863 double __x0 = j0f(__x);
864 double __x1 = j1f(__x);
865 for (int __i = 1; __i < __n; ++__i) {
866 double __x2 = (2 * __i) / __x * __x1 - __x0;
867 __x0 = __x1;
868 __x1 = __x2;
869 }
870
871 return __x1;
872}
873
874// BEGIN INTRINSICS
875#if defined OCML_BASIC_ROUNDED_OPERATIONS
877inline double __dadd_rd(double __x, double __y) {
878 return __ocml_add_rtn_f64(__x, __y);
879}
880#endif
882inline double __dadd_rn(double __x, double __y) { return __x + __y; }
883#if defined OCML_BASIC_ROUNDED_OPERATIONS
885inline double __dadd_ru(double __x, double __y) {
886 return __ocml_add_rtp_f64(__x, __y);
887}
889inline double __dadd_rz(double __x, double __y) {
890 return __ocml_add_rtz_f64(__x, __y);
891}
893inline double __ddiv_rd(double __x, double __y) {
894 return __ocml_div_rtn_f64(__x, __y);
895}
896#endif
898inline double __ddiv_rn(double __x, double __y) { return __x / __y; }
899#if defined OCML_BASIC_ROUNDED_OPERATIONS
901inline double __ddiv_ru(double __x, double __y) {
902 return __ocml_div_rtp_f64(__x, __y);
903}
905inline double __ddiv_rz(double __x, double __y) {
906 return __ocml_div_rtz_f64(__x, __y);
907}
909inline double __dmul_rd(double __x, double __y) {
910 return __ocml_mul_rtn_f64(__x, __y);
911}
912#endif
914inline double __dmul_rn(double __x, double __y) { return __x * __y; }
915#if defined OCML_BASIC_ROUNDED_OPERATIONS
917inline double __dmul_ru(double __x, double __y) {
918 return __ocml_mul_rtp_f64(__x, __y);
919}
921inline double __dmul_rz(double __x, double __y) {
922 return __ocml_mul_rtz_f64(__x, __y);
923}
925inline double __drcp_rd(double __x) { return __llvm_amdgcn_rcp_f64(__x); }
926#endif
928inline double __drcp_rn(double __x) { return __llvm_amdgcn_rcp_f64(__x); }
929#if defined OCML_BASIC_ROUNDED_OPERATIONS
931inline double __drcp_ru(double __x) { return __llvm_amdgcn_rcp_f64(__x); }
933inline double __drcp_rz(double __x) { return __llvm_amdgcn_rcp_f64(__x); }
935inline double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); }
936#endif
938inline double __dsqrt_rn(double __x) { return __ocml_sqrt_f64(__x); }
939#if defined OCML_BASIC_ROUNDED_OPERATIONS
941inline double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); }
943inline double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); }
945inline double __dsub_rd(double __x, double __y) {
946 return __ocml_sub_rtn_f64(__x, __y);
947}
948#endif
950inline double __dsub_rn(double __x, double __y) { return __x - __y; }
951#if defined OCML_BASIC_ROUNDED_OPERATIONS
953inline double __dsub_ru(double __x, double __y) {
954 return __ocml_sub_rtp_f64(__x, __y);
955}
957inline double __dsub_rz(double __x, double __y) {
958 return __ocml_sub_rtz_f64(__x, __y);
959}
961inline double __fma_rd(double __x, double __y, double __z) {
962 return __ocml_fma_rtn_f64(__x, __y, __z);
963}
964#endif
966inline double __fma_rn(double __x, double __y, double __z) {
967 return __ocml_fma_f64(__x, __y, __z);
968}
969#if defined OCML_BASIC_ROUNDED_OPERATIONS
971inline double __fma_ru(double __x, double __y, double __z) {
972 return __ocml_fma_rtp_f64(__x, __y, __z);
973}
975inline double __fma_rz(double __x, double __y, double __z) {
976 return __ocml_fma_rtz_f64(__x, __y, __z);
977}
978#endif
979// END INTRINSICS
980// END DOUBLE
981
982// BEGIN INTEGER
984inline int abs(int __x) {
985 int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1);
986 return (__x ^ __sgn) - __sgn;
987}
989inline long labs(long __x) {
990 long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1);
991 return (__x ^ __sgn) - __sgn;
992}
994inline long long llabs(long long __x) {
995 long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
996 return (__x ^ __sgn) - __sgn;
997}
998
999#if defined(__cplusplus)
1001inline long abs(long __x) { return labs(__x); }
1003inline long long abs(long long __x) { return llabs(__x); }
1004#endif
1005// END INTEGER
1006
1009 return __ocml_fma_f16(__x, __y, __z);
1010}
1011
1013inline float fma(float __x, float __y, float __z) {
1014 return fmaf(__x, __y, __z);
1015}
1016
1017#pragma push_macro("__DEF_FUN1")
1018#pragma push_macro("__DEF_FUN2")
1019#pragma push_macro("__DEF_FUNI")
1020#pragma push_macro("__DEF_FLOAT_FUN2I")
1021#pragma push_macro("__HIP_OVERLOAD1")
1022#pragma push_macro("__HIP_OVERLOAD2")
1023
1024// __hip_enable_if::type is a type function which returns __T if __B is true.
1025template <bool __B, class __T = void> struct __hip_enable_if {};
1026
1027template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
1028
1029// __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
1030// avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
1031// floor(double).
1032#define __HIP_OVERLOAD1(__retty, __fn) \
1033 template <typename __T> \
1034 __DEVICE__ typename __hip_enable_if<std::numeric_limits<__T>::is_integer, \
1035 __retty>::type \
1036 __fn(__T __x) { \
1037 return ::__fn((double)__x); \
1038 }
1039
1040// __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double
1041// or integer argument to avoid compilation error due to ambibuity. e.g.
1042// max(5.0f, 6.0) is resolved with max(double, double).
1043#define __HIP_OVERLOAD2(__retty, __fn) \
1044 template <typename __T1, typename __T2> \
1045 __DEVICE__ \
1046 typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && \
1047 std::numeric_limits<__T2>::is_specialized, \
1048 __retty>::type \
1049 __fn(__T1 __x, __T2 __y) { \
1050 return __fn((double)__x, (double)__y); \
1051 }
1052
1053// Define cmath functions with float argument and returns float.
1054#define __DEF_FUN1(__retty, __func) \
1055 __DEVICE__ \
1056 inline float __func(float __x) { return __func##f(__x); } \
1057 __HIP_OVERLOAD1(__retty, __func)
1058
1059// Define cmath functions with float argument and returns __retty.
1060#define __DEF_FUNI(__retty, __func) \
1061 __DEVICE__ \
1062 inline __retty __func(float __x) { return __func##f(__x); } \
1063 __HIP_OVERLOAD1(__retty, __func)
1064
1065// define cmath functions with two float arguments.
1066#define __DEF_FUN2(__retty, __func) \
1067 __DEVICE__ \
1068 inline float __func(float __x, float __y) { return __func##f(__x, __y); } \
1069 __HIP_OVERLOAD2(__retty, __func)
1070
1071__DEF_FUN1(double, acos)
1072__DEF_FUN1(double, acosh)
1073__DEF_FUN1(double, asin)
1074__DEF_FUN1(double, asinh)
1075__DEF_FUN1(double, atan)
1077__DEF_FUN1(double, atanh)
1078__DEF_FUN1(double, cbrt)
1079__DEF_FUN1(double, ceil)
1081__DEF_FUN1(double, cos)
1082__DEF_FUN1(double, cosh)
1083__DEF_FUN1(double, erf)
1084__DEF_FUN1(double, erfc)
1085__DEF_FUN1(double, exp)
1086__DEF_FUN1(double, exp2)
1087__DEF_FUN1(double, expm1)
1088__DEF_FUN1(double, fabs)
1090__DEF_FUN1(double, floor)
1094//__HIP_OVERLOAD1(int, fpclassify)
1096__DEF_FUNI(int, ilogb)
1105//__HIP_OVERLOAD1(bool, isnormal)
1107__DEF_FUN1(double, lgamma)
1108__DEF_FUN1(double, log)
1109__DEF_FUN1(double, log10)
1110__DEF_FUN1(double, log1p)
1111__DEF_FUN1(double, log2)
1112__DEF_FUN1(double, logb)
1113__DEF_FUNI(long long, llrint)
1114__DEF_FUNI(long long, llround)
1115__DEF_FUNI(long, lrint)
1116__DEF_FUNI(long, lround)
1124__DEF_FUN1(double, sin)
1125__DEF_FUN1(double, sinh)
1126__DEF_FUN1(double, sqrt)
1127__DEF_FUN1(double, tan)
1128__DEF_FUN1(double, tanh)
1129__DEF_FUN1(double, tgamma)
1131
1132// define cmath functions with a float and an integer argument.
1133#define __DEF_FLOAT_FUN2I(__func) \
1134 __DEVICE__ \
1135 inline float __func(float __x, int __y) { return __func##f(__x, __y); }
1137
1138template <class T> __DEVICE__ inline T min(T __arg1, T __arg2) {
1139 return (__arg1 < __arg2) ? __arg1 : __arg2;
1140}
1141
1142template <class T> __DEVICE__ inline T max(T __arg1, T __arg2) {
1143 return (__arg1 > __arg2) ? __arg1 : __arg2;
1144}
1145
1146__DEVICE__ inline int min(int __arg1, int __arg2) {
1147 return (__arg1 < __arg2) ? __arg1 : __arg2;
1148}
1149__DEVICE__ inline int max(int __arg1, int __arg2) {
1150 return (__arg1 > __arg2) ? __arg1 : __arg2;
1151}
1152
1154inline float max(float __x, float __y) { return fmaxf(__x, __y); }
1155
1157inline double max(double __x, double __y) { return fmax(__x, __y); }
1158
1160inline float min(float __x, float __y) { return fminf(__x, __y); }
1161
1163inline double min(double __x, double __y) { return fmin(__x, __y); }
1164
1165__HIP_OVERLOAD2(double, max)
1166__HIP_OVERLOAD2(double, min)
1167
1168__host__ inline static int min(int __arg1, int __arg2) {
1169 return std::min(__arg1, __arg2);
1170}
1171
1172__host__ inline static int max(int __arg1, int __arg2) {
1173 return std::max(__arg1, __arg2);
1174}
1175
1176#pragma pop_macro("__DEF_FUN1")
1177#pragma pop_macro("__DEF_FUN2")
1178#pragma pop_macro("__DEF_FUNI")
1179#pragma pop_macro("__DEF_FLOAT_FUN2I")
1180#pragma pop_macro("__HIP_OVERLOAD1")
1181#pragma pop_macro("__HIP_OVERLOAD2")
1182#pragma pop_macro("__DEVICE__")
1183#pragma pop_macro("__RETURN_TYPE")
1184
1185#endif // __CLANG_HIP_MATH_H__
__DEVICE__ bool isunordered(float __x, float __y)
Test if arguments are unordered.
__DEVICE__ bool isgreater(float __x, float __y)
Returns the component-wise compare of x > y.
__DEVICE__ bool islessgreater(float __x, float __y)
Returns the component-wise compare of (x < y) || (x > y) .
__DEVICE__ bool islessequal(float __x, float __y)
Returns the component-wise compare of x <= y.
__DEVICE__ bool isless(float __x, float __y)
Returns the component-wise compare of x < y.
__DEVICE__ bool isgreaterequal(float __x, float __y)
Returns the component-wise compare of x >= y.
__DEVICE__ float __fsqrt_rd(float __a)
__DEVICE__ float __fdiv_rd(float __a, float __b)
__DEVICE__ double __dsub_ru(double __a, double __b)
__DEVICE__ double __drcp_ru(double __a)
__DEVICE__ float __frcp_rz(float __a)
__DEVICE__ float __fmul_ru(float __a, float __b)
__DEVICE__ double __dsub_rd(double __a, double __b)
__DEVICE__ float __frcp_ru(float __a)
__DEVICE__ float __frcp_rd(float __a)
__DEVICE__ double __dmul_ru(double __a, double __b)
__DEVICE__ float __fmaf_ru(float __a, float __b, float __c)
__DEVICE__ double __fma_rz(double __a, double __b, double __c)
__DEVICE__ double __fma_rd(double __a, double __b, double __c)
__DEVICE__ double __dmul_rd(double __a, double __b)
__DEVICE__ double __ddiv_ru(double __a, double __b)
__DEVICE__ double __ddiv_rd(double __a, double __b)
__DEVICE__ double __dadd_ru(double __a, double __b)
__DEVICE__ float __fmul_rd(float __a, float __b)
__DEVICE__ float __fsub_rd(float __a, float __b)
__DEVICE__ float __fsub_rz(float __a, float __b)
__DEVICE__ double __fma_ru(double __a, double __b, double __c)
__DEVICE__ double __dsqrt_ru(double __a)
__DEVICE__ float __fsqrt_rz(float __a)
__DEVICE__ double __dsub_rz(double __a, double __b)
__DEVICE__ float __fadd_rd(float __a, float __b)
__DEVICE__ float __fmul_rz(float __a, float __b)
__DEVICE__ float __fadd_rz(float __a, float __b)
__DEVICE__ double __dsqrt_rd(double __a)
__DEVICE__ float __fmaf_rd(float __a, float __b, float __c)
__DEVICE__ double __dadd_rd(double __a, double __b)
__DEVICE__ double __dsqrt_rz(double __a)
__DEVICE__ double __drcp_rd(double __a)
__DEVICE__ float __fdiv_rz(float __a, float __b)
__DEVICE__ float __fmaf_rz(float __a, float __b, float __c)
__DEVICE__ double __drcp_rz(double __a)
__DEVICE__ float __fsub_ru(float __a, float __b)
__DEVICE__ double __dmul_rz(double __a, double __b)
__DEVICE__ float __fsqrt_ru(float __a)
__DEVICE__ float __fadd_ru(float __a, float __b)
__DEVICE__ float __fdiv_ru(float __a, float __b)
__DEVICE__ double __ddiv_rz(double __a, double __b)
__DEVICE__ double __dadd_rz(double __a, double __b)
__device__ double __ocml_i0_f64(double)
__device__ float __ocml_j1_f32(float)
__device__ double __ocml_remquo_f64(double, double, __attribute__((address_space(5))) int *)
__device__ double __ocml_modf_f64(double, __attribute__((address_space(5))) double *)
__device__ float __ocml_cospi_f32(float)
__device__ float __ocml_i0_f32(float)
__device__ double __ocml_lgamma_f64(double)
__device__ float __ocml_sincos_f32(float, __attribute__((address_space(5))) float *)
__device__ float __ocml_remquo_f32(float, float, __attribute__((address_space(5))) int *)
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
__device__ float __ocml_y0_f32(float)
__device__ float __ocml_modf_f32(float, __attribute__((address_space(5))) float *)
__device__ float __ocml_i1_f32(float)
__device__ float __ocml_lgamma_f32(float)
__device__ float __ocml_frexp_f32(float, __attribute__((address_space(5))) int *)
__device__ double __ocml_sinpi_f64(double)
__device__ double
__device__ double __ocml_cospi_f64(double)
__device__ double __ocml_frexp_f64(double, __attribute__((address_space(5))) int *)
__device__ double __ocml_tgamma_f64(double)
__device__ float __ocml_tan_f32(float)
__device__ float __ocml_tgamma_f32(float)
__device__ double __ocml_sincos_f64(double, __attribute__((address_space(5))) double *)
__device__ float __ocml_sinpi_f32(float)
__device__ double __ocml_j1_f64(double)
__device__ double __ocml_y1_f64(double)
__device__ double __ocml_j0_f64(double)
__device__ float __ocml_cos_f32(float)
__device__ float __ocml_y1_f32(float)
__device__ float __ocml_j0_f32(float)
__device__ double __ocml_cos_f64(double)
__device__ double __ocml_i1_f64(double)
__device__ double __ocml_sin_f64(double)
__device__ double __ocml_sincospi_f64(double, __attribute__((address_space(5))) double *)
__device__ _Float16
__device__ float __ocml_sin_f32(float)
__device__ float
__device__ float __ocml_sincospi_f32(float, __attribute__((address_space(5))) float *)
__device__ float __ocml_native_sin_f32(float)
__device__ float __ocml_native_cos_f32(float)
__device__ double __ocml_y0_f64(double)
__device__ double __ocml_tan_f64(double)
__DEVICE__ double __dsub_rn(double __x, double __y)
__DEVICE__ float sinpif(float __x)
__DEVICE__ float tanf(float __x)
__DEVICE__ float log2f(float __x)
__DEVICE__ float y0f(float __x)
__DEVICE__ float tanhf(float __x)
__DEVICE__ float coshf(float __x)
__DEVICE__ float log10f(float __x)
#define __HIP_OVERLOAD2(__retty, __fn)
__DEVICE__ float j1f(float __x)
__DEVICE__ float ldexpf(float __x, int __e)
__DEVICE__ long long int llroundf(float __x)
__DEVICE__ double rhypot(double __x, double __y)
__DEVICE__ double normcdfinv(double __x)
__DEVICE__ double norm3d(double __x, double __y, double __z)
__DEVICE__ float truncf(float __x)
__DEVICE__ float remainderf(float __x, float __y)
__DEVICE__ float fabsf(float __x)
__DEVICE__ float __fdiv_rn(float __x, float __y)
__DEVICE__ float scalbnf(float __x, int __n)
__DEVICE__ float cyl_bessel_i0f(float __x)
__DEVICE__ float lgammaf(float __x)
__DEVICE__ float cospif(float __x)
#define __DEF_FUNI(__retty, __func)
__DEVICE__ double __dsqrt_rn(double __x)
__DEVICE__ float frexpf(float __x, int *__nptr)
__DEVICE__ float tgammaf(float __x)
__DEVICE__ float __sinf(float __x)
__DEVICE__ float erfinvf(float __x)
__DEVICE__ float modff(float __x, float *__iptr)
__DEVICE__ double erfinv(double __x)
__DEVICE__ float expm1f(float __x)
#define __DEF_FUN2(__retty, __func)
__DEVICE__ float sinhf(float __x)
__DEVICE__ double j0(double __x)
__DEVICE__ float y1f(float __x)
__DEVICE__ float acosf(float __x)
__DEVICE__ float fmaf(float __x, float __y, float __z)
__DEVICE__ float cyl_bessel_i1f(float __x)
__DEVICE__ float fmodf(float __x, float __y)
__DEVICE__ float log1pf(float __x)
__DEVICE__ float atan2f(float __x, float __y)
__DEVICE__ float copysignf(float __x, float __y)
__DEVICE__ double j1(double __x)
__DEVICE__ long labs(long __x)
__DEVICE__ float rnormf(int __dim, const float *__a)
#define __DEF_FLOAT_FUN2I(__func)
__DEVICE__ float rnorm4df(float __x, float __y, float __z, float __w)
__DEVICE__ float __cosf(float __x)
__DEVICE__ float erff(float __x)
__DEVICE__ float atanf(float __x)
__DEVICE__ float rnorm3df(float __x, float __y, float __z)
__DEVICE__ double norm(int __dim, const double *__a)
__DEVICE__ float erfcxf(float __x)
__DEVICE__ float erfcinvf(float __x)
__DEVICE__ float asinf(float __x)
__DEVICE__ float abs(float __x)
__DEVICE__ long int lroundf(float __x)
__DEVICE__ float __fdividef(float __x, float __y)
__DEVICE__ float __frsqrt_rn(float __x)
__DEVICE__ float __log2f(float __x)
__DEVICE__ float norm4df(float __x, float __y, float __z, float __w)
__DEVICE__ double jn(int __n, double __x)
__DEVICE__ __RETURN_TYPE signbit(float __x)
Test for sign bit.
#define __HIP_OVERLOAD1(__retty, __fn)
__DEVICE__ float __exp10f(float __x)
__DEVICE__ float __frcp_rn(float __x)
#define __DEVICE__
__DEVICE__ float ynf(int __n, float __x)
__DEVICE__ float powf(float __x, float __y)
__DEVICE__ float __fsub_rn(float __x, float __y)
__DEVICE__ double __dadd_rn(double __x, double __y)
__DEVICE__ float sinf(float __x)
__DEVICE__ float __tanf(float __x)
__DEVICE__ float remquof(float __x, float __y, int *__quo)
__DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp)
__DEVICE__ double normcdf(double __x)
__DEVICE__ float __fsqrt_rn(float __x)
__DEVICE__ float hypotf(float __x, float __y)
__DEVICE__ float __fmaf_rn(float __x, float __y, float __z)
__DEVICE__ void sincosf(float __x, float *__sinptr, float *__cosptr)
__DEVICE__ float exp10f(float __x)
__DEVICE__ double y1(double __x)
__DEVICE__ float fmaxf(float __x, float __y)
__DEVICE__ float fminf(float __x, float __y)
__DEVICE__ double erfcinv(double __x)
__DEVICE__ float logf(float __x)
__DEVICE__ float __fadd_rn(float __x, float __y)
__DEVICE__ double cospi(double __x)
__DEVICE__ uint64_t __make_mantissa(const char *__tagp)
__DEVICE__ double rsqrt(double __x)
__DEVICE__ float erfcf(float __x)
__DEVICE__ float atanhf(float __x)
__DEVICE__ float asinhf(float __x)
__DEVICE__ float __expf(float __x)
__DEVICE__ double norm4d(double __x, double __y, double __z, double __w)
#define __RETURN_TYPE
__DEVICE__ float __logf(float __x)
__DEVICE__ double __fma_rn(double __x, double __y, double __z)
__DEVICE__ double nan(const char *__tagp)
__DEVICE__ double rnorm(int __dim, const double *__a)
__DEVICE__ float j0f(float __x)
__DEVICE__ float rsqrtf(float __x)
__DEVICE__ T max(T __arg1, T __arg2)
__DEVICE__ float jnf(int __n, float __x)
__DEVICE__ double sinpi(double __x)
__DEVICE__ float logbf(float __x)
__DEVICE__ double y0(double __x)
__DEVICE__ T min(T __arg1, T __arg2)
__DEVICE__ __RETURN_TYPE isinf(float __x)
Test for infinity value (+ve or -ve) .
__DEVICE__ void __sincosf(float __x, float *__sinptr, float *__cosptr)
__DEVICE__ double yn(int __n, double __x)
__DEVICE__ float rhypotf(float __x, float __y)
__DEVICE__ float exp2f(float __x)
__DEVICE__ double cyl_bessel_i0(double __x)
__DEVICE__ double __ddiv_rn(double __x, double __y)
__DEVICE__ double cyl_bessel_i1(double __x)
__DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp)
__DEVICE__ float ceilf(float __x)
__DEVICE__ double rcbrt(double __x)
__DEVICE__ double rnorm3d(double __x, double __y, double __z)
__DEVICE__ float normcdfinvf(float __x)
__DEVICE__ float norm3df(float __x, float __y, float __z)
__DEVICE__ void sincos(double __x, double *__sinptr, double *__cosptr)
__DEVICE__ __RETURN_TYPE isnan(float __x)
Test for a NaN.
__DEVICE__ float fdimf(float __x, float __y)
__DEVICE__ double __dmul_rn(double __x, double __y)
#define __DEF_FUN1(__retty, __func)
__DEVICE__ float normf(int __dim, const float *__a)
__DEVICE__ float nearbyintf(float __x)
__DEVICE__ int ilogbf(float __x)
__DEVICE__ float floorf(float __x)
__DEVICE__ long long llabs(long long __x)
__DEVICE__ float sqrtf(float __x)
__DEVICE__ float roundf(float __x)
__DEVICE__ void sincospif(float __x, float *__sinptr, float *__cosptr)
__DEVICE__ double __drcp_rn(double __x)
__DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp)
__DEVICE__ long int lrintf(float __x)
__DEVICE__ float acoshf(float __x)
__DEVICE__ double modf(double __x, double *__iptr)
__DEVICE__ __RETURN_TYPE isfinite(float __x)
Test for finite value.
__DEVICE__ float cosf(float __x)
__DEVICE__ float expf(float __x)
__DEVICE__ float nextafterf(float __x, float __y)
__DEVICE__ double rnorm4d(double __x, double __y, double __z, double __w)
__DEVICE__ long long int llrintf(float __x)
__DEVICE__ double erfcx(double __x)
__DEVICE__ float fdividef(float __x, float __y)
__DEVICE__ float rcbrtf(float __x)
__DEVICE__ double exp10(double __x)
__DEVICE__ float __log10f(float __x)
__DEVICE__ float cbrtf(float __x)
__DEVICE__ float nanf(const char *__tagp)
__DEVICE__ float __fmul_rn(float __x, float __y)
__DEVICE__ void sincospi(double __x, double *__sinptr, double *__cosptr)
__DEVICE__ float scalblnf(float __x, long int __n)
__DEVICE__ float rintf(float __x)
__DEVICE__ float normcdff(float __x)
__DEVICE__ float __saturatef(float __x)
__DEVICE__ float __powf(float __x, float __y)
static __inline unsigned char unsigned int __x
Definition adxintrin.h:22
static __inline unsigned char unsigned int unsigned int __y
Definition adxintrin.h:22
static __inline__ uint32_t
Definition arm_cde.h:25
static __inline__ uint64_t
Definition arm_cde.h:31
static __inline__ void int __a
Definition emmintrin.h:4185
#define CHAR_BIT
Definition limits.h:63
#define INT_MAX
Definition limits.h:46
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.
#define true
Definition stdbool.h:16
#define sinh(__x)
Definition tgmath.h:373
#define asin(__x)
Definition tgmath.h:112
#define scalbln(__x, __y)
Definition tgmath.h:1182
#define sqrt(__x)
Definition tgmath.h:520
#define acos(__x)
Definition tgmath.h:83
#define fmin(__x, __y)
Definition tgmath.h:780
#define exp(__x)
Definition tgmath.h:431
#define ilogb(__x)
Definition tgmath.h:851
#define copysign(__x, __y)
Definition tgmath.h:618
#define erf(__x)
Definition tgmath.h:636
#define atanh(__x)
Definition tgmath.h:228
#define remquo(__x, __y, __z)
Definition tgmath.h:1111
#define nextafter(__x, __y)
Definition tgmath.h:1055
#define frexp(__x, __y)
Definition tgmath.h:816
#define asinh(__x)
Definition tgmath.h:199
#define erfc(__x)
Definition tgmath.h:653
#define atan2(__x, __y)
Definition tgmath.h:566
#define hypot(__x, __y)
Definition tgmath.h:833
#define exp2(__x)
Definition tgmath.h:670
#define sin(__x)
Definition tgmath.h:286
#define cbrt(__x)
Definition tgmath.h:584
#define log2(__x)
Definition tgmath.h:970
#define llround(__x)
Definition tgmath.h:919
#define cosh(__x)
Definition tgmath.h:344
#define trunc(__x)
Definition tgmath.h:1216
#define fmax(__x, __y)
Definition tgmath.h:762
#define ldexp(__x, __y)
Definition tgmath.h:868
#define acosh(__x)
Definition tgmath.h:170
#define tgamma(__x)
Definition tgmath.h:1199
#define scalbn(__x, __y)
Definition tgmath.h:1165
#define round(__x)
Definition tgmath.h:1148
#define fmod(__x, __y)
Definition tgmath.h:798
#define llrint(__x)
Definition tgmath.h:902
#define tan(__x)
Definition tgmath.h:315
#define cos(__x)
Definition tgmath.h:257
#define log10(__x)
Definition tgmath.h:936
#define fabs(__x)
Definition tgmath.h:549
#define pow(__x, __y)
Definition tgmath.h:490
#define log1p(__x)
Definition tgmath.h:953
#define rint(__x)
Definition tgmath.h:1131
#define expm1(__x)
Definition tgmath.h:687
#define remainder(__x, __y)
Definition tgmath.h:1090
#define fdim(__x, __y)
Definition tgmath.h:704
#define lgamma(__x)
Definition tgmath.h:885
#define tanh(__x)
Definition tgmath.h:402
#define lrint(__x)
Definition tgmath.h:1004
#define atan(__x)
Definition tgmath.h:141
#define floor(__x)
Definition tgmath.h:722
#define ceil(__x)
Definition tgmath.h:601
#define log(__x)
Definition tgmath.h:460
#define logb(__x)
Definition tgmath.h:987
#define nearbyint(__x)
Definition tgmath.h:1038
#define lround(__x)
Definition tgmath.h:1021
#define fma(__x, __y, __z)
Definition tgmath.h:742