clang  6.0.0svn
__clang_cuda_cmath.h
Go to the documentation of this file.
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 
29 #include <limits>
30 
31 // CUDA lets us use various std math functions on the device side. This file
32 // works in concert with __clang_cuda_math_forward_declares.h to make this work.
33 //
34 // Specifically, the forward-declares header declares __device__ overloads for
35 // these functions in the global namespace, then pulls them into namespace std
36 // with 'using' statements. Then this file implements those functions, after
37 // their implementations have been pulled in.
38 //
39 // It's important that we declare the functions in the global namespace and pull
40 // them into namespace std with using statements, as opposed to simply declaring
41 // these functions in namespace std, because our device functions need to
42 // overload the standard library functions, which may be declared in the global
43 // namespace or in std, depending on the degree of conformance of the stdlib
44 // implementation. Declaring in the global namespace and pulling into namespace
45 // std covers all of the known knowns.
46 
47 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
48 
49 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
50 __DEVICE__ long abs(long __n) { return ::labs(__n); }
51 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
52 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
53 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
54 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
55 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
56 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
57 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
58 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
59 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
60 __DEVICE__ float exp(float __x) { return ::expf(__x); }
61 __DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
62 __DEVICE__ float floor(float __x) { return ::floorf(__x); }
63 __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
65  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
66  FP_ZERO, __x);
67 }
68 __DEVICE__ int fpclassify(double __x) {
69  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
70  FP_ZERO, __x);
71 }
72 __DEVICE__ float frexp(float __arg, int *__exp) {
73  return ::frexpf(__arg, __exp);
74 }
75 
76 // For inscrutable reasons, the CUDA headers define these functions for us on
77 // Windows.
78 #ifndef _MSC_VER
79 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
80 __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
81 __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
82 // For inscrutable reasons, __finite(), the double-precision version of
83 // __finitef, does not exist when compiling for MacOS. __isfinited is available
84 // everywhere and is just as good.
85 __DEVICE__ bool isfinite(double __x) { return ::__isfinited(__x); }
86 __DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
87 __DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
88 #endif
89 
90 __DEVICE__ bool isgreater(float __x, float __y) {
91  return __builtin_isgreater(__x, __y);
92 }
93 __DEVICE__ bool isgreater(double __x, double __y) {
94  return __builtin_isgreater(__x, __y);
95 }
96 __DEVICE__ bool isgreaterequal(float __x, float __y) {
97  return __builtin_isgreaterequal(__x, __y);
98 }
99 __DEVICE__ bool isgreaterequal(double __x, double __y) {
100  return __builtin_isgreaterequal(__x, __y);
101 }
102 __DEVICE__ bool isless(float __x, float __y) {
103  return __builtin_isless(__x, __y);
104 }
105 __DEVICE__ bool isless(double __x, double __y) {
106  return __builtin_isless(__x, __y);
107 }
108 __DEVICE__ bool islessequal(float __x, float __y) {
109  return __builtin_islessequal(__x, __y);
110 }
111 __DEVICE__ bool islessequal(double __x, double __y) {
112  return __builtin_islessequal(__x, __y);
113 }
114 __DEVICE__ bool islessgreater(float __x, float __y) {
115  return __builtin_islessgreater(__x, __y);
116 }
117 __DEVICE__ bool islessgreater(double __x, double __y) {
118  return __builtin_islessgreater(__x, __y);
119 }
120 __DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
121 __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
122 __DEVICE__ bool isunordered(float __x, float __y) {
123  return __builtin_isunordered(__x, __y);
124 }
125 __DEVICE__ bool isunordered(double __x, double __y) {
126  return __builtin_isunordered(__x, __y);
127 }
128 __DEVICE__ float ldexp(float __arg, int __exp) {
129  return ::ldexpf(__arg, __exp);
130 }
131 __DEVICE__ float log(float __x) { return ::logf(__x); }
132 __DEVICE__ float log10(float __x) { return ::log10f(__x); }
133 __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
134 __DEVICE__ float pow(float __base, float __exp) {
135  return ::powf(__base, __exp);
136 }
137 __DEVICE__ float pow(float __base, int __iexp) {
138  return ::powif(__base, __iexp);
139 }
140 __DEVICE__ double pow(double __base, int __iexp) {
141  return ::powi(__base, __iexp);
142 }
143 __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
144 __DEVICE__ bool signbit(double __x) { return ::__signbitd(__x); }
145 __DEVICE__ float sin(float __x) { return ::sinf(__x); }
146 __DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
147 __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
148 __DEVICE__ float tan(float __x) { return ::tanf(__x); }
149 __DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
150 
151 // Notably missing above is nexttoward. We omit it because
152 // libdevice doesn't provide an implementation, and we don't want to be in the
153 // business of implementing tricky libm functions in this header.
154 
155 // Now we've defined everything we promised we'd define in
156 // __clang_cuda_math_forward_declares.h. We need to do two additional things to
157 // fix up our math functions.
158 //
159 // 1) Define __device__ overloads for e.g. sin(int). The CUDA headers define
160 // only sin(float) and sin(double), which means that e.g. sin(0) is
161 // ambiguous.
162 //
163 // 2) Pull the __device__ overloads of "foobarf" math functions into namespace
164 // std. These are defined in the CUDA headers in the global namespace,
165 // independent of everything else we've done here.
166 
167 // We can't use std::enable_if, because we want to be pre-C++11 compatible. But
168 // we go ahead and unconditionally define functions that are only available when
169 // compiling for C++11 to match the behavior of the CUDA headers.
170 template<bool __B, class __T = void>
172 
173 template <class __T> struct __clang_cuda_enable_if<true, __T> {
174  typedef __T type;
175 };
176 
177 // Defines an overload of __fn that accepts one integral argument, calls
178 // __fn((double)x), and returns __retty.
179 #define __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(__retty, __fn) \
180  template <typename __T> \
181  __DEVICE__ \
182  typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer, \
183  __retty>::type \
184  __fn(__T __x) { \
185  return ::__fn((double)__x); \
186  }
187 
188 // Defines an overload of __fn that accepts one two arithmetic arguments, calls
189 // __fn((double)x, (double)y), and returns a double.
190 //
191 // Note this is different from OVERLOAD_1, which generates an overload that
192 // accepts only *integral* arguments.
193 #define __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(__retty, __fn) \
194  template <typename __T1, typename __T2> \
195  __DEVICE__ typename __clang_cuda_enable_if< \
196  std::numeric_limits<__T1>::is_specialized && \
197  std::numeric_limits<__T2>::is_specialized, \
198  __retty>::type \
199  __fn(__T1 __x, __T2 __y) { \
200  return __fn((double)__x, (double)__y); \
201  }
202 
263 
264 #undef __CUDA_CLANG_FN_INTEGER_OVERLOAD_1
265 #undef __CUDA_CLANG_FN_INTEGER_OVERLOAD_2
266 
267 // Overloads for functions that don't match the patterns expected by
268 // __CUDA_CLANG_FN_INTEGER_OVERLOAD_{1,2}.
269 template <typename __T1, typename __T2, typename __T3>
271  std::numeric_limits<__T1>::is_specialized &&
272  std::numeric_limits<__T2>::is_specialized &&
273  std::numeric_limits<__T3>::is_specialized,
274  double>::type
275 fma(__T1 __x, __T2 __y, __T3 __z) {
276  return std::fma((double)__x, (double)__y, (double)__z);
277 }
278 
279 template <typename __T>
281  double>::type
282 frexp(__T __x, int *__exp) {
283  return std::frexp((double)__x, __exp);
284 }
285 
286 template <typename __T>
288  double>::type
289 ldexp(__T __x, int __exp) {
290  return std::ldexp((double)__x, __exp);
291 }
292 
293 template <typename __T1, typename __T2>
295  std::numeric_limits<__T1>::is_specialized &&
296  std::numeric_limits<__T2>::is_specialized,
297  double>::type
298 remquo(__T1 __x, __T2 __y, int *__quo) {
299  return std::remquo((double)__x, (double)__y, __quo);
300 }
301 
302 template <typename __T>
304  double>::type
305 scalbln(__T __x, long __exp) {
306  return std::scalbln((double)__x, __exp);
307 }
308 
309 template <typename __T>
311  double>::type
312 scalbn(__T __x, int __exp) {
313  return std::scalbn((double)__x, __exp);
314 }
315 
316 // We need to define these overloads in exactly the namespace our standard
317 // library uses (including the right inline namespace), otherwise they won't be
318 // picked up by other functions in the standard library (e.g. functions in
319 // <complex>). Thus the ugliness below.
320 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
321 _LIBCPP_BEGIN_NAMESPACE_STD
322 #else
323 namespace std {
324 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
325 _GLIBCXX_BEGIN_NAMESPACE_VERSION
326 #endif
327 #endif
328 
329 // Pull the new overloads we defined above into namespace std.
394 
395 // Well this is fun: We need to pull these symbols in for libc++, but we can't
396 // pull them in with libstdc++, because its ::isinf and ::isnan are different
397 // than its std::isinf and std::isnan.
398 #ifndef __GLIBCXX__
401 #endif
402 
403 // Finally, pull the "foobarf" functions that CUDA defines in its headers into
404 // namespace std.
405 using ::acosf;
406 using ::acoshf;
407 using ::asinf;
408 using ::asinhf;
409 using ::atan2f;
410 using ::atanf;
411 using ::atanhf;
412 using ::cbrtf;
413 using ::ceilf;
414 using ::copysignf;
415 using ::cosf;
416 using ::coshf;
417 using ::erfcf;
418 using ::erff;
419 using ::exp2f;
420 using ::expf;
421 using ::expm1f;
422 using ::fabsf;
423 using ::fdimf;
424 using ::floorf;
425 using ::fmaf;
426 using ::fmaxf;
427 using ::fminf;
428 using ::fmodf;
429 using ::frexpf;
430 using ::hypotf;
431 using ::ilogbf;
432 using ::ldexpf;
433 using ::lgammaf;
434 using ::llrintf;
435 using ::llroundf;
436 using ::log10f;
437 using ::log1pf;
438 using ::log2f;
439 using ::logbf;
440 using ::logf;
441 using ::lrintf;
442 using ::lroundf;
443 using ::modff;
444 using ::nearbyintf;
445 using ::nextafterf;
446 using ::powf;
447 using ::remainderf;
448 using ::remquof;
449 using ::rintf;
450 using ::roundf;
451 using ::scalblnf;
452 using ::scalbnf;
453 using ::sinf;
454 using ::sinhf;
455 using ::sqrtf;
456 using ::tanf;
457 using ::tanhf;
458 using ::tgammaf;
459 using ::truncf;
460 
461 #ifdef _LIBCPP_END_NAMESPACE_STD
462 _LIBCPP_END_NAMESPACE_STD
463 #else
464 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
465 _GLIBCXX_END_NAMESPACE_VERSION
466 #endif
467 } // namespace std
468 #endif
469 
470 #undef __DEVICE__
471 
472 #endif
static __inline unsigned char unsigned int unsigned int __y
Definition: adxintrin.h:36
__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__ bool signbit(float __x)
Test for sign bit.
__DEVICE__ float sinh(float __x)
Compute hyperbolic sine.
#define log1p(__x)
Definition: tgmath.h:969
__DEVICE__ float atan2(float __x, float __y)
Arc tangent of y / x.
#define trunc(__x)
Definition: tgmath.h:1232
__DEVICE__ long long abs(long long __n)
#define atanh(__x)
Definition: tgmath.h:244
#define log2(__x)
Definition: tgmath.h:986
__DEVICE__ float floor(float __x)
Round to integral value using the round to -ve infinity rounding mode.
#define erfc(__x)
Definition: tgmath.h:669
__DEVICE__ float sqrt(float __x)
Compute square root.
__DEVICE__ float modf(float __x, float *__iptr)
Decompose a floating-point number.
__DEVICE__ float asin(float __x)
Arc sine function.
__DEVICE__ float ceil(float __x)
Round to integral value using the round to positive infinity rounding mode.
#define fmax(__x, __y)
Definition: tgmath.h:778
__DEVICE__ bool isinf(float __x)
Test for infinity value (+ve or -ve) .
#define __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(__retty, __fn)
#define __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(__retty, __fn)
__DEVICE__ float acos(float __x)
Arc cosine function.
__DEVICE__ float log(float __x)
Compute natural logarithm.
Definition: Format.h:1900
#define fmin(__x, __y)
Definition: tgmath.h:796
#define logb(__x)
Definition: tgmath.h:1003
__DEVICE__ int fpclassify(float __x)
__DEVICE__ float sin(float __x)
Compute sine.
#define __DEVICE__
__DEVICE__ long labs(long)
#define asinh(__x)
Definition: tgmath.h:215
#define remainder(__x, __y)
Definition: tgmath.h:1106
__DEVICE__ bool isfinite(float __x)
Test for finite value.
#define tgamma(__x)
Definition: tgmath.h:1215
__DEVICE__ float fmod(float __x, float __y)
Modulus.
__DEVICE__ float tan(float __x)
Compute tangent.
#define lrint(__x)
Definition: tgmath.h:1020
__DEVICE__ bool islessequal(float __x, float __y)
Returns the component-wise compare of x <= y.
#define cbrt(__x)
Definition: tgmath.h:600
__DEVICE__ float ldexp(float __arg, int __exp)
Multiply x by 2 to the power n.
#define copysign(__x, __y)
Definition: tgmath.h:634
__DEVICE__ float fabs(float __x)
Compute absolute value of a floating-point number.
__DEVICE__ bool isunordered(float __x, float __y)
Test if arguments are unordered.
__DEVICE__ float frexp(float __arg, int *__exp)
Extract mantissa and exponent from x.
#define lgamma(__x)
Definition: tgmath.h:901
#define expm1(__x)
Definition: tgmath.h:703
#define acosh(__x)
Definition: tgmath.h:186
static __inline unsigned char unsigned int __x
Definition: adxintrin.h:36
__DEVICE__ __clang_cuda_enable_if< std::numeric_limits< __T >::is_integer, double >::type scalbln(__T __x, long __exp)
__DEVICE__ bool isgreaterequal(float __x, float __y)
Returns the component-wise compare of x >= y.
#define rint(__x)
Definition: tgmath.h:1147
#define fdim(__x, __y)
Definition: tgmath.h:720
#define round(__x)
Definition: tgmath.h:1164
#define exp2(__x)
Definition: tgmath.h:686
#define ilogb(__x)
Definition: tgmath.h:867
#define lround(__x)
Definition: tgmath.h:1037
__DEVICE__ float cos(float __x)
Compute cosine.
__DEVICE__ float log10(float __x)
Compute a base 10 logarithm.
__DEVICE__ bool isnan(float __x)
Test for a NaN.
__DEVICE__ bool isgreater(float __x, float __y)
Returns the component-wise compare of x > y.
__DEVICE__ float atan(float __x)
Arc tangent function.
__DEVICE__ bool islessgreater(float __x, float __y)
Returns the component-wise compare of (x < y) || (x > y) .
__DEVICE__ long long llround(float)
__DEVICE__ long long llabs(long long)
__DEVICE__ __clang_cuda_enable_if< std::numeric_limits< __T >::is_integer, double >::type scalbn(__T __x, int __exp)
#define erf(__x)
Definition: tgmath.h:652
#define llrint(__x)
Definition: tgmath.h:918
#define nextafter(__x, __y)
Definition: tgmath.h:1071
#define nearbyint(__x)
Definition: tgmath.h:1054
__DEVICE__ float cosh(float __x)
Compute hyperbolic cosine.
__DEVICE__ float pow(float __base, float __exp)
Compute x to the power y.
__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__ bool isless(float __x, float __y)
Returns the component-wise compare of x < y.
__DEVICE__ float tanh(float __x)
Compute hyperbolic tangent.
#define true
Definition: stdbool.h:32
__DEVICE__ float exp(float __x)
Compute the base e exponential function of x.
#define hypot(__x, __y)
Definition: tgmath.h:849
__DEVICE__ bool isnormal(float __x)
Test for a normal value.