1 /*===---- __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
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.
31 //
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.
36 //
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.
44
45 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
46
abs(long long __n)47 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
abs(long __n)48 __DEVICE__ long abs(long __n) { return ::labs(__n); }
abs(float __x)49 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
abs(double __x)50 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
acos(float __x)51 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
asin(float __x)52 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
atan(float __x)53 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
atan2(float __x,float __y)54 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
ceil(float __x)55 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
cos(float __x)56 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
cosh(float __x)57 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
exp(float __x)58 __DEVICE__ float exp(float __x) { return ::expf(__x); }
fabs(float __x)59 __DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
floor(float __x)60 __DEVICE__ float floor(float __x) { return ::floorf(__x); }
fmod(float __x,float __y)61 __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
fpclassify(float __x)62 __DEVICE__ int fpclassify(float __x) {
63 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
64 FP_ZERO, __x);
65 }
fpclassify(double __x)66 __DEVICE__ int fpclassify(double __x) {
67 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
68 FP_ZERO, __x);
69 }
frexp(float __arg,int * __exp)70 __DEVICE__ float frexp(float __arg, int *__exp) {
71 return ::frexpf(__arg, __exp);
72 }
isinf(float __x)73 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
isinf(double __x)74 __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
isfinite(float __x)75 __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
isfinite(double __x)76 __DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
isgreater(float __x,float __y)77 __DEVICE__ bool isgreater(float __x, float __y) {
78 return __builtin_isgreater(__x, __y);
79 }
isgreater(double __x,double __y)80 __DEVICE__ bool isgreater(double __x, double __y) {
81 return __builtin_isgreater(__x, __y);
82 }
isgreaterequal(float __x,float __y)83 __DEVICE__ bool isgreaterequal(float __x, float __y) {
84 return __builtin_isgreaterequal(__x, __y);
85 }
isgreaterequal(double __x,double __y)86 __DEVICE__ bool isgreaterequal(double __x, double __y) {
87 return __builtin_isgreaterequal(__x, __y);
88 }
isless(float __x,float __y)89 __DEVICE__ bool isless(float __x, float __y) {
90 return __builtin_isless(__x, __y);
91 }
isless(double __x,double __y)92 __DEVICE__ bool isless(double __x, double __y) {
93 return __builtin_isless(__x, __y);
94 }
islessequal(float __x,float __y)95 __DEVICE__ bool islessequal(float __x, float __y) {
96 return __builtin_islessequal(__x, __y);
97 }
islessequal(double __x,double __y)98 __DEVICE__ bool islessequal(double __x, double __y) {
99 return __builtin_islessequal(__x, __y);
100 }
islessgreater(float __x,float __y)101 __DEVICE__ bool islessgreater(float __x, float __y) {
102 return __builtin_islessgreater(__x, __y);
103 }
islessgreater(double __x,double __y)104 __DEVICE__ bool islessgreater(double __x, double __y) {
105 return __builtin_islessgreater(__x, __y);
106 }
isnan(float __x)107 __DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
isnan(double __x)108 __DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
isnormal(float __x)109 __DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
isnormal(double __x)110 __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
isunordered(float __x,float __y)111 __DEVICE__ bool isunordered(float __x, float __y) {
112 return __builtin_isunordered(__x, __y);
113 }
isunordered(double __x,double __y)114 __DEVICE__ bool isunordered(double __x, double __y) {
115 return __builtin_isunordered(__x, __y);
116 }
ldexp(float __arg,int __exp)117 __DEVICE__ float ldexp(float __arg, int __exp) {
118 return ::ldexpf(__arg, __exp);
119 }
log(float __x)120 __DEVICE__ float log(float __x) { return ::logf(__x); }
log10(float __x)121 __DEVICE__ float log10(float __x) { return ::log10f(__x); }
modf(float __x,float * __iptr)122 __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
nexttoward(float __from,float __to)123 __DEVICE__ float nexttoward(float __from, float __to) {
124 return __builtin_nexttowardf(__from, __to);
125 }
nexttoward(double __from,double __to)126 __DEVICE__ double nexttoward(double __from, double __to) {
127 return __builtin_nexttoward(__from, __to);
128 }
pow(float __base,float __exp)129 __DEVICE__ float pow(float __base, float __exp) {
130 return ::powf(__base, __exp);
131 }
pow(float __base,int __iexp)132 __DEVICE__ float pow(float __base, int __iexp) {
133 return ::powif(__base, __iexp);
134 }
pow(double __base,int __iexp)135 __DEVICE__ double pow(double __base, int __iexp) {
136 return ::powi(__base, __iexp);
137 }
signbit(float __x)138 __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
signbit(double __x)139 __DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
sin(float __x)140 __DEVICE__ float sin(float __x) { return ::sinf(__x); }
sinh(float __x)141 __DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
sqrt(float __x)142 __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
tan(float __x)143 __DEVICE__ float tan(float __x) { return ::tanf(__x); }
tanh(float __x)144 __DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
145
146 #undef __DEVICE__
147
148 #endif
149