[CUDA] Declare our __device__ math functions in the same inline namespace as our standard library.

Summary:
Currently we declare our inline __device__ math functions in namespace
std.  But libstdc++ and libc++ declare these functions in an inline
namespace inside namespace std.  We need to match this because, in a
later patch, we want to get e.g. <complex> to use our device overloads,
and it only will if those overloads are in the right inline namespace.

Reviewers: tra

Subscribers: cfe-commits, jhen

Differential Revision: https://reviews.llvm.org/D24977

llvm-svn: 283678

1 files changed, 21 insertions, 1 deletions

diff --git a/clang/lib/Headers/__clang_cuda_cmath.h b/clang/lib/Headers/__clang_cuda_cmath.h
index 01aaa2f88ae..569374584e4 100644
--- a/clang/lib/Headers/__clang_cuda_cmath.h
+++ b/clang/lib/Headers/__clang_cuda_cmath.h
@@ -316,7 +316,19 @@ scalbn(__T __x, int __exp) {
   return std::scalbn((double)__x, __exp);
 }
 
+// We need to define these overloads in exactly the namespace our standard
+// library uses (including the right inline namespace), otherwise they won't be
+// picked up by other functions in the standard library (e.g. functions in
+// <complex>).  Thus the ugliness below.
+#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
+_LIBCPP_BEGIN_NAMESPACE_STD
+#else
 namespace std {
+#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
+_GLIBCXX_BEGIN_NAMESPACE_VERSION
+#endif
+#endif
+
 // Pull the new overloads we defined above into namespace std.
 using ::acos;
 using ::acosh;
@@ -451,7 +463,15 @@ using ::tanf;
 using ::tanhf;
 using ::tgammaf;
 using ::truncf;
-}
+
+#ifdef _LIBCPP_END_NAMESPACE_STD
+_LIBCPP_END_NAMESPACE_STD
+#else
+#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
+_GLIBCXX_END_NAMESPACE_VERSION
+#endif
+} // namespace std
+#endif
 
 #undef __DEVICE__