Index: cfe/trunk/lib/Headers/__clang_cuda_cmath.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_cmath.h +++ cfe/trunk/lib/Headers/__clang_cuda_cmath.h @@ -26,54 +26,39 @@ #error "This file is for CUDA compilation only." #endif -// CUDA allows using math functions form std:: on device side. This -// file provides __device__ overloads for math functions that map to -// appropriate math functions provided by CUDA headers or to compiler -// builtins if CUDA does not provide a suitable function. +// CUDA lets us use various std math functions on the device side. This file +// works in concert with __clang_cuda_math_forward_declares.h to make this work. +// +// Specifically, the forward-declares header declares __device__ overloads for +// these functions in the global namespace, then pulls them into namespace std +// with 'using' statements. Then this file implements those functions, after +// the implementations have been pulled in. +// +// It's important that we declare the functions in the global namespace and pull +// them into namespace std with using statements, as opposed to simply declaring +// these functions in namespace std, because our device functions need to +// overload the standard library functions, which may be declared in the global +// namespace or in std, depending on the degree of conformance of the stdlib +// implementation. Declaring in the global namespace and pulling into namespace +// std covers all of the known knowns. #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) -namespace std { __DEVICE__ long long abs(long long __n) { return ::llabs(__n); } __DEVICE__ long abs(long __n) { return ::labs(__n); } -using ::abs; __DEVICE__ float abs(float __x) { return ::fabsf(__x); } __DEVICE__ double abs(double __x) { return ::fabs(__x); } __DEVICE__ float acos(float __x) { return ::acosf(__x); } -using ::acos; -using ::acosh; __DEVICE__ float asin(float __x) { return ::asinf(__x); } -using ::asin; -using ::asinh; __DEVICE__ float atan(float __x) { return ::atanf(__x); } -using ::atan; __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); } -using ::atan2; -using ::atanh; -using ::cbrt; __DEVICE__ float ceil(float __x) { return ::ceilf(__x); } -using ::ceil; -using ::copysign; __DEVICE__ float cos(float __x) { return ::cosf(__x); } -using ::cos; __DEVICE__ float cosh(float __x) { return ::coshf(__x); } -using ::cosh; -using ::erf; -using ::erfc; __DEVICE__ float exp(float __x) { return ::expf(__x); } -using ::exp; -using ::exp2; -using ::expm1; __DEVICE__ float fabs(float __x) { return ::fabsf(__x); } -using ::fabs; -using ::fdim; __DEVICE__ float floor(float __x) { return ::floorf(__x); } -using ::floor; -using ::fma; -using ::fmax; -using ::fmin; __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); } -using ::fmod; __DEVICE__ int fpclassify(float __x) { return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, FP_ZERO, __x); @@ -85,9 +70,8 @@ __DEVICE__ float frexp(float __arg, int *__exp) { return ::frexpf(__arg, __exp); } -using ::frexp; -using ::hypot; -using ::ilogb; +__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } +__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } __DEVICE__ bool isfinite(double __x) { return ::__finite(__x); } __DEVICE__ bool isgreater(float __x, float __y) { @@ -102,8 +86,6 @@ __DEVICE__ bool isgreaterequal(double __x, double __y) { return __builtin_isgreaterequal(__x, __y); } -__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } -__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } __DEVICE__ bool isless(float __x, float __y) { return __builtin_isless(__x, __y); } @@ -132,36 +114,18 @@ __DEVICE__ bool isunordered(double __x, double __y) { return __builtin_isunordered(__x, __y); } -using ::labs; __DEVICE__ float ldexp(float __arg, int __exp) { return ::ldexpf(__arg, __exp); } -using ::ldexp; -using ::lgamma; -using ::llabs; -using ::llrint; __DEVICE__ float log(float __x) { return ::logf(__x); } -using ::log; __DEVICE__ float log10(float __x) { return ::log10f(__x); } -using ::log10; -using ::log1p; -using ::log2; -using ::logb; -using ::lrint; -using ::lround; __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } -using ::modf; -using ::nan; -using ::nanf; -using ::nearbyint; -using ::nextafter; __DEVICE__ float nexttoward(float __from, float __to) { return __builtin_nexttowardf(__from, __to); } __DEVICE__ double nexttoward(double __from, double __to) { return __builtin_nexttoward(__from, __to); } -using ::pow; __DEVICE__ float pow(float __base, float __exp) { return ::powf(__base, __exp); } @@ -171,28 +135,13 @@ __DEVICE__ double pow(double __base, int __iexp) { return ::powi(__base, __iexp); } -using ::remainder; -using ::remquo; -using ::rint; -using ::round; -using ::scalbln; -using ::scalbn; __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); } __DEVICE__ bool signbit(double __x) { return ::__signbit(__x); } __DEVICE__ float sin(float __x) { return ::sinf(__x); } -using ::sin; __DEVICE__ float sinh(float __x) { return ::sinhf(__x); } -using ::sinh; __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); } -using ::sqrt; __DEVICE__ float tan(float __x) { return ::tanf(__x); } -using ::tan; __DEVICE__ float tanh(float __x) { return ::tanhf(__x); } -using ::tanh; -using ::tgamma; -using ::trunc; - -} // namespace std #undef __DEVICE__ Index: cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h +++ cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h @@ -37,153 +37,225 @@ #define __DEVICE__ \ static __inline__ __attribute__((always_inline)) __attribute__((device)) +__DEVICE__ double abs(double); +__DEVICE__ float abs(float); __DEVICE__ int abs(int); +__DEVICE__ long abs(long); +__DEVICE__ long long abs(long long); __DEVICE__ double acos(double); -__DEVICE__ float acosh(float); +__DEVICE__ float acos(float); __DEVICE__ double acosh(double); +__DEVICE__ float acosh(float); __DEVICE__ double asin(double); -__DEVICE__ float asinh(float); +__DEVICE__ float asin(float); __DEVICE__ double asinh(double); -__DEVICE__ double atan(double); +__DEVICE__ float asinh(float); __DEVICE__ double atan2(double, double); -__DEVICE__ float atanh(float); +__DEVICE__ float atan2(float, float); +__DEVICE__ double atan(double); +__DEVICE__ float atan(float); __DEVICE__ double atanh(double); -__DEVICE__ float cbrt(float); +__DEVICE__ float atanh(float); __DEVICE__ double cbrt(double); +__DEVICE__ float cbrt(float); __DEVICE__ double ceil(double); -__DEVICE__ float copysign(float, float); +__DEVICE__ float ceil(float); __DEVICE__ double copysign(double, double); +__DEVICE__ float copysign(float, float); __DEVICE__ double cos(double); +__DEVICE__ float cos(float); __DEVICE__ double cosh(double); -__DEVICE__ float erf(float); -__DEVICE__ double erf(double); -__DEVICE__ float erfc(float); +__DEVICE__ float cosh(float); __DEVICE__ double erfc(double); -__DEVICE__ double exp(double); -__DEVICE__ float exp2(float); +__DEVICE__ float erfc(float); +__DEVICE__ double erf(double); +__DEVICE__ float erf(float); __DEVICE__ double exp2(double); -__DEVICE__ float expm1(float); +__DEVICE__ float exp2(float); +__DEVICE__ double exp(double); +__DEVICE__ float exp(float); __DEVICE__ double expm1(double); +__DEVICE__ float expm1(float); __DEVICE__ double fabs(double); -__DEVICE__ float fdim(float, float); +__DEVICE__ float fabs(float); __DEVICE__ double fdim(double, double); +__DEVICE__ float fdim(float, float); __DEVICE__ double floor(double); -__DEVICE__ float fma(float, float, float); +__DEVICE__ float floor(float); __DEVICE__ double fma(double, double, double); -__DEVICE__ float fmax(float, float); +__DEVICE__ float fma(float, float, float); __DEVICE__ double fmax(double, double); -__DEVICE__ float fmin(float, float); +__DEVICE__ float fmax(float, float); __DEVICE__ double fmin(double, double); +__DEVICE__ float fmin(float, float); __DEVICE__ double fmod(double, double); +__DEVICE__ float fmod(float, float); +__DEVICE__ int fpclassify(double); +__DEVICE__ int fpclassify(float); __DEVICE__ double frexp(double, int *); -__DEVICE__ float hypot(float, float); +__DEVICE__ float frexp(float, int *); __DEVICE__ double hypot(double, double); -__DEVICE__ int ilogb(float); +__DEVICE__ float hypot(float, float); __DEVICE__ int ilogb(double); +__DEVICE__ int ilogb(float); +__DEVICE__ bool isfinite(double); +__DEVICE__ bool isfinite(float); +__DEVICE__ bool isgreater(double, double); +__DEVICE__ bool isgreaterequal(double, double); +__DEVICE__ bool isgreaterequal(float, float); +__DEVICE__ bool isgreater(float, float); +__DEVICE__ bool isinf(double); +__DEVICE__ bool isinf(float); +__DEVICE__ bool isless(double, double); +__DEVICE__ bool islessequal(double, double); +__DEVICE__ bool islessequal(float, float); +__DEVICE__ bool isless(float, float); +__DEVICE__ bool islessgreater(double, double); +__DEVICE__ bool islessgreater(float, float); +__DEVICE__ bool isnan(double); +__DEVICE__ bool isnan(float); +__DEVICE__ bool isnormal(double); +__DEVICE__ bool isnormal(float); +__DEVICE__ bool isunordered(double, double); +__DEVICE__ bool isunordered(float, float); __DEVICE__ long labs(long); __DEVICE__ double ldexp(double, int); -__DEVICE__ float lgamma(float); +__DEVICE__ float ldexp(float, int); __DEVICE__ double lgamma(double); +__DEVICE__ float lgamma(float); __DEVICE__ long long llabs(long long); -__DEVICE__ long long llrint(float); __DEVICE__ long long llrint(double); -__DEVICE__ double log(double); +__DEVICE__ long long llrint(float); __DEVICE__ double log10(double); -__DEVICE__ float log1p(float); +__DEVICE__ float log10(float); __DEVICE__ double log1p(double); -__DEVICE__ float log2(float); +__DEVICE__ float log1p(float); __DEVICE__ double log2(double); -__DEVICE__ float logb(float); +__DEVICE__ float log2(float); __DEVICE__ double logb(double); -__DEVICE__ long lrint(float); +__DEVICE__ float logb(float); +__DEVICE__ double log(double); +__DEVICE__ float log(float); __DEVICE__ long lrint(double); -__DEVICE__ long lround(float); +__DEVICE__ long lrint(float); __DEVICE__ long lround(double); +__DEVICE__ long lround(float); __DEVICE__ double modf(double, double *); +__DEVICE__ float modf(float, float *); __DEVICE__ double nan(const char *); __DEVICE__ float nanf(const char *); -__DEVICE__ float nearbyint(float); __DEVICE__ double nearbyint(double); -__DEVICE__ float nextafter(float, float); +__DEVICE__ float nearbyint(float); __DEVICE__ double nextafter(double, double); +__DEVICE__ float nextafter(float, float); +__DEVICE__ double nexttoward(double, double); +__DEVICE__ float nexttoward(float, float); __DEVICE__ double pow(double, double); -__DEVICE__ float remainder(float, float); +__DEVICE__ double pow(double, int); +__DEVICE__ float pow(float, float); +__DEVICE__ float pow(float, int); __DEVICE__ double remainder(double, double); -__DEVICE__ float remquo(float, float, int *); +__DEVICE__ float remainder(float, float); __DEVICE__ double remquo(double, double, int *); -__DEVICE__ float rint(float); +__DEVICE__ float remquo(float, float, int *); __DEVICE__ double rint(double); -__DEVICE__ float round(float); +__DEVICE__ float rint(float); __DEVICE__ double round(double); -__DEVICE__ float scalbln(float, long); +__DEVICE__ float round(float); __DEVICE__ double scalbln(double, long); -__DEVICE__ float scalbn(float, int); +__DEVICE__ float scalbln(float, long); __DEVICE__ double scalbn(double, int); +__DEVICE__ float scalbn(float, int); +__DEVICE__ bool signbit(double); +__DEVICE__ bool signbit(float); __DEVICE__ double sin(double); +__DEVICE__ float sin(float); __DEVICE__ double sinh(double); +__DEVICE__ float sinh(float); __DEVICE__ double sqrt(double); +__DEVICE__ float sqrt(float); __DEVICE__ double tan(double); +__DEVICE__ float tan(float); __DEVICE__ double tanh(double); -__DEVICE__ float tgamma(float); +__DEVICE__ float tanh(float); __DEVICE__ double tgamma(double); -__DEVICE__ float trunc(float); +__DEVICE__ float tgamma(float); __DEVICE__ double trunc(double); +__DEVICE__ float trunc(float); namespace std { -__DEVICE__ long long abs(long long); -__DEVICE__ long abs(long); -__DEVICE__ float abs(float); -__DEVICE__ double abs(double); -__DEVICE__ float acos(float); -__DEVICE__ float asin(float); -__DEVICE__ float atan(float); -__DEVICE__ float atan2(float, float); -__DEVICE__ float ceil(float); -__DEVICE__ float cos(float); -__DEVICE__ float cosh(float); -__DEVICE__ float exp(float); -__DEVICE__ float fabs(float); -__DEVICE__ float floor(float); -__DEVICE__ float fmod(float, float); -__DEVICE__ int fpclassify(float); -__DEVICE__ int fpclassify(double); -__DEVICE__ float frexp(float, int *); -__DEVICE__ bool isfinite(float); -__DEVICE__ bool isfinite(double); -__DEVICE__ bool isgreater(float, float); -__DEVICE__ bool isgreater(double, double); -__DEVICE__ bool isgreaterequal(float, float); -__DEVICE__ bool isgreaterequal(double, double); -__DEVICE__ bool isinf(float); -__DEVICE__ bool isinf(double); -__DEVICE__ bool isless(float, float); -__DEVICE__ bool isless(double, double); -__DEVICE__ bool islessequal(float, float); -__DEVICE__ bool islessequal(double, double); -__DEVICE__ bool islessgreater(float, float); -__DEVICE__ bool islessgreater(double, double); -__DEVICE__ bool isnan(float); -__DEVICE__ bool isnan(double); -__DEVICE__ bool isnormal(float); -__DEVICE__ bool isnormal(double); -__DEVICE__ bool isunordered(float, float); -__DEVICE__ bool isunordered(double, double); -__DEVICE__ float ldexp(float, int); -__DEVICE__ float log(float); -__DEVICE__ float log10(float); -__DEVICE__ float modf(float, float *); -__DEVICE__ float nexttoward(float, float); -__DEVICE__ double nexttoward(double, double); -__DEVICE__ float pow(float, float); -__DEVICE__ float pow(float, int); -__DEVICE__ double pow(double, int); -__DEVICE__ bool signbit(float); -__DEVICE__ bool signbit(double); -__DEVICE__ float sin(float); -__DEVICE__ float sinh(float); -__DEVICE__ float sqrt(float); -__DEVICE__ float tan(float); -__DEVICE__ float tanh(float); +using ::abs; +using ::acos; +using ::acosh; +using ::asin; +using ::asinh; +using ::atan; +using ::atan2; +using ::atanh; +using ::cbrt; +using ::ceil; +using ::copysign; +using ::cos; +using ::cosh; +using ::erf; +using ::erfc; +using ::exp; +using ::exp2; +using ::expm1; +using ::fabs; +using ::fdim; +using ::floor; +using ::fma; +using ::fmax; +using ::fmin; +using ::fmod; +using ::fpclassify; +using ::frexp; +using ::hypot; +using ::ilogb; +using ::isfinite; +using ::isgreater; +using ::isgreaterequal; +using ::isinf; +using ::isless; +using ::islessequal; +using ::islessgreater; +using ::isnan; +using ::isnormal; +using ::isunordered; +using ::labs; +using ::ldexp; +using ::lgamma; +using ::llabs; +using ::llrint; +using ::log; +using ::log10; +using ::log1p; +using ::log2; +using ::logb; +using ::lrint; +using ::lround; +using ::modf; +using ::nan; +using ::nanf; +using ::nearbyint; +using ::nextafter; +using ::nexttoward; +using ::pow; +using ::remainder; +using ::remquo; +using ::rint; +using ::round; +using ::scalbln; +using ::scalbn; +using ::signbit; +using ::sin; +using ::sinh; +using ::sqrt; +using ::tan; +using ::tanh; +using ::tgamma; +using ::trunc; } // namespace std #pragma pop_macro("__DEVICE__") Index: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h +++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -189,8 +189,21 @@ // we have to include it and it will in turn include .hpp #include "sm_30_intrinsics.h" #include "sm_32_intrinsics.hpp" + #undef __MATH_FUNCTIONS_HPP__ + +// math_functions.hpp defines ::signbit as a __host__ __device__ function. This +// conflicts with libstdc++'s constexpr ::signbit, so we have to rename +// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's +// conditional on __GNUC__. :) +#pragma push_macro("signbit") +#pragma push_macro("__GNUC__") +#undef __GNUC__ +#define signbit __ignored_cuda_signbit #include "math_functions.hpp" +#pragma pop_macro("__GNUC__") +#pragma pop_macro("signbit") + #pragma pop_macro("__host__") #include "texture_indirect_functions.h"