849 lines · c
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-exception6 *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#endif16 17#if !defined(__HIPCC_RTC__)18#if defined(__cplusplus)19#include <limits>20#include <type_traits>21#include <utility>22#endif23#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__ constexpr32#else33#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#endif44__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 not53// 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-conforming70// `isinf(float)` and `isnan(float)` implementations that return an `int`. We do71// this by providing two versions of these functions, differing only in the72// return type. To avoid conflicting definitions we disable implicit base73// function generation. That means we will end up with two specializations, one74// 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 variants80// (note that we do not create implicit base functions here). To avoid81// this clash we add a new trait to some of them that is always true82// (this is LLVM after all ;)). It will only influence the mangled name83// 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 variant94#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 variant105#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 because168// ocml doesn't provide an implementation, and we don't want to be in the169// 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_OVERLOAD182 183// BEGIN DEF_FUN184 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_FUN261 262// BEGIN HIP_OVERLOAD263 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 >= 201103L382// __hip_promote383template <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 407template <> struct __numeric_type<void> { static const bool value = true; };408 409template <class _A1, class _A2 = void, class _A3 = void,410 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value411 &&__numeric_type<_A3>::value>412class __promote_imp {413public:414 static const bool value = false;415};416 417template <class _A1, class _A2, class _A3>418class __promote_imp<_A1, _A2, _A3, true> {419private: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 424public:425 typedef decltype(__type1() + __type2() + __type3()) type;426 static const bool value = true;427};428 429template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {430private:431 typedef typename __promote_imp<_A1>::type __type1;432 typedef typename __promote_imp<_A2>::type __type2;433 434public:435 typedef decltype(__type1() + __type2()) type;436 static const bool value = true;437};438 439template <class _A1> class __promote_imp<_A1, void, void, true> {440public:441 typedef typename __numeric_type<_A1>::type type;442 static const bool value = true;443};444 445template <class _A1, class _A2 = void, class _A3 = void>446class __promote : public __promote_imp<_A1, _A2, _A3> {};447#endif //__cplusplus >= 201103L448} // namespace __hip449 450// __HIP_OVERLOAD1 is used to resolve function calls with integer argument to451// avoid compilation error due to ambiguity. e.g. floor(5) is resolved with452// 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/double462// or integer argument to avoid compilation error due to ambiguity. e.g.463// max(5.0f, 6.0) is resolved with max(double, double).464#if __cplusplus >= 201103L465#define __HIP_OVERLOAD2(__retty, __fn) \466 template <typename __T1, typename __T2> \467 __DEVICE__ __CONSTEXPR__ \468 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \469 __hip::is_arithmetic<__T2>::value, \470 __retty>::type \471 __fn(__T1 __x, __T2 __y) { \472 typedef typename __hip::__promote<__T1, __T2>::type __arg_type; \473 return __fn((__arg_type)__x, (__arg_type)__y); \474 }475#else476#define __HIP_OVERLOAD2(__retty, __fn) \477 template <typename __T1, typename __T2> \478 __DEVICE__ __CONSTEXPR__ \479 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \480 __hip::is_arithmetic<__T2>::value, \481 __retty>::type \482 __fn(__T1 __x, __T2 __y) { \483 return __fn((double)__x, (double)__y); \484 }485#endif486 487__HIP_OVERLOAD1(double, acos)488__HIP_OVERLOAD1(double, acosh)489__HIP_OVERLOAD1(double, asin)490__HIP_OVERLOAD1(double, asinh)491__HIP_OVERLOAD1(double, atan)492__HIP_OVERLOAD2(double, atan2)493__HIP_OVERLOAD1(double, atanh)494__HIP_OVERLOAD1(double, cbrt)495__HIP_OVERLOAD1(double, ceil)496__HIP_OVERLOAD2(double, copysign)497__HIP_OVERLOAD1(double, cos)498__HIP_OVERLOAD1(double, cosh)499__HIP_OVERLOAD1(double, erf)500__HIP_OVERLOAD1(double, erfc)501__HIP_OVERLOAD1(double, exp)502__HIP_OVERLOAD1(double, exp2)503__HIP_OVERLOAD1(double, expm1)504__HIP_OVERLOAD1(double, fabs)505__HIP_OVERLOAD2(double, fdim)506__HIP_OVERLOAD1(double, floor)507__HIP_OVERLOAD2(double, fmax)508__HIP_OVERLOAD2(double, fmin)509__HIP_OVERLOAD2(double, fmod)510#if !defined(__HIPCC_RTC__)511__HIP_OVERLOAD1(int, fpclassify)512#endif // !defined(__HIPCC_RTC__)513__HIP_OVERLOAD2(double, hypot)514__HIP_OVERLOAD1(int, ilogb)515__HIP_OVERLOAD1(bool, isfinite)516__HIP_OVERLOAD2(bool, isgreater)517__HIP_OVERLOAD2(bool, isgreaterequal)518__HIP_OVERLOAD1(bool, isinf)519__HIP_OVERLOAD2(bool, isless)520__HIP_OVERLOAD2(bool, islessequal)521__HIP_OVERLOAD2(bool, islessgreater)522__HIP_OVERLOAD1(bool, isnan)523__HIP_OVERLOAD1(bool, isnormal)524__HIP_OVERLOAD2(bool, isunordered)525__HIP_OVERLOAD1(double, lgamma)526__HIP_OVERLOAD1(double, log)527__HIP_OVERLOAD1(double, log10)528__HIP_OVERLOAD1(double, log1p)529__HIP_OVERLOAD1(double, log2)530__HIP_OVERLOAD1(double, logb)531__HIP_OVERLOAD1(long long, llrint)532__HIP_OVERLOAD1(long long, llround)533__HIP_OVERLOAD1(long, lrint)534__HIP_OVERLOAD1(long, lround)535__HIP_OVERLOAD1(double, nearbyint)536__HIP_OVERLOAD2(double, nextafter)537__HIP_OVERLOAD2(double, pow)538__HIP_OVERLOAD2(double, remainder)539__HIP_OVERLOAD1(double, rint)540__HIP_OVERLOAD1(double, round)541__HIP_OVERLOAD1(bool, signbit)542__HIP_OVERLOAD1(double, sin)543__HIP_OVERLOAD1(double, sinh)544__HIP_OVERLOAD1(double, sqrt)545__HIP_OVERLOAD1(double, tan)546__HIP_OVERLOAD1(double, tanh)547__HIP_OVERLOAD1(double, tgamma)548__HIP_OVERLOAD1(double, trunc)549 550// Overload these but don't add them to std, they are not part of cmath.551__HIP_OVERLOAD2(double, max)552__HIP_OVERLOAD2(double, min)553 554// Additional Overloads that don't quite match HIP_OVERLOAD.555#if __cplusplus >= 201103L556template <typename __T1, typename __T2, typename __T3>557__DEVICE__ __CONSTEXPR__ typename __hip_enable_if<558 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&559 __hip::is_arithmetic<__T3>::value,560 typename __hip::__promote<__T1, __T2, __T3>::type>::type561fma(__T1 __x, __T2 __y, __T3 __z) {562 typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;563 return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);564}565#else566template <typename __T1, typename __T2, typename __T3>567__DEVICE__ __CONSTEXPR__568 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&569 __hip::is_arithmetic<__T2>::value &&570 __hip::is_arithmetic<__T3>::value,571 double>::type572 fma(__T1 __x, __T2 __y, __T3 __z) {573 return ::fma((double)__x, (double)__y, (double)__z);574}575#endif576 577template <typename __T>578__DEVICE__ __CONSTEXPR__579 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type580 frexp(__T __x, int *__exp) {581 return ::frexp((double)__x, __exp);582}583 584template <typename __T>585__DEVICE__ __CONSTEXPR__586 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type587 ldexp(__T __x, int __exp) {588 return ::ldexp((double)__x, __exp);589}590 591template <typename __T>592__DEVICE__ __CONSTEXPR__593 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type594 modf(__T __x, double *__exp) {595 return ::modf((double)__x, __exp);596}597 598#if __cplusplus >= 201103L599template <typename __T1, typename __T2>600__DEVICE__ __CONSTEXPR__601 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&602 __hip::is_arithmetic<__T2>::value,603 typename __hip::__promote<__T1, __T2>::type>::type604 remquo(__T1 __x, __T2 __y, int *__quo) {605 typedef typename __hip::__promote<__T1, __T2>::type __result_type;606 return ::remquo((__result_type)__x, (__result_type)__y, __quo);607}608#else609template <typename __T1, typename __T2>610__DEVICE__ __CONSTEXPR__611 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&612 __hip::is_arithmetic<__T2>::value,613 double>::type614 remquo(__T1 __x, __T2 __y, int *__quo) {615 return ::remquo((double)__x, (double)__y, __quo);616}617#endif618 619template <typename __T>620__DEVICE__ __CONSTEXPR__621 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type622 scalbln(__T __x, long int __exp) {623 return ::scalbln((double)__x, __exp);624}625 626template <typename __T>627__DEVICE__ __CONSTEXPR__628 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type629 scalbn(__T __x, int __exp) {630 return ::scalbn((double)__x, __exp);631}632 633#pragma pop_macro("__HIP_OVERLOAD1")634#pragma pop_macro("__HIP_OVERLOAD2")635 636// END HIP_OVERLOAD637 638// END DEF_FUN and HIP_OVERLOAD639 640#endif // ifndef __OPENMP_AMDGCN__641#endif // defined(__cplusplus)642 643#ifndef __OPENMP_AMDGCN__644// Define these overloads inside the namespace our standard library uses.645#if !defined(__HIPCC_RTC__)646#ifdef _LIBCPP_BEGIN_NAMESPACE_STD647_LIBCPP_BEGIN_NAMESPACE_STD648#else649namespace std {650#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION651_GLIBCXX_BEGIN_NAMESPACE_VERSION652#endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION653#endif // _LIBCPP_BEGIN_NAMESPACE_STD654 655// Pull the new overloads we defined above into namespace std.656// using ::abs; - This may be considered for C++.657using ::acos;658using ::acosh;659using ::asin;660using ::asinh;661using ::atan;662using ::atan2;663using ::atanh;664using ::cbrt;665using ::ceil;666using ::copysign;667using ::cos;668using ::cosh;669using ::erf;670using ::erfc;671using ::exp;672using ::exp2;673using ::expm1;674using ::fabs;675using ::fdim;676using ::floor;677using ::fma;678using ::fmax;679using ::fmin;680using ::fmod;681using ::fpclassify;682using ::frexp;683using ::hypot;684using ::ilogb;685using ::isfinite;686using ::isgreater;687using ::isgreaterequal;688using ::isless;689using ::islessequal;690using ::islessgreater;691using ::isnormal;692using ::isunordered;693using ::ldexp;694using ::lgamma;695using ::llrint;696using ::llround;697using ::log;698using ::log10;699using ::log1p;700using ::log2;701using ::logb;702using ::lrint;703using ::lround;704using ::modf;705// using ::nan; - This may be considered for C++.706// using ::nanf; - This may be considered for C++.707// using ::nanl; - This is not yet defined.708using ::nearbyint;709using ::nextafter;710// using ::nexttoward; - Omit this since we do not have a definition.711using ::pow;712using ::remainder;713using ::remquo;714using ::rint;715using ::round;716using ::scalbln;717using ::scalbn;718using ::signbit;719using ::sin;720using ::sinh;721using ::sqrt;722using ::tan;723using ::tanh;724using ::tgamma;725using ::trunc;726 727// Well this is fun: We need to pull these symbols in for libc++, but we can't728// pull them in with libstdc++, because its ::isinf and ::isnan are different729// than its std::isinf and std::isnan.730#ifndef __GLIBCXX__731using ::isinf;732using ::isnan;733#endif734 735// Finally, pull the "foobarf" functions that HIP defines into std.736using ::acosf;737using ::acoshf;738using ::asinf;739using ::asinhf;740using ::atan2f;741using ::atanf;742using ::atanhf;743using ::cbrtf;744using ::ceilf;745using ::copysignf;746using ::cosf;747using ::coshf;748using ::erfcf;749using ::erff;750using ::exp2f;751using ::expf;752using ::expm1f;753using ::fabsf;754using ::fdimf;755using ::floorf;756using ::fmaf;757using ::fmaxf;758using ::fminf;759using ::fmodf;760using ::frexpf;761using ::hypotf;762using ::ilogbf;763using ::ldexpf;764using ::lgammaf;765using ::llrintf;766using ::llroundf;767using ::log10f;768using ::log1pf;769using ::log2f;770using ::logbf;771using ::logf;772using ::lrintf;773using ::lroundf;774using ::modff;775using ::nearbyintf;776using ::nextafterf;777// using ::nexttowardf; - Omit this since we do not have a definition.778using ::powf;779using ::remainderf;780using ::remquof;781using ::rintf;782using ::roundf;783using ::scalblnf;784using ::scalbnf;785using ::sinf;786using ::sinhf;787using ::sqrtf;788using ::tanf;789using ::tanhf;790using ::tgammaf;791using ::truncf;792 793#ifdef _LIBCPP_END_NAMESPACE_STD794_LIBCPP_END_NAMESPACE_STD795#else796#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION797_GLIBCXX_END_NAMESPACE_VERSION798#endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION799} // namespace std800#endif // _LIBCPP_END_NAMESPACE_STD801#endif // !defined(__HIPCC_RTC__)802 803// Define device-side math functions from <ymath.h> on MSVC.804#if !defined(__HIPCC_RTC__)805#if defined(_MSC_VER)806 807// Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers.808// But, from VS2019, it's only included in `<complex>`. Need to include809// `<ymath.h>` here to ensure C functions declared there won't be markded as810// `__host__` and `__device__` through `<complex>` wrapper.811#include <ymath.h>812 813#if defined(__cplusplus)814extern "C" {815#endif // defined(__cplusplus)816__DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Cosh(double x,817 double y) {818 return cosh(x) * y;819}820__DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FCosh(float x,821 float y) {822 return coshf(x) * y;823}824__DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _Dtest(double *p) {825 return fpclassify(*p);826}827__DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _FDtest(float *p) {828 return fpclassify(*p);829}830__DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Sinh(double x,831 double y) {832 return sinh(x) * y;833}834__DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FSinh(float x,835 float y) {836 return sinhf(x) * y;837}838#if defined(__cplusplus)839}840#endif // defined(__cplusplus)841#endif // defined(_MSC_VER)842#endif // !defined(__HIPCC_RTC__)843#endif // ifndef __OPENMP_AMDGCN__844 845#pragma pop_macro("__DEVICE__")846#pragma pop_macro("__CONSTEXPR__")847 848#endif // __CLANG_HIP_CMATH_H__849