blob: 9fe3f712c7e17826b3d46c2daad7201440b81c50 [file] [log] [blame]
Ben Murdoch097c5b22016-05-18 11:27:45 +01001/*===---- __clang_cuda_cmath.h - Device-side CUDA cmath support ------------===
2 *
3 * Permission is hereby granted, free of charge, to any person obtaining a copy
4 * of this software and associated documentation files (the "Software"), to deal
5 * in the Software without restriction, including without limitation the rights
6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 * copies of the Software, and to permit persons to whom the Software is
8 * furnished to do so, subject to the following conditions:
9 *
10 * The above copyright notice and this permission notice shall be included in
11 * all copies or substantial portions of the Software.
12 *
13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 * THE SOFTWARE.
20 *
21 *===-----------------------------------------------------------------------===
22 */
23#ifndef __CLANG_CUDA_CMATH_H__
24#define __CLANG_CUDA_CMATH_H__
25#ifndef __CUDA__
26#error "This file is for CUDA compilation only."
27#endif
28
Ben Murdochc5610432016-08-08 18:44:38 +010029// CUDA allows using math functions form std:: on device side. This
30// file provides __device__ overloads for math functions that map to
31// appropriate math functions provided by CUDA headers or to compiler
32// builtins if CUDA does not provide a suitable function.
Ben Murdoch097c5b22016-05-18 11:27:45 +010033
34#define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
35
Ben Murdochc5610432016-08-08 18:44:38 +010036namespace std {
Ben Murdoch097c5b22016-05-18 11:27:45 +010037__DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
38__DEVICE__ long abs(long __n) { return ::labs(__n); }
Ben Murdochc5610432016-08-08 18:44:38 +010039using ::abs;
Ben Murdoch097c5b22016-05-18 11:27:45 +010040__DEVICE__ float abs(float __x) { return ::fabsf(__x); }
41__DEVICE__ double abs(double __x) { return ::fabs(__x); }
42__DEVICE__ float acos(float __x) { return ::acosf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010043using ::acos;
44using ::acosh;
Ben Murdoch097c5b22016-05-18 11:27:45 +010045__DEVICE__ float asin(float __x) { return ::asinf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010046using ::asin;
47using ::asinh;
Ben Murdoch097c5b22016-05-18 11:27:45 +010048__DEVICE__ float atan(float __x) { return ::atanf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010049using ::atan;
Ben Murdoch097c5b22016-05-18 11:27:45 +010050__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
Ben Murdochc5610432016-08-08 18:44:38 +010051using ::atan2;
52using ::atanh;
53using ::cbrt;
Ben Murdoch097c5b22016-05-18 11:27:45 +010054__DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010055using ::ceil;
56using ::copysign;
Ben Murdoch097c5b22016-05-18 11:27:45 +010057__DEVICE__ float cos(float __x) { return ::cosf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010058using ::cos;
Ben Murdoch097c5b22016-05-18 11:27:45 +010059__DEVICE__ float cosh(float __x) { return ::coshf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010060using ::cosh;
61using ::erf;
62using ::erfc;
Ben Murdoch097c5b22016-05-18 11:27:45 +010063__DEVICE__ float exp(float __x) { return ::expf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010064using ::exp;
65using ::exp2;
66using ::expm1;
Ben Murdoch097c5b22016-05-18 11:27:45 +010067__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010068using ::fabs;
69using ::fdim;
Ben Murdoch097c5b22016-05-18 11:27:45 +010070__DEVICE__ float floor(float __x) { return ::floorf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +010071using ::floor;
72using ::fma;
73using ::fmax;
74using ::fmin;
Ben Murdoch097c5b22016-05-18 11:27:45 +010075__DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
Ben Murdochc5610432016-08-08 18:44:38 +010076using ::fmod;
Ben Murdoch097c5b22016-05-18 11:27:45 +010077__DEVICE__ int fpclassify(float __x) {
78 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
79 FP_ZERO, __x);
80}
81__DEVICE__ int fpclassify(double __x) {
82 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
83 FP_ZERO, __x);
84}
85__DEVICE__ float frexp(float __arg, int *__exp) {
86 return ::frexpf(__arg, __exp);
87}
Ben Murdochc5610432016-08-08 18:44:38 +010088using ::frexp;
89using ::hypot;
90using ::ilogb;
Ben Murdoch097c5b22016-05-18 11:27:45 +010091__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
92__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
93__DEVICE__ bool isgreater(float __x, float __y) {
94 return __builtin_isgreater(__x, __y);
95}
96__DEVICE__ bool isgreater(double __x, double __y) {
97 return __builtin_isgreater(__x, __y);
98}
99__DEVICE__ bool isgreaterequal(float __x, float __y) {
100 return __builtin_isgreaterequal(__x, __y);
101}
102__DEVICE__ bool isgreaterequal(double __x, double __y) {
103 return __builtin_isgreaterequal(__x, __y);
104}
Ben Murdochc5610432016-08-08 18:44:38 +0100105__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
106__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
Ben Murdoch097c5b22016-05-18 11:27:45 +0100107__DEVICE__ bool isless(float __x, float __y) {
108 return __builtin_isless(__x, __y);
109}
110__DEVICE__ bool isless(double __x, double __y) {
111 return __builtin_isless(__x, __y);
112}
113__DEVICE__ bool islessequal(float __x, float __y) {
114 return __builtin_islessequal(__x, __y);
115}
116__DEVICE__ bool islessequal(double __x, double __y) {
117 return __builtin_islessequal(__x, __y);
118}
119__DEVICE__ bool islessgreater(float __x, float __y) {
120 return __builtin_islessgreater(__x, __y);
121}
122__DEVICE__ bool islessgreater(double __x, double __y) {
123 return __builtin_islessgreater(__x, __y);
124}
125__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
126__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
127__DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
128__DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
129__DEVICE__ bool isunordered(float __x, float __y) {
130 return __builtin_isunordered(__x, __y);
131}
132__DEVICE__ bool isunordered(double __x, double __y) {
133 return __builtin_isunordered(__x, __y);
134}
Ben Murdochc5610432016-08-08 18:44:38 +0100135using ::labs;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100136__DEVICE__ float ldexp(float __arg, int __exp) {
137 return ::ldexpf(__arg, __exp);
138}
Ben Murdochc5610432016-08-08 18:44:38 +0100139using ::ldexp;
140using ::lgamma;
141using ::llabs;
142using ::llrint;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100143__DEVICE__ float log(float __x) { return ::logf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100144using ::log;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100145__DEVICE__ float log10(float __x) { return ::log10f(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100146using ::log10;
147using ::log1p;
148using ::log2;
149using ::logb;
150using ::lrint;
151using ::lround;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100152__DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
Ben Murdochc5610432016-08-08 18:44:38 +0100153using ::modf;
154using ::nan;
155using ::nanf;
156using ::nearbyint;
157using ::nextafter;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100158__DEVICE__ float nexttoward(float __from, float __to) {
159 return __builtin_nexttowardf(__from, __to);
160}
161__DEVICE__ double nexttoward(double __from, double __to) {
162 return __builtin_nexttoward(__from, __to);
163}
Ben Murdochc5610432016-08-08 18:44:38 +0100164using ::pow;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100165__DEVICE__ float pow(float __base, float __exp) {
166 return ::powf(__base, __exp);
167}
168__DEVICE__ float pow(float __base, int __iexp) {
169 return ::powif(__base, __iexp);
170}
171__DEVICE__ double pow(double __base, int __iexp) {
172 return ::powi(__base, __iexp);
173}
Ben Murdochc5610432016-08-08 18:44:38 +0100174using ::remainder;
175using ::remquo;
176using ::rint;
177using ::round;
178using ::scalbln;
179using ::scalbn;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100180__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
181__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
182__DEVICE__ float sin(float __x) { return ::sinf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100183using ::sin;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100184__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100185using ::sinh;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100186__DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100187using ::sqrt;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100188__DEVICE__ float tan(float __x) { return ::tanf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100189using ::tan;
Ben Murdoch097c5b22016-05-18 11:27:45 +0100190__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
Ben Murdochc5610432016-08-08 18:44:38 +0100191using ::tanh;
192using ::tgamma;
193using ::trunc;
194
195} // namespace std
Ben Murdoch097c5b22016-05-18 11:27:45 +0100196
197#undef __DEVICE__
198
199#endif