Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 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 | |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 29 | // 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 Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 33 | |
| 34 | #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) |
| 35 | |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 36 | namespace std { |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 37 | __DEVICE__ long long abs(long long __n) { return ::llabs(__n); } |
| 38 | __DEVICE__ long abs(long __n) { return ::labs(__n); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 39 | using ::abs; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 40 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 43 | using ::acos; |
| 44 | using ::acosh; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 45 | __DEVICE__ float asin(float __x) { return ::asinf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 46 | using ::asin; |
| 47 | using ::asinh; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 48 | __DEVICE__ float atan(float __x) { return ::atanf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 49 | using ::atan; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 50 | __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 51 | using ::atan2; |
| 52 | using ::atanh; |
| 53 | using ::cbrt; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 54 | __DEVICE__ float ceil(float __x) { return ::ceilf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 55 | using ::ceil; |
| 56 | using ::copysign; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 57 | __DEVICE__ float cos(float __x) { return ::cosf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 58 | using ::cos; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 59 | __DEVICE__ float cosh(float __x) { return ::coshf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 60 | using ::cosh; |
| 61 | using ::erf; |
| 62 | using ::erfc; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 63 | __DEVICE__ float exp(float __x) { return ::expf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 64 | using ::exp; |
| 65 | using ::exp2; |
| 66 | using ::expm1; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 67 | __DEVICE__ float fabs(float __x) { return ::fabsf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 68 | using ::fabs; |
| 69 | using ::fdim; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 70 | __DEVICE__ float floor(float __x) { return ::floorf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 71 | using ::floor; |
| 72 | using ::fma; |
| 73 | using ::fmax; |
| 74 | using ::fmin; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 75 | __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 76 | using ::fmod; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 77 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 88 | using ::frexp; |
| 89 | using ::hypot; |
| 90 | using ::ilogb; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 91 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 105 | __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } |
| 106 | __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 107 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 135 | using ::labs; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 136 | __DEVICE__ float ldexp(float __arg, int __exp) { |
| 137 | return ::ldexpf(__arg, __exp); |
| 138 | } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 139 | using ::ldexp; |
| 140 | using ::lgamma; |
| 141 | using ::llabs; |
| 142 | using ::llrint; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 143 | __DEVICE__ float log(float __x) { return ::logf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 144 | using ::log; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 145 | __DEVICE__ float log10(float __x) { return ::log10f(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 146 | using ::log10; |
| 147 | using ::log1p; |
| 148 | using ::log2; |
| 149 | using ::logb; |
| 150 | using ::lrint; |
| 151 | using ::lround; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 152 | __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 153 | using ::modf; |
| 154 | using ::nan; |
| 155 | using ::nanf; |
| 156 | using ::nearbyint; |
| 157 | using ::nextafter; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 158 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 164 | using ::pow; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 165 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 174 | using ::remainder; |
| 175 | using ::remquo; |
| 176 | using ::rint; |
| 177 | using ::round; |
| 178 | using ::scalbln; |
| 179 | using ::scalbn; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 180 | __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 Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 183 | using ::sin; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 184 | __DEVICE__ float sinh(float __x) { return ::sinhf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 185 | using ::sinh; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 186 | __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 187 | using ::sqrt; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 188 | __DEVICE__ float tan(float __x) { return ::tanf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 189 | using ::tan; |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 190 | __DEVICE__ float tanh(float __x) { return ::tanhf(__x); } |
Ben Murdoch | c561043 | 2016-08-08 18:44:38 +0100 | [diff] [blame^] | 191 | using ::tanh; |
| 192 | using ::tgamma; |
| 193 | using ::trunc; |
| 194 | |
| 195 | } // namespace std |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 196 | |
| 197 | #undef __DEVICE__ |
| 198 | |
| 199 | #endif |