16bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker/*===---- __clang_cuda_cmath.h - Device-side CUDA cmath support ------------=== 26bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * 36bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * Permission is hereby granted, free of charge, to any person obtaining a copy 46bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * of this software and associated documentation files (the "Software"), to deal 56bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * in the Software without restriction, including without limitation the rights 66bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 76bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * copies of the Software, and to permit persons to whom the Software is 86bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * furnished to do so, subject to the following conditions: 96bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * 106bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * The above copyright notice and this permission notice shall be included in 116bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * all copies or substantial portions of the Software. 126bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * 136bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 146bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 156bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 166bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 176bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 186bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 196bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * THE SOFTWARE. 206bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker * 216bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker *===-----------------------------------------------------------------------=== 226bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker */ 236bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#ifndef __CLANG_CUDA_CMATH_H__ 246bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#define __CLANG_CUDA_CMATH_H__ 256bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#ifndef __CUDA__ 266bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#error "This file is for CUDA compilation only." 276bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#endif 286bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker 296bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// CUDA lets us use various std math functions on the device side. This file 306bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// works in concert with __clang_cuda_math_forward_declares.h to make this work. 316bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// 326bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// Specifically, the forward-declares header declares __device__ overloads for 336bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// these functions in the global namespace, then pulls them into namespace std 346bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// with 'using' statements. Then this file implements those functions, after 356bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// the implementations have been pulled in. 366bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// 376bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// It's important that we declare the functions in the global namespace and pull 386bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// them into namespace std with using statements, as opposed to simply declaring 396bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// these functions in namespace std, because our device functions need to 406bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// overload the standard library functions, which may be declared in the global 416bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// namespace or in std, depending on the degree of conformance of the stdlib 426bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// implementation. Declaring in the global namespace and pulling into namespace 436bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker// std covers all of the known knowns. 446bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker 456bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) 466bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker 476bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ long long abs(long long __n) { return ::llabs(__n); } 486bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ long abs(long __n) { return ::labs(__n); } 496bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float abs(float __x) { return ::fabsf(__x); } 506bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ double abs(double __x) { return ::fabs(__x); } 516bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float acos(float __x) { return ::acosf(__x); } 526bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float asin(float __x) { return ::asinf(__x); } 536bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float atan(float __x) { return ::atanf(__x); } 546bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); } 556bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float ceil(float __x) { return ::ceilf(__x); } 566bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float cos(float __x) { return ::cosf(__x); } 576bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float cosh(float __x) { return ::coshf(__x); } 586bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float exp(float __x) { return ::expf(__x); } 596bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float fabs(float __x) { return ::fabsf(__x); } 606bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float floor(float __x) { return ::floorf(__x); } 616bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); } 626bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ int fpclassify(float __x) { 636bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 646bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker FP_ZERO, __x); 656bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 666bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ int fpclassify(double __x) { 676bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 686bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker FP_ZERO, __x); 696bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 706bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float frexp(float __arg, int *__exp) { 716bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return ::frexpf(__arg, __exp); 726bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 736bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } 746bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } 756bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } 766bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); } 776bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isgreater(float __x, float __y) { 786bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isgreater(__x, __y); 796bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 806bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isgreater(double __x, double __y) { 816bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isgreater(__x, __y); 826bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 836bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isgreaterequal(float __x, float __y) { 846bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isgreaterequal(__x, __y); 856bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 866bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isgreaterequal(double __x, double __y) { 876bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isgreaterequal(__x, __y); 886bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 896bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isless(float __x, float __y) { 906bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isless(__x, __y); 916bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 926bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isless(double __x, double __y) { 936bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isless(__x, __y); 946bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 956bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool islessequal(float __x, float __y) { 966bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_islessequal(__x, __y); 976bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 986bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool islessequal(double __x, double __y) { 996bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_islessequal(__x, __y); 1006bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1016bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool islessgreater(float __x, float __y) { 1026bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_islessgreater(__x, __y); 1036bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1046bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool islessgreater(double __x, double __y) { 1056bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_islessgreater(__x, __y); 1066bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1076bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); } 1086bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); } 1096bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); } 1106bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); } 1116bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isunordered(float __x, float __y) { 1126bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isunordered(__x, __y); 1136bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1146bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool isunordered(double __x, double __y) { 1156bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_isunordered(__x, __y); 1166bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1176bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float ldexp(float __arg, int __exp) { 1186bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return ::ldexpf(__arg, __exp); 1196bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1206bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float log(float __x) { return ::logf(__x); } 1216bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float log10(float __x) { return ::log10f(__x); } 1226bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } 1236bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float nexttoward(float __from, float __to) { 1246bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_nexttowardf(__from, __to); 1256bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1266bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ double nexttoward(double __from, double __to) { 1276bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return __builtin_nexttoward(__from, __to); 1286bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1296bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float pow(float __base, float __exp) { 1306bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return ::powf(__base, __exp); 1316bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1326bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float pow(float __base, int __iexp) { 1336bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return ::powif(__base, __iexp); 1346bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1356bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ double pow(double __base, int __iexp) { 1366bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker return ::powi(__base, __iexp); 1376bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker} 1386bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); } 1396bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); } 1406bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float sin(float __x) { return ::sinf(__x); } 1416bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float sinh(float __x) { return ::sinhf(__x); } 1426bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); } 1436bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float tan(float __x) { return ::tanf(__x); } 1446bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker__DEVICE__ float tanh(float __x) { return ::tanhf(__x); } 1456bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker 1466bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#undef __DEVICE__ 1476bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker 1486bdbd720989797e8a53237ef3ef213c4114f869gitbuildkicker#endif 149