11188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker/*===---- __clang_cuda_cmath.h - Device-side CUDA cmath support ------------=== 21188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * 31188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * Permission is hereby granted, free of charge, to any person obtaining a copy 41188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * of this software and associated documentation files (the "Software"), to deal 51188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * in the Software without restriction, including without limitation the rights 61188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 71188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * copies of the Software, and to permit persons to whom the Software is 81188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * furnished to do so, subject to the following conditions: 91188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * 101188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * The above copyright notice and this permission notice shall be included in 111188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * all copies or substantial portions of the Software. 121188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * 131188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 141188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 151188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 161188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 171188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 181188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 191188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * THE SOFTWARE. 201188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker * 211188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker *===-----------------------------------------------------------------------=== 221188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker */ 231188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#ifndef __CLANG_CUDA_CMATH_H__ 241188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#define __CLANG_CUDA_CMATH_H__ 251188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#ifndef __CUDA__ 261188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#error "This file is for CUDA compilation only." 271188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#endif 281188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker 291188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// CUDA lets us use various std math functions on the device side. This file 301188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// works in concert with __clang_cuda_math_forward_declares.h to make this work. 311188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// 321188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// Specifically, the forward-declares header declares __device__ overloads for 331188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// these functions in the global namespace, then pulls them into namespace std 341188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// with 'using' statements. Then this file implements those functions, after 351188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// the implementations have been pulled in. 361188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// 371188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// It's important that we declare the functions in the global namespace and pull 381188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// them into namespace std with using statements, as opposed to simply declaring 391188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// these functions in namespace std, because our device functions need to 401188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// overload the standard library functions, which may be declared in the global 411188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// namespace or in std, depending on the degree of conformance of the stdlib 421188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// implementation. Declaring in the global namespace and pulling into namespace 431188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker// std covers all of the known knowns. 441188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker 451188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) 461188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker 471188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ long long abs(long long __n) { return ::llabs(__n); } 481188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ long abs(long __n) { return ::labs(__n); } 491188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float abs(float __x) { return ::fabsf(__x); } 501188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ double abs(double __x) { return ::fabs(__x); } 511188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float acos(float __x) { return ::acosf(__x); } 521188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float asin(float __x) { return ::asinf(__x); } 531188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float atan(float __x) { return ::atanf(__x); } 541188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); } 551188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float ceil(float __x) { return ::ceilf(__x); } 561188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float cos(float __x) { return ::cosf(__x); } 571188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float cosh(float __x) { return ::coshf(__x); } 581188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float exp(float __x) { return ::expf(__x); } 591188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float fabs(float __x) { return ::fabsf(__x); } 601188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float floor(float __x) { return ::floorf(__x); } 611188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); } 621188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ int fpclassify(float __x) { 631188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 641188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker FP_ZERO, __x); 651188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 661188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ int fpclassify(double __x) { 671188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, 681188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker FP_ZERO, __x); 691188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 701188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float frexp(float __arg, int *__exp) { 711188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return ::frexpf(__arg, __exp); 721188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 731188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } 741188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } 751188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } 761188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); } 771188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isgreater(float __x, float __y) { 781188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isgreater(__x, __y); 791188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 801188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isgreater(double __x, double __y) { 811188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isgreater(__x, __y); 821188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 831188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isgreaterequal(float __x, float __y) { 841188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isgreaterequal(__x, __y); 851188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 861188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isgreaterequal(double __x, double __y) { 871188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isgreaterequal(__x, __y); 881188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 891188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isless(float __x, float __y) { 901188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isless(__x, __y); 911188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 921188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isless(double __x, double __y) { 931188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isless(__x, __y); 941188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 951188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool islessequal(float __x, float __y) { 961188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_islessequal(__x, __y); 971188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 981188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool islessequal(double __x, double __y) { 991188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_islessequal(__x, __y); 1001188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1011188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool islessgreater(float __x, float __y) { 1021188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_islessgreater(__x, __y); 1031188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1041188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool islessgreater(double __x, double __y) { 1051188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_islessgreater(__x, __y); 1061188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1071188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); } 1081188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); } 1091188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); } 1101188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); } 1111188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isunordered(float __x, float __y) { 1121188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isunordered(__x, __y); 1131188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1141188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool isunordered(double __x, double __y) { 1151188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_isunordered(__x, __y); 1161188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1171188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float ldexp(float __arg, int __exp) { 1181188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return ::ldexpf(__arg, __exp); 1191188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1201188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float log(float __x) { return ::logf(__x); } 1211188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float log10(float __x) { return ::log10f(__x); } 1221188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } 1231188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float nexttoward(float __from, float __to) { 1241188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_nexttowardf(__from, __to); 1251188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1261188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ double nexttoward(double __from, double __to) { 1271188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return __builtin_nexttoward(__from, __to); 1281188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1291188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float pow(float __base, float __exp) { 1301188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return ::powf(__base, __exp); 1311188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1321188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float pow(float __base, int __iexp) { 1331188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return ::powif(__base, __iexp); 1341188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1351188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ double pow(double __base, int __iexp) { 1361188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker return ::powi(__base, __iexp); 1371188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker} 1381188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); } 1391188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); } 1401188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float sin(float __x) { return ::sinf(__x); } 1411188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float sinh(float __x) { return ::sinhf(__x); } 1421188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); } 1431188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float tan(float __x) { return ::tanf(__x); } 1441188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker__DEVICE__ float tanh(float __x) { return ::tanhf(__x); } 1451188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker 1461188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#undef __DEVICE__ 1471188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker 1481188dcf30923cb444143ffa4b83dc951037e76agitbuildkicker#endif 149