diff options
| -rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 1 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaDeclAttr.cpp | 16 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/bad-attributes.cu | 2 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/extern-shared.cu | 14 |
4 files changed, 31 insertions, 2 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a65cf97ca60..66a5bce86c3 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -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 " diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 250380b6a1c..ff63d7b1ce6 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3696,6 +3696,19 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, 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, Scope *scope, Decl *D, 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); diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu index 4cb43e240b0..f3e1688462d 100644 --- a/clang/test/SemaCUDA/bad-attributes.cu +++ b/clang/test/SemaCUDA/bad-attributes.cu @@ -42,6 +42,8 @@ __constant__ __shared__ int z8; // expected-error {{attributes are not compatib __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}} diff --git a/clang/test/SemaCUDA/extern-shared.cu b/clang/test/SemaCUDA/extern-shared.cu new file mode 100644 index 00000000000..9450bbfa5ad --- /dev/null +++ b/clang/test/SemaCUDA/extern-shared.cu @@ -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'}} |
