[CUDA] Work around a new quirk in CUDA9 headers.

[tra]

Summary

In CUDA-9 some of device-side math functions that we need are conditionally
defined within '#if _GLIBCXX_MATH_H'. We need to temporarily undo the guard
around inclusion of math_functions.hpp

[jlebar]

This is a bit of a Chesterton's Fence -- do we know why they're doing this?

I guess it's probably going to be OK because our overriding semantics will make it OK, and our test-suite tests (should) exercise all of math.h. But I'm still a little worried about it.

[tra]

I don't think we really care why they do it for nvcc.
My understanding is that nvcc needs to avoid name clashes between their implementation of functions and the ones that come from the host headers and that's why they have to tread really carefully around host includes.

We just need to get those definitions of the device-side functions and we don't care whether host-side functions with the same name are visible at the same time. I wish CUDA would just put all device-side function definitions into a separate file/namespace and let us cherry-pick them at will. Alas, we have to work around nvcc's work arounds...

BTW, this change essentially augments the job that the "#undef GNUC" above used to do in older CUDA versions. CUDA9 replaced GNUC with _GLIBCXX_MATH_H in CUDA-9 in some places.

[jlebar]

BTW, this change essentially augments the job that the "#undef GNUC" above used to do in older CUDA versions. CUDA9 replaced GNUC with _GLIBCXX_MATH_H in CUDA-9 in some places.

Ah, that's right. Okay then. :)

cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h

 // 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
+
+// CUDA-9 omits device-side definitions of some math functions if it sees
+// include guard from math.h wrapper from libstdc++. We have to undo the header
+// guard temporarily to get the definitions we need.
+#pragma push_macro("_GLIBCXX_MATH_H")
+#if CUDA_VERSION >= 9000
+#undef _GLIBCXX_MATH_H
+#endif
+
 #include "math_functions.hpp"
+#pragma pop_macro("_GLIBCXX_MATH_H")
 #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.