diff options
| -rw-r--r-- | clang/lib/Sema/SemaDeclAttr.cpp | 4 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/extern-shared.cu | 11 |
2 files changed, 13 insertions, 2 deletions
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 21cace432c4..e06f110f0e4 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3714,7 +3714,9 @@ static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) { Attr.getName())) return; auto *VD = cast<VarDecl>(D); - if (VD->hasExternalStorage()) { + // extern __shared__ is only allowed on arrays with no length (e.g. + // "int x[]"). + if (VD->hasExternalStorage() && !isa<IncompleteArrayType>(VD->getType())) { S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD; return; } diff --git a/clang/test/SemaCUDA/extern-shared.cu b/clang/test/SemaCUDA/extern-shared.cu index 9450bbfa5ad..9fd7bbba593 100644 --- a/clang/test/SemaCUDA/extern-shared.cu +++ b/clang/test/SemaCUDA/extern-shared.cu @@ -5,10 +5,19 @@ __device__ void foo() { extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}} + extern __shared__ int arr[]; // ok + extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}} + extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}} + extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}} } __host__ __device__ void bar() { - extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}} + extern __shared__ int arr[]; // ok + extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}} + extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}} + extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}} } extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}} +extern __shared__ int global_arr[]; // ok +extern __shared__ int global_arr1[1]; // expected-error {{__shared__ variable 'global_arr1' cannot be 'extern'}} |
