[CUDA] Tweak math forward declares so we're compatible with libstdc++…

…4.9.

Summary:
See comments in patch; we were assuming that some stdlib math functions
would be defined in namespace std, when in fact the spec says they
should be defined in the global namespace.  libstdc++4.9 became more
conforming and broke us.

This new implementation seems to cover the known knowns.

Reviewers: rsmith

Subscribers: cfe-commits, tra

Differential Revision: http://reviews.llvm.org/D18882

llvm-svn: 265751

clang/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__ int fpclassify(double __x) {
 __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(float __x, float __y) {
 __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(float __x, float __y) {
 __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__ float pow(float __base, int __iexp) {
 __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__