Warning: This file is not a C or C++ file. It does not have highlighting.

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__
40__DEVICE__ __CONSTEXPR__ float fabs(float __x) { return ::fabsf(__x); }
41__DEVICE__ __CONSTEXPR__ float sin(float __x) { return ::sinf(__x); }
42__DEVICE__ __CONSTEXPR__ float cos(float __x) { return ::cosf(__x); }
43#endif
44__DEVICE__ __CONSTEXPR__ double abs(double __x) { return ::fabs(__x); }
45__DEVICE__ __CONSTEXPR__ float abs(float __x) { return ::fabsf(__x); }
46__DEVICE__ __CONSTEXPR__ long long abs(long long __n) { return ::llabs(__n); }
47__DEVICE__ __CONSTEXPR__ long abs(long __n) { return ::labs(__n); }
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.
54__DEVICE__ __CONSTEXPR__ int fpclassify(float __x) {
55 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
56 FP_ZERO, __x);
57}
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
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
86__DEVICE__ __CONSTEXPR__ int isinf(float __x) { return ::__isinff(__x); }
87__DEVICE__ __CONSTEXPR__ int isinf(double __x) { return ::__isinf(__x); }
88__DEVICE__ __CONSTEXPR__ int isfinite(float __x) { return ::__finitef(__x); }
89__DEVICE__ __CONSTEXPR__ int isfinite(double __x) { return ::__finite(__x); }
90__DEVICE__ __CONSTEXPR__ int isnan(float __x) { return ::__isnanf(__x); }
91__DEVICE__ __CONSTEXPR__ int isnan(double __x) { return ::__isnan(__x); }
92
93#pragma omp end declare variant
94#endif // defined(__OPENMP_AMDGCN__)
95
96__DEVICE__ __CONSTEXPR__ bool isinf(float __x) { return ::__isinff(__x); }
97__DEVICE__ __CONSTEXPR__ bool isinf(double __x) { return ::__isinf(__x); }
98__DEVICE__ __CONSTEXPR__ bool isfinite(float __x) { return ::__finitef(__x); }
99__DEVICE__ __CONSTEXPR__ bool isfinite(double __x) { return ::__finite(__x); }
100__DEVICE__ __CONSTEXPR__ bool isnan(float __x) { return ::__isnanf(__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
107__DEVICE__ __CONSTEXPR__ bool isgreater(float __x, float __y) {
108 return __builtin_isgreater(__x, __y);
109}
110__DEVICE__ __CONSTEXPR__ bool isgreater(double __x, double __y) {
111 return __builtin_isgreater(__x, __y);
112}
113__DEVICE__ __CONSTEXPR__ bool isgreaterequal(float __x, float __y) {
114 return __builtin_isgreaterequal(__x, __y);
115}
116__DEVICE__ __CONSTEXPR__ bool isgreaterequal(double __x, double __y) {
117 return __builtin_isgreaterequal(__x, __y);
118}
119__DEVICE__ __CONSTEXPR__ bool isless(float __x, float __y) {
120 return __builtin_isless(__x, __y);
121}
122__DEVICE__ __CONSTEXPR__ bool isless(double __x, double __y) {
123 return __builtin_isless(__x, __y);
124}
125__DEVICE__ __CONSTEXPR__ bool islessequal(float __x, float __y) {
126 return __builtin_islessequal(__x, __y);
127}
128__DEVICE__ __CONSTEXPR__ bool islessequal(double __x, double __y) {
129 return __builtin_islessequal(__x, __y);
130}
131__DEVICE__ __CONSTEXPR__ bool islessgreater(float __x, float __y) {
132 return __builtin_islessgreater(__x, __y);
133}
134__DEVICE__ __CONSTEXPR__ bool islessgreater(double __x, double __y) {
135 return __builtin_islessgreater(__x, __y);
136}
137__DEVICE__ __CONSTEXPR__ bool isnormal(float __x) {
138 return __builtin_isnormal(__x);
139}
140__DEVICE__ __CONSTEXPR__ bool isnormal(double __x) {
141 return __builtin_isnormal(__x);
142}
143__DEVICE__ __CONSTEXPR__ bool isunordered(float __x, float __y) {
144 return __builtin_isunordered(__x, __y);
145}
146__DEVICE__ __CONSTEXPR__ bool isunordered(double __x, double __y) {
147 return __builtin_isunordered(__x, __y);
148}
149__DEVICE__ __CONSTEXPR__ float modf(float __x, float *__iptr) {
150 return ::modff(__x, __iptr);
151}
152__DEVICE__ __CONSTEXPR__ float pow(float __base, int __iexp) {
153 return ::powif(__base, __iexp);
154}
155__DEVICE__ __CONSTEXPR__ double pow(double __base, int __iexp) {
156 return ::powi(__base, __iexp);
157}
158__DEVICE__ __CONSTEXPR__ float remquo(float __x, float __y, int *__quo) {
159 return ::remquof(__x, __y, __quo);
160}
161__DEVICE__ __CONSTEXPR__ float scalbln(float __x, long int __n) {
162 return ::scalblnf(__x, __n);
163}
164__DEVICE__ __CONSTEXPR__ bool signbit(float __x) { return ::__signbitf(__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.
172__DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y,
173 _Float16 __z) {
174 return __builtin_fmaf16(__x, __y, __z);
175}
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.
268template <bool __B, class __T = void> struct __hip_enable_if {};
269
270template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
271
272namespace __hip {
273template <class _Tp> struct is_integral {
274 enum { value = 0 };
275};
276template <> struct is_integral<bool> {
277 enum { value = 1 };
278};
279template <> struct is_integral<char> {
280 enum { value = 1 };
281};
282template <> struct is_integral<signed char> {
283 enum { value = 1 };
284};
285template <> struct is_integral<unsigned char> {
286 enum { value = 1 };
287};
288template <> struct is_integral<wchar_t> {
289 enum { value = 1 };
290};
291template <> struct is_integral<short> {
292 enum { value = 1 };
293};
294template <> struct is_integral<unsigned short> {
295 enum { value = 1 };
296};
297template <> struct is_integral<int> {
298 enum { value = 1 };
299};
300template <> struct is_integral<unsigned int> {
301 enum { value = 1 };
302};
303template <> struct is_integral<long> {
304 enum { value = 1 };
305};
306template <> struct is_integral<unsigned long> {
307 enum { value = 1 };
308};
309template <> struct is_integral<long long> {
310 enum { value = 1 };
311};
312template <> struct is_integral<unsigned long long> {
313 enum { value = 1 };
314};
315
316// ToDo: specializes is_arithmetic<_Float16>
317template <class _Tp> struct is_arithmetic {
318 enum { value = 0 };
319};
320template <> struct is_arithmetic<bool> {
321 enum { value = 1 };
322};
323template <> struct is_arithmetic<char> {
324 enum { value = 1 };
325};
326template <> struct is_arithmetic<signed char> {
327 enum { value = 1 };
328};
329template <> struct is_arithmetic<unsigned char> {
330 enum { value = 1 };
331};
332template <> struct is_arithmetic<wchar_t> {
333 enum { value = 1 };
334};
335template <> struct is_arithmetic<short> {
336 enum { value = 1 };
337};
338template <> struct is_arithmetic<unsigned short> {
339 enum { value = 1 };
340};
341template <> struct is_arithmetic<int> {
342 enum { value = 1 };
343};
344template <> struct is_arithmetic<unsigned int> {
345 enum { value = 1 };
346};
347template <> struct is_arithmetic<long> {
348 enum { value = 1 };
349};
350template <> struct is_arithmetic<unsigned long> {
351 enum { value = 1 };
352};
353template <> struct is_arithmetic<long long> {
354 enum { value = 1 };
355};
356template <> struct is_arithmetic<unsigned long long> {
357 enum { value = 1 };
358};
359template <> struct is_arithmetic<float> {
360 enum { value = 1 };
361};
362template <> struct is_arithmetic<double> {
363 enum { value = 1 };
364};
365
366struct true_type {
367 static const __constant__ bool value = true;
368};
369struct false_type {
370 static const __constant__ bool value = false;
371};
372
373template <typename __T, typename __U> struct is_same : public false_type {};
374template <typename __T> struct is_same<__T, __T> : public true_type {};
375
376template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
377
378template <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
383template <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
402template <> struct __numeric_type<void> { static const bool value = true; };
403
404template <class _A1, class _A2 = void, class _A3 = void,
405 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
406 &&__numeric_type<_A3>::value>
407class __promote_imp {
408public:
409 static const bool value = false;
410};
411
412template <class _A1, class _A2, class _A3>
413class __promote_imp<_A1, _A2, _A3, true> {
414private:
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
419public:
420 typedef decltype(__type1() + __type2() + __type3()) type;
421 static const bool value = true;
422};
423
424template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
425private:
426 typedef typename __promote_imp<_A1>::type __type1;
427 typedef typename __promote_imp<_A2>::type __type2;
428
429public:
430 typedef decltype(__type1() + __type2()) type;
431 static const bool value = true;
432};
433
434template <class _A1> class __promote_imp<_A1, void, void, true> {
435public:
436 typedef typename __numeric_type<_A1>::type type;
437 static const bool value = true;
438};
439
440template <class _A1, class _A2 = void, class _A3 = void>
441class __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
550template <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
555fma(__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
560template <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
571template <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
578template <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
585template <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
593template <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
603template <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
613template <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
620template <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
643namespace 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++.
651using ::acos;
652using ::acosh;
653using ::asin;
654using ::asinh;
655using ::atan;
656using ::atan2;
657using ::atanh;
658using ::cbrt;
659using ::ceil;
660using ::copysign;
661using ::cos;
662using ::cosh;
663using ::erf;
664using ::erfc;
665using ::exp;
666using ::exp2;
667using ::expm1;
668using ::fabs;
669using ::fdim;
670using ::floor;
671using ::fma;
672using ::fmax;
673using ::fmin;
674using ::fmod;
675using ::fpclassify;
676using ::frexp;
677using ::hypot;
678using ::ilogb;
679using ::isfinite;
680using ::isgreater;
681using ::isgreaterequal;
682using ::isless;
683using ::islessequal;
684using ::islessgreater;
685using ::isnormal;
686using ::isunordered;
687using ::ldexp;
688using ::lgamma;
689using ::llrint;
690using ::llround;
691using ::log;
692using ::log10;
693using ::log1p;
694using ::log2;
695using ::logb;
696using ::lrint;
697using ::lround;
698using ::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.
702using ::nearbyint;
703using ::nextafter;
704// using ::nexttoward; - Omit this since we do not have a definition.
705using ::pow;
706using ::remainder;
707using ::remquo;
708using ::rint;
709using ::round;
710using ::scalbln;
711using ::scalbn;
712using ::signbit;
713using ::sin;
714using ::sinh;
715using ::sqrt;
716using ::tan;
717using ::tanh;
718using ::tgamma;
719using ::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__
725using ::isinf;
726using ::isnan;
727#endif
728
729// Finally, pull the "foobarf" functions that HIP defines into std.
730using ::acosf;
731using ::acoshf;
732using ::asinf;
733using ::asinhf;
734using ::atan2f;
735using ::atanf;
736using ::atanhf;
737using ::cbrtf;
738using ::ceilf;
739using ::copysignf;
740using ::cosf;
741using ::coshf;
742using ::erfcf;
743using ::erff;
744using ::exp2f;
745using ::expf;
746using ::expm1f;
747using ::fabsf;
748using ::fdimf;
749using ::floorf;
750using ::fmaf;
751using ::fmaxf;
752using ::fminf;
753using ::fmodf;
754using ::frexpf;
755using ::hypotf;
756using ::ilogbf;
757using ::ldexpf;
758using ::lgammaf;
759using ::llrintf;
760using ::llroundf;
761using ::log10f;
762using ::log1pf;
763using ::log2f;
764using ::logbf;
765using ::logf;
766using ::lrintf;
767using ::lroundf;
768using ::modff;
769using ::nearbyintf;
770using ::nextafterf;
771// using ::nexttowardf; - Omit this since we do not have a definition.
772using ::powf;
773using ::remainderf;
774using ::remquof;
775using ::rintf;
776using ::roundf;
777using ::scalblnf;
778using ::scalbnf;
779using ::sinf;
780using ::sinhf;
781using ::sqrtf;
782using ::tanf;
783using ::tanhf;
784using ::tgammaf;
785using ::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)
808extern "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

Warning: This file is not a C or C++ file. It does not have highlighting.

source code of clang/lib/Headers/__clang_hip_cmath.h