From: Justin Lebar <jlebar@google.com>
Date: Sun, 2 Oct 2016 15:24:50 +0000 (+0000)
Subject: [CUDA] Allow extern __shared__ on empty-length arrays.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=90917bdf3bb2a971530383aec34966ecff7e16f3;p=clang

[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 21cace432c..e06f110f0e 100644
--- 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 9450bbfa5a..9fd7bbba59 100644
--- 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'}}