1 | |
2 | |
3 | |
4 | |
5 | |
6 | |
7 | |
8 | |
9 | |
10 | |
11 | |
12 | |
13 | |
14 | |
15 | |
16 | |
17 | |
18 | |
19 | |
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 | |
32 | |
33 | |
34 | |
35 | |
36 | |
37 | |
38 | |
39 | |
40 | |
41 | |
42 | |
43 | |
44 | |
45 | |
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); } |
64 | __DEVICE__ int fpclassify(float __x) { |
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 | |
77 | |
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 | |
83 | |
84 | |
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 | |
152 | |
153 | |
154 | |
155 | |
156 | |
157 | |
158 | |
159 | |
160 | |
161 | |
162 | |
163 | |
164 | |
165 | |
166 | |
167 | |
168 | |
169 | |
170 | template<bool __B, class __T = void> |
171 | struct __clang_cuda_enable_if {}; |
172 | |
173 | template <class __T> struct __clang_cuda_enable_if<true, __T> { |
174 | typedef __T type; |
175 | }; |
176 | |
177 | |
178 | |
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 | |
189 | |
190 | |
191 | |
192 | |
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 | |
203 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, acos) |
204 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, acosh) |
205 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, asin) |
206 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, asinh) |
207 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, atan) |
208 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, atan2); |
209 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, atanh) |
210 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, cbrt) |
211 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, ceil) |
212 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, copysign); |
213 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, cos) |
214 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, cosh) |
215 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, erf) |
216 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, erfc) |
217 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, exp) |
218 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, exp2) |
219 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, expm1) |
220 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, fabs) |
221 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fdim); |
222 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, floor) |
223 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fmax); |
224 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fmin); |
225 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fmod); |
226 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(int, fpclassify) |
227 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, hypot); |
228 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(int, ilogb) |
229 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, isfinite) |
230 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isgreater); |
231 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isgreaterequal); |
232 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, isinf); |
233 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isless); |
234 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, islessequal); |
235 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, islessgreater); |
236 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, isnan); |
237 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, isnormal) |
238 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isunordered); |
239 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, lgamma) |
240 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, log) |
241 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, log10) |
242 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, log1p) |
243 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, log2) |
244 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, logb) |
245 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long long, llrint) |
246 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long long, llround) |
247 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long, lrint) |
248 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(long, lround) |
249 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, nearbyint); |
250 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, nextafter); |
251 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, pow); |
252 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, remainder); |
253 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, rint); |
254 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, round); |
255 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, signbit) |
256 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, sin) |
257 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, sinh) |
258 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, sqrt) |
259 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, tan) |
260 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, tanh) |
261 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, tgamma) |
262 | __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, trunc); |
263 | |
264 | #undef __CUDA_CLANG_FN_INTEGER_OVERLOAD_1 |
265 | #undef __CUDA_CLANG_FN_INTEGER_OVERLOAD_2 |
266 | |
267 | |
268 | |
269 | template <typename __T1, typename __T2, typename __T3> |
270 | __DEVICE__ typename __clang_cuda_enable_if< |
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> |
280 | __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer, |
281 | double>::type |
282 | frexp(__T __x, int *__exp) { |
283 | return std::frexp((double)__x, __exp); |
284 | } |
285 | |
286 | template <typename __T> |
287 | __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer, |
288 | double>::type |
289 | ldexp(__T __x, int __exp) { |
290 | return std::ldexp((double)__x, __exp); |
291 | } |
292 | |
293 | template <typename __T1, typename __T2> |
294 | __DEVICE__ typename __clang_cuda_enable_if< |
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> |
303 | __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer, |
304 | double>::type |
305 | scalbln(__T __x, long __exp) { |
306 | return std::scalbln((double)__x, __exp); |
307 | } |
308 | |
309 | template <typename __T> |
310 | __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer, |
311 | double>::type |
312 | scalbn(__T __x, int __exp) { |
313 | return std::scalbn((double)__x, __exp); |
314 | } |
315 | |
316 | |
317 | |
318 | |
319 | |
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 | |
330 | using ::acos; |
331 | using ::acosh; |
332 | using ::asin; |
333 | using ::asinh; |
334 | using ::atan; |
335 | using ::atan2; |
336 | using ::atanh; |
337 | using ::cbrt; |
338 | using ::ceil; |
339 | using ::copysign; |
340 | using ::cos; |
341 | using ::cosh; |
342 | using ::erf; |
343 | using ::erfc; |
344 | using ::exp; |
345 | using ::exp2; |
346 | using ::expm1; |
347 | using ::fabs; |
348 | using ::fdim; |
349 | using ::floor; |
350 | using ::fma; |
351 | using ::fmax; |
352 | using ::fmin; |
353 | using ::fmod; |
354 | using ::fpclassify; |
355 | using ::frexp; |
356 | using ::hypot; |
357 | using ::ilogb; |
358 | using ::isfinite; |
359 | using ::isgreater; |
360 | using ::isgreaterequal; |
361 | using ::isless; |
362 | using ::islessequal; |
363 | using ::islessgreater; |
364 | using ::isnormal; |
365 | using ::isunordered; |
366 | using ::ldexp; |
367 | using ::lgamma; |
368 | using ::llrint; |
369 | using ::llround; |
370 | using ::log; |
371 | using ::log10; |
372 | using ::log1p; |
373 | using ::log2; |
374 | using ::logb; |
375 | using ::lrint; |
376 | using ::lround; |
377 | using ::nearbyint; |
378 | using ::nextafter; |
379 | using ::pow; |
380 | using ::remainder; |
381 | using ::remquo; |
382 | using ::rint; |
383 | using ::round; |
384 | using ::scalbln; |
385 | using ::scalbn; |
386 | using ::signbit; |
387 | using ::sin; |
388 | using ::sinh; |
389 | using ::sqrt; |
390 | using ::tan; |
391 | using ::tanh; |
392 | using ::tgamma; |
393 | using ::trunc; |
394 | |
395 | |
396 | |
397 | |
398 | #ifndef __GLIBCXX__ |
399 | using ::isinf; |
400 | using ::isnan; |
401 | #endif |
402 | |
403 | |
404 | |
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 | } |
468 | #endif |
469 | |
470 | #undef __DEVICE__ |
471 | |
472 | #endif |
473 | |