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)
32__DEVICE__ double abs(double __x) { return ::fabs(__x); }
33__DEVICE__ float abs(float __x) { return ::fabsf(__x); }
34__DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
35__DEVICE__ long abs(long __n) { return ::labs(__n); }
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.
42__DEVICE__ int fpclassify(float __x) {
43  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
44                              FP_ZERO, __x);
45}
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
52__DEVICE__ float frexp(float __arg, int *__exp) {
53  return ::frexpf(__arg, __exp);
54}
55__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
56__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
57__DEVICE__ bool isgreater(float __x, float __y) {
58  return __builtin_isgreater(__x, __y);
59}
60__DEVICE__ bool isgreater(double __x, double __y) {
61  return __builtin_isgreater(__x, __y);
62}
63__DEVICE__ bool isgreaterequal(float __x, float __y) {
64  return __builtin_isgreaterequal(__x, __y);
65}
66__DEVICE__ bool isgreaterequal(double __x, double __y) {
67  return __builtin_isgreaterequal(__x, __y);
68}
69__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
70__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
71__DEVICE__ bool isless(float __x, float __y) {
72  return __builtin_isless(__x, __y);
73}
74__DEVICE__ bool isless(double __x, double __y) {
75  return __builtin_isless(__x, __y);
76}
77__DEVICE__ bool islessequal(float __x, float __y) {
78  return __builtin_islessequal(__x, __y);
79}
80__DEVICE__ bool islessequal(double __x, double __y) {
81  return __builtin_islessequal(__x, __y);
82}
83__DEVICE__ bool islessgreater(float __x, float __y) {
84  return __builtin_islessgreater(__x, __y);
85}
86__DEVICE__ bool islessgreater(double __x, double __y) {
87  return __builtin_islessgreater(__x, __y);
88}
89__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
90__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
91__DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
92__DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
93__DEVICE__ bool isunordered(float __x, float __y) {
94  return __builtin_isunordered(__x, __y);
95}
96__DEVICE__ bool isunordered(double __x, double __y) {
97  return __builtin_isunordered(__x, __y);
98}
99__DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
100__DEVICE__ float pow(float __base, int __iexp) {
101  return ::powif(__base, __iexp);
102}
103__DEVICE__ double pow(double __base, int __iexp) {
104  return ::powi(__base, __iexp);
105}
106__DEVICE__ float remquo(float __x, float __y, int *__quo) {
107  return ::remquof(__x, __y, __quo);
108}
109__DEVICE__ float scalbln(float __x, long int __n) {
110  return ::scalblnf(__x, __n);
111}
112__DEVICE__ bool signbit(float __x) { return ::__signbitf(__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.
120__DEVICE__ _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) {
121  return __ocml_fma_f16(__x, __y, __z);
122}
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.
213template <bool __B, class __T = void> struct __hip_enable_if {};
214
215template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
216
217namespace __hip {
218template <class _Tp> struct is_integral {
219  enum { value = 0 };
220};
221template <> struct is_integral<bool> {
222  enum { value = 1 };
223};
224template <> struct is_integral<char> {
225  enum { value = 1 };
226};
227template <> struct is_integral<signed char> {
228  enum { value = 1 };
229};
230template <> struct is_integral<unsigned char> {
231  enum { value = 1 };
232};
233template <> struct is_integral<wchar_t> {
234  enum { value = 1 };
235};
236template <> struct is_integral<short> {
237  enum { value = 1 };
238};
239template <> struct is_integral<unsigned short> {
240  enum { value = 1 };
241};
242template <> struct is_integral<int> {
243  enum { value = 1 };
244};
245template <> struct is_integral<unsigned int> {
246  enum { value = 1 };
247};
248template <> struct is_integral<long> {
249  enum { value = 1 };
250};
251template <> struct is_integral<unsigned long> {
252  enum { value = 1 };
253};
254template <> struct is_integral<long long> {
255  enum { value = 1 };
256};
257template <> struct is_integral<unsigned long long> {
258  enum { value = 1 };
259};
260
261// ToDo: specializes is_arithmetic<_Float16>
262template <class _Tp> struct is_arithmetic {
263  enum { value = 0 };
264};
265template <> struct is_arithmetic<bool> {
266  enum { value = 1 };
267};
268template <> struct is_arithmetic<char> {
269  enum { value = 1 };
270};
271template <> struct is_arithmetic<signed char> {
272  enum { value = 1 };
273};
274template <> struct is_arithmetic<unsigned char> {
275  enum { value = 1 };
276};
277template <> struct is_arithmetic<wchar_t> {
278  enum { value = 1 };
279};
280template <> struct is_arithmetic<short> {
281  enum { value = 1 };
282};
283template <> struct is_arithmetic<unsigned short> {
284  enum { value = 1 };
285};
286template <> struct is_arithmetic<int> {
287  enum { value = 1 };
288};
289template <> struct is_arithmetic<unsigned int> {
290  enum { value = 1 };
291};
292template <> struct is_arithmetic<long> {
293  enum { value = 1 };
294};
295template <> struct is_arithmetic<unsigned long> {
296  enum { value = 1 };
297};
298template <> struct is_arithmetic<long long> {
299  enum { value = 1 };
300};
301template <> struct is_arithmetic<unsigned long long> {
302  enum { value = 1 };
303};
304template <> struct is_arithmetic<float> {
305  enum { value = 1 };
306};
307template <> struct is_arithmetic<double> {
308  enum { value = 1 };
309};
310
311struct true_type {
312  static const __constant__ bool value = true;
313};
314struct false_type {
315  static const __constant__ bool value = false;
316};
317
318template <typename __T, typename __U> struct is_same : public false_type {};
319template <typename __T> struct is_same<__T, __T> : public true_type {};
320
321template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
322
323template <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
328template <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
347template <> struct __numeric_type<void> { static const bool value = true; };
348
349template <class _A1, class _A2 = void, class _A3 = void,
350          bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
351              &&__numeric_type<_A3>::value>
352class __promote_imp {
353public:
354  static const bool value = false;
355};
356
357template <class _A1, class _A2, class _A3>
358class __promote_imp<_A1, _A2, _A3, true> {
359private:
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
364public:
365  typedef decltype(__type1() + __type2() + __type3()) type;
366  static const bool value = true;
367};
368
369template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
370private:
371  typedef typename __promote_imp<_A1>::type __type1;
372  typedef typename __promote_imp<_A2>::type __type2;
373
374public:
375  typedef decltype(__type1() + __type2()) type;
376  static const bool value = true;
377};
378
379template <class _A1> class __promote_imp<_A1, void, void, true> {
380public:
381  typedef typename __numeric_type<_A1>::type type;
382  static const bool value = true;
383};
384
385template <class _A1, class _A2 = void, class _A3 = void>
386class __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
494template <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
499fma(__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
504template <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
509fma(__T1 __x, __T2 __y, __T3 __z) {
510  return ::fma((double)__x, (double)__y, (double)__z);
511}
512#endif
513
514template <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
521template <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
528template <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
536template <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
546template <typename __T1, typename __T2>
547__DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
548                                        __hip::is_arithmetic<__T2>::value,
549                                    double>::type
550remquo(__T1 __x, __T2 __y, int *__quo) {
551  return ::remquo((double)__x, (double)__y, __quo);
552}
553#endif
554
555template <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
562template <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
583namespace 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++.
591using ::acos;
592using ::acosh;
593using ::asin;
594using ::asinh;
595using ::atan;
596using ::atan2;
597using ::atanh;
598using ::cbrt;
599using ::ceil;
600using ::copysign;
601using ::cos;
602using ::cosh;
603using ::erf;
604using ::erfc;
605using ::exp;
606using ::exp2;
607using ::expm1;
608using ::fabs;
609using ::fdim;
610using ::floor;
611using ::fma;
612using ::fmax;
613using ::fmin;
614using ::fmod;
615using ::fpclassify;
616using ::frexp;
617using ::hypot;
618using ::ilogb;
619using ::isfinite;
620using ::isgreater;
621using ::isgreaterequal;
622using ::isless;
623using ::islessequal;
624using ::islessgreater;
625using ::isnormal;
626using ::isunordered;
627using ::ldexp;
628using ::lgamma;
629using ::llrint;
630using ::llround;
631using ::log;
632using ::log10;
633using ::log1p;
634using ::log2;
635using ::logb;
636using ::lrint;
637using ::lround;
638using ::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.
642using ::nearbyint;
643using ::nextafter;
644// using ::nexttoward; - Omit this since we do not have a definition.
645using ::pow;
646using ::remainder;
647using ::remquo;
648using ::rint;
649using ::round;
650using ::scalbln;
651using ::scalbn;
652using ::signbit;
653using ::sin;
654using ::sinh;
655using ::sqrt;
656using ::tan;
657using ::tanh;
658using ::tgamma;
659using ::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__
665using ::isinf;
666using ::isnan;
667#endif
668
669// Finally, pull the "foobarf" functions that HIP defines into std.
670using ::acosf;
671using ::acoshf;
672using ::asinf;
673using ::asinhf;
674using ::atan2f;
675using ::atanf;
676using ::atanhf;
677using ::cbrtf;
678using ::ceilf;
679using ::copysignf;
680using ::cosf;
681using ::coshf;
682using ::erfcf;
683using ::erff;
684using ::exp2f;
685using ::expf;
686using ::expm1f;
687using ::fabsf;
688using ::fdimf;
689using ::floorf;
690using ::fmaf;
691using ::fmaxf;
692using ::fminf;
693using ::fmodf;
694using ::frexpf;
695using ::hypotf;
696using ::ilogbf;
697using ::ldexpf;
698using ::lgammaf;
699using ::llrintf;
700using ::llroundf;
701using ::log10f;
702using ::log1pf;
703using ::log2f;
704using ::logbf;
705using ::logf;
706using ::lrintf;
707using ::lroundf;
708using ::modff;
709using ::nearbyintf;
710using ::nextafterf;
711// using ::nexttowardf; - Omit this since we do not have a definition.
712using ::powf;
713using ::remainderf;
714using ::remquof;
715using ::rintf;
716using ::roundf;
717using ::scalblnf;
718using ::scalbnf;
719using ::sinf;
720using ::sinhf;
721using ::sqrtf;
722using ::tanf;
723using ::tanhf;
724using ::tgammaf;
725using ::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)
748extern "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