Author: jlebar Date: Fri Sep 30 18:57:30 2016 New Revision: 282985 URL: http://llvm.org/viewvc/llvm-project?rev=282985&view=rev Log: [CUDA] Disallow 'extern __shared__' variables.
Also add a test that we disallow __constant__ __shared__ int x; because it's possible to break this without breaking __shared__ __constant__ int x; Reviewers: rnk Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D25125 Added: cfe/trunk/test/SemaCUDA/extern-shared.cu 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=282985&r1=282984&r2=282985&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Fri Sep 30 18:57:30 2016 @@ -6722,6 +6722,7 @@ def err_device_static_local_var : Error< 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 warn_non_pod_vararg_with_format_string : Warning< "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic " Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=282985&r1=282984&r2=282985&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original) +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Fri Sep 30 18:57:30 2016 @@ -3696,6 +3696,19 @@ static void handleOptimizeNoneAttr(Sema D->addAttr(Optnone); } +static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(), + Attr.getName())) + return; + auto *VD = cast<VarDecl>(D); + if (VD->hasExternalStorage()) { + S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD; + return; + } + D->addAttr(::new (S.Context) CUDASharedAttr( + Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); +} + static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) { if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(), Attr.getName()) || @@ -5639,8 +5652,7 @@ static void ProcessDeclAttribute(Sema &S handleSimpleAttribute<NoThrowAttr>(S, D, Attr); break; case AttributeList::AT_CUDAShared: - handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D, - Attr); + handleSharedAttr(S, D, Attr); break; case AttributeList::AT_VecReturn: handleVecReturnAttr(S, D, Attr); Modified: cfe/trunk/test/SemaCUDA/bad-attributes.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-attributes.cu?rev=282985&r1=282984&r2=282985&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/bad-attributes.cu (original) +++ cfe/trunk/test/SemaCUDA/bad-attributes.cu Fri Sep 30 18:57:30 2016 @@ -42,6 +42,8 @@ __constant__ __shared__ int z8; // expe __shared__ __device__ int z9; __shared__ __constant__ int z10; // expected-error {{attributes are not compatible}} // expected-note@-1 {{conflicting attribute is here}} +__constant__ __shared__ int z10a; // expected-error {{attributes are not compatible}} +// expected-note@-1 {{conflicting attribute is here}} __global__ __device__ void z11(); // expected-error {{attributes are not compatible}} // expected-note@-1 {{conflicting attribute is here}} Added: cfe/trunk/test/SemaCUDA/extern-shared.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/extern-shared.cu?rev=282985&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/extern-shared.cu (added) +++ cfe/trunk/test/SemaCUDA/extern-shared.cu Fri Sep 30 18:57:30 2016 @@ -0,0 +1,14 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s + +#include "Inputs/cuda.h" + +__device__ void foo() { + extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}} +} + +__host__ __device__ void bar() { + extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}} +} + +extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
