[CUDA] Disallow __shared__ variables in host functions.

[jlebar]

Update to new CUDADiagIfHostCode API.

[rnk]

lgtm

cfe/trunk/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 err_cuda_host_shared : Error<
+    "__shared__ local variables not allowed in "
+    "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
 def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
 
 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<

cfe/trunk/lib/Sema/SemaDeclAttr.cpp

if (checkAttrMutualExclusion<CUDAConstantAttr>(S, D, Attr.getRange(),
     return;
   auto *VD = cast<VarDecl>(D);
   // 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;
   }
+  if (S.getLangOpts().CUDA && VD->hasLocalStorage() &&
+      S.CUDADiagIfHostCode(Attr.getLoc(), diag::err_cuda_host_shared)
+          << S.CurrentCUDATarget())
+    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(),

cfe/trunk/test/SemaCUDA/bad-attributes.cu

 __global__ static inline void foobar() {};
 #ifdef EXPECT_INLINE_WARNING
 // expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}}
 #endif
 
 __constant__ int global_constant;
 void host_fn() {
   __constant__ int c; // expected-error {{__constant__ variables must be global}}
+  __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
 }
 __device__ void device_fn() {
   __constant__ int c; // expected-error {{__constant__ variables must be global}}
 }