diff options
| author | Justin Lebar <jlebar@google.com> | 2016-10-02 15:24:50 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-10-02 15:24:50 +0000 |
| commit | 281ce2af178e4f751c22418ea33ff056c376ad7c (patch) | |
| tree | ff557d9d3a2baa283ec9dd269fa015faca34d9b3 | |
| parent | bbce7e21e50f937ab32acdf3c84c9834e4d1bce1 (diff) | |
| download | bcm5719-llvm-281ce2af178e4f751c22418ea33ff056c376ad7c.tar.gz bcm5719-llvm-281ce2af178e4f751c22418ea33ff056c376ad7c.zip | |
[CUDA] Allow extern __shared__ on empty-length arrays.
"extern __shared__ int x[]" is OK.
llvm-svn: 283068
| -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'}} |
