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 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, 64 FP_ZERO, __x); 65 } 66 __DEVICE__ int fpclassify(double __x) { 67 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 68 FP_ZERO, __x); 69 } 70 __DEVICE__ float frexp(float __arg, int *__exp) { 71 return ::frexpf(__arg, __exp); 72 } 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); 79 } 80 __DEVICE__ bool isgreater(double __x, double __y) { 81 return __builtin_isgreater(__x, __y); 82 } 83 __DEVICE__ bool isgreaterequal(float __x, float __y) { 84 return __builtin_isgreaterequal(__x, __y); 85 } 86 __DEVICE__ bool isgreaterequal(double __x, double __y) { 87 return __builtin_isgreaterequal(__x, __y); 88 } 89 __DEVICE__ bool isless(float __x, float __y) { 90 return __builtin_isless(__x, __y); 91 } 92 __DEVICE__ bool isless(double __x, double __y) { 93 return __builtin_isless(__x, __y); 94 } 95 __DEVICE__ bool islessequal(float __x, float __y) { 96 return __builtin_islessequal(__x, __y); 97 } 98 __DEVICE__ bool islessequal(double __x, double __y) { 99 return __builtin_islessequal(__x, __y); 100 } 101 __DEVICE__ bool islessgreater(float __x, float __y) { 102 return __builtin_islessgreater(__x, __y); 103 } 104 __DEVICE__ bool islessgreater(double __x, double __y) { 105 return __builtin_islessgreater(__x, __y); 106 } 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); 113 } 114 __DEVICE__ bool isunordered(double __x, double __y) { 115 return __builtin_isunordered(__x, __y); 116 } 117 __DEVICE__ float ldexp(float __arg, int __exp) { 118 return ::ldexpf(__arg, __exp); 119 } 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); 125 } 126 __DEVICE__ double nexttoward(double __from, double __to) { 127 return __builtin_nexttoward(__from, __to); 128 } 129 __DEVICE__ float pow(float __base, float __exp) { 130 return ::powf(__base, __exp); 131 } 132 __DEVICE__ float pow(float __base, int __iexp) { 133 return ::powif(__base, __iexp); 134 } 135 __DEVICE__ double pow(double __base, int __iexp) { 136 return ::powi(__base, __iexp); 137 } 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); } 145 146 #undef __DEVICE__ 147 148 #endif 149
