Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 1 | /*===- __clang_math_forward_declares.h - Prototypes of __device__ math fns --=== |
| 2 | * |
Logan Chien | df4f766 | 2019-09-04 16:45:23 -0700 | [diff] [blame] | 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 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 6 | * |
| 7 | *===-----------------------------------------------------------------------=== |
| 8 | */ |
| 9 | #ifndef __CLANG__CUDA_MATH_FORWARD_DECLARES_H__ |
| 10 | #define __CLANG__CUDA_MATH_FORWARD_DECLARES_H__ |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 11 | #if !defined(__CUDA__) && !__HIP__ |
| 12 | #error "This file is for CUDA/HIP compilation only." |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 13 | #endif |
| 14 | |
| 15 | // This file forward-declares of some math functions we (or the CUDA headers) |
| 16 | // will define later. We need to do this, and do it before cmath is included, |
| 17 | // because the standard library may have constexpr math functions. In the |
| 18 | // absence of a prior __device__ decl, those constexpr functions may become |
| 19 | // implicitly host+device. host+device functions can't be overloaded, so that |
| 20 | // would preclude the use of our own __device__ overloads for these functions. |
| 21 | |
| 22 | #pragma push_macro("__DEVICE__") |
| 23 | #define __DEVICE__ \ |
| 24 | static __inline__ __attribute__((always_inline)) __attribute__((device)) |
| 25 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 26 | __DEVICE__ long abs(long); |
| 27 | __DEVICE__ long long abs(long long); |
Logan Chien | df4f766 | 2019-09-04 16:45:23 -0700 | [diff] [blame] | 28 | __DEVICE__ double abs(double); |
| 29 | __DEVICE__ float abs(float); |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 30 | __DEVICE__ int abs(int); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 31 | __DEVICE__ double acos(double); |
| 32 | __DEVICE__ float acos(float); |
| 33 | __DEVICE__ double acosh(double); |
| 34 | __DEVICE__ float acosh(float); |
| 35 | __DEVICE__ double asin(double); |
| 36 | __DEVICE__ float asin(float); |
| 37 | __DEVICE__ double asinh(double); |
| 38 | __DEVICE__ float asinh(float); |
| 39 | __DEVICE__ double atan2(double, double); |
| 40 | __DEVICE__ float atan2(float, float); |
| 41 | __DEVICE__ double atan(double); |
| 42 | __DEVICE__ float atan(float); |
| 43 | __DEVICE__ double atanh(double); |
| 44 | __DEVICE__ float atanh(float); |
| 45 | __DEVICE__ double cbrt(double); |
| 46 | __DEVICE__ float cbrt(float); |
| 47 | __DEVICE__ double ceil(double); |
| 48 | __DEVICE__ float ceil(float); |
| 49 | __DEVICE__ double copysign(double, double); |
| 50 | __DEVICE__ float copysign(float, float); |
| 51 | __DEVICE__ double cos(double); |
| 52 | __DEVICE__ float cos(float); |
| 53 | __DEVICE__ double cosh(double); |
| 54 | __DEVICE__ float cosh(float); |
| 55 | __DEVICE__ double erfc(double); |
| 56 | __DEVICE__ float erfc(float); |
| 57 | __DEVICE__ double erf(double); |
| 58 | __DEVICE__ float erf(float); |
| 59 | __DEVICE__ double exp2(double); |
| 60 | __DEVICE__ float exp2(float); |
| 61 | __DEVICE__ double exp(double); |
| 62 | __DEVICE__ float exp(float); |
| 63 | __DEVICE__ double expm1(double); |
| 64 | __DEVICE__ float expm1(float); |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 65 | __DEVICE__ double fabs(double); |
| 66 | __DEVICE__ float fabs(float); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 67 | __DEVICE__ double fdim(double, double); |
| 68 | __DEVICE__ float fdim(float, float); |
| 69 | __DEVICE__ double floor(double); |
| 70 | __DEVICE__ float floor(float); |
| 71 | __DEVICE__ double fma(double, double, double); |
| 72 | __DEVICE__ float fma(float, float, float); |
| 73 | __DEVICE__ double fmax(double, double); |
| 74 | __DEVICE__ float fmax(float, float); |
| 75 | __DEVICE__ double fmin(double, double); |
| 76 | __DEVICE__ float fmin(float, float); |
| 77 | __DEVICE__ double fmod(double, double); |
| 78 | __DEVICE__ float fmod(float, float); |
| 79 | __DEVICE__ int fpclassify(double); |
| 80 | __DEVICE__ int fpclassify(float); |
| 81 | __DEVICE__ double frexp(double, int *); |
| 82 | __DEVICE__ float frexp(float, int *); |
| 83 | __DEVICE__ double hypot(double, double); |
| 84 | __DEVICE__ float hypot(float, float); |
| 85 | __DEVICE__ int ilogb(double); |
| 86 | __DEVICE__ int ilogb(float); |
Logan Chien | df4f766 | 2019-09-04 16:45:23 -0700 | [diff] [blame] | 87 | #ifdef _MSC_VER |
| 88 | __DEVICE__ bool isfinite(long double); |
| 89 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 90 | __DEVICE__ bool isfinite(double); |
| 91 | __DEVICE__ bool isfinite(float); |
| 92 | __DEVICE__ bool isgreater(double, double); |
| 93 | __DEVICE__ bool isgreaterequal(double, double); |
| 94 | __DEVICE__ bool isgreaterequal(float, float); |
| 95 | __DEVICE__ bool isgreater(float, float); |
Logan Chien | df4f766 | 2019-09-04 16:45:23 -0700 | [diff] [blame] | 96 | #ifdef _MSC_VER |
| 97 | __DEVICE__ bool isinf(long double); |
| 98 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 99 | __DEVICE__ bool isinf(double); |
| 100 | __DEVICE__ bool isinf(float); |
| 101 | __DEVICE__ bool isless(double, double); |
| 102 | __DEVICE__ bool islessequal(double, double); |
| 103 | __DEVICE__ bool islessequal(float, float); |
| 104 | __DEVICE__ bool isless(float, float); |
| 105 | __DEVICE__ bool islessgreater(double, double); |
| 106 | __DEVICE__ bool islessgreater(float, float); |
Logan Chien | df4f766 | 2019-09-04 16:45:23 -0700 | [diff] [blame] | 107 | #ifdef _MSC_VER |
| 108 | __DEVICE__ bool isnan(long double); |
| 109 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 110 | __DEVICE__ bool isnan(double); |
| 111 | __DEVICE__ bool isnan(float); |
| 112 | __DEVICE__ bool isnormal(double); |
| 113 | __DEVICE__ bool isnormal(float); |
| 114 | __DEVICE__ bool isunordered(double, double); |
| 115 | __DEVICE__ bool isunordered(float, float); |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 116 | __DEVICE__ long labs(long); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 117 | __DEVICE__ double ldexp(double, int); |
| 118 | __DEVICE__ float ldexp(float, int); |
| 119 | __DEVICE__ double lgamma(double); |
| 120 | __DEVICE__ float lgamma(float); |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 121 | __DEVICE__ long long llabs(long long); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 122 | __DEVICE__ long long llrint(double); |
| 123 | __DEVICE__ long long llrint(float); |
| 124 | __DEVICE__ double log10(double); |
| 125 | __DEVICE__ float log10(float); |
| 126 | __DEVICE__ double log1p(double); |
| 127 | __DEVICE__ float log1p(float); |
| 128 | __DEVICE__ double log2(double); |
| 129 | __DEVICE__ float log2(float); |
| 130 | __DEVICE__ double logb(double); |
| 131 | __DEVICE__ float logb(float); |
| 132 | __DEVICE__ double log(double); |
| 133 | __DEVICE__ float log(float); |
| 134 | __DEVICE__ long lrint(double); |
| 135 | __DEVICE__ long lrint(float); |
| 136 | __DEVICE__ long lround(double); |
| 137 | __DEVICE__ long lround(float); |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 138 | __DEVICE__ long long llround(float); // No llround(double). |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 139 | __DEVICE__ double modf(double, double *); |
| 140 | __DEVICE__ float modf(float, float *); |
| 141 | __DEVICE__ double nan(const char *); |
| 142 | __DEVICE__ float nanf(const char *); |
| 143 | __DEVICE__ double nearbyint(double); |
| 144 | __DEVICE__ float nearbyint(float); |
| 145 | __DEVICE__ double nextafter(double, double); |
| 146 | __DEVICE__ float nextafter(float, float); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 147 | __DEVICE__ double pow(double, double); |
| 148 | __DEVICE__ double pow(double, int); |
| 149 | __DEVICE__ float pow(float, float); |
| 150 | __DEVICE__ float pow(float, int); |
| 151 | __DEVICE__ double remainder(double, double); |
| 152 | __DEVICE__ float remainder(float, float); |
| 153 | __DEVICE__ double remquo(double, double, int *); |
| 154 | __DEVICE__ float remquo(float, float, int *); |
| 155 | __DEVICE__ double rint(double); |
| 156 | __DEVICE__ float rint(float); |
| 157 | __DEVICE__ double round(double); |
| 158 | __DEVICE__ float round(float); |
| 159 | __DEVICE__ double scalbln(double, long); |
| 160 | __DEVICE__ float scalbln(float, long); |
| 161 | __DEVICE__ double scalbn(double, int); |
| 162 | __DEVICE__ float scalbn(float, int); |
Pirama Arumuga Nainar | 986b880 | 2021-06-03 16:00:34 -0700 | [diff] [blame] | 163 | #ifdef _MSC_VER |
| 164 | __DEVICE__ bool signbit(long double); |
| 165 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 166 | __DEVICE__ bool signbit(double); |
| 167 | __DEVICE__ bool signbit(float); |
| 168 | __DEVICE__ double sin(double); |
| 169 | __DEVICE__ float sin(float); |
| 170 | __DEVICE__ double sinh(double); |
| 171 | __DEVICE__ float sinh(float); |
| 172 | __DEVICE__ double sqrt(double); |
| 173 | __DEVICE__ float sqrt(float); |
| 174 | __DEVICE__ double tan(double); |
| 175 | __DEVICE__ float tan(float); |
| 176 | __DEVICE__ double tanh(double); |
| 177 | __DEVICE__ float tanh(float); |
| 178 | __DEVICE__ double tgamma(double); |
| 179 | __DEVICE__ float tgamma(float); |
| 180 | __DEVICE__ double trunc(double); |
| 181 | __DEVICE__ float trunc(float); |
| 182 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 183 | // Notably missing above is nexttoward, which we don't define on |
| 184 | // the device side because libdevice doesn't give us an implementation, and we |
| 185 | // don't want to be in the business of writing one ourselves. |
| 186 | |
| 187 | // We need to define these overloads in exactly the namespace our standard |
| 188 | // library uses (including the right inline namespace), otherwise they won't be |
| 189 | // picked up by other functions in the standard library (e.g. functions in |
| 190 | // <complex>). Thus the ugliness below. |
| 191 | #ifdef _LIBCPP_BEGIN_NAMESPACE_STD |
| 192 | _LIBCPP_BEGIN_NAMESPACE_STD |
| 193 | #else |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 194 | namespace std { |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 195 | #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION |
| 196 | _GLIBCXX_BEGIN_NAMESPACE_VERSION |
| 197 | #endif |
| 198 | #endif |
| 199 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 200 | using ::abs; |
| 201 | using ::acos; |
| 202 | using ::acosh; |
| 203 | using ::asin; |
| 204 | using ::asinh; |
| 205 | using ::atan; |
| 206 | using ::atan2; |
| 207 | using ::atanh; |
| 208 | using ::cbrt; |
| 209 | using ::ceil; |
| 210 | using ::copysign; |
| 211 | using ::cos; |
| 212 | using ::cosh; |
| 213 | using ::erf; |
| 214 | using ::erfc; |
| 215 | using ::exp; |
| 216 | using ::exp2; |
| 217 | using ::expm1; |
| 218 | using ::fabs; |
| 219 | using ::fdim; |
| 220 | using ::floor; |
| 221 | using ::fma; |
| 222 | using ::fmax; |
| 223 | using ::fmin; |
| 224 | using ::fmod; |
| 225 | using ::fpclassify; |
| 226 | using ::frexp; |
| 227 | using ::hypot; |
| 228 | using ::ilogb; |
| 229 | using ::isfinite; |
| 230 | using ::isgreater; |
| 231 | using ::isgreaterequal; |
| 232 | using ::isinf; |
| 233 | using ::isless; |
| 234 | using ::islessequal; |
| 235 | using ::islessgreater; |
| 236 | using ::isnan; |
| 237 | using ::isnormal; |
| 238 | using ::isunordered; |
| 239 | using ::labs; |
| 240 | using ::ldexp; |
| 241 | using ::lgamma; |
| 242 | using ::llabs; |
| 243 | using ::llrint; |
| 244 | using ::log; |
| 245 | using ::log10; |
| 246 | using ::log1p; |
| 247 | using ::log2; |
| 248 | using ::logb; |
| 249 | using ::lrint; |
| 250 | using ::lround; |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 251 | using ::llround; |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 252 | using ::modf; |
| 253 | using ::nan; |
| 254 | using ::nanf; |
| 255 | using ::nearbyint; |
| 256 | using ::nextafter; |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 257 | using ::pow; |
| 258 | using ::remainder; |
| 259 | using ::remquo; |
| 260 | using ::rint; |
| 261 | using ::round; |
| 262 | using ::scalbln; |
| 263 | using ::scalbn; |
| 264 | using ::signbit; |
| 265 | using ::sin; |
| 266 | using ::sinh; |
| 267 | using ::sqrt; |
| 268 | using ::tan; |
| 269 | using ::tanh; |
| 270 | using ::tgamma; |
| 271 | using ::trunc; |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 272 | |
| 273 | #ifdef _LIBCPP_END_NAMESPACE_STD |
| 274 | _LIBCPP_END_NAMESPACE_STD |
| 275 | #else |
| 276 | #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION |
| 277 | _GLIBCXX_END_NAMESPACE_VERSION |
| 278 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 279 | } // namespace std |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 280 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 281 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 282 | #pragma pop_macro("__DEVICE__") |
| 283 | |
| 284 | #endif |