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