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'}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits