Subversion Repositories QNX 8.QNX8 LLVM/Clang compiler suite

Rev

Blame | Last modification | View Log | Download | RSS feed

  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 __ocml_fma_f16(__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.   typedef decltype(__test(declval<_Tp>())) type;
  399.   static const bool value = !is_same<type, void>::value;
  400. };
  401.  
  402. template <> struct __numeric_type<void> { static const bool value = true; };
  403.  
  404. template <class _A1, class _A2 = void, class _A3 = void,
  405.           bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
  406.               &&__numeric_type<_A3>::value>
  407. class __promote_imp {
  408. public:
  409.   static const bool value = false;
  410. };
  411.  
  412. template <class _A1, class _A2, class _A3>
  413. class __promote_imp<_A1, _A2, _A3, true> {
  414. private:
  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.  
  419. public:
  420.   typedef decltype(__type1() + __type2() + __type3()) type;
  421.   static const bool value = true;
  422. };
  423.  
  424. template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
  425. private:
  426.   typedef typename __promote_imp<_A1>::type __type1;
  427.   typedef typename __promote_imp<_A2>::type __type2;
  428.  
  429. public:
  430.   typedef decltype(__type1() + __type2()) type;
  431.   static const bool value = true;
  432. };
  433.  
  434. template <class _A1> class __promote_imp<_A1, void, void, true> {
  435. public:
  436.   typedef typename __numeric_type<_A1>::type type;
  437.   static const bool value = true;
  438. };
  439.  
  440. template <class _A1, class _A2 = void, class _A3 = void>
  441. class __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
  550. template <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
  555. fma(__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
  560. template <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.  
  571. template <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.  
  578. template <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.  
  585. template <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
  593. template <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
  603. template <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.  
  613. template <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.  
  620. template <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
  643. namespace 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++.
  651. using ::acos;
  652. using ::acosh;
  653. using ::asin;
  654. using ::asinh;
  655. using ::atan;
  656. using ::atan2;
  657. using ::atanh;
  658. using ::cbrt;
  659. using ::ceil;
  660. using ::copysign;
  661. using ::cos;
  662. using ::cosh;
  663. using ::erf;
  664. using ::erfc;
  665. using ::exp;
  666. using ::exp2;
  667. using ::expm1;
  668. using ::fabs;
  669. using ::fdim;
  670. using ::floor;
  671. using ::fma;
  672. using ::fmax;
  673. using ::fmin;
  674. using ::fmod;
  675. using ::fpclassify;
  676. using ::frexp;
  677. using ::hypot;
  678. using ::ilogb;
  679. using ::isfinite;
  680. using ::isgreater;
  681. using ::isgreaterequal;
  682. using ::isless;
  683. using ::islessequal;
  684. using ::islessgreater;
  685. using ::isnormal;
  686. using ::isunordered;
  687. using ::ldexp;
  688. using ::lgamma;
  689. using ::llrint;
  690. using ::llround;
  691. using ::log;
  692. using ::log10;
  693. using ::log1p;
  694. using ::log2;
  695. using ::logb;
  696. using ::lrint;
  697. using ::lround;
  698. using ::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.
  702. using ::nearbyint;
  703. using ::nextafter;
  704. // using ::nexttoward; - Omit this since we do not have a definition.
  705. using ::pow;
  706. using ::remainder;
  707. using ::remquo;
  708. using ::rint;
  709. using ::round;
  710. using ::scalbln;
  711. using ::scalbn;
  712. using ::signbit;
  713. using ::sin;
  714. using ::sinh;
  715. using ::sqrt;
  716. using ::tan;
  717. using ::tanh;
  718. using ::tgamma;
  719. using ::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__
  725. using ::isinf;
  726. using ::isnan;
  727. #endif
  728.  
  729. // Finally, pull the "foobarf" functions that HIP defines into std.
  730. using ::acosf;
  731. using ::acoshf;
  732. using ::asinf;
  733. using ::asinhf;
  734. using ::atan2f;
  735. using ::atanf;
  736. using ::atanhf;
  737. using ::cbrtf;
  738. using ::ceilf;
  739. using ::copysignf;
  740. using ::cosf;
  741. using ::coshf;
  742. using ::erfcf;
  743. using ::erff;
  744. using ::exp2f;
  745. using ::expf;
  746. using ::expm1f;
  747. using ::fabsf;
  748. using ::fdimf;
  749. using ::floorf;
  750. using ::fmaf;
  751. using ::fmaxf;
  752. using ::fminf;
  753. using ::fmodf;
  754. using ::frexpf;
  755. using ::hypotf;
  756. using ::ilogbf;
  757. using ::ldexpf;
  758. using ::lgammaf;
  759. using ::llrintf;
  760. using ::llroundf;
  761. using ::log10f;
  762. using ::log1pf;
  763. using ::log2f;
  764. using ::logbf;
  765. using ::logf;
  766. using ::lrintf;
  767. using ::lroundf;
  768. using ::modff;
  769. using ::nearbyintf;
  770. using ::nextafterf;
  771. // using ::nexttowardf; - Omit this since we do not have a definition.
  772. using ::powf;
  773. using ::remainderf;
  774. using ::remquof;
  775. using ::rintf;
  776. using ::roundf;
  777. using ::scalblnf;
  778. using ::scalbnf;
  779. using ::sinf;
  780. using ::sinhf;
  781. using ::sqrtf;
  782. using ::tanf;
  783. using ::tanhf;
  784. using ::tgammaf;
  785. using ::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)
  808. extern "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.