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