r261776 - [CUDA] Add hack so code which includes "curand.h" doesn't break.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 24 13:49:32 PST 2016
Date: Wed Feb 24 15:49:31 2016
New Revision: 261776
[CUDA] Add hack so code which includes "curand.h" doesn't break.
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.
Subscribers: echristo, cfe-commits
Differential Revision: http://reviews.llvm.org/D17562
--- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Wed Feb 24 15:49:31 2016
@@ -247,5 +247,19 @@ __device__ static inline void *malloc(si
+// 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.)
+#define dim3 __cuda_builtin_blockDim_t
+#define uint3 __cuda_builtin_threadIdx_t
#endif // __CUDA__
#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
More information about the cfe-commits