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