172169Sphantom/*===- __clang_math_forward_declares.h - Prototypes of __device__ math fns --===
2116875Sphantom *
372169Sphantom * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
472169Sphantom * See https://llvm.org/LICENSE.txt for license information.
5235785Stheraven * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6235785Stheraven *
7235785Stheraven *===-----------------------------------------------------------------------===
8235785Stheraven */
9235785Stheraven#ifndef __CLANG__CUDA_MATH_FORWARD_DECLARES_H__
1072169Sphantom#define __CLANG__CUDA_MATH_FORWARD_DECLARES_H__
1172169Sphantom#ifndef __CUDA__
1272169Sphantom#error "This file is for CUDA compilation only."
1372169Sphantom#endif
1472169Sphantom
1572169Sphantom// This file forward-declares of some math functions we (or the CUDA headers)
1672169Sphantom// will define later.  We need to do this, and do it before cmath is included,
1772169Sphantom// because the standard library may have constexpr math functions.  In the
1872169Sphantom// absence of a prior __device__ decl, those constexpr functions may become
1972169Sphantom// implicitly host+device.  host+device functions can't be overloaded, so that
2072169Sphantom// would preclude the use of our own __device__ overloads for these functions.
2172169Sphantom
2272169Sphantom#pragma push_macro("__DEVICE__")
2372169Sphantom#ifdef _OPENMP
2472169Sphantom#define __DEVICE__ static __inline__ __attribute__((always_inline))
2572169Sphantom#else
2672169Sphantom#define __DEVICE__                                                             \
2772169Sphantom  static __inline__ __attribute__((always_inline)) __attribute__((device))
2872169Sphantom#endif
2972169Sphantom
3072169Sphantom// For C++ 17 we need to include noexcept attribute to be compatible
3172169Sphantom// with the header-defined version. This may be removed once
3292986Sobrien// variant is supported.
3392986Sobrien#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
3492986Sobrien#define __NOEXCEPT noexcept
3572263Sache#else
3672561Sache#define __NOEXCEPT
37116875Sphantom#endif
3872561Sache
3972263Sache#if !(defined(_OPENMP) && defined(__cplusplus))
4087658Sphantom__DEVICE__ long abs(long);
4172169Sphantom__DEVICE__ long long abs(long long);
42116875Sphantom__DEVICE__ double abs(double);
4372169Sphantom__DEVICE__ float abs(float);
44116875Sphantom#endif
4572169Sphantom// While providing the CUDA declarations and definitions for math functions,
4672169Sphantom// we may manually define additional functions.
4772169Sphantom// TODO: Once variant is supported the additional functions will have
4872169Sphantom// to be removed.
49235785Stheraven#if defined(_OPENMP) && defined(__cplusplus)
50116875Sphantom__DEVICE__ const double abs(const double);
51235785Stheraven__DEVICE__ const float abs(const float);
52235785Stheraven#endif
53235785Stheraven__DEVICE__ int abs(int) __NOEXCEPT;
5472169Sphantom__DEVICE__ double acos(double);
5572169Sphantom__DEVICE__ float acos(float);
5672169Sphantom__DEVICE__ double acosh(double);
5772263Sache__DEVICE__ float acosh(float);
58235785Stheraven__DEVICE__ double asin(double);
59116134Sache__DEVICE__ float asin(float);
6072263Sache__DEVICE__ double asinh(double);
61116134Sache__DEVICE__ float asinh(float);
62116134Sache__DEVICE__ double atan2(double, double);
6372263Sache__DEVICE__ float atan2(float, float);
6472263Sache__DEVICE__ double atan(double);
6572169Sphantom__DEVICE__ float atan(float);
6672169Sphantom__DEVICE__ double atanh(double);
67235785Stheraven__DEVICE__ float atanh(float);
6872169Sphantom__DEVICE__ double cbrt(double);
6972169Sphantom__DEVICE__ float cbrt(float);
70235785Stheraven__DEVICE__ double ceil(double);
7172169Sphantom__DEVICE__ float ceil(float);
7272169Sphantom__DEVICE__ double copysign(double, double);
73235785Stheraven__DEVICE__ float copysign(float, float);
7472169Sphantom__DEVICE__ double cos(double);
7572169Sphantom__DEVICE__ float cos(float);
76235785Stheraven__DEVICE__ double cosh(double);
7772169Sphantom__DEVICE__ float cosh(float);
7872169Sphantom__DEVICE__ double erfc(double);
79235785Stheraven__DEVICE__ float erfc(float);
8072169Sphantom__DEVICE__ double erf(double);
8172169Sphantom__DEVICE__ float erf(float);
82235785Stheraven__DEVICE__ double exp2(double);
8372169Sphantom__DEVICE__ float exp2(float);
8472169Sphantom__DEVICE__ double exp(double);
8572169Sphantom__DEVICE__ float exp(float);
86235785Stheraven__DEVICE__ double expm1(double);
8772169Sphantom__DEVICE__ float expm1(float);
8872169Sphantom__DEVICE__ double fabs(double) __NOEXCEPT;
8972169Sphantom__DEVICE__ float fabs(float) __NOEXCEPT;
90235785Stheraven__DEVICE__ double fdim(double, double);
9172169Sphantom__DEVICE__ float fdim(float, float);
9272169Sphantom__DEVICE__ double floor(double);
9372169Sphantom__DEVICE__ float floor(float);
9472169Sphantom__DEVICE__ double fma(double, double, double);
95235785Stheraven__DEVICE__ float fma(float, float, float);
9672169Sphantom__DEVICE__ double fmax(double, double);
9772169Sphantom__DEVICE__ float fmax(float, float);
9872169Sphantom__DEVICE__ double fmin(double, double);
9972169Sphantom__DEVICE__ float fmin(float, float);
100235785Stheraven__DEVICE__ double fmod(double, double);
10172169Sphantom__DEVICE__ float fmod(float, float);
102197765Sedwin__DEVICE__ int fpclassify(double);
103197765Sedwin__DEVICE__ int fpclassify(float);
104197765Sedwin__DEVICE__ double frexp(double, int *);
105197765Sedwin__DEVICE__ float frexp(float, int *);
106235785Stheraven__DEVICE__ double hypot(double, double);
107197765Sedwin__DEVICE__ float hypot(float, float);
10872169Sphantom__DEVICE__ int ilogb(double);
10972706Sphantom__DEVICE__ int ilogb(float);
11072169Sphantom#ifdef _MSC_VER
11172169Sphantom__DEVICE__ bool isfinite(long double);
11272169Sphantom#endif
11372706Sphantom__DEVICE__ bool isfinite(double);
11472169Sphantom__DEVICE__ bool isfinite(float);
11572169Sphantom__DEVICE__ bool isgreater(double, double);
11672169Sphantom__DEVICE__ bool isgreaterequal(double, double);
11772706Sphantom__DEVICE__ bool isgreaterequal(float, float);
11872169Sphantom__DEVICE__ bool isgreater(float, float);
11972169Sphantom#ifdef _MSC_VER
12072169Sphantom__DEVICE__ bool isinf(long double);
12172706Sphantom#endif
12272169Sphantom__DEVICE__ bool isinf(double);
12372169Sphantom__DEVICE__ bool isinf(float);
12472169Sphantom__DEVICE__ bool isless(double, double);
12572706Sphantom__DEVICE__ bool islessequal(double, double);
12672169Sphantom__DEVICE__ bool islessequal(float, float);
12772169Sphantom__DEVICE__ bool isless(float, float);
12872706Sphantom__DEVICE__ bool islessgreater(double, double);
129235785Stheraven__DEVICE__ bool islessgreater(float, float);
13072169Sphantom#ifdef _MSC_VER
13172706Sphantom__DEVICE__ bool isnan(long double);
132235785Stheraven#endif
13372169Sphantom__DEVICE__ bool isnan(double);
13472169Sphantom__DEVICE__ bool isnan(float);
135235785Stheraven__DEVICE__ bool isnormal(double);
13672169Sphantom__DEVICE__ bool isnormal(float);
13772169Sphantom__DEVICE__ bool isunordered(double, double);
138235785Stheraven__DEVICE__ bool isunordered(float, float);
13972169Sphantom__DEVICE__ long labs(long) __NOEXCEPT;
14087658Sphantom__DEVICE__ double ldexp(double, int);
141116875Sphantom__DEVICE__ float ldexp(float, int);
142116875Sphantom__DEVICE__ double lgamma(double);
143116875Sphantom__DEVICE__ float lgamma(float);
14472706Sphantom__DEVICE__ long long llabs(long long) __NOEXCEPT;
14572706Sphantom__DEVICE__ long long llrint(double);
146235785Stheraven__DEVICE__ long long llrint(float);
14772169Sphantom__DEVICE__ double log10(double);
14872706Sphantom__DEVICE__ float log10(float);
149235785Stheraven__DEVICE__ double log1p(double);
15072169Sphantom__DEVICE__ float log1p(float);
151116875Sphantom__DEVICE__ double log2(double);
152116875Sphantom__DEVICE__ float log2(float);
153116875Sphantom__DEVICE__ double logb(double);
15472706Sphantom__DEVICE__ float logb(float);
15572245Sache#if defined(_OPENMP) && defined(__cplusplus)
156235785Stheraven__DEVICE__ long double log(long double);
15772561Sache#endif
158235785Stheraven__DEVICE__ double log(double);
15972561Sache__DEVICE__ float log(float);
160235785Stheraven__DEVICE__ long lrint(double);
16172568Sache__DEVICE__ long lrint(float);
16272561Sache__DEVICE__ long lround(double);
16372568Sache__DEVICE__ long lround(float);
164235785Stheraven__DEVICE__ long long llround(float); // No llround(double).
16572568Sache__DEVICE__ double modf(double, double *);
16672568Sache__DEVICE__ float modf(float, float *);
16772568Sache__DEVICE__ double nan(const char *);
16872568Sache__DEVICE__ float nanf(const char *);
16972568Sache__DEVICE__ double nearbyint(double);
17072568Sache__DEVICE__ float nearbyint(float);
171235785Stheraven__DEVICE__ double nextafter(double, double);
172235785Stheraven__DEVICE__ float nextafter(float, float);
173235785Stheraven__DEVICE__ double pow(double, double);
174235785Stheraven__DEVICE__ double pow(double, int);
17572568Sache__DEVICE__ float pow(float, float);
17672561Sache__DEVICE__ float pow(float, int);
17772561Sache__DEVICE__ double remainder(double, double);
17872561Sache__DEVICE__ float remainder(float, float);
17972243Sache__DEVICE__ double remquo(double, double, int *);
180116875Sphantom__DEVICE__ float remquo(float, float, int *);
181235785Stheraven__DEVICE__ double rint(double);
18274459Sache__DEVICE__ float rint(float);
18372169Sphantom__DEVICE__ double round(double);
18472169Sphantom__DEVICE__ float round(float);
18572169Sphantom__DEVICE__ double scalbln(double, long);
18672169Sphantom__DEVICE__ float scalbln(float, long);
18772169Sphantom__DEVICE__ double scalbn(double, int);
188235785Stheraven__DEVICE__ float scalbn(float, int);
189235785Stheraven__DEVICE__ bool signbit(double);
190235785Stheraven__DEVICE__ bool signbit(float);
191235785Stheraven__DEVICE__ double sin(double);
192235785Stheraven__DEVICE__ float sin(float);
193235785Stheraven__DEVICE__ double sinh(double);
194__DEVICE__ float sinh(float);
195__DEVICE__ double sqrt(double);
196__DEVICE__ float sqrt(float);
197__DEVICE__ double tan(double);
198__DEVICE__ float tan(float);
199__DEVICE__ double tanh(double);
200__DEVICE__ float tanh(float);
201__DEVICE__ double tgamma(double);
202__DEVICE__ float tgamma(float);
203__DEVICE__ double trunc(double);
204__DEVICE__ float trunc(float);
205
206// Notably missing above is nexttoward, which we don't define on
207// the device side because libdevice doesn't give us an implementation, and we
208// don't want to be in the business of writing one ourselves.
209
210// We need to define these overloads in exactly the namespace our standard
211// library uses (including the right inline namespace), otherwise they won't be
212// picked up by other functions in the standard library (e.g. functions in
213// <complex>).  Thus the ugliness below.
214#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
215_LIBCPP_BEGIN_NAMESPACE_STD
216#else
217namespace std {
218#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
219_GLIBCXX_BEGIN_NAMESPACE_VERSION
220#endif
221#endif
222
223using ::abs;
224using ::acos;
225using ::acosh;
226using ::asin;
227using ::asinh;
228using ::atan;
229using ::atan2;
230using ::atanh;
231using ::cbrt;
232using ::ceil;
233using ::copysign;
234using ::cos;
235using ::cosh;
236using ::erf;
237using ::erfc;
238using ::exp;
239using ::exp2;
240using ::expm1;
241using ::fabs;
242using ::fdim;
243using ::floor;
244using ::fma;
245using ::fmax;
246using ::fmin;
247using ::fmod;
248using ::fpclassify;
249using ::frexp;
250using ::hypot;
251using ::ilogb;
252using ::isfinite;
253using ::isgreater;
254using ::isgreaterequal;
255using ::isinf;
256using ::isless;
257using ::islessequal;
258using ::islessgreater;
259using ::isnan;
260using ::isnormal;
261using ::isunordered;
262using ::labs;
263using ::ldexp;
264using ::lgamma;
265using ::llabs;
266using ::llrint;
267using ::log;
268using ::log10;
269using ::log1p;
270using ::log2;
271using ::logb;
272using ::lrint;
273using ::lround;
274using ::llround;
275using ::modf;
276using ::nan;
277using ::nanf;
278using ::nearbyint;
279using ::nextafter;
280using ::pow;
281using ::remainder;
282using ::remquo;
283using ::rint;
284using ::round;
285using ::scalbln;
286using ::scalbn;
287using ::signbit;
288using ::sin;
289using ::sinh;
290using ::sqrt;
291using ::tan;
292using ::tanh;
293using ::tgamma;
294using ::trunc;
295
296#ifdef _LIBCPP_END_NAMESPACE_STD
297_LIBCPP_END_NAMESPACE_STD
298#else
299#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
300_GLIBCXX_END_NAMESPACE_VERSION
301#endif
302} // namespace std
303#endif
304
305#undef __NOEXCEPT
306#pragma pop_macro("__DEVICE__")
307
308#endif
309