9 #ifndef __CLANG_CUDA_CMATH_H__
10 #define __CLANG_CUDA_CMATH_H__
12 #error "This file is for CUDA compilation only."
15 #ifndef __OPENMP_NVPTX__
35 #ifdef __OPENMP_NVPTX__
36 #define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
38 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
57 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
61 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
70 #if !defined(_MSC_VER) || defined(__OPENMP_NVPTX__)
78 #if defined(__OPENMP_NVPTX__)
79 #pragma omp begin declare variant match( \
80 implementation = {extension(disable_implicit_base)})
88 #pragma omp begin declare variant match(implementation = {vendor(llvm)})
97 #pragma omp end declare variant
111 #if defined(__OPENMP_NVPTX__)
112 #pragma omp end declare variant
118 return __builtin_isgreater(__x,
__y);
121 return __builtin_isgreater(__x,
__y);
124 return __builtin_isgreaterequal(__x,
__y);
127 return __builtin_isgreaterequal(__x,
__y);
130 return __builtin_isless(__x,
__y);
133 return __builtin_isless(__x,
__y);
136 return __builtin_islessequal(__x,
__y);
139 return __builtin_islessequal(__x,
__y);
142 return __builtin_islessgreater(__x,
__y);
145 return __builtin_islessgreater(__x,
__y);
150 return __builtin_isunordered(__x,
__y);
153 return __builtin_isunordered(__x,
__y);
181 #ifdef __OPENMP_NVPTX__
191 #ifndef __OPENMP_NVPTX__
208 template<
bool __B,
class __T =
void>
217 #define __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(__retty, __fn) \
218 template <typename __T> \
220 typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer, \
223 return ::__fn((double)__x); \
231 #define __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(__retty, __fn) \
232 template <typename __T1, typename __T2> \
233 __DEVICE__ typename __clang_cuda_enable_if< \
234 std::numeric_limits<__T1>::is_specialized && \
235 std::numeric_limits<__T2>::is_specialized, \
237 __fn(__T1 __x, __T2 __y) { \
238 return __fn((double)__x, (double)__y); \
302 #undef __CUDA_CLANG_FN_INTEGER_OVERLOAD_1
303 #undef __CUDA_CLANG_FN_INTEGER_OVERLOAD_2
307 template <
typename __T1,
typename __T2,
typename __T3>
309 std::numeric_limits<__T1>::is_specialized &&
310 std::numeric_limits<__T2>::is_specialized &&
311 std::numeric_limits<__T3>::is_specialized,
314 return std::fma((
double)__x, (
double)
__y, (
double)__z);
317 template <
typename __T>
324 template <
typename __T>
331 template <
typename __T1,
typename __T2>
333 std::numeric_limits<__T1>::is_specialized &&
334 std::numeric_limits<__T2>::is_specialized,
340 template <
typename __T>
347 template <
typename __T>
358 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
359 _LIBCPP_BEGIN_NAMESPACE_STD
362 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
363 _GLIBCXX_BEGIN_NAMESPACE_VERSION
499 #ifdef _LIBCPP_END_NAMESPACE_STD
500 _LIBCPP_END_NAMESPACE_STD
502 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
503 _GLIBCXX_END_NAMESPACE_VERSION
__DEVICE__ float sinh(float __x)
Compute hyperbolic sine.
__DEVICE__ __clang_cuda_enable_if< std::numeric_limits< __T1 >::is_specialized &&std::numeric_limits< __T2 >::is_specialized &&std::numeric_limits< __T3 >::is_specialized, double >::type fma(__T1 __x, __T2 __y, __T3 __z)
__DEVICE__ __clang_cuda_enable_if< std::numeric_limits< __T1 >::is_specialized &&std::numeric_limits< __T2 >::is_specialized, double >::type remquo(__T1 __x, __T2 __y, int *__quo)
__DEVICE__ float log10(float __x)
Compute a base 10 logarithm.
__DEVICE__ bool isunordered(float __x, float __y)
Test if arguments are unordered.
__DEVICE__ __clang_cuda_enable_if< std::numeric_limits< __T >::is_integer, double >::type scalbln(__T __x, long __exp)
__DEVICE__ bool isgreater(float __x, float __y)
Returns the component-wise compare of x > y.
__DEVICE__ float tanh(float __x)
Compute hyperbolic tangent.
__DEVICE__ bool islessgreater(float __x, float __y)
Returns the component-wise compare of (x < y) || (x > y) .
__DEVICE__ bool isnan(float __x)
Test for a NaN.
__DEVICE__ int fpclassify(float __x)
__DEVICE__ float sin(float __x)
Compute sine.
__DEVICE__ __clang_cuda_enable_if< std::numeric_limits< __T >::is_integer, double >::type scalbn(__T __x, int __exp)
__DEVICE__ bool isfinite(float __x)
Test for finite value.
__DEVICE__ float cos(float __x)
Compute cosine.
__DEVICE__ float floor(float __x)
Round to integral value using the round to -ve infinity rounding mode.
__DEVICE__ bool signbit(float __x)
Test for sign bit.
__DEVICE__ bool isinf(float __x)
Test for infinity value (+ve or -ve) .
__DEVICE__ float asin(float __x)
Arc sine function.
__DEVICE__ float fmod(float __x, float __y)
Modulus.
__DEVICE__ float acos(float __x)
Arc cosine function.
__DEVICE__ float pow(float __base, float __exp)
Compute x to the power y.
__DEVICE__ float ceil(float __x)
Round to integral value using the round to positive infinity rounding mode.
__DEVICE__ float tan(float __x)
Compute tangent.
__DEVICE__ float sqrt(float __x)
Compute square root.
__DEVICE__ float log(float __x)
Compute natural logarithm.
__DEVICE__ float modf(float __x, float *__iptr)
#define __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(__retty, __fn)
__DEVICE__ bool islessequal(float __x, float __y)
Returns the component-wise compare of x <= y.
__DEVICE__ float frexp(float __arg, int *__exp)
#define __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(__retty, __fn)
__DEVICE__ float exp(float __x)
Compute the base e exponential function of x.
__DEVICE__ long long abs(long long __n)
__DEVICE__ bool isless(float __x, float __y)
Returns the component-wise compare of x < y.
__DEVICE__ float fabs(float __x)
Compute absolute value of a floating-point number.
__DEVICE__ bool isnormal(float __x)
Test for a normal value.
__DEVICE__ float atan2(float __x, float __y)
Arc tangent of y / x.
__DEVICE__ float atan(float __x)
Arc tangent function.
__DEVICE__ float cosh(float __x)
Compute hyperbolic cosine.
__DEVICE__ float ldexp(float __arg, int __exp)
Multiply x by 2 to the power n.
__DEVICE__ bool isgreaterequal(float __x, float __y)
Returns the component-wise compare of x >= y.
__DEVICE__ int __isinff(float __a)
__DEVICE__ int __isnan(double __a)
__DEVICE__ int __isinf(double __a)
__DEVICE__ int __signbitf(float __a)
__DEVICE__ int __finitef(float __a)
__DEVICE__ int __signbitd(double __a)
__DEVICE__ int __isnanf(float __a)
__DEVICE__ int __isfinited(double __a)
__DEVICE__ double powi(double __a, int __b)
__DEVICE__ float fabsf(float __a)
__DEVICE__ float fmodf(float __a, float __b)
__DEVICE__ float remainderf(float __a, float __b)
__DEVICE__ float exp2f(float __a)
__DEVICE__ float acosf(float __a)
__DEVICE__ float fmaf(float __a, float __b, float __c)
__DEVICE__ float cbrtf(float __a)
__DEVICE__ float remquof(float __a, float __b, int *__c)
__DEVICE__ float tanf(float __a)
__DEVICE__ long labs(long __a)
__DEVICE__ float nextafterf(float __a, float __b)
__DEVICE__ float fmaxf(float __a, float __b)
__DEVICE__ long long llabs(long long __a)
__DEVICE__ float fminf(float __a, float __b)
__DEVICE__ float log2f(float __a)
__DEVICE__ float copysignf(float __a, float __b)
__DEVICE__ float truncf(float __a)
__DEVICE__ float fdimf(float __a, float __b)
__DEVICE__ long lrintf(float __a)
__DEVICE__ long long llrintf(float __a)
__DEVICE__ float cosf(float __a)
__DEVICE__ float sinf(float __a)
__DEVICE__ float logf(float __a)
__DEVICE__ float erff(float __a)
__DEVICE__ float floorf(float __f)
__DEVICE__ float ceilf(float __a)
__DEVICE__ float nearbyintf(float __a)
__DEVICE__ float atanf(float __a)
__DEVICE__ float atanhf(float __a)
__DEVICE__ float tanhf(float __a)
__DEVICE__ float rintf(float __a)
__DEVICE__ float atan2f(float __a, float __b)
__DEVICE__ float sinhf(float __a)
__DEVICE__ float acoshf(float __a)
__DEVICE__ float log10f(float __a)
__DEVICE__ float ldexpf(float __a, int __b)
__DEVICE__ float modff(float __a, float *__b)
__DEVICE__ float logbf(float __a)
__DEVICE__ float powif(float __a, int __b)
__DEVICE__ float coshf(float __a)
__DEVICE__ float asinhf(float __a)
__DEVICE__ float roundf(float __a)
__DEVICE__ long lroundf(float __a)
__DEVICE__ float scalbnf(float __a, int __b)
__DEVICE__ float erfcf(float __a)
__DEVICE__ int ilogbf(float __a)
__DEVICE__ float powf(float __a, float __b)
__DEVICE__ float frexpf(float __a, int *__b)
__DEVICE__ float sqrtf(float __a)
__DEVICE__ float expf(float __a)
__DEVICE__ float expm1f(float __a)
__DEVICE__ float scalblnf(float __a, long __b)
__DEVICE__ float tgammaf(float __a)
__DEVICE__ float log1pf(float __a)
__DEVICE__ float lgammaf(float __a)
__DEVICE__ long long llroundf(float __a)
__DEVICE__ float hypotf(float __a, float __b)
__DEVICE__ float asinf(float __a)
static __inline__ uint32_t uint32_t __y
const internal::VariadicAllOfMatcher< Type > type
Matches Types in the clang AST.
#define copysign(__x, __y)
#define nextafter(__x, __y)
#define remainder(__x, __y)