blob: 0afe4db556dbc00843246fd1d551bbded71be993 [file] [log] [blame]
Logan Chien2833ffb2018-10-09 10:03:24 +08001/*===- __clang_math_forward_declares.h - Prototypes of __device__ math fns --===
2 *
Logan Chiendf4f7662019-09-04 16:45:23 -07003 * 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
Logan Chien2833ffb2018-10-09 10:03:24 +08006 *
7 *===-----------------------------------------------------------------------===
8 */
9#ifndef __CLANG__CUDA_MATH_FORWARD_DECLARES_H__
10#define __CLANG__CUDA_MATH_FORWARD_DECLARES_H__
11#ifndef __CUDA__
12#error "This file is for CUDA compilation only."
13#endif
14
15// This file forward-declares of some math functions we (or the CUDA headers)
16// will define later. We need to do this, and do it before cmath is included,
17// because the standard library may have constexpr math functions. In the
18// absence of a prior __device__ decl, those constexpr functions may become
19// implicitly host+device. host+device functions can't be overloaded, so that
20// would preclude the use of our own __device__ overloads for these functions.
21
22#pragma push_macro("__DEVICE__")
Logan Chiendf4f7662019-09-04 16:45:23 -070023#ifdef _OPENMP
24#define __DEVICE__ static __inline__ __attribute__((always_inline))
25#else
Logan Chien2833ffb2018-10-09 10:03:24 +080026#define __DEVICE__ \
27 static __inline__ __attribute__((always_inline)) __attribute__((device))
Logan Chiendf4f7662019-09-04 16:45:23 -070028#endif
Logan Chien2833ffb2018-10-09 10:03:24 +080029
Logan Chiendf4f7662019-09-04 16:45:23 -070030// For C++ 17 we need to include noexcept attribute to be compatible
31// with the header-defined version. This may be removed once
32// variant is supported.
33#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
34#define __NOEXCEPT noexcept
35#else
36#define __NOEXCEPT
37#endif
38
39#if !(defined(_OPENMP) && defined(__cplusplus))
Logan Chien2833ffb2018-10-09 10:03:24 +080040__DEVICE__ long abs(long);
41__DEVICE__ long long abs(long long);
Logan Chiendf4f7662019-09-04 16:45:23 -070042__DEVICE__ double abs(double);
43__DEVICE__ float abs(float);
44#endif
45// While providing the CUDA declarations and definitions for math functions,
46// we may manually define additional functions.
47// TODO: Once variant is supported the additional functions will have
48// to be removed.
49#if defined(_OPENMP) && defined(__cplusplus)
50__DEVICE__ const double abs(const double);
51__DEVICE__ const float abs(const float);
52#endif
53__DEVICE__ int abs(int) __NOEXCEPT;
Logan Chien2833ffb2018-10-09 10:03:24 +080054__DEVICE__ double acos(double);
55__DEVICE__ float acos(float);
56__DEVICE__ double acosh(double);
57__DEVICE__ float acosh(float);
58__DEVICE__ double asin(double);
59__DEVICE__ float asin(float);
60__DEVICE__ double asinh(double);
61__DEVICE__ float asinh(float);
62__DEVICE__ double atan2(double, double);
63__DEVICE__ float atan2(float, float);
64__DEVICE__ double atan(double);
65__DEVICE__ float atan(float);
66__DEVICE__ double atanh(double);
67__DEVICE__ float atanh(float);
68__DEVICE__ double cbrt(double);
69__DEVICE__ float cbrt(float);
70__DEVICE__ double ceil(double);
71__DEVICE__ float ceil(float);
72__DEVICE__ double copysign(double, double);
73__DEVICE__ float copysign(float, float);
74__DEVICE__ double cos(double);
75__DEVICE__ float cos(float);
76__DEVICE__ double cosh(double);
77__DEVICE__ float cosh(float);
78__DEVICE__ double erfc(double);
79__DEVICE__ float erfc(float);
80__DEVICE__ double erf(double);
81__DEVICE__ float erf(float);
82__DEVICE__ double exp2(double);
83__DEVICE__ float exp2(float);
84__DEVICE__ double exp(double);
85__DEVICE__ float exp(float);
86__DEVICE__ double expm1(double);
87__DEVICE__ float expm1(float);
Logan Chiendf4f7662019-09-04 16:45:23 -070088__DEVICE__ double fabs(double) __NOEXCEPT;
89__DEVICE__ float fabs(float) __NOEXCEPT;
Logan Chien2833ffb2018-10-09 10:03:24 +080090__DEVICE__ double fdim(double, double);
91__DEVICE__ float fdim(float, float);
92__DEVICE__ double floor(double);
93__DEVICE__ float floor(float);
94__DEVICE__ double fma(double, double, double);
95__DEVICE__ float fma(float, float, float);
96__DEVICE__ double fmax(double, double);
97__DEVICE__ float fmax(float, float);
98__DEVICE__ double fmin(double, double);
99__DEVICE__ float fmin(float, float);
100__DEVICE__ double fmod(double, double);
101__DEVICE__ float fmod(float, float);
102__DEVICE__ int fpclassify(double);
103__DEVICE__ int fpclassify(float);
104__DEVICE__ double frexp(double, int *);
105__DEVICE__ float frexp(float, int *);
106__DEVICE__ double hypot(double, double);
107__DEVICE__ float hypot(float, float);
108__DEVICE__ int ilogb(double);
109__DEVICE__ int ilogb(float);
Logan Chiendf4f7662019-09-04 16:45:23 -0700110#ifdef _MSC_VER
111__DEVICE__ bool isfinite(long double);
112#endif
Logan Chien2833ffb2018-10-09 10:03:24 +0800113__DEVICE__ bool isfinite(double);
114__DEVICE__ bool isfinite(float);
115__DEVICE__ bool isgreater(double, double);
116__DEVICE__ bool isgreaterequal(double, double);
117__DEVICE__ bool isgreaterequal(float, float);
118__DEVICE__ bool isgreater(float, float);
Logan Chiendf4f7662019-09-04 16:45:23 -0700119#ifdef _MSC_VER
120__DEVICE__ bool isinf(long double);
121#endif
Logan Chien2833ffb2018-10-09 10:03:24 +0800122__DEVICE__ bool isinf(double);
123__DEVICE__ bool isinf(float);
124__DEVICE__ bool isless(double, double);
125__DEVICE__ bool islessequal(double, double);
126__DEVICE__ bool islessequal(float, float);
127__DEVICE__ bool isless(float, float);
128__DEVICE__ bool islessgreater(double, double);
129__DEVICE__ bool islessgreater(float, float);
Logan Chiendf4f7662019-09-04 16:45:23 -0700130#ifdef _MSC_VER
131__DEVICE__ bool isnan(long double);
132#endif
Logan Chien2833ffb2018-10-09 10:03:24 +0800133__DEVICE__ bool isnan(double);
134__DEVICE__ bool isnan(float);
135__DEVICE__ bool isnormal(double);
136__DEVICE__ bool isnormal(float);
137__DEVICE__ bool isunordered(double, double);
138__DEVICE__ bool isunordered(float, float);
Logan Chiendf4f7662019-09-04 16:45:23 -0700139__DEVICE__ long labs(long) __NOEXCEPT;
Logan Chien2833ffb2018-10-09 10:03:24 +0800140__DEVICE__ double ldexp(double, int);
141__DEVICE__ float ldexp(float, int);
142__DEVICE__ double lgamma(double);
143__DEVICE__ float lgamma(float);
Logan Chiendf4f7662019-09-04 16:45:23 -0700144__DEVICE__ long long llabs(long long) __NOEXCEPT;
Logan Chien2833ffb2018-10-09 10:03:24 +0800145__DEVICE__ long long llrint(double);
146__DEVICE__ long long llrint(float);
147__DEVICE__ double log10(double);
148__DEVICE__ float log10(float);
149__DEVICE__ double log1p(double);
150__DEVICE__ float log1p(float);
151__DEVICE__ double log2(double);
152__DEVICE__ float log2(float);
153__DEVICE__ double logb(double);
154__DEVICE__ float logb(float);
Logan Chiendf4f7662019-09-04 16:45:23 -0700155#if defined(_OPENMP) && defined(__cplusplus)
156__DEVICE__ long double log(long double);
157#endif
Logan Chien2833ffb2018-10-09 10:03:24 +0800158__DEVICE__ double log(double);
159__DEVICE__ float log(float);
160__DEVICE__ long lrint(double);
161__DEVICE__ long lrint(float);
162__DEVICE__ long lround(double);
163__DEVICE__ long lround(float);
Logan Chien55afb0a2018-10-15 10:42:14 +0800164__DEVICE__ long long llround(float); // No llround(double).
Logan Chien2833ffb2018-10-09 10:03:24 +0800165__DEVICE__ double modf(double, double *);
166__DEVICE__ float modf(float, float *);
167__DEVICE__ double nan(const char *);
168__DEVICE__ float nanf(const char *);
169__DEVICE__ double nearbyint(double);
170__DEVICE__ float nearbyint(float);
171__DEVICE__ double nextafter(double, double);
172__DEVICE__ float nextafter(float, float);
Logan Chien2833ffb2018-10-09 10:03:24 +0800173__DEVICE__ double pow(double, double);
174__DEVICE__ double pow(double, int);
175__DEVICE__ float pow(float, float);
176__DEVICE__ float pow(float, int);
177__DEVICE__ double remainder(double, double);
178__DEVICE__ float remainder(float, float);
179__DEVICE__ double remquo(double, double, int *);
180__DEVICE__ float remquo(float, float, int *);
181__DEVICE__ double rint(double);
182__DEVICE__ float rint(float);
183__DEVICE__ double round(double);
184__DEVICE__ float round(float);
185__DEVICE__ double scalbln(double, long);
186__DEVICE__ float scalbln(float, long);
187__DEVICE__ double scalbn(double, int);
188__DEVICE__ float scalbn(float, int);
189__DEVICE__ bool signbit(double);
190__DEVICE__ bool signbit(float);
191__DEVICE__ double sin(double);
192__DEVICE__ float sin(float);
193__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
Logan Chien55afb0a2018-10-15 10:42:14 +0800206// 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
Logan Chien2833ffb2018-10-09 10:03:24 +0800217namespace std {
Logan Chien55afb0a2018-10-15 10:42:14 +0800218#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
219_GLIBCXX_BEGIN_NAMESPACE_VERSION
220#endif
221#endif
222
Logan Chien2833ffb2018-10-09 10:03:24 +0800223using ::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;
Logan Chien55afb0a2018-10-15 10:42:14 +0800274using ::llround;
Logan Chien2833ffb2018-10-09 10:03:24 +0800275using ::modf;
276using ::nan;
277using ::nanf;
278using ::nearbyint;
279using ::nextafter;
Logan Chien2833ffb2018-10-09 10:03:24 +0800280using ::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;
Logan Chien55afb0a2018-10-15 10:42:14 +0800295
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
Logan Chien2833ffb2018-10-09 10:03:24 +0800302} // namespace std
Logan Chien55afb0a2018-10-15 10:42:14 +0800303#endif
Logan Chien2833ffb2018-10-09 10:03:24 +0800304
Logan Chiendf4f7662019-09-04 16:45:23 -0700305#undef __NOEXCEPT
Logan Chien2833ffb2018-10-09 10:03:24 +0800306#pragma pop_macro("__DEVICE__")
307
308#endif