From: Justin Lebar Date: Wed, 24 Feb 2016 21:49:31 +0000 (+0000) Subject: [CUDA] Add hack so code which includes "curand.h" doesn't break. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=5a68fbd86b5d46737cb456194701fb25474ee87d;p=clang [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 47cb6a7e6a..fb527dc9ab 100644 --- 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__