Author: jlebar Date: Thu Oct 13 13:45:13 2016 New Revision: 284144 URL: http://llvm.org/viewvc/llvm-project?rev=284144&view=rev Log: [CUDA] Disallow __shared__ variables in host functions.
Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25143 Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/lib/Sema/SemaDeclAttr.cpp cfe/trunk/test/SemaCUDA/bad-attributes.cu Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=284144&r1=284143&r2=284144&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Oct 13 13:45:13 2016 @@ -6747,6 +6747,9 @@ def err_cuda_vla : Error< "cannot use variable-length arrays in " "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">; +def err_cuda_host_shared : Error< + "__shared__ local variables not allowed in " + "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">; def warn_non_pod_vararg_with_format_string : Warning< Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=284144&r1=284143&r2=284144&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original) +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Thu Oct 13 13:45:13 2016 @@ -3724,6 +3724,10 @@ static void handleSharedAttr(Sema &S, De S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD; return; } + if (S.getLangOpts().CUDA && VD->hasLocalStorage() && + S.CUDADiagIfHostCode(Attr.getLoc(), diag::err_cuda_host_shared) + << S.CurrentCUDATarget()) + return; D->addAttr(::new (S.Context) CUDASharedAttr( Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); } Modified: cfe/trunk/test/SemaCUDA/bad-attributes.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-attributes.cu?rev=284144&r1=284143&r2=284144&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/bad-attributes.cu (original) +++ cfe/trunk/test/SemaCUDA/bad-attributes.cu Thu Oct 13 13:45:13 2016 @@ -65,6 +65,7 @@ __global__ static inline void foobar() { __constant__ int global_constant; void host_fn() { __constant__ int c; // expected-error {{__constant__ variables must be global}} + __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}} } __device__ void device_fn() { __constant__ int c; // expected-error {{__constant__ variables must be global}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
