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