Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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 | |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 13 | #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 14 | #error "This file is for HIP and OpenMP AMDGCN device compilation only." |
| 15 | #endif |
| 16 | |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 17 | #if !defined(__HIPCC_RTC__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 18 | #if defined(__cplusplus) |
| 19 | #include <limits> |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 20 | #include <type_traits> |
| 21 | #include <utility> |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 22 | #endif |
| 23 | #include <limits.h> |
| 24 | #include <stdint.h> |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 25 | #endif // !defined(__HIPCC_RTC__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 26 | |
| 27 | #pragma push_macro("__DEVICE__") |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 28 | #pragma push_macro("__CONSTEXPR__") |
| 29 | #ifdef __OPENMP_AMDGCN__ |
| 30 | #define __DEVICE__ static __attribute__((always_inline, nothrow)) |
| 31 | #define __CONSTEXPR__ constexpr |
| 32 | #else |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 33 | #define __DEVICE__ static __device__ inline __attribute__((always_inline)) |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 34 | #define __CONSTEXPR__ |
| 35 | #endif // __OPENMP_AMDGCN__ |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 36 | |
| 37 | // Start with functions that cannot be defined by DEF macros below. |
| 38 | #if defined(__cplusplus) |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 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) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 49 | return ::fmaf(__x, __y, __z); |
| 50 | } |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 51 | #if !defined(__HIPCC_RTC__) |
| 52 | // The value returned by fpclassify is platform dependent, therefore it is not |
| 53 | // supported by hipRTC. |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 54 | __DEVICE__ __CONSTEXPR__ int fpclassify(float __x) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 55 | return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, |
| 56 | FP_ZERO, __x); |
| 57 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 58 | __DEVICE__ __CONSTEXPR__ int fpclassify(double __x) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 59 | return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, |
| 60 | FP_ZERO, __x); |
| 61 | } |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 62 | #endif // !defined(__HIPCC_RTC__) |
| 63 | |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 64 | __DEVICE__ __CONSTEXPR__ float frexp(float __arg, int *__exp) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 65 | return ::frexpf(__arg, __exp); |
| 66 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 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) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 108 | return __builtin_isgreater(__x, __y); |
| 109 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 110 | __DEVICE__ __CONSTEXPR__ bool isgreater(double __x, double __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 111 | return __builtin_isgreater(__x, __y); |
| 112 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 113 | __DEVICE__ __CONSTEXPR__ bool isgreaterequal(float __x, float __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 114 | return __builtin_isgreaterequal(__x, __y); |
| 115 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 116 | __DEVICE__ __CONSTEXPR__ bool isgreaterequal(double __x, double __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 117 | return __builtin_isgreaterequal(__x, __y); |
| 118 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 119 | __DEVICE__ __CONSTEXPR__ bool isless(float __x, float __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 120 | return __builtin_isless(__x, __y); |
| 121 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 122 | __DEVICE__ __CONSTEXPR__ bool isless(double __x, double __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 123 | return __builtin_isless(__x, __y); |
| 124 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 125 | __DEVICE__ __CONSTEXPR__ bool islessequal(float __x, float __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 126 | return __builtin_islessequal(__x, __y); |
| 127 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 128 | __DEVICE__ __CONSTEXPR__ bool islessequal(double __x, double __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 129 | return __builtin_islessequal(__x, __y); |
| 130 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 131 | __DEVICE__ __CONSTEXPR__ bool islessgreater(float __x, float __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 132 | return __builtin_islessgreater(__x, __y); |
| 133 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 134 | __DEVICE__ __CONSTEXPR__ bool islessgreater(double __x, double __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 135 | return __builtin_islessgreater(__x, __y); |
| 136 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 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) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 144 | return __builtin_isunordered(__x, __y); |
| 145 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 146 | __DEVICE__ __CONSTEXPR__ bool isunordered(double __x, double __y) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 147 | return __builtin_isunordered(__x, __y); |
| 148 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 149 | __DEVICE__ __CONSTEXPR__ float modf(float __x, float *__iptr) { |
| 150 | return ::modff(__x, __iptr); |
| 151 | } |
| 152 | __DEVICE__ __CONSTEXPR__ float pow(float __base, int __iexp) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 153 | return ::powif(__base, __iexp); |
| 154 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 155 | __DEVICE__ __CONSTEXPR__ double pow(double __base, int __iexp) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 156 | return ::powi(__base, __iexp); |
| 157 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 158 | __DEVICE__ __CONSTEXPR__ float remquo(float __x, float __y, int *__quo) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 159 | return ::remquof(__x, __y, __quo); |
| 160 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 161 | __DEVICE__ __CONSTEXPR__ float scalbln(float __x, long int __n) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 162 | return ::scalblnf(__x, __n); |
| 163 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 164 | __DEVICE__ __CONSTEXPR__ bool signbit(float __x) { return ::__signbitf(__x); } |
| 165 | __DEVICE__ __CONSTEXPR__ bool signbit(double __x) { return ::__signbit(__x); } |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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. |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 172 | __DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y, |
| 173 | _Float16 __z) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 174 | return __ocml_fma_f16(__x, __y, __z); |
| 175 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 176 | __DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 177 | return __ocml_pown_f16(__base, __iexp); |
| 178 | } |
| 179 | |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 180 | #ifndef __OPENMP_AMDGCN__ |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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) \ |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 191 | __DEVICE__ __CONSTEXPR__ __retty __func(float __x) { return __func##f(__x); } |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 192 | |
| 193 | // Define cmath functions with two float arguments and returns __retty. |
| 194 | #define __DEF_FUN2(__retty, __func) \ |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 195 | __DEVICE__ __CONSTEXPR__ __retty __func(float __x, float __y) { \ |
| 196 | return __func##f(__x, __y); \ |
| 197 | } |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 198 | |
| 199 | // Define cmath functions with a float and an int argument and returns __retty. |
| 200 | #define __DEF_FUN2_FI(__retty, __func) \ |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 201 | __DEVICE__ __CONSTEXPR__ __retty __func(float __x, int __y) { \ |
| 202 | return __func##f(__x, __y); \ |
| 203 | } |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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 | |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 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 | |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 380 | // decltype is only available in C++11 and above. |
| 381 | #if __cplusplus >= 201103L |
| 382 | // __hip_promote |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 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 | |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 398 | typedef decltype(__test(declval<_Tp>())) type; |
| 399 | static const bool value = !is_same<type, void>::value; |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 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> {}; |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 442 | #endif //__cplusplus >= 201103L |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 443 | } // namespace __hip |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 444 | |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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> \ |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 450 | __DEVICE__ __CONSTEXPR__ \ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 451 | typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type \ |
| 452 | __fn(__T __x) { \ |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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). |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 459 | #if __cplusplus >= 201103L |
| 460 | #define __HIP_OVERLOAD2(__retty, __fn) \ |
| 461 | template <typename __T1, typename __T2> \ |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 462 | __DEVICE__ __CONSTEXPR__ typename __hip_enable_if< \ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 463 | __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, \ |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 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 |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 470 | #define __HIP_OVERLOAD2(__retty, __fn) \ |
| 471 | template <typename __T1, typename __T2> \ |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 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) { \ |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 477 | return __fn((double)__x, (double)__y); \ |
| 478 | } |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 479 | #endif |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 480 | |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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) |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 504 | #if !defined(__HIPCC_RTC__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 505 | __HIP_OVERLOAD1(int, fpclassify) |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 506 | #endif // !defined(__HIPCC_RTC__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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. |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 549 | #if __cplusplus >= 201103L |
| 550 | template <typename __T1, typename __T2, typename __T3> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 551 | __DEVICE__ __CONSTEXPR__ typename __hip_enable_if< |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 552 | __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value && |
| 553 | __hip::is_arithmetic<__T3>::value, |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 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 |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 560 | template <typename __T1, typename __T2, typename __T3> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 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) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 567 | return ::fma((double)__x, (double)__y, (double)__z); |
| 568 | } |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 569 | #endif |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 570 | |
| 571 | template <typename __T> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 572 | __DEVICE__ __CONSTEXPR__ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 573 | typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 574 | frexp(__T __x, int *__exp) { |
| 575 | return ::frexp((double)__x, __exp); |
| 576 | } |
| 577 | |
| 578 | template <typename __T> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 579 | __DEVICE__ __CONSTEXPR__ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 580 | typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 581 | ldexp(__T __x, int __exp) { |
| 582 | return ::ldexp((double)__x, __exp); |
| 583 | } |
| 584 | |
| 585 | template <typename __T> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 586 | __DEVICE__ __CONSTEXPR__ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 587 | typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 588 | modf(__T __x, double *__exp) { |
| 589 | return ::modf((double)__x, __exp); |
| 590 | } |
| 591 | |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 592 | #if __cplusplus >= 201103L |
| 593 | template <typename __T1, typename __T2> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 594 | __DEVICE__ __CONSTEXPR__ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 595 | typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && |
| 596 | __hip::is_arithmetic<__T2>::value, |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 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 |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 603 | template <typename __T1, typename __T2> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 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) { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 609 | return ::remquo((double)__x, (double)__y, __quo); |
| 610 | } |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 611 | #endif |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 612 | |
| 613 | template <typename __T> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 614 | __DEVICE__ __CONSTEXPR__ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 615 | typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 616 | scalbln(__T __x, long int __exp) { |
| 617 | return ::scalbln((double)__x, __exp); |
| 618 | } |
| 619 | |
| 620 | template <typename __T> |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 621 | __DEVICE__ __CONSTEXPR__ |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 622 | typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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 | |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 634 | #endif // ifndef __OPENMP_AMDGCN__ |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 635 | #endif // defined(__cplusplus) |
| 636 | |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 637 | #ifndef __OPENMP_AMDGCN__ |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 638 | // Define these overloads inside the namespace our standard library uses. |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 639 | #if !defined(__HIPCC_RTC__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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 |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 646 | #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION |
| 647 | #endif // _LIBCPP_BEGIN_NAMESPACE_STD |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 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 |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 792 | #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 793 | } // namespace std |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 794 | #endif // _LIBCPP_END_NAMESPACE_STD |
| 795 | #endif // !defined(__HIPCC_RTC__) |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 796 | |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 797 | // Define device-side math functions from <ymath.h> on MSVC. |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 798 | #if !defined(__HIPCC_RTC__) |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 799 | #if defined(_MSC_VER) |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 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 | |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 807 | #if defined(__cplusplus) |
| 808 | extern "C" { |
| 809 | #endif // defined(__cplusplus) |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 810 | __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Cosh(double x, |
| 811 | double y) { |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 812 | return cosh(x) * y; |
| 813 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 814 | __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FCosh(float x, |
| 815 | float y) { |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 816 | return coshf(x) * y; |
| 817 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 818 | __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _Dtest(double *p) { |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 819 | return fpclassify(*p); |
| 820 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 821 | __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _FDtest(float *p) { |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 822 | return fpclassify(*p); |
| 823 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 824 | __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Sinh(double x, |
| 825 | double y) { |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 826 | return sinh(x) * y; |
| 827 | } |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 828 | __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FSinh(float x, |
| 829 | float y) { |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 830 | return sinhf(x) * y; |
| 831 | } |
| 832 | #if defined(__cplusplus) |
| 833 | } |
| 834 | #endif // defined(__cplusplus) |
| 835 | #endif // defined(_MSC_VER) |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 836 | #endif // !defined(__HIPCC_RTC__) |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 837 | #endif // ifndef __OPENMP_AMDGCN__ |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 838 | |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 839 | #pragma pop_macro("__DEVICE__") |
Pirama Arumuga Nainar | 494f645 | 2021-12-02 10:42:14 -0800 | [diff] [blame] | 840 | #pragma pop_macro("__CONSTEXPR__") |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 841 | |
| 842 | #endif // __CLANG_HIP_CMATH_H__ |