r265751 - [CUDA] Tweak math forward declares so we're compatible with libstdc++4.9.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 7 16:55:53 PDT 2016
Author: jlebar
Date: Thu Apr 7 18:55:53 2016
New Revision: 265751
URL: http://llvm.org/viewvc/llvm-project?rev=265751&view=rev
Log:
[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
Modified:
cfe/trunk/lib/Headers/__clang_cuda_cmath.h
cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h
cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
Modified: cfe/trunk/lib/Headers/__clang_cuda_cmath.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_cmath.h?rev=265751&r1=265750&r2=265751&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_cmath.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_cmath.h Thu Apr 7 18:55:53 2016
@@ -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
__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, f
__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 _
__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__
Modified: cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h?rev=265751&r1=265750&r2=265751&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h Thu Apr 7 18:55:53 2016
@@ -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__")
Modified: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h?rev=265751&r1=265750&r2=265751&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Thu Apr 7 18:55:53 2016
@@ -189,8 +189,21 @@ static inline __device__ void __brkpt(in
// 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"
More information about the cfe-commits
mailing list