[CUDA] Disallow 'extern __shared__' variables.

[jlebar]

Fix typo (and add a test to catch it).

[tra]

LGTM.

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_shared_var_init : Error<
     "initialization is not supported for __shared__ variables.">;
 def err_device_static_local_var : Error<
     "Within a __device__/__global__ function, "
     "only __shared__ variables may be marked \"static\"">;
 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 "
   "%select{function|block|method|constructor}2; expected type from format "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;
 // The arguments to this diagnostic should match the warning above.
 def err_cannot_pass_objc_interface_to_vararg_format : Error<
   "cannot pass object with interface type %1 by value to variadic "

clang/lib/Sema/SemaDeclAttr.cpp

 
 static void handleOptimizeNoneAttr(Sema &S, Decl *D,
                                    const AttributeList &Attr) {
   if (OptimizeNoneAttr *Optnone = S.mergeOptimizeNoneAttr(
           D, Attr.getRange(), Attr.getAttributeSpellingListIndex()))
     D->addAttr(Optnone);
 }
 
+static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+  if (checkAttrMutualExclusion<CUDASharedAttr>(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()) ||
       checkAttrMutualExclusion<CUDAHostAttr>(S, D, Attr.getRange(),
                                              Attr.getName())) {
     return;
   }
   FunctionDecl *FD = cast<FunctionDecl>(D);
case AttributeList::AT_Naked:
     break;
   case AttributeList::AT_NoReturn:
     handleNoReturnAttr(S, D, Attr);
     break;
   case AttributeList::AT_NoThrow:
     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);
     break;
   case AttributeList::AT_ObjCOwnership:
     handleObjCOwnershipAttr(S, D, Attr);
     break;
   case AttributeList::AT_ObjCPreciseLifetime:

clang/test/SemaCUDA/extern-shared.cu

+// 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'}}