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

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Sun Oct  2 10:24:50 2016
New Revision: 283068

URL: http://llvm.org/viewvc/llvm-project?rev=283068&view=rev
Log:
[CUDA] Allow extern __shared__ on empty-length arrays.

"extern __shared__ int x[]" is OK.

Modified:
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/test/SemaCUDA/extern-shared.cu

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=283068&r1=283067&r2=283068&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Sun Oct  2 10:24:50 2016
@@ -3714,7 +3714,9 @@ static void handleSharedAttr(Sema &S, De
                                                  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;
   }

Modified: cfe/trunk/test/SemaCUDA/extern-shared.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/extern-shared.cu?rev=283068&r1=283067&r2=283068&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/extern-shared.cu (original)
+++ cfe/trunk/test/SemaCUDA/extern-shared.cu Sun Oct  2 10:24:50 2016
@@ -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'}}