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