[CUDA] Add hack so code which includes "curand.h" doesn't break.

Summary:
curand.h includes curand_mtgp32_kernel.h.  In host mode, this header
redefines threadIdx and blockDim, giving them their "proper" types of
uint3 and dim3, respectively.

clang has its own plan for these variables -- their types are magic
builtin classes.  So these redefinitions are incompatible.

As a hack, we force-include the offending CUDA header and use #defines
to get the right types for threadIdx and blockDim.

Reviewers: tra

Subscribers: echristo, cfe-commits

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@261776 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Headers/__clang_cuda_runtime_wrapper.h b/lib/Headers/__clang_cuda_runtime_wrapper.h
index 47cb6a7e6a6a92ac8426294ddef7f0e13411f5cc..fb527dc9abc1fde5c1cee561fd4f46c03ae7d0db 100644 (file)
--- a/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ b/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -247,5 +247,19 @@ __device__ static inline void *malloc(size_t __size) {
 
 #include <__clang_cuda_cmath.h>
 
+// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
+// mode, giving them their "proper" types of dim3 and uint3.  This is
+// incompatible with the types we give in cuda_builtin_vars.h.  As as hack,
+// force-include the header (nvcc doesn't include it by default) but redefine
+// dim3 and uint3 to our builtin types.  (Thankfully dim3 and uint3 are only
+// used here for the redeclarations of blockDim and threadIdx.)
+#pragma push_macro("dim3")
+#pragma push_macro("uint3")
+#define dim3 __cuda_builtin_blockDim_t
+#define uint3 __cuda_builtin_threadIdx_t
+#include "curand_mtgp32_kernel.h"
+#pragma pop_macro("dim3")
+#pragma pop_macro("uint3")
+
 #endif // __CUDA__
 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__