| /*===---- __clang_cuda_cmath.h - Device-side CUDA cmath support ------------=== |
| * |
| * Permission is hereby granted, free of charge, to any person obtaining a copy |
| * of this software and associated documentation files (the "Software"), to deal |
| * in the Software without restriction, including without limitation the rights |
| * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell |
| * copies of the Software, and to permit persons to whom the Software is |
| * furnished to do so, subject to the following conditions: |
| * |
| * The above copyright notice and this permission notice shall be included in |
| * all copies or substantial portions of the Software. |
| * |
| * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR |
| * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, |
| * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE |
| * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER |
| * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, |
| * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN |
| * THE SOFTWARE. |
| * |
| *===-----------------------------------------------------------------------=== |
| */ |
| #ifndef __CLANG_CUDA_CMATH_H__ |
| #define __CLANG_CUDA_CMATH_H__ |
| #ifndef __CUDA__ |
| #error "This file is for CUDA compilation only." |
| #endif |
| |
| // CUDA allows using math functions form std:: on device side. This |
| // file provides __device__ overloads for math functions that map to |
| // appropriate math functions provided by CUDA headers or to compiler |
| // builtins if CUDA does not provide a suitable function. |
| |
| #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) |
| |
| namespace std { |
| __DEVICE__ long long abs(long long __n) { return ::llabs(__n); } |
| __DEVICE__ long abs(long __n) { return ::labs(__n); } |
| using ::abs; |
| __DEVICE__ float abs(float __x) { return ::fabsf(__x); } |
| __DEVICE__ double abs(double __x) { return ::fabs(__x); } |
| __DEVICE__ float acos(float __x) { return ::acosf(__x); } |
| using ::acos; |
| using ::acosh; |
| __DEVICE__ float asin(float __x) { return ::asinf(__x); } |
| using ::asin; |
| using ::asinh; |
| __DEVICE__ float atan(float __x) { return ::atanf(__x); } |
| using ::atan; |
| __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); } |
| using ::atan2; |
| using ::atanh; |
| using ::cbrt; |
| __DEVICE__ float ceil(float __x) { return ::ceilf(__x); } |
| using ::ceil; |
| using ::copysign; |
| __DEVICE__ float cos(float __x) { return ::cosf(__x); } |
| using ::cos; |
| __DEVICE__ float cosh(float __x) { return ::coshf(__x); } |
| using ::cosh; |
| using ::erf; |
| using ::erfc; |
| __DEVICE__ float exp(float __x) { return ::expf(__x); } |
| using ::exp; |
| using ::exp2; |
| using ::expm1; |
| __DEVICE__ float fabs(float __x) { return ::fabsf(__x); } |
| using ::fabs; |
| using ::fdim; |
| __DEVICE__ float floor(float __x) { return ::floorf(__x); } |
| using ::floor; |
| using ::fma; |
| using ::fmax; |
| using ::fmin; |
| __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); } |
| using ::fmod; |
| __DEVICE__ int fpclassify(float __x) { |
| return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, |
| FP_ZERO, __x); |
| } |
| __DEVICE__ int fpclassify(double __x) { |
| return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, |
| FP_ZERO, __x); |
| } |
| __DEVICE__ float frexp(float __arg, int *__exp) { |
| return ::frexpf(__arg, __exp); |
| } |
| using ::frexp; |
| using ::hypot; |
| using ::ilogb; |
| __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } |
| __DEVICE__ bool isfinite(double __x) { return ::__finite(__x); } |
| __DEVICE__ bool isgreater(float __x, float __y) { |
| return __builtin_isgreater(__x, __y); |
| } |
| __DEVICE__ bool isgreater(double __x, double __y) { |
| return __builtin_isgreater(__x, __y); |
| } |
| __DEVICE__ bool isgreaterequal(float __x, float __y) { |
| return __builtin_isgreaterequal(__x, __y); |
| } |
| __DEVICE__ bool isgreaterequal(double __x, double __y) { |
| return __builtin_isgreaterequal(__x, __y); |
| } |
| __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } |
| __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } |
| __DEVICE__ bool isless(float __x, float __y) { |
| return __builtin_isless(__x, __y); |
| } |
| __DEVICE__ bool isless(double __x, double __y) { |
| return __builtin_isless(__x, __y); |
| } |
| __DEVICE__ bool islessequal(float __x, float __y) { |
| return __builtin_islessequal(__x, __y); |
| } |
| __DEVICE__ bool islessequal(double __x, double __y) { |
| return __builtin_islessequal(__x, __y); |
| } |
| __DEVICE__ bool islessgreater(float __x, float __y) { |
| return __builtin_islessgreater(__x, __y); |
| } |
| __DEVICE__ bool islessgreater(double __x, double __y) { |
| return __builtin_islessgreater(__x, __y); |
| } |
| __DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); } |
| __DEVICE__ bool isnan(double __x) { return ::__isnan(__x); } |
| __DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); } |
| __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); } |
| __DEVICE__ bool isunordered(float __x, float __y) { |
| return __builtin_isunordered(__x, __y); |
| } |
| __DEVICE__ bool isunordered(double __x, double __y) { |
| return __builtin_isunordered(__x, __y); |
| } |
| using ::labs; |
| __DEVICE__ float ldexp(float __arg, int __exp) { |
| return ::ldexpf(__arg, __exp); |
| } |
| using ::ldexp; |
| using ::lgamma; |
| using ::llabs; |
| using ::llrint; |
| __DEVICE__ float log(float __x) { return ::logf(__x); } |
| using ::log; |
| __DEVICE__ float log10(float __x) { return ::log10f(__x); } |
| using ::log10; |
| using ::log1p; |
| using ::log2; |
| using ::logb; |
| using ::lrint; |
| using ::lround; |
| __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } |
| using ::modf; |
| using ::nan; |
| using ::nanf; |
| using ::nearbyint; |
| using ::nextafter; |
| __DEVICE__ float nexttoward(float __from, float __to) { |
| return __builtin_nexttowardf(__from, __to); |
| } |
| __DEVICE__ double nexttoward(double __from, double __to) { |
| return __builtin_nexttoward(__from, __to); |
| } |
| using ::pow; |
| __DEVICE__ float pow(float __base, float __exp) { |
| return ::powf(__base, __exp); |
| } |
| __DEVICE__ float pow(float __base, int __iexp) { |
| return ::powif(__base, __iexp); |
| } |
| __DEVICE__ double pow(double __base, int __iexp) { |
| return ::powi(__base, __iexp); |
| } |
| using ::remainder; |
| using ::remquo; |
| using ::rint; |
| using ::round; |
| using ::scalbln; |
| using ::scalbn; |
| __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); } |
| __DEVICE__ bool signbit(double __x) { return ::__signbit(__x); } |
| __DEVICE__ float sin(float __x) { return ::sinf(__x); } |
| using ::sin; |
| __DEVICE__ float sinh(float __x) { return ::sinhf(__x); } |
| using ::sinh; |
| __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); } |
| using ::sqrt; |
| __DEVICE__ float tan(float __x) { return ::tanf(__x); } |
| using ::tan; |
| __DEVICE__ float tanh(float __x) { return ::tanhf(__x); } |
| using ::tanh; |
| using ::tgamma; |
| using ::trunc; |
| |
| } // namespace std |
| |
| #undef __DEVICE__ |
| |
| #endif |