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

[jlebar]

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.

[jlebar]

Thank you for the review!

cfe/trunk/lib/Headers/__clang_cuda_cmath.h

  *===-----------------------------------------------------------------------===
  */
 #ifndef __CLANG_CUDA_CMATH_H__
 #define __CLANG_CUDA_CMATH_H__
 #ifndef __CUDA__
 #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);
 }
 __DEVICE__ int fpclassify(double __x) {
   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
                               FP_ZERO, __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) {
   return __builtin_isgreater(__x, __y);
 }
 __DEVICE__ bool isgreater(double __x, double __y) {
   return __builtin_isgreater(__x, __y);
 }
 __DEVICE__ bool isgreaterequal(float __x, float __y) {
   return __builtin_isgreaterequal(__x, __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);
 }
 __DEVICE__ bool isless(double __x, double __y) {
   return __builtin_isless(__x, __y);
 }
 __DEVICE__ bool islessequal(float __x, float __y) {
   return __builtin_islessequal(__x, __y);

 __DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
 __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
 __DEVICE__ bool isunordered(float __x, float __y) {
   return __builtin_isunordered(__x, __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);
 }
 __DEVICE__ float pow(float __base, int __iexp) {
   return ::powif(__base, __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__
 
 #endif

cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h

 // absence of a prior __device__ decl, those constexpr functions may become
 // implicitly host+device.  host+device functions can't be overloaded, so that
 // would preclude the use of our own __device__ overloads for these functions.
 
 #pragma push_macro("__DEVICE__")
 #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__")
 
 #endif

cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h

 #include "device_functions.hpp"
 #include "sm_20_atomic_functions.hpp"
 #include "sm_20_intrinsics.hpp"
 #include "sm_32_atomic_functions.hpp"
 // sm_30_intrinsics.h has declarations that use default argument, so
 // 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"
 
 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
 #pragma pop_macro("__CUDA_ARCH__")
 #pragma pop_macro("__THROW")