1 /*===---- __clang_hip_cmath.h - HIP cmath 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_CMATH_H__ 11 #define __CLANG_HIP_CMATH_H__ 12 13 #if !defined(__HIP__) 14 #error "This file is for HIP and OpenMP AMDGCN device compilation only." 15 #endif 16 17 #if defined(__cplusplus) 18 #include <limits> 19 #include <type_traits> 20 #include <utility> 21 #endif 22 #include <limits.h> 23 #include <stdint.h> 24 25 #pragma push_macro("__DEVICE__") 26 #define __DEVICE__ static __device__ inline __attribute__((always_inline)) 27 28 // Start with functions that cannot be defined by DEF macros below. 29 #if defined(__cplusplus) 30 __DEVICE__ double abs(double __x) { return ::fabs(__x); } 31 __DEVICE__ float abs(float __x) { return ::fabsf(__x); } 32 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); } 33 __DEVICE__ long abs(long __n) { return ::labs(__n); } 34 __DEVICE__ float fma(float __x, float __y, float __z) { 35 return ::fmaf(__x, __y, __z); 36 } 37 __DEVICE__ int fpclassify(float __x) { 38 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 39 FP_ZERO, __x); 40 } 41 __DEVICE__ int fpclassify(double __x) { 42 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 43 FP_ZERO, __x); 44 } 45 __DEVICE__ float frexp(float __arg, int *__exp) { 46 return ::frexpf(__arg, __exp); 47 } 48 __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } 49 __DEVICE__ bool isfinite(double __x) { return ::__finite(__x); } 50 __DEVICE__ bool isgreater(float __x, float __y) { 51 return __builtin_isgreater(__x, __y); 52 } 53 __DEVICE__ bool isgreater(double __x, double __y) { 54 return __builtin_isgreater(__x, __y); 55 } 56 __DEVICE__ bool isgreaterequal(float __x, float __y) { 57 return __builtin_isgreaterequal(__x, __y); 58 } 59 __DEVICE__ bool isgreaterequal(double __x, double __y) { 60 return __builtin_isgreaterequal(__x, __y); 61 } 62 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } 63 __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } 64 __DEVICE__ bool isless(float __x, float __y) { 65 return __builtin_isless(__x, __y); 66 } 67 __DEVICE__ bool isless(double __x, double __y) { 68 return __builtin_isless(__x, __y); 69 } 70 __DEVICE__ bool islessequal(float __x, float __y) { 71 return __builtin_islessequal(__x, __y); 72 } 73 __DEVICE__ bool islessequal(double __x, double __y) { 74 return __builtin_islessequal(__x, __y); 75 } 76 __DEVICE__ bool islessgreater(float __x, float __y) { 77 return __builtin_islessgreater(__x, __y); 78 } 79 __DEVICE__ bool islessgreater(double __x, double __y) { 80 return __builtin_islessgreater(__x, __y); 81 } 82 __DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); } 83 __DEVICE__ bool isnan(double __x) { return ::__isnan(__x); } 84 __DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); } 85 __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); } 86 __DEVICE__ bool isunordered(float __x, float __y) { 87 return __builtin_isunordered(__x, __y); 88 } 89 __DEVICE__ bool isunordered(double __x, double __y) { 90 return __builtin_isunordered(__x, __y); 91 } 92 __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } 93 __DEVICE__ float pow(float __base, int __iexp) { 94 return ::powif(__base, __iexp); 95 } 96 __DEVICE__ double pow(double __base, int __iexp) { 97 return ::powi(__base, __iexp); 98 } 99 __DEVICE__ float remquo(float __x, float __y, int *__quo) { 100 return ::remquof(__x, __y, __quo); 101 } 102 __DEVICE__ float scalbln(float __x, long int __n) { 103 return ::scalblnf(__x, __n); 104 } 105 __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); } 106 __DEVICE__ bool signbit(double __x) { return ::__signbit(__x); } 107 108 // Notably missing above is nexttoward. We omit it because 109 // ocml doesn't provide an implementation, and we don't want to be in the 110 // business of implementing tricky libm functions in this header. 111 112 // Other functions. 113 __DEVICE__ _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) { 114 return __ocml_fma_f16(__x, __y, __z); 115 } 116 __DEVICE__ _Float16 pow(_Float16 __base, int __iexp) { 117 return __ocml_pown_f16(__base, __iexp); 118 } 119 120 // BEGIN DEF_FUN and HIP_OVERLOAD 121 122 // BEGIN DEF_FUN 123 124 #pragma push_macro("__DEF_FUN1") 125 #pragma push_macro("__DEF_FUN2") 126 #pragma push_macro("__DEF_FUN2_FI") 127 128 // Define cmath functions with float argument and returns __retty. 129 #define __DEF_FUN1(__retty, __func) \ 130 __DEVICE__ \ 131 __retty __func(float __x) { return __func##f(__x); } 132 133 // Define cmath functions with two float arguments and returns __retty. 134 #define __DEF_FUN2(__retty, __func) \ 135 __DEVICE__ \ 136 __retty __func(float __x, float __y) { return __func##f(__x, __y); } 137 138 // Define cmath functions with a float and an int argument and returns __retty. 139 #define __DEF_FUN2_FI(__retty, __func) \ 140 __DEVICE__ \ 141 __retty __func(float __x, int __y) { return __func##f(__x, __y); } 142 143 __DEF_FUN1(float, acos) 144 __DEF_FUN1(float, acosh) 145 __DEF_FUN1(float, asin) 146 __DEF_FUN1(float, asinh) 147 __DEF_FUN1(float, atan) 148 __DEF_FUN2(float, atan2) 149 __DEF_FUN1(float, atanh) 150 __DEF_FUN1(float, cbrt) 151 __DEF_FUN1(float, ceil) 152 __DEF_FUN2(float, copysign) 153 __DEF_FUN1(float, cos) 154 __DEF_FUN1(float, cosh) 155 __DEF_FUN1(float, erf) 156 __DEF_FUN1(float, erfc) 157 __DEF_FUN1(float, exp) 158 __DEF_FUN1(float, exp2) 159 __DEF_FUN1(float, expm1) 160 __DEF_FUN1(float, fabs) 161 __DEF_FUN2(float, fdim) 162 __DEF_FUN1(float, floor) 163 __DEF_FUN2(float, fmax) 164 __DEF_FUN2(float, fmin) 165 __DEF_FUN2(float, fmod) 166 __DEF_FUN2(float, hypot) 167 __DEF_FUN1(int, ilogb) 168 __DEF_FUN2_FI(float, ldexp) 169 __DEF_FUN1(float, lgamma) 170 __DEF_FUN1(float, log) 171 __DEF_FUN1(float, log10) 172 __DEF_FUN1(float, log1p) 173 __DEF_FUN1(float, log2) 174 __DEF_FUN1(float, logb) 175 __DEF_FUN1(long long, llrint) 176 __DEF_FUN1(long long, llround) 177 __DEF_FUN1(long, lrint) 178 __DEF_FUN1(long, lround) 179 __DEF_FUN1(float, nearbyint) 180 __DEF_FUN2(float, nextafter) 181 __DEF_FUN2(float, pow) 182 __DEF_FUN2(float, remainder) 183 __DEF_FUN1(float, rint) 184 __DEF_FUN1(float, round) 185 __DEF_FUN2_FI(float, scalbn) 186 __DEF_FUN1(float, sin) 187 __DEF_FUN1(float, sinh) 188 __DEF_FUN1(float, sqrt) 189 __DEF_FUN1(float, tan) 190 __DEF_FUN1(float, tanh) 191 __DEF_FUN1(float, tgamma) 192 __DEF_FUN1(float, trunc) 193 194 #pragma pop_macro("__DEF_FUN1") 195 #pragma pop_macro("__DEF_FUN2") 196 #pragma pop_macro("__DEF_FUN2_FI") 197 198 // END DEF_FUN 199 200 // BEGIN HIP_OVERLOAD 201 202 #pragma push_macro("__HIP_OVERLOAD1") 203 #pragma push_macro("__HIP_OVERLOAD2") 204 205 // __hip_enable_if::type is a type function which returns __T if __B is true. 206 template <bool __B, class __T = void> struct __hip_enable_if {}; 207 208 template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; }; 209 210 // decltype is only available in C++11 and above. 211 #if __cplusplus >= 201103L 212 // __hip_promote 213 namespace __hip { 214 215 template <class _Tp> struct __numeric_type { 216 static void __test(...); 217 static _Float16 __test(_Float16); 218 static float __test(float); 219 static double __test(char); 220 static double __test(int); 221 static double __test(unsigned); 222 static double __test(long); 223 static double __test(unsigned long); 224 static double __test(long long); 225 static double __test(unsigned long long); 226 static double __test(double); 227 // No support for long double, use double instead. 228 static double __test(long double); 229 230 typedef decltype(__test(std::declval<_Tp>())) type; 231 static const bool value = !std::is_same<type, void>::value; 232 }; 233 234 template <> struct __numeric_type<void> { static const bool value = true; }; 235 236 template <class _A1, class _A2 = void, class _A3 = void, 237 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value 238 &&__numeric_type<_A3>::value> 239 class __promote_imp { 240 public: 241 static const bool value = false; 242 }; 243 244 template <class _A1, class _A2, class _A3> 245 class __promote_imp<_A1, _A2, _A3, true> { 246 private: 247 typedef typename __promote_imp<_A1>::type __type1; 248 typedef typename __promote_imp<_A2>::type __type2; 249 typedef typename __promote_imp<_A3>::type __type3; 250 251 public: 252 typedef decltype(__type1() + __type2() + __type3()) type; 253 static const bool value = true; 254 }; 255 256 template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> { 257 private: 258 typedef typename __promote_imp<_A1>::type __type1; 259 typedef typename __promote_imp<_A2>::type __type2; 260 261 public: 262 typedef decltype(__type1() + __type2()) type; 263 static const bool value = true; 264 }; 265 266 template <class _A1> class __promote_imp<_A1, void, void, true> { 267 public: 268 typedef typename __numeric_type<_A1>::type type; 269 static const bool value = true; 270 }; 271 272 template <class _A1, class _A2 = void, class _A3 = void> 273 class __promote : public __promote_imp<_A1, _A2, _A3> {}; 274 275 } // namespace __hip 276 #endif //__cplusplus >= 201103L 277 278 // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to 279 // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with 280 // floor(double). 281 #define __HIP_OVERLOAD1(__retty, __fn) \ 282 template <typename __T> \ 283 __DEVICE__ typename __hip_enable_if<std::numeric_limits<__T>::is_integer, \ 284 __retty>::type \ 285 __fn(__T __x) { \ 286 return ::__fn((double)__x); \ 287 } 288 289 // __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double 290 // or integer argument to avoid compilation error due to ambibuity. e.g. 291 // max(5.0f, 6.0) is resolved with max(double, double). 292 #if __cplusplus >= 201103L 293 #define __HIP_OVERLOAD2(__retty, __fn) \ 294 template <typename __T1, typename __T2> \ 295 __DEVICE__ typename __hip_enable_if< \ 296 std::numeric_limits<__T1>::is_specialized && \ 297 std::numeric_limits<__T2>::is_specialized, \ 298 typename __hip::__promote<__T1, __T2>::type>::type \ 299 __fn(__T1 __x, __T2 __y) { \ 300 typedef typename __hip::__promote<__T1, __T2>::type __result_type; \ 301 return __fn((__result_type)__x, (__result_type)__y); \ 302 } 303 #else 304 #define __HIP_OVERLOAD2(__retty, __fn) \ 305 template <typename __T1, typename __T2> \ 306 __DEVICE__ \ 307 typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && \ 308 std::numeric_limits<__T2>::is_specialized, \ 309 __retty>::type \ 310 __fn(__T1 __x, __T2 __y) { \ 311 return __fn((double)__x, (double)__y); \ 312 } 313 #endif 314 315 __HIP_OVERLOAD1(double, abs) 316 __HIP_OVERLOAD1(double, acos) 317 __HIP_OVERLOAD1(double, acosh) 318 __HIP_OVERLOAD1(double, asin) 319 __HIP_OVERLOAD1(double, asinh) 320 __HIP_OVERLOAD1(double, atan) 321 __HIP_OVERLOAD2(double, atan2) 322 __HIP_OVERLOAD1(double, atanh) 323 __HIP_OVERLOAD1(double, cbrt) 324 __HIP_OVERLOAD1(double, ceil) 325 __HIP_OVERLOAD2(double, copysign) 326 __HIP_OVERLOAD1(double, cos) 327 __HIP_OVERLOAD1(double, cosh) 328 __HIP_OVERLOAD1(double, erf) 329 __HIP_OVERLOAD1(double, erfc) 330 __HIP_OVERLOAD1(double, exp) 331 __HIP_OVERLOAD1(double, exp2) 332 __HIP_OVERLOAD1(double, expm1) 333 __HIP_OVERLOAD1(double, fabs) 334 __HIP_OVERLOAD2(double, fdim) 335 __HIP_OVERLOAD1(double, floor) 336 __HIP_OVERLOAD2(double, fmax) 337 __HIP_OVERLOAD2(double, fmin) 338 __HIP_OVERLOAD2(double, fmod) 339 __HIP_OVERLOAD1(int, fpclassify) 340 __HIP_OVERLOAD2(double, hypot) 341 __HIP_OVERLOAD1(int, ilogb) 342 __HIP_OVERLOAD1(bool, isfinite) 343 __HIP_OVERLOAD2(bool, isgreater) 344 __HIP_OVERLOAD2(bool, isgreaterequal) 345 __HIP_OVERLOAD1(bool, isinf) 346 __HIP_OVERLOAD2(bool, isless) 347 __HIP_OVERLOAD2(bool, islessequal) 348 __HIP_OVERLOAD2(bool, islessgreater) 349 __HIP_OVERLOAD1(bool, isnan) 350 __HIP_OVERLOAD1(bool, isnormal) 351 __HIP_OVERLOAD2(bool, isunordered) 352 __HIP_OVERLOAD1(double, lgamma) 353 __HIP_OVERLOAD1(double, log) 354 __HIP_OVERLOAD1(double, log10) 355 __HIP_OVERLOAD1(double, log1p) 356 __HIP_OVERLOAD1(double, log2) 357 __HIP_OVERLOAD1(double, logb) 358 __HIP_OVERLOAD1(long long, llrint) 359 __HIP_OVERLOAD1(long long, llround) 360 __HIP_OVERLOAD1(long, lrint) 361 __HIP_OVERLOAD1(long, lround) 362 __HIP_OVERLOAD1(double, nearbyint) 363 __HIP_OVERLOAD2(double, nextafter) 364 __HIP_OVERLOAD2(double, pow) 365 __HIP_OVERLOAD2(double, remainder) 366 __HIP_OVERLOAD1(double, rint) 367 __HIP_OVERLOAD1(double, round) 368 __HIP_OVERLOAD1(bool, signbit) 369 __HIP_OVERLOAD1(double, sin) 370 __HIP_OVERLOAD1(double, sinh) 371 __HIP_OVERLOAD1(double, sqrt) 372 __HIP_OVERLOAD1(double, tan) 373 __HIP_OVERLOAD1(double, tanh) 374 __HIP_OVERLOAD1(double, tgamma) 375 __HIP_OVERLOAD1(double, trunc) 376 377 // Overload these but don't add them to std, they are not part of cmath. 378 __HIP_OVERLOAD2(double, max) 379 __HIP_OVERLOAD2(double, min) 380 381 // Additional Overloads that don't quite match HIP_OVERLOAD. 382 #if __cplusplus >= 201103L 383 template <typename __T1, typename __T2, typename __T3> 384 __DEVICE__ typename __hip_enable_if< 385 std::numeric_limits<__T1>::is_specialized && 386 std::numeric_limits<__T2>::is_specialized && 387 std::numeric_limits<__T3>::is_specialized, 388 typename __hip::__promote<__T1, __T2, __T3>::type>::type 389 fma(__T1 __x, __T2 __y, __T3 __z) { 390 typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type; 391 return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z); 392 } 393 #else 394 template <typename __T1, typename __T2, typename __T3> 395 __DEVICE__ 396 typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && 397 std::numeric_limits<__T2>::is_specialized && 398 std::numeric_limits<__T3>::is_specialized, 399 double>::type 400 fma(__T1 __x, __T2 __y, __T3 __z) { 401 return ::fma((double)__x, (double)__y, (double)__z); 402 } 403 #endif 404 405 template <typename __T> 406 __DEVICE__ 407 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type 408 frexp(__T __x, int *__exp) { 409 return ::frexp((double)__x, __exp); 410 } 411 412 template <typename __T> 413 __DEVICE__ 414 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type 415 ldexp(__T __x, int __exp) { 416 return ::ldexp((double)__x, __exp); 417 } 418 419 template <typename __T> 420 __DEVICE__ 421 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type 422 modf(__T __x, double *__exp) { 423 return ::modf((double)__x, __exp); 424 } 425 426 #if __cplusplus >= 201103L 427 template <typename __T1, typename __T2> 428 __DEVICE__ 429 typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && 430 std::numeric_limits<__T2>::is_specialized, 431 typename __hip::__promote<__T1, __T2>::type>::type 432 remquo(__T1 __x, __T2 __y, int *__quo) { 433 typedef typename __hip::__promote<__T1, __T2>::type __result_type; 434 return ::remquo((__result_type)__x, (__result_type)__y, __quo); 435 } 436 #else 437 template <typename __T1, typename __T2> 438 __DEVICE__ 439 typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && 440 std::numeric_limits<__T2>::is_specialized, 441 double>::type 442 remquo(__T1 __x, __T2 __y, int *__quo) { 443 return ::remquo((double)__x, (double)__y, __quo); 444 } 445 #endif 446 447 template <typename __T> 448 __DEVICE__ 449 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type 450 scalbln(__T __x, long int __exp) { 451 return ::scalbln((double)__x, __exp); 452 } 453 454 template <typename __T> 455 __DEVICE__ 456 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type 457 scalbn(__T __x, int __exp) { 458 return ::scalbn((double)__x, __exp); 459 } 460 461 #pragma pop_macro("__HIP_OVERLOAD1") 462 #pragma pop_macro("__HIP_OVERLOAD2") 463 464 // END HIP_OVERLOAD 465 466 // END DEF_FUN and HIP_OVERLOAD 467 468 #endif // defined(__cplusplus) 469 470 // Define these overloads inside the namespace our standard library uses. 471 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD 472 _LIBCPP_BEGIN_NAMESPACE_STD 473 #else 474 namespace std { 475 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION 476 _GLIBCXX_BEGIN_NAMESPACE_VERSION 477 #endif 478 #endif 479 480 // Pull the new overloads we defined above into namespace std. 481 // using ::abs; - This may be considered for C++. 482 using ::acos; 483 using ::acosh; 484 using ::asin; 485 using ::asinh; 486 using ::atan; 487 using ::atan2; 488 using ::atanh; 489 using ::cbrt; 490 using ::ceil; 491 using ::copysign; 492 using ::cos; 493 using ::cosh; 494 using ::erf; 495 using ::erfc; 496 using ::exp; 497 using ::exp2; 498 using ::expm1; 499 using ::fabs; 500 using ::fdim; 501 using ::floor; 502 using ::fma; 503 using ::fmax; 504 using ::fmin; 505 using ::fmod; 506 using ::fpclassify; 507 using ::frexp; 508 using ::hypot; 509 using ::ilogb; 510 using ::isfinite; 511 using ::isgreater; 512 using ::isgreaterequal; 513 using ::isless; 514 using ::islessequal; 515 using ::islessgreater; 516 using ::isnormal; 517 using ::isunordered; 518 using ::ldexp; 519 using ::lgamma; 520 using ::llrint; 521 using ::llround; 522 using ::log; 523 using ::log10; 524 using ::log1p; 525 using ::log2; 526 using ::logb; 527 using ::lrint; 528 using ::lround; 529 using ::modf; 530 // using ::nan; - This may be considered for C++. 531 // using ::nanf; - This may be considered for C++. 532 // using ::nanl; - This is not yet defined. 533 using ::nearbyint; 534 using ::nextafter; 535 // using ::nexttoward; - Omit this since we do not have a definition. 536 using ::pow; 537 using ::remainder; 538 using ::remquo; 539 using ::rint; 540 using ::round; 541 using ::scalbln; 542 using ::scalbn; 543 using ::signbit; 544 using ::sin; 545 using ::sinh; 546 using ::sqrt; 547 using ::tan; 548 using ::tanh; 549 using ::tgamma; 550 using ::trunc; 551 552 // Well this is fun: We need to pull these symbols in for libc++, but we can't 553 // pull them in with libstdc++, because its ::isinf and ::isnan are different 554 // than its std::isinf and std::isnan. 555 #ifndef __GLIBCXX__ 556 using ::isinf; 557 using ::isnan; 558 #endif 559 560 // Finally, pull the "foobarf" functions that HIP defines into std. 561 using ::acosf; 562 using ::acoshf; 563 using ::asinf; 564 using ::asinhf; 565 using ::atan2f; 566 using ::atanf; 567 using ::atanhf; 568 using ::cbrtf; 569 using ::ceilf; 570 using ::copysignf; 571 using ::cosf; 572 using ::coshf; 573 using ::erfcf; 574 using ::erff; 575 using ::exp2f; 576 using ::expf; 577 using ::expm1f; 578 using ::fabsf; 579 using ::fdimf; 580 using ::floorf; 581 using ::fmaf; 582 using ::fmaxf; 583 using ::fminf; 584 using ::fmodf; 585 using ::frexpf; 586 using ::hypotf; 587 using ::ilogbf; 588 using ::ldexpf; 589 using ::lgammaf; 590 using ::llrintf; 591 using ::llroundf; 592 using ::log10f; 593 using ::log1pf; 594 using ::log2f; 595 using ::logbf; 596 using ::logf; 597 using ::lrintf; 598 using ::lroundf; 599 using ::modff; 600 using ::nearbyintf; 601 using ::nextafterf; 602 // using ::nexttowardf; - Omit this since we do not have a definition. 603 using ::powf; 604 using ::remainderf; 605 using ::remquof; 606 using ::rintf; 607 using ::roundf; 608 using ::scalblnf; 609 using ::scalbnf; 610 using ::sinf; 611 using ::sinhf; 612 using ::sqrtf; 613 using ::tanf; 614 using ::tanhf; 615 using ::tgammaf; 616 using ::truncf; 617 618 #ifdef _LIBCPP_END_NAMESPACE_STD 619 _LIBCPP_END_NAMESPACE_STD 620 #else 621 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION 622 _GLIBCXX_END_NAMESPACE_VERSION 623 #endif 624 } // namespace std 625 #endif 626 627 // Define device-side math functions from <ymath.h> on MSVC. 628 #if defined(_MSC_VER) 629 630 // Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers. 631 // But, from VS2019, it's only included in `<complex>`. Need to include 632 // `<ymath.h>` here to ensure C functions declared there won't be markded as 633 // `__host__` and `__device__` through `<complex>` wrapper. 634 #include <ymath.h> 635 636 #if defined(__cplusplus) 637 extern "C" { 638 #endif // defined(__cplusplus) 639 __DEVICE__ __attribute__((overloadable)) double _Cosh(double x, double y) { 640 return cosh(x) * y; 641 } 642 __DEVICE__ __attribute__((overloadable)) float _FCosh(float x, float y) { 643 return coshf(x) * y; 644 } 645 __DEVICE__ __attribute__((overloadable)) short _Dtest(double *p) { 646 return fpclassify(*p); 647 } 648 __DEVICE__ __attribute__((overloadable)) short _FDtest(float *p) { 649 return fpclassify(*p); 650 } 651 __DEVICE__ __attribute__((overloadable)) double _Sinh(double x, double y) { 652 return sinh(x) * y; 653 } 654 __DEVICE__ __attribute__((overloadable)) float _FSinh(float x, float y) { 655 return sinhf(x) * y; 656 } 657 #if defined(__cplusplus) 658 } 659 #endif // defined(__cplusplus) 660 #endif // defined(_MSC_VER) 661 662 #pragma pop_macro("__DEVICE__") 663 664 #endif // __CLANG_HIP_CMATH_H__ 665