1 /*===---- __clang_cuda_cmath.h - Device-side CUDA cmath support ------------===
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:
10 * The above copyright notice and this permission notice shall be included in
11 * all copies or substantial portions of the Software.
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
21 *===-----------------------------------------------------------------------===
23 #ifndef __CLANG_CUDA_CMATH_H__
24 #define __CLANG_CUDA_CMATH_H__
26 #error "This file is for CUDA compilation only."
29 // CUDA lets us use various std math functions on the device side. This file
30 // works in concert with __clang_cuda_math_forward_declares.h to make this work.
32 // Specifically, the forward-declares header declares __device__ overloads for
33 // these functions in the global namespace, then pulls them into namespace std
34 // with 'using' statements. Then this file implements those functions, after
35 // the implementations have been pulled in.
37 // It's important that we declare the functions in the global namespace and pull
38 // them into namespace std with using statements, as opposed to simply declaring
39 // these functions in namespace std, because our device functions need to
40 // overload the standard library functions, which may be declared in the global
41 // namespace or in std, depending on the degree of conformance of the stdlib
42 // implementation. Declaring in the global namespace and pulling into namespace
43 // std covers all of the known knowns.
45 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
47 __DEVICE__
long long abs(long long __n
) { return ::llabs(__n
); }
48 __DEVICE__
long abs(long __n
) { return ::labs(__n
); }
49 __DEVICE__
float abs(float __x
) { return ::fabsf(__x
); }
50 __DEVICE__
double abs(double __x
) { return ::fabs(__x
); }
51 __DEVICE__
float acos(float __x
) { return ::acosf(__x
); }
52 __DEVICE__
float asin(float __x
) { return ::asinf(__x
); }
53 __DEVICE__
float atan(float __x
) { return ::atanf(__x
); }
54 __DEVICE__
float atan2(float __x
, float __y
) { return ::atan2f(__x
, __y
); }
55 __DEVICE__
float ceil(float __x
) { return ::ceilf(__x
); }
56 __DEVICE__
float cos(float __x
) { return ::cosf(__x
); }
57 __DEVICE__
float cosh(float __x
) { return ::coshf(__x
); }
58 __DEVICE__
float exp(float __x
) { return ::expf(__x
); }
59 __DEVICE__
float fabs(float __x
) { return ::fabsf(__x
); }
60 __DEVICE__
float floor(float __x
) { return ::floorf(__x
); }
61 __DEVICE__
float fmod(float __x
, float __y
) { return ::fmodf(__x
, __y
); }
62 __DEVICE__
int fpclassify(float __x
) {
63 return __builtin_fpclassify(FP_NAN
, FP_INFINITE
, FP_NORMAL
, FP_SUBNORMAL
,
66 __DEVICE__
int fpclassify(double __x
) {
67 return __builtin_fpclassify(FP_NAN
, FP_INFINITE
, FP_NORMAL
, FP_SUBNORMAL
,
70 __DEVICE__
float frexp(float __arg
, int *__exp
) {
71 return ::frexpf(__arg
, __exp
);
73 __DEVICE__
bool isinf(float __x
) { return ::__isinff(__x
); }
74 __DEVICE__
bool isinf(double __x
) { return ::__isinf(__x
); }
75 __DEVICE__
bool isfinite(float __x
) { return ::__finitef(__x
); }
76 __DEVICE__
bool isfinite(double __x
) { return ::__finite(__x
); }
77 __DEVICE__
bool isgreater(float __x
, float __y
) {
78 return __builtin_isgreater(__x
, __y
);
80 __DEVICE__
bool isgreater(double __x
, double __y
) {
81 return __builtin_isgreater(__x
, __y
);
83 __DEVICE__
bool isgreaterequal(float __x
, float __y
) {
84 return __builtin_isgreaterequal(__x
, __y
);
86 __DEVICE__
bool isgreaterequal(double __x
, double __y
) {
87 return __builtin_isgreaterequal(__x
, __y
);
89 __DEVICE__
bool isless(float __x
, float __y
) {
90 return __builtin_isless(__x
, __y
);
92 __DEVICE__
bool isless(double __x
, double __y
) {
93 return __builtin_isless(__x
, __y
);
95 __DEVICE__
bool islessequal(float __x
, float __y
) {
96 return __builtin_islessequal(__x
, __y
);
98 __DEVICE__
bool islessequal(double __x
, double __y
) {
99 return __builtin_islessequal(__x
, __y
);
101 __DEVICE__
bool islessgreater(float __x
, float __y
) {
102 return __builtin_islessgreater(__x
, __y
);
104 __DEVICE__
bool islessgreater(double __x
, double __y
) {
105 return __builtin_islessgreater(__x
, __y
);
107 __DEVICE__
bool isnan(float __x
) { return ::__isnanf(__x
); }
108 __DEVICE__
bool isnan(double __x
) { return ::__isnan(__x
); }
109 __DEVICE__
bool isnormal(float __x
) { return __builtin_isnormal(__x
); }
110 __DEVICE__
bool isnormal(double __x
) { return __builtin_isnormal(__x
); }
111 __DEVICE__
bool isunordered(float __x
, float __y
) {
112 return __builtin_isunordered(__x
, __y
);
114 __DEVICE__
bool isunordered(double __x
, double __y
) {
115 return __builtin_isunordered(__x
, __y
);
117 __DEVICE__
float ldexp(float __arg
, int __exp
) {
118 return ::ldexpf(__arg
, __exp
);
120 __DEVICE__
float log(float __x
) { return ::logf(__x
); }
121 __DEVICE__
float log10(float __x
) { return ::log10f(__x
); }
122 __DEVICE__
float modf(float __x
, float *__iptr
) { return ::modff(__x
, __iptr
); }
123 __DEVICE__
float nexttoward(float __from
, float __to
) {
124 return __builtin_nexttowardf(__from
, __to
);
126 __DEVICE__
double nexttoward(double __from
, double __to
) {
127 return __builtin_nexttoward(__from
, __to
);
129 __DEVICE__
float pow(float __base
, float __exp
) {
130 return ::powf(__base
, __exp
);
132 __DEVICE__
float pow(float __base
, int __iexp
) {
133 return ::powif(__base
, __iexp
);
135 __DEVICE__
double pow(double __base
, int __iexp
) {
136 return ::powi(__base
, __iexp
);
138 __DEVICE__
bool signbit(float __x
) { return ::__signbitf(__x
); }
139 __DEVICE__
bool signbit(double __x
) { return ::__signbit(__x
); }
140 __DEVICE__
float sin(float __x
) { return ::sinf(__x
); }
141 __DEVICE__
float sinh(float __x
) { return ::sinhf(__x
); }
142 __DEVICE__
float sqrt(float __x
) { return ::sqrtf(__x
); }
143 __DEVICE__
float tan(float __x
) { return ::tanf(__x
); }
144 __DEVICE__
float tanh(float __x
) { return ::tanhf(__x
); }