r282985 - [CUDA] Disallow 'extern __shared__' variables.

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Fri Sep 30 18:57:30 2016
New Revision: 282985

URL: http://llvm.org/viewvc/llvm-project?rev=282985&view=rev
Log:
[CUDA] Disallow 'extern __shared__' variables.

Also add a test that we disallow

  __constant__ __shared__ int x;

because it's possible to break this without breaking

  __shared__ __constant__ int x;

Reviewers: rnk

Subscribers: cfe-commits, tra

Differential Revision: https://reviews.llvm.org/D25125

Added:
    cfe/trunk/test/SemaCUDA/extern-shared.cu
Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/test/SemaCUDA/bad-attributes.cu

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=282985&r1=282984&r2=282985&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Fri Sep 30 18:57:30 2016
@@ -6722,6 +6722,7 @@ def err_device_static_local_var : Error<
 def err_cuda_vla : Error<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
+def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
 
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=282985&r1=282984&r2=282985&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Fri Sep 30 18:57:30 2016
@@ -3696,6 +3696,19 @@ static void handleOptimizeNoneAttr(Sema
     D->addAttr(Optnone);
 }
 
+static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+  if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(),
+                                                 Attr.getName()))
+    return;
+  auto *VD = cast<VarDecl>(D);
+  if (VD->hasExternalStorage()) {
+    S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
+    return;
+  }
+  D->addAttr(::new (S.Context) CUDASharedAttr(
+      Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex()));
+}
+
 static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(),
                                                Attr.getName()) ||
@@ -5639,8 +5652,7 @@ static void ProcessDeclAttribute(Sema &S
     handleSimpleAttribute<NoThrowAttr>(S, D, Attr);
     break;
   case AttributeList::AT_CUDAShared:
-    handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D,
-                                                                          Attr);
+    handleSharedAttr(S, D, Attr);
     break;
   case AttributeList::AT_VecReturn:
     handleVecReturnAttr(S, D, Attr);

Modified: cfe/trunk/test/SemaCUDA/bad-attributes.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-attributes.cu?rev=282985&r1=282984&r2=282985&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu (original)
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu Fri Sep 30 18:57:30 2016
@@ -42,6 +42,8 @@ __constant__ __shared__ int z8;  // expe
 __shared__ __device__ int z9;
 __shared__ __constant__ int z10;  // expected-error {{attributes are not compatible}}
 // expected-note at -1 {{conflicting attribute is here}}
+__constant__ __shared__ int z10a;  // expected-error {{attributes are not compatible}}
+// expected-note at -1 {{conflicting attribute is here}}
 
 __global__ __device__ void z11();  // expected-error {{attributes are not compatible}}
 // expected-note at -1 {{conflicting attribute is here}}

Added: cfe/trunk/test/SemaCUDA/extern-shared.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/extern-shared.cu?rev=282985&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/extern-shared.cu (added)
+++ cfe/trunk/test/SemaCUDA/extern-shared.cu Fri Sep 30 18:57:30 2016
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+__device__ void foo() {
+  extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
+}
+
+__host__ __device__ void bar() {
+  extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
+}
+
+extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}}