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