[CUDA] Disallow __constant__ local variables.

[tra]

LGTM.

clang/include/clang/Basic/DiagnosticSemaKinds.td

6727

Nit: Technically they are allowed in namespace scope.

[jlebar]

clang/include/clang/Basic/DiagnosticSemaKinds.td

6727

That's still a "global variable"? Or do you think calling it such will be confusing?

[tra]

clang/include/clang/Basic/DiagnosticSemaKinds.td

6727

It's not clear whether you mean global storage class or global namespace.
The code checks for global storage, but error message is could be interpreted either way, IMO.

I'll leave phrasing up to you.

[jlebar]

clang/include/clang/Basic/DiagnosticSemaKinds.td

6727

It's not clear whether you mean global storage class or global namespace.

So there's actually no such thing as "global" storage class. It's *static* storage (so helpful), meaning, a global variable or a static variable inside of a function.

But __constant__ symbols must actually be global variables -- nvcc doesn't allow them to appear inside function bodies.

I think "global" is the right way to describe what we're going after. Until a Language Lawyer corrects me, anyway. :)

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 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<
   "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 handleConstantAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+  if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, Attr.getRange(),
+                                               Attr.getName()))
+    return;
+  auto *VD = cast<VarDecl>(D);
+  if (!VD->hasGlobalStorage()) {
+    S.Diag(Attr.getLoc(), diag::err_cuda_nonglobal_constant);
+    return;
+  }
+  D->addAttr(::new (S.Context) CUDAConstantAttr(
+      Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex()));
+}
+
 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;
case AttributeList::AT_Availability:
     break;
   case AttributeList::AT_CarriesDependency:
     handleDependencyAttr(S, scope, D, Attr);
     break;
   case AttributeList::AT_Common:
     handleCommonAttr(S, D, Attr);
     break;
   case AttributeList::AT_CUDAConstant:
-    handleSimpleAttributeWithExclusions<CUDAConstantAttr, CUDASharedAttr>(S, D,
-                                                                          Attr);
+    handleConstantAttr(S, D, Attr);
     break;
   case AttributeList::AT_PassObjectSize:
     handlePassObjectSizeAttr(S, D, Attr);
     break;
   case AttributeList::AT_Constructor:
     handleConstructorAttr(S, D, Attr);
     break;
   case AttributeList::AT_CXX11NoReturn:

clang/test/SemaCUDA/bad-attributes.cu

struct S {
   // Although this is implicitly inline, we shouldn't warn.
   __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}}
 };
 
 __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}}
+}
+__device__ void device_fn() {
+  __constant__ int c; // expected-error {{__constant__ variables must be global}}
+}