[CUDA] Allow extern __shared__ on empty-length arrays.

"extern __shared__ int x[]" is OK.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283068 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 21cace432c48beddc5f80281787de9de2397cee0..e06f110f0e4917ae5f39d8b0a3c8dc8c9a071c7d 100644 (file)
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/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/test/SemaCUDA/extern-shared.cu b/test/SemaCUDA/extern-shared.cu
index 9450bbfa5ad6842d148159cb15bd915267e20d14..9fd7bbba593c959b9712f68edef3eb6ed28e436a 100644 (file)
--- a/test/SemaCUDA/extern-shared.cu
+++ b/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'}}