1 /*===---- __clang_hip_cmath.h - HIP cmath decls -----------------------------===
2 *
3 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 * See https://llvm.org/LICENSE.txt for license information.
5 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 *
7 *===-----------------------------------------------------------------------===
8 */
9
10 #ifndef __CLANG_HIP_CMATH_H__
11 #define __CLANG_HIP_CMATH_H__
12
13 #if !defined(__HIP__)
14 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
15 #endif
16
17 #if !defined(__HIPCC_RTC__)
18 #if defined(__cplusplus)
19 #include <limits>
20 #include <type_traits>
21 #include <utility>
22 #endif
23 #include <limits.h>
24 #include <stdint.h>
25 #endif // !defined(__HIPCC_RTC__)
26
27 #pragma push_macro("__DEVICE__")
28 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
29
30 // Start with functions that cannot be defined by DEF macros below.
31 #if defined(__cplusplus)
abs(double __x)32 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
abs(float __x)33 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
abs(long long __n)34 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
abs(long __n)35 __DEVICE__ long abs(long __n) { return ::labs(__n); }
fma(float __x,float __y,float __z)36 __DEVICE__ float fma(float __x, float __y, float __z) {
37 return ::fmaf(__x, __y, __z);
38 }
39 #if !defined(__HIPCC_RTC__)
40 // The value returned by fpclassify is platform dependent, therefore it is not
41 // supported by hipRTC.
fpclassify(float __x)42 __DEVICE__ int fpclassify(float __x) {
43 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
44 FP_ZERO, __x);
45 }
fpclassify(double __x)46 __DEVICE__ int fpclassify(double __x) {
47 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
48 FP_ZERO, __x);
49 }
50 #endif // !defined(__HIPCC_RTC__)
51
frexp(float __arg,int * __exp)52 __DEVICE__ float frexp(float __arg, int *__exp) {
53 return ::frexpf(__arg, __exp);
54 }
isfinite(float __x)55 __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
isfinite(double __x)56 __DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
isgreater(float __x,float __y)57 __DEVICE__ bool isgreater(float __x, float __y) {
58 return __builtin_isgreater(__x, __y);
59 }
isgreater(double __x,double __y)60 __DEVICE__ bool isgreater(double __x, double __y) {
61 return __builtin_isgreater(__x, __y);
62 }
isgreaterequal(float __x,float __y)63 __DEVICE__ bool isgreaterequal(float __x, float __y) {
64 return __builtin_isgreaterequal(__x, __y);
65 }
isgreaterequal(double __x,double __y)66 __DEVICE__ bool isgreaterequal(double __x, double __y) {
67 return __builtin_isgreaterequal(__x, __y);
68 }
isinf(float __x)69 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
isinf(double __x)70 __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
isless(float __x,float __y)71 __DEVICE__ bool isless(float __x, float __y) {
72 return __builtin_isless(__x, __y);
73 }
isless(double __x,double __y)74 __DEVICE__ bool isless(double __x, double __y) {
75 return __builtin_isless(__x, __y);
76 }
islessequal(float __x,float __y)77 __DEVICE__ bool islessequal(float __x, float __y) {
78 return __builtin_islessequal(__x, __y);
79 }
islessequal(double __x,double __y)80 __DEVICE__ bool islessequal(double __x, double __y) {
81 return __builtin_islessequal(__x, __y);
82 }
islessgreater(float __x,float __y)83 __DEVICE__ bool islessgreater(float __x, float __y) {
84 return __builtin_islessgreater(__x, __y);
85 }
islessgreater(double __x,double __y)86 __DEVICE__ bool islessgreater(double __x, double __y) {
87 return __builtin_islessgreater(__x, __y);
88 }
isnan(float __x)89 __DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
isnan(double __x)90 __DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
isnormal(float __x)91 __DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
isnormal(double __x)92 __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
isunordered(float __x,float __y)93 __DEVICE__ bool isunordered(float __x, float __y) {
94 return __builtin_isunordered(__x, __y);
95 }
isunordered(double __x,double __y)96 __DEVICE__ bool isunordered(double __x, double __y) {
97 return __builtin_isunordered(__x, __y);
98 }
modf(float __x,float * __iptr)99 __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
pow(float __base,int __iexp)100 __DEVICE__ float pow(float __base, int __iexp) {
101 return ::powif(__base, __iexp);
102 }
pow(double __base,int __iexp)103 __DEVICE__ double pow(double __base, int __iexp) {
104 return ::powi(__base, __iexp);
105 }
remquo(float __x,float __y,int * __quo)106 __DEVICE__ float remquo(float __x, float __y, int *__quo) {
107 return ::remquof(__x, __y, __quo);
108 }
scalbln(float __x,long int __n)109 __DEVICE__ float scalbln(float __x, long int __n) {
110 return ::scalblnf(__x, __n);
111 }
signbit(float __x)112 __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
signbit(double __x)113 __DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
114
115 // Notably missing above is nexttoward. We omit it because
116 // ocml doesn't provide an implementation, and we don't want to be in the
117 // business of implementing tricky libm functions in this header.
118
119 // Other functions.
fma(_Float16 __x,_Float16 __y,_Float16 __z)120 __DEVICE__ _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) {
121 return __ocml_fma_f16(__x, __y, __z);
122 }
pow(_Float16 __base,int __iexp)123 __DEVICE__ _Float16 pow(_Float16 __base, int __iexp) {
124 return __ocml_pown_f16(__base, __iexp);
125 }
126
127 // BEGIN DEF_FUN and HIP_OVERLOAD
128
129 // BEGIN DEF_FUN
130
131 #pragma push_macro("__DEF_FUN1")
132 #pragma push_macro("__DEF_FUN2")
133 #pragma push_macro("__DEF_FUN2_FI")
134
135 // Define cmath functions with float argument and returns __retty.
136 #define __DEF_FUN1(__retty, __func) \
137 __DEVICE__ \
138 __retty __func(float __x) { return __func##f(__x); }
139
140 // Define cmath functions with two float arguments and returns __retty.
141 #define __DEF_FUN2(__retty, __func) \
142 __DEVICE__ \
143 __retty __func(float __x, float __y) { return __func##f(__x, __y); }
144
145 // Define cmath functions with a float and an int argument and returns __retty.
146 #define __DEF_FUN2_FI(__retty, __func) \
147 __DEVICE__ \
148 __retty __func(float __x, int __y) { return __func##f(__x, __y); }
149
150 __DEF_FUN1(float, acos)
151 __DEF_FUN1(float, acosh)
152 __DEF_FUN1(float, asin)
153 __DEF_FUN1(float, asinh)
154 __DEF_FUN1(float, atan)
155 __DEF_FUN2(float, atan2)
156 __DEF_FUN1(float, atanh)
157 __DEF_FUN1(float, cbrt)
158 __DEF_FUN1(float, ceil)
159 __DEF_FUN2(float, copysign)
160 __DEF_FUN1(float, cos)
161 __DEF_FUN1(float, cosh)
162 __DEF_FUN1(float, erf)
163 __DEF_FUN1(float, erfc)
164 __DEF_FUN1(float, exp)
165 __DEF_FUN1(float, exp2)
166 __DEF_FUN1(float, expm1)
167 __DEF_FUN1(float, fabs)
168 __DEF_FUN2(float, fdim)
169 __DEF_FUN1(float, floor)
170 __DEF_FUN2(float, fmax)
171 __DEF_FUN2(float, fmin)
172 __DEF_FUN2(float, fmod)
173 __DEF_FUN2(float, hypot)
174 __DEF_FUN1(int, ilogb)
175 __DEF_FUN2_FI(float, ldexp)
176 __DEF_FUN1(float, lgamma)
177 __DEF_FUN1(float, log)
178 __DEF_FUN1(float, log10)
179 __DEF_FUN1(float, log1p)
180 __DEF_FUN1(float, log2)
181 __DEF_FUN1(float, logb)
182 __DEF_FUN1(long long, llrint)
183 __DEF_FUN1(long long, llround)
184 __DEF_FUN1(long, lrint)
185 __DEF_FUN1(long, lround)
186 __DEF_FUN1(float, nearbyint)
187 __DEF_FUN2(float, nextafter)
188 __DEF_FUN2(float, pow)
189 __DEF_FUN2(float, remainder)
190 __DEF_FUN1(float, rint)
191 __DEF_FUN1(float, round)
192 __DEF_FUN2_FI(float, scalbn)
193 __DEF_FUN1(float, sin)
194 __DEF_FUN1(float, sinh)
195 __DEF_FUN1(float, sqrt)
196 __DEF_FUN1(float, tan)
197 __DEF_FUN1(float, tanh)
198 __DEF_FUN1(float, tgamma)
199 __DEF_FUN1(float, trunc)
200
201 #pragma pop_macro("__DEF_FUN1")
202 #pragma pop_macro("__DEF_FUN2")
203 #pragma pop_macro("__DEF_FUN2_FI")
204
205 // END DEF_FUN
206
207 // BEGIN HIP_OVERLOAD
208
209 #pragma push_macro("__HIP_OVERLOAD1")
210 #pragma push_macro("__HIP_OVERLOAD2")
211
212 // __hip_enable_if::type is a type function which returns __T if __B is true.
213 template <bool __B, class __T = void> struct __hip_enable_if {};
214
215 template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
216
217 namespace __hip {
218 template <class _Tp> struct is_integral {
219 enum { value = 0 };
220 };
221 template <> struct is_integral<bool> {
222 enum { value = 1 };
223 };
224 template <> struct is_integral<char> {
225 enum { value = 1 };
226 };
227 template <> struct is_integral<signed char> {
228 enum { value = 1 };
229 };
230 template <> struct is_integral<unsigned char> {
231 enum { value = 1 };
232 };
233 template <> struct is_integral<wchar_t> {
234 enum { value = 1 };
235 };
236 template <> struct is_integral<short> {
237 enum { value = 1 };
238 };
239 template <> struct is_integral<unsigned short> {
240 enum { value = 1 };
241 };
242 template <> struct is_integral<int> {
243 enum { value = 1 };
244 };
245 template <> struct is_integral<unsigned int> {
246 enum { value = 1 };
247 };
248 template <> struct is_integral<long> {
249 enum { value = 1 };
250 };
251 template <> struct is_integral<unsigned long> {
252 enum { value = 1 };
253 };
254 template <> struct is_integral<long long> {
255 enum { value = 1 };
256 };
257 template <> struct is_integral<unsigned long long> {
258 enum { value = 1 };
259 };
260
261 // ToDo: specializes is_arithmetic<_Float16>
262 template <class _Tp> struct is_arithmetic {
263 enum { value = 0 };
264 };
265 template <> struct is_arithmetic<bool> {
266 enum { value = 1 };
267 };
268 template <> struct is_arithmetic<char> {
269 enum { value = 1 };
270 };
271 template <> struct is_arithmetic<signed char> {
272 enum { value = 1 };
273 };
274 template <> struct is_arithmetic<unsigned char> {
275 enum { value = 1 };
276 };
277 template <> struct is_arithmetic<wchar_t> {
278 enum { value = 1 };
279 };
280 template <> struct is_arithmetic<short> {
281 enum { value = 1 };
282 };
283 template <> struct is_arithmetic<unsigned short> {
284 enum { value = 1 };
285 };
286 template <> struct is_arithmetic<int> {
287 enum { value = 1 };
288 };
289 template <> struct is_arithmetic<unsigned int> {
290 enum { value = 1 };
291 };
292 template <> struct is_arithmetic<long> {
293 enum { value = 1 };
294 };
295 template <> struct is_arithmetic<unsigned long> {
296 enum { value = 1 };
297 };
298 template <> struct is_arithmetic<long long> {
299 enum { value = 1 };
300 };
301 template <> struct is_arithmetic<unsigned long long> {
302 enum { value = 1 };
303 };
304 template <> struct is_arithmetic<float> {
305 enum { value = 1 };
306 };
307 template <> struct is_arithmetic<double> {
308 enum { value = 1 };
309 };
310
311 struct true_type {
312 static const __constant__ bool value = true;
313 };
314 struct false_type {
315 static const __constant__ bool value = false;
316 };
317
318 template <typename __T, typename __U> struct is_same : public false_type {};
319 template <typename __T> struct is_same<__T, __T> : public true_type {};
320
321 template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
322
323 template <typename __T> typename add_rvalue_reference<__T>::type declval();
324
325 // decltype is only available in C++11 and above.
326 #if __cplusplus >= 201103L
327 // __hip_promote
328 template <class _Tp> struct __numeric_type {
329 static void __test(...);
330 static _Float16 __test(_Float16);
331 static float __test(float);
332 static double __test(char);
333 static double __test(int);
334 static double __test(unsigned);
335 static double __test(long);
336 static double __test(unsigned long);
337 static double __test(long long);
338 static double __test(unsigned long long);
339 static double __test(double);
340 // No support for long double, use double instead.
341 static double __test(long double);
342
343 typedef decltype(__test(declval<_Tp>())) type;
344 static const bool value = !is_same<type, void>::value;
345 };
346
347 template <> struct __numeric_type<void> { static const bool value = true; };
348
349 template <class _A1, class _A2 = void, class _A3 = void,
350 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
351 &&__numeric_type<_A3>::value>
352 class __promote_imp {
353 public:
354 static const bool value = false;
355 };
356
357 template <class _A1, class _A2, class _A3>
358 class __promote_imp<_A1, _A2, _A3, true> {
359 private:
360 typedef typename __promote_imp<_A1>::type __type1;
361 typedef typename __promote_imp<_A2>::type __type2;
362 typedef typename __promote_imp<_A3>::type __type3;
363
364 public:
365 typedef decltype(__type1() + __type2() + __type3()) type;
366 static const bool value = true;
367 };
368
369 template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
370 private:
371 typedef typename __promote_imp<_A1>::type __type1;
372 typedef typename __promote_imp<_A2>::type __type2;
373
374 public:
375 typedef decltype(__type1() + __type2()) type;
376 static const bool value = true;
377 };
378
379 template <class _A1> class __promote_imp<_A1, void, void, true> {
380 public:
381 typedef typename __numeric_type<_A1>::type type;
382 static const bool value = true;
383 };
384
385 template <class _A1, class _A2 = void, class _A3 = void>
386 class __promote : public __promote_imp<_A1, _A2, _A3> {};
387 #endif //__cplusplus >= 201103L
388 } // namespace __hip
389
390 // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
391 // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
392 // floor(double).
393 #define __HIP_OVERLOAD1(__retty, __fn) \
394 template <typename __T> \
395 __DEVICE__ \
396 typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type \
397 __fn(__T __x) { \
398 return ::__fn((double)__x); \
399 }
400
401 // __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double
402 // or integer argument to avoid compilation error due to ambibuity. e.g.
403 // max(5.0f, 6.0) is resolved with max(double, double).
404 #if __cplusplus >= 201103L
405 #define __HIP_OVERLOAD2(__retty, __fn) \
406 template <typename __T1, typename __T2> \
407 __DEVICE__ typename __hip_enable_if< \
408 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, \
409 typename __hip::__promote<__T1, __T2>::type>::type \
410 __fn(__T1 __x, __T2 __y) { \
411 typedef typename __hip::__promote<__T1, __T2>::type __result_type; \
412 return __fn((__result_type)__x, (__result_type)__y); \
413 }
414 #else
415 #define __HIP_OVERLOAD2(__retty, __fn) \
416 template <typename __T1, typename __T2> \
417 __DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \
418 __hip::is_arithmetic<__T2>::value, \
419 __retty>::type \
420 __fn(__T1 __x, __T2 __y) { \
421 return __fn((double)__x, (double)__y); \
422 }
423 #endif
424
425 __HIP_OVERLOAD1(double, acos)
426 __HIP_OVERLOAD1(double, acosh)
427 __HIP_OVERLOAD1(double, asin)
428 __HIP_OVERLOAD1(double, asinh)
429 __HIP_OVERLOAD1(double, atan)
430 __HIP_OVERLOAD2(double, atan2)
431 __HIP_OVERLOAD1(double, atanh)
432 __HIP_OVERLOAD1(double, cbrt)
433 __HIP_OVERLOAD1(double, ceil)
434 __HIP_OVERLOAD2(double, copysign)
435 __HIP_OVERLOAD1(double, cos)
436 __HIP_OVERLOAD1(double, cosh)
437 __HIP_OVERLOAD1(double, erf)
438 __HIP_OVERLOAD1(double, erfc)
439 __HIP_OVERLOAD1(double, exp)
440 __HIP_OVERLOAD1(double, exp2)
441 __HIP_OVERLOAD1(double, expm1)
442 __HIP_OVERLOAD1(double, fabs)
443 __HIP_OVERLOAD2(double, fdim)
444 __HIP_OVERLOAD1(double, floor)
445 __HIP_OVERLOAD2(double, fmax)
446 __HIP_OVERLOAD2(double, fmin)
447 __HIP_OVERLOAD2(double, fmod)
448 #if !defined(__HIPCC_RTC__)
449 __HIP_OVERLOAD1(int, fpclassify)
450 #endif // !defined(__HIPCC_RTC__)
451 __HIP_OVERLOAD2(double, hypot)
452 __HIP_OVERLOAD1(int, ilogb)
453 __HIP_OVERLOAD1(bool, isfinite)
454 __HIP_OVERLOAD2(bool, isgreater)
455 __HIP_OVERLOAD2(bool, isgreaterequal)
456 __HIP_OVERLOAD1(bool, isinf)
457 __HIP_OVERLOAD2(bool, isless)
458 __HIP_OVERLOAD2(bool, islessequal)
459 __HIP_OVERLOAD2(bool, islessgreater)
460 __HIP_OVERLOAD1(bool, isnan)
461 __HIP_OVERLOAD1(bool, isnormal)
462 __HIP_OVERLOAD2(bool, isunordered)
463 __HIP_OVERLOAD1(double, lgamma)
464 __HIP_OVERLOAD1(double, log)
465 __HIP_OVERLOAD1(double, log10)
466 __HIP_OVERLOAD1(double, log1p)
467 __HIP_OVERLOAD1(double, log2)
468 __HIP_OVERLOAD1(double, logb)
469 __HIP_OVERLOAD1(long long, llrint)
470 __HIP_OVERLOAD1(long long, llround)
471 __HIP_OVERLOAD1(long, lrint)
472 __HIP_OVERLOAD1(long, lround)
473 __HIP_OVERLOAD1(double, nearbyint)
474 __HIP_OVERLOAD2(double, nextafter)
475 __HIP_OVERLOAD2(double, pow)
476 __HIP_OVERLOAD2(double, remainder)
477 __HIP_OVERLOAD1(double, rint)
478 __HIP_OVERLOAD1(double, round)
479 __HIP_OVERLOAD1(bool, signbit)
480 __HIP_OVERLOAD1(double, sin)
481 __HIP_OVERLOAD1(double, sinh)
482 __HIP_OVERLOAD1(double, sqrt)
483 __HIP_OVERLOAD1(double, tan)
484 __HIP_OVERLOAD1(double, tanh)
485 __HIP_OVERLOAD1(double, tgamma)
486 __HIP_OVERLOAD1(double, trunc)
487
488 // Overload these but don't add them to std, they are not part of cmath.
489 __HIP_OVERLOAD2(double, max)
490 __HIP_OVERLOAD2(double, min)
491
492 // Additional Overloads that don't quite match HIP_OVERLOAD.
493 #if __cplusplus >= 201103L
494 template <typename __T1, typename __T2, typename __T3>
495 __DEVICE__ typename __hip_enable_if<
496 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
497 __hip::is_arithmetic<__T3>::value,
498 typename __hip::__promote<__T1, __T2, __T3>::type>::type
499 fma(__T1 __x, __T2 __y, __T3 __z) {
500 typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
501 return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);
502 }
503 #else
504 template <typename __T1, typename __T2, typename __T3>
505 __DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
506 __hip::is_arithmetic<__T2>::value &&
507 __hip::is_arithmetic<__T3>::value,
508 double>::type
509 fma(__T1 __x, __T2 __y, __T3 __z) {
510 return ::fma((double)__x, (double)__y, (double)__z);
511 }
512 #endif
513
514 template <typename __T>
515 __DEVICE__
516 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
517 frexp(__T __x, int *__exp) {
518 return ::frexp((double)__x, __exp);
519 }
520
521 template <typename __T>
522 __DEVICE__
523 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
524 ldexp(__T __x, int __exp) {
525 return ::ldexp((double)__x, __exp);
526 }
527
528 template <typename __T>
529 __DEVICE__
530 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
531 modf(__T __x, double *__exp) {
532 return ::modf((double)__x, __exp);
533 }
534
535 #if __cplusplus >= 201103L
536 template <typename __T1, typename __T2>
537 __DEVICE__
538 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
539 __hip::is_arithmetic<__T2>::value,
540 typename __hip::__promote<__T1, __T2>::type>::type
541 remquo(__T1 __x, __T2 __y, int *__quo) {
542 typedef typename __hip::__promote<__T1, __T2>::type __result_type;
543 return ::remquo((__result_type)__x, (__result_type)__y, __quo);
544 }
545 #else
546 template <typename __T1, typename __T2>
547 __DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
548 __hip::is_arithmetic<__T2>::value,
549 double>::type
550 remquo(__T1 __x, __T2 __y, int *__quo) {
551 return ::remquo((double)__x, (double)__y, __quo);
552 }
553 #endif
554
555 template <typename __T>
556 __DEVICE__
557 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
558 scalbln(__T __x, long int __exp) {
559 return ::scalbln((double)__x, __exp);
560 }
561
562 template <typename __T>
563 __DEVICE__
564 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
565 scalbn(__T __x, int __exp) {
566 return ::scalbn((double)__x, __exp);
567 }
568
569 #pragma pop_macro("__HIP_OVERLOAD1")
570 #pragma pop_macro("__HIP_OVERLOAD2")
571
572 // END HIP_OVERLOAD
573
574 // END DEF_FUN and HIP_OVERLOAD
575
576 #endif // defined(__cplusplus)
577
578 // Define these overloads inside the namespace our standard library uses.
579 #if !defined(__HIPCC_RTC__)
580 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
581 _LIBCPP_BEGIN_NAMESPACE_STD
582 #else
583 namespace std {
584 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
585 _GLIBCXX_BEGIN_NAMESPACE_VERSION
586 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
587 #endif // _LIBCPP_BEGIN_NAMESPACE_STD
588
589 // Pull the new overloads we defined above into namespace std.
590 // using ::abs; - This may be considered for C++.
591 using ::acos;
592 using ::acosh;
593 using ::asin;
594 using ::asinh;
595 using ::atan;
596 using ::atan2;
597 using ::atanh;
598 using ::cbrt;
599 using ::ceil;
600 using ::copysign;
601 using ::cos;
602 using ::cosh;
603 using ::erf;
604 using ::erfc;
605 using ::exp;
606 using ::exp2;
607 using ::expm1;
608 using ::fabs;
609 using ::fdim;
610 using ::floor;
611 using ::fma;
612 using ::fmax;
613 using ::fmin;
614 using ::fmod;
615 using ::fpclassify;
616 using ::frexp;
617 using ::hypot;
618 using ::ilogb;
619 using ::isfinite;
620 using ::isgreater;
621 using ::isgreaterequal;
622 using ::isless;
623 using ::islessequal;
624 using ::islessgreater;
625 using ::isnormal;
626 using ::isunordered;
627 using ::ldexp;
628 using ::lgamma;
629 using ::llrint;
630 using ::llround;
631 using ::log;
632 using ::log10;
633 using ::log1p;
634 using ::log2;
635 using ::logb;
636 using ::lrint;
637 using ::lround;
638 using ::modf;
639 // using ::nan; - This may be considered for C++.
640 // using ::nanf; - This may be considered for C++.
641 // using ::nanl; - This is not yet defined.
642 using ::nearbyint;
643 using ::nextafter;
644 // using ::nexttoward; - Omit this since we do not have a definition.
645 using ::pow;
646 using ::remainder;
647 using ::remquo;
648 using ::rint;
649 using ::round;
650 using ::scalbln;
651 using ::scalbn;
652 using ::signbit;
653 using ::sin;
654 using ::sinh;
655 using ::sqrt;
656 using ::tan;
657 using ::tanh;
658 using ::tgamma;
659 using ::trunc;
660
661 // Well this is fun: We need to pull these symbols in for libc++, but we can't
662 // pull them in with libstdc++, because its ::isinf and ::isnan are different
663 // than its std::isinf and std::isnan.
664 #ifndef __GLIBCXX__
665 using ::isinf;
666 using ::isnan;
667 #endif
668
669 // Finally, pull the "foobarf" functions that HIP defines into std.
670 using ::acosf;
671 using ::acoshf;
672 using ::asinf;
673 using ::asinhf;
674 using ::atan2f;
675 using ::atanf;
676 using ::atanhf;
677 using ::cbrtf;
678 using ::ceilf;
679 using ::copysignf;
680 using ::cosf;
681 using ::coshf;
682 using ::erfcf;
683 using ::erff;
684 using ::exp2f;
685 using ::expf;
686 using ::expm1f;
687 using ::fabsf;
688 using ::fdimf;
689 using ::floorf;
690 using ::fmaf;
691 using ::fmaxf;
692 using ::fminf;
693 using ::fmodf;
694 using ::frexpf;
695 using ::hypotf;
696 using ::ilogbf;
697 using ::ldexpf;
698 using ::lgammaf;
699 using ::llrintf;
700 using ::llroundf;
701 using ::log10f;
702 using ::log1pf;
703 using ::log2f;
704 using ::logbf;
705 using ::logf;
706 using ::lrintf;
707 using ::lroundf;
708 using ::modff;
709 using ::nearbyintf;
710 using ::nextafterf;
711 // using ::nexttowardf; - Omit this since we do not have a definition.
712 using ::powf;
713 using ::remainderf;
714 using ::remquof;
715 using ::rintf;
716 using ::roundf;
717 using ::scalblnf;
718 using ::scalbnf;
719 using ::sinf;
720 using ::sinhf;
721 using ::sqrtf;
722 using ::tanf;
723 using ::tanhf;
724 using ::tgammaf;
725 using ::truncf;
726
727 #ifdef _LIBCPP_END_NAMESPACE_STD
728 _LIBCPP_END_NAMESPACE_STD
729 #else
730 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
731 _GLIBCXX_END_NAMESPACE_VERSION
732 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
733 } // namespace std
734 #endif // _LIBCPP_END_NAMESPACE_STD
735 #endif // !defined(__HIPCC_RTC__)
736
737 // Define device-side math functions from <ymath.h> on MSVC.
738 #if !defined(__HIPCC_RTC__)
739 #if defined(_MSC_VER)
740
741 // Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers.
742 // But, from VS2019, it's only included in `<complex>`. Need to include
743 // `<ymath.h>` here to ensure C functions declared there won't be markded as
744 // `__host__` and `__device__` through `<complex>` wrapper.
745 #include <ymath.h>
746
747 #if defined(__cplusplus)
748 extern "C" {
749 #endif // defined(__cplusplus)
750 __DEVICE__ __attribute__((overloadable)) double _Cosh(double x, double y) {
751 return cosh(x) * y;
752 }
753 __DEVICE__ __attribute__((overloadable)) float _FCosh(float x, float y) {
754 return coshf(x) * y;
755 }
756 __DEVICE__ __attribute__((overloadable)) short _Dtest(double *p) {
757 return fpclassify(*p);
758 }
759 __DEVICE__ __attribute__((overloadable)) short _FDtest(float *p) {
760 return fpclassify(*p);
761 }
762 __DEVICE__ __attribute__((overloadable)) double _Sinh(double x, double y) {
763 return sinh(x) * y;
764 }
765 __DEVICE__ __attribute__((overloadable)) float _FSinh(float x, float y) {
766 return sinhf(x) * y;
767 }
768 #if defined(__cplusplus)
769 }
770 #endif // defined(__cplusplus)
771 #endif // defined(_MSC_VER)
772 #endif // !defined(__HIPCC_RTC__)
773
774 #pragma pop_macro("__DEVICE__")
775
776 #endif // __CLANG_HIP_CMATH_H__
777