clang/lib/Headers/__clang_hip_cmath.h

   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.
 268 template <bool __B, class __T = void> struct __hip_enable_if {};
 269
 270 template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
 271
 272 namespace __hip {
 273 template <class _Tp> struct is_integral {
 274   enum { value = 0 };
 275 };
 276 template <> struct is_integral<bool> {
 277   enum { value = 1 };
 278 };
 279 template <> struct is_integral<char> {
 280   enum { value = 1 };
 281 };
 282 template <> struct is_integral<signed char> {
 283   enum { value = 1 };
 284 };
 285 template <> struct is_integral<unsigned char> {
 286   enum { value = 1 };
 287 };
 288 template <> struct is_integral<wchar_t> {
 289   enum { value = 1 };
 290 };
 291 template <> struct is_integral<short> {
 292   enum { value = 1 };
 293 };
 294 template <> struct is_integral<unsigned short> {
 295   enum { value = 1 };
 296 };
 297 template <> struct is_integral<int> {
 298   enum { value = 1 };
 299 };
 300 template <> struct is_integral<unsigned int> {
 301   enum { value = 1 };
 302 };
 303 template <> struct is_integral<long> {
 304   enum { value = 1 };
 305 };
 306 template <> struct is_integral<unsigned long> {
 307   enum { value = 1 };
 308 };
 309 template <> struct is_integral<long long> {
 310   enum { value = 1 };
 311 };
 312 template <> struct is_integral<unsigned long long> {
 313   enum { value = 1 };
 314 };
 315
 316 // ToDo: specializes is_arithmetic<_Float16>
 317 template <class _Tp> struct is_arithmetic {
 318   enum { value = 0 };
 319 };
 320 template <> struct is_arithmetic<bool> {
 321   enum { value = 1 };
 322 };
 323 template <> struct is_arithmetic<char> {
 324   enum { value = 1 };
 325 };
 326 template <> struct is_arithmetic<signed char> {
 327   enum { value = 1 };
 328 };
 329 template <> struct is_arithmetic<unsigned char> {
 330   enum { value = 1 };
 331 };
 332 template <> struct is_arithmetic<wchar_t> {
 333   enum { value = 1 };
 334 };
 335 template <> struct is_arithmetic<short> {
 336   enum { value = 1 };
 337 };
 338 template <> struct is_arithmetic<unsigned short> {
 339   enum { value = 1 };
 340 };
 341 template <> struct is_arithmetic<int> {
 342   enum { value = 1 };
 343 };
 344 template <> struct is_arithmetic<unsigned int> {
 345   enum { value = 1 };
 346 };
 347 template <> struct is_arithmetic<long> {
 348   enum { value = 1 };
 349 };
 350 template <> struct is_arithmetic<unsigned long> {
 351   enum { value = 1 };
 352 };
 353 template <> struct is_arithmetic<long long> {
 354   enum { value = 1 };
 355 };
 356 template <> struct is_arithmetic<unsigned long long> {
 357   enum { value = 1 };
 358 };
 359 template <> struct is_arithmetic<float> {
 360   enum { value = 1 };
 361 };
 362 template <> struct is_arithmetic<double> {
 363   enum { value = 1 };
 364 };
 365
 366 struct true_type {
 367   static const __constant__ bool value = true;
 368 };
 369 struct false_type {
 370   static const __constant__ bool value = false;
 371 };
 372
 373 template <typename __T, typename __U> struct is_same : public false_type {};
 374 template <typename __T> struct is_same<__T, __T> : public true_type {};
 375
 376 template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
 377
 378 template <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
 383 template <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   template <typename _U>
 399   static auto __test_impl(int) -> decltype(__test(declval<_U>()));
 400
 401   template <typename _U> static void __test_impl(...);
 402
 403   typedef decltype(__test_impl<_Tp>(0)) type;
 404   static const bool value = !is_same<type, void>::value;
 405 };
 406
 407 template <> struct __numeric_type<void> { static const bool value = true; };
 408
 409 template <class _A1, class _A2 = void, class _A3 = void,
 410           bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
 411               &&__numeric_type<_A3>::value>
 412 class __promote_imp {
 413 public:
 414   static const bool value = false;
 415 };
 416
 417 template <class _A1, class _A2, class _A3>
 418 class __promote_imp<_A1, _A2, _A3, true> {
 419 private:
 420   typedef typename __promote_imp<_A1>::type __type1;
 421   typedef typename __promote_imp<_A2>::type __type2;
 422   typedef typename __promote_imp<_A3>::type __type3;
 423
 424 public:
 425   typedef decltype(__type1() + __type2() + __type3()) type;
 426   static const bool value = true;
 427 };
 428
 429 template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
 430 private:
 431   typedef typename __promote_imp<_A1>::type __type1;
 432   typedef typename __promote_imp<_A2>::type __type2;
 433
 434 public:
 435   typedef decltype(__type1() + __type2()) type;
 436   static const bool value = true;
 437 };
 438
 439 template <class _A1> class __promote_imp<_A1, void, void, true> {
 440 public:
 441   typedef typename __numeric_type<_A1>::type type;
 442   static const bool value = true;
 443 };
 444
 445 template <class _A1, class _A2 = void, class _A3 = void>
 446 class __promote : public __promote_imp<_A1, _A2, _A3> {};
 447 #endif //__cplusplus >= 201103L
 448 } // namespace __hip
 449
 450 // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
 451 // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
 452 // floor(double).
 453 #define __HIP_OVERLOAD1(__retty, __fn)                                         \
 454   template <typename __T>                                                      \
 455   __DEVICE__ __CONSTEXPR__                                                     \
 456       typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type  \
 457       __fn(__T __x) {                                                          \
 458     return ::__fn((double)__x);                                                \
 459   }
 460
 461 // __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double
 462 // or integer argument to avoid compilation error due to ambibuity. e.g.
 463 // max(5.0f, 6.0) is resolved with max(double, double).
 464 #if __cplusplus >= 201103L
 465 #define __HIP_OVERLOAD2(__retty, __fn)                                         \
 466   template <typename __T1, typename __T2>                                      \
 467   __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<                           \
 468       __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value,  \
 469       typename __hip::__promote<__T1, __T2>::type>::type                       \
 470   __fn(__T1 __x, __T2 __y) {                                                   \
 471     typedef typename __hip::__promote<__T1, __T2>::type __result_type;         \
 472     return __fn((__result_type)__x, (__result_type)__y);                       \
 473   }
 474 #else
 475 #define __HIP_OVERLOAD2(__retty, __fn)                                         \
 476   template <typename __T1, typename __T2>                                      \
 477   __DEVICE__ __CONSTEXPR__                                                     \
 478       typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&            \
 479                                    __hip::is_arithmetic<__T2>::value,          \
 480                                __retty>::type                                  \
 481       __fn(__T1 __x, __T2 __y) {                                               \
 482     return __fn((double)__x, (double)__y);                                     \
 483   }
 484 #endif
 485
 486 __HIP_OVERLOAD1(double, acos)
 487 __HIP_OVERLOAD1(double, acosh)
 488 __HIP_OVERLOAD1(double, asin)
 489 __HIP_OVERLOAD1(double, asinh)
 490 __HIP_OVERLOAD1(double, atan)
 491 __HIP_OVERLOAD2(double, atan2)
 492 __HIP_OVERLOAD1(double, atanh)
 493 __HIP_OVERLOAD1(double, cbrt)
 494 __HIP_OVERLOAD1(double, ceil)
 495 __HIP_OVERLOAD2(double, copysign)
 496 __HIP_OVERLOAD1(double, cos)
 497 __HIP_OVERLOAD1(double, cosh)
 498 __HIP_OVERLOAD1(double, erf)
 499 __HIP_OVERLOAD1(double, erfc)
 500 __HIP_OVERLOAD1(double, exp)
 501 __HIP_OVERLOAD1(double, exp2)
 502 __HIP_OVERLOAD1(double, expm1)
 503 __HIP_OVERLOAD1(double, fabs)
 504 __HIP_OVERLOAD2(double, fdim)
 505 __HIP_OVERLOAD1(double, floor)
 506 __HIP_OVERLOAD2(double, fmax)
 507 __HIP_OVERLOAD2(double, fmin)
 508 __HIP_OVERLOAD2(double, fmod)
 509 #if !defined(__HIPCC_RTC__)
 510 __HIP_OVERLOAD1(int, fpclassify)
 511 #endif // !defined(__HIPCC_RTC__)
 512 __HIP_OVERLOAD2(double, hypot)
 513 __HIP_OVERLOAD1(int, ilogb)
 514 __HIP_OVERLOAD1(bool, isfinite)
 515 __HIP_OVERLOAD2(bool, isgreater)
 516 __HIP_OVERLOAD2(bool, isgreaterequal)
 517 __HIP_OVERLOAD1(bool, isinf)
 518 __HIP_OVERLOAD2(bool, isless)
 519 __HIP_OVERLOAD2(bool, islessequal)
 520 __HIP_OVERLOAD2(bool, islessgreater)
 521 __HIP_OVERLOAD1(bool, isnan)
 522 __HIP_OVERLOAD1(bool, isnormal)
 523 __HIP_OVERLOAD2(bool, isunordered)
 524 __HIP_OVERLOAD1(double, lgamma)
 525 __HIP_OVERLOAD1(double, log)
 526 __HIP_OVERLOAD1(double, log10)
 527 __HIP_OVERLOAD1(double, log1p)
 528 __HIP_OVERLOAD1(double, log2)
 529 __HIP_OVERLOAD1(double, logb)
 530 __HIP_OVERLOAD1(long long, llrint)
 531 __HIP_OVERLOAD1(long long, llround)
 532 __HIP_OVERLOAD1(long, lrint)
 533 __HIP_OVERLOAD1(long, lround)
 534 __HIP_OVERLOAD1(double, nearbyint)
 535 __HIP_OVERLOAD2(double, nextafter)
 536 __HIP_OVERLOAD2(double, pow)
 537 __HIP_OVERLOAD2(double, remainder)
 538 __HIP_OVERLOAD1(double, rint)
 539 __HIP_OVERLOAD1(double, round)
 540 __HIP_OVERLOAD1(bool, signbit)
 541 __HIP_OVERLOAD1(double, sin)
 542 __HIP_OVERLOAD1(double, sinh)
 543 __HIP_OVERLOAD1(double, sqrt)
 544 __HIP_OVERLOAD1(double, tan)
 545 __HIP_OVERLOAD1(double, tanh)
 546 __HIP_OVERLOAD1(double, tgamma)
 547 __HIP_OVERLOAD1(double, trunc)
 548
 549 // Overload these but don't add them to std, they are not part of cmath.
 550 __HIP_OVERLOAD2(double, max)
 551 __HIP_OVERLOAD2(double, min)
 552
 553 // Additional Overloads that don't quite match HIP_OVERLOAD.
 554 #if __cplusplus >= 201103L
 555 template <typename __T1, typename __T2, typename __T3>
 556 __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<
 557     __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
 558         __hip::is_arithmetic<__T3>::value,
 559     typename __hip::__promote<__T1, __T2, __T3>::type>::type
 560 fma(__T1 __x, __T2 __y, __T3 __z) {
 561   typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
 562   return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);
 563 }
 564 #else
 565 template <typename __T1, typename __T2, typename __T3>
 566 __DEVICE__ __CONSTEXPR__
 567     typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
 568                                  __hip::is_arithmetic<__T2>::value &&
 569                                  __hip::is_arithmetic<__T3>::value,
 570                              double>::type
 571     fma(__T1 __x, __T2 __y, __T3 __z) {
 572   return ::fma((double)__x, (double)__y, (double)__z);
 573 }
 574 #endif
 575
 576 template <typename __T>
 577 __DEVICE__ __CONSTEXPR__
 578     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
 579     frexp(__T __x, int *__exp) {
 580   return ::frexp((double)__x, __exp);
 581 }
 582
 583 template <typename __T>
 584 __DEVICE__ __CONSTEXPR__
 585     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
 586     ldexp(__T __x, int __exp) {
 587   return ::ldexp((double)__x, __exp);
 588 }
 589
 590 template <typename __T>
 591 __DEVICE__ __CONSTEXPR__
 592     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
 593     modf(__T __x, double *__exp) {
 594   return ::modf((double)__x, __exp);
 595 }
 596
 597 #if __cplusplus >= 201103L
 598 template <typename __T1, typename __T2>
 599 __DEVICE__ __CONSTEXPR__
 600     typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
 601                                  __hip::is_arithmetic<__T2>::value,
 602                              typename __hip::__promote<__T1, __T2>::type>::type
 603     remquo(__T1 __x, __T2 __y, int *__quo) {
 604   typedef typename __hip::__promote<__T1, __T2>::type __result_type;
 605   return ::remquo((__result_type)__x, (__result_type)__y, __quo);
 606 }
 607 #else
 608 template <typename __T1, typename __T2>
 609 __DEVICE__ __CONSTEXPR__
 610     typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
 611                                  __hip::is_arithmetic<__T2>::value,
 612                              double>::type
 613     remquo(__T1 __x, __T2 __y, int *__quo) {
 614   return ::remquo((double)__x, (double)__y, __quo);
 615 }
 616 #endif
 617
 618 template <typename __T>
 619 __DEVICE__ __CONSTEXPR__
 620     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
 621     scalbln(__T __x, long int __exp) {
 622   return ::scalbln((double)__x, __exp);
 623 }
 624
 625 template <typename __T>
 626 __DEVICE__ __CONSTEXPR__
 627     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
 628     scalbn(__T __x, int __exp) {
 629   return ::scalbn((double)__x, __exp);
 630 }
 631
 632 #pragma pop_macro("__HIP_OVERLOAD1")
 633 #pragma pop_macro("__HIP_OVERLOAD2")
 634
 635 // END HIP_OVERLOAD
 636
 637 // END DEF_FUN and HIP_OVERLOAD
 638
 639 #endif // ifndef __OPENMP_AMDGCN__
 640 #endif // defined(__cplusplus)
 641
 642 #ifndef __OPENMP_AMDGCN__
 643 // Define these overloads inside the namespace our standard library uses.
 644 #if !defined(__HIPCC_RTC__)
 645 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
 646 _LIBCPP_BEGIN_NAMESPACE_STD
 647 #else
 648 namespace std {
 649 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 650 _GLIBCXX_BEGIN_NAMESPACE_VERSION
 651 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
 652 #endif // _LIBCPP_BEGIN_NAMESPACE_STD
 653
 654 // Pull the new overloads we defined above into namespace std.
 655 // using ::abs; - This may be considered for C++.
 656 using ::acos;
 657 using ::acosh;
 658 using ::asin;
 659 using ::asinh;
 660 using ::atan;
 661 using ::atan2;
 662 using ::atanh;
 663 using ::cbrt;
 664 using ::ceil;
 665 using ::copysign;
 666 using ::cos;
 667 using ::cosh;
 668 using ::erf;
 669 using ::erfc;
 670 using ::exp;
 671 using ::exp2;
 672 using ::expm1;
 673 using ::fabs;
 674 using ::fdim;
 675 using ::floor;
 676 using ::fma;
 677 using ::fmax;
 678 using ::fmin;
 679 using ::fmod;
 680 using ::fpclassify;
 681 using ::frexp;
 682 using ::hypot;
 683 using ::ilogb;
 684 using ::isfinite;
 685 using ::isgreater;
 686 using ::isgreaterequal;
 687 using ::isless;
 688 using ::islessequal;
 689 using ::islessgreater;
 690 using ::isnormal;
 691 using ::isunordered;
 692 using ::ldexp;
 693 using ::lgamma;
 694 using ::llrint;
 695 using ::llround;
 696 using ::log;
 697 using ::log10;
 698 using ::log1p;
 699 using ::log2;
 700 using ::logb;
 701 using ::lrint;
 702 using ::lround;
 703 using ::modf;
 704 // using ::nan; - This may be considered for C++.
 705 // using ::nanf; - This may be considered for C++.
 706 // using ::nanl; - This is not yet defined.
 707 using ::nearbyint;
 708 using ::nextafter;
 709 // using ::nexttoward; - Omit this since we do not have a definition.
 710 using ::pow;
 711 using ::remainder;
 712 using ::remquo;
 713 using ::rint;
 714 using ::round;
 715 using ::scalbln;
 716 using ::scalbn;
 717 using ::signbit;
 718 using ::sin;
 719 using ::sinh;
 720 using ::sqrt;
 721 using ::tan;
 722 using ::tanh;
 723 using ::tgamma;
 724 using ::trunc;
 725
 726 // Well this is fun: We need to pull these symbols in for libc++, but we can't
 727 // pull them in with libstdc++, because its ::isinf and ::isnan are different
 728 // than its std::isinf and std::isnan.
 729 #ifndef __GLIBCXX__
 730 using ::isinf;
 731 using ::isnan;
 732 #endif
 733
 734 // Finally, pull the "foobarf" functions that HIP defines into std.
 735 using ::acosf;
 736 using ::acoshf;
 737 using ::asinf;
 738 using ::asinhf;
 739 using ::atan2f;
 740 using ::atanf;
 741 using ::atanhf;
 742 using ::cbrtf;
 743 using ::ceilf;
 744 using ::copysignf;
 745 using ::cosf;
 746 using ::coshf;
 747 using ::erfcf;
 748 using ::erff;
 749 using ::exp2f;
 750 using ::expf;
 751 using ::expm1f;
 752 using ::fabsf;
 753 using ::fdimf;
 754 using ::floorf;
 755 using ::fmaf;
 756 using ::fmaxf;
 757 using ::fminf;
 758 using ::fmodf;
 759 using ::frexpf;
 760 using ::hypotf;
 761 using ::ilogbf;
 762 using ::ldexpf;
 763 using ::lgammaf;
 764 using ::llrintf;
 765 using ::llroundf;
 766 using ::log10f;
 767 using ::log1pf;
 768 using ::log2f;
 769 using ::logbf;
 770 using ::logf;
 771 using ::lrintf;
 772 using ::lroundf;
 773 using ::modff;
 774 using ::nearbyintf;
 775 using ::nextafterf;
 776 // using ::nexttowardf; - Omit this since we do not have a definition.
 777 using ::powf;
 778 using ::remainderf;
 779 using ::remquof;
 780 using ::rintf;
 781 using ::roundf;
 782 using ::scalblnf;
 783 using ::scalbnf;
 784 using ::sinf;
 785 using ::sinhf;
 786 using ::sqrtf;
 787 using ::tanf;
 788 using ::tanhf;
 789 using ::tgammaf;
 790 using ::truncf;
 791
 792 #ifdef _LIBCPP_END_NAMESPACE_STD
 793 _LIBCPP_END_NAMESPACE_STD
 794 #else
 795 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 796 _GLIBCXX_END_NAMESPACE_VERSION
 797 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
 798 } // namespace std
 799 #endif // _LIBCPP_END_NAMESPACE_STD
 800 #endif // !defined(__HIPCC_RTC__)
 801
 802 // Define device-side math functions from <ymath.h> on MSVC.
 803 #if !defined(__HIPCC_RTC__)
 804 #if defined(_MSC_VER)
 805
 806 // Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers.
 807 // But, from VS2019, it's only included in `<complex>`. Need to include
 808 // `<ymath.h>` here to ensure C functions declared there won't be markded as
 809 // `__host__` and `__device__` through `<complex>` wrapper.
 810 #include <ymath.h>
 811
 812 #if defined(__cplusplus)
 813 extern "C" {
 814 #endif // defined(__cplusplus)
 815 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Cosh(double x,
 816                                                                     double y) {
 817   return cosh(x) * y;
 818 }
 819 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FCosh(float x,
 820                                                                     float y) {
 821   return coshf(x) * y;
 822 }
 823 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _Dtest(double *p) {
 824   return fpclassify(*p);
 825 }
 826 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _FDtest(float *p) {
 827   return fpclassify(*p);
 828 }
 829 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Sinh(double x,
 830                                                                     double y) {
 831   return sinh(x) * y;
 832 }
 833 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FSinh(float x,
 834                                                                     float y) {
 835   return sinhf(x) * y;
 836 }
 837 #if defined(__cplusplus)
 838 }
 839 #endif // defined(__cplusplus)
 840 #endif // defined(_MSC_VER)
 841 #endif // !defined(__HIPCC_RTC__)
 842 #endif // ifndef __OPENMP_AMDGCN__
 843
 844 #pragma pop_macro("__DEVICE__")
 845 #pragma pop_macro("__CONSTEXPR__")
 846
 847 #endif // __CLANG_HIP_CMATH_H__