[CUDA] Only __shared__ variables can be static local on device side.

[tra]

Summary

According to CUDA programming guide (v7.5):

E.2.9.4: Within the body of a device or global function, only shared variables may be declared with static storage class.

[jlebar]

What are we supposed to do if we encounter a static shared variable in an HD function? Presumably that also should be an error if we invoke the HD function from the device?

[tra]

In D20034#423945, @jlebar wrote:

What are we supposed to do if we encounter a static shared variable in an HD function? Presumably that also should be an error if we invoke the HD function from the device?

nvcc produces an error only of such HD function is used from a host function:

error: a function scope variable cannot be declared with "shared" inside a host function

Considering that each block gets its own instance of that shared variable, there's no way to access it from host in principle so the error makes sense. I'll add it in a separate patch.

[tra]

Updated tests in CodeGenCUDA/address-spaces.cu

[jlebar]

This patch regresses Eigen, because it raises an error even on host+device functions.

include/clang/Basic/DiagnosticSemaKinds.td

def err_cuda_unattributed_constexpr_cannot_overload_device : Error<
   "attribute, or build with -fno-cuda-host-device-constexpr.">;
 def note_cuda_conflicting_device_function_declared_here : Note<
   "conflicting __device__ function declared here">;
 def err_dynamic_var_init : Error<
     "dynamic initialization is not supported for "
     "__device__, __constant__, and __shared__ variables.">;
 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 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 "
   "%select{function|block|method|constructor}2; expected type from format "

lib/Sema/SemaDecl.cpp

if (VD->getTLSKind() && !hasDependentAlignment(VD)) {
       if (Context.getDeclAlign(VD) > MaxAlignChars) {
         Diag(VD->getLocation(), diag::err_tls_var_aligned_over_maximum)
           << (unsigned)Context.getDeclAlign(VD).getQuantity() << VD
           << (unsigned)MaxAlignChars.getQuantity();
       }
     }
   }
 
-  // Static locals inherit dll attributes from their function.
   if (VD->isStaticLocal()) {
     if (FunctionDecl *FD =
             dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
+      // Static locals inherit dll attributes from their function.
       if (Attr *A = getDLLAttr(FD)) {
         auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext()));
         NewAttr->setInherited(true);
         VD->addAttr(NewAttr);
       }
+      // CUDA E.2.9.4: Within the body of a __device__ or __global__
+      // function, only __shared__ variables may be declared with
+      // static storage class.
+      if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+          (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>()) &&
+          !VD->hasAttr<CUDASharedAttr>()) {
+        Diag(VD->getLocation(), diag::err_device_static_local_var);
+        VD->setInvalidDecl();
+      }
     }
   }
 
   // Perform check for initializers of device-side global variables.
   // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
   // 7.5). CUDA also allows constant initializers for __constant__ and
   // __device__ variables.
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {

test/CodeGenCUDA/address-spaces.cu

__device__ void foo() {
   i++;
 
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
   j++;
 
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
   k++;
 
-  static int li;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
-  li++;
-
-  __constant__ int lj;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
-  lj++;
-
   __shared__ int lk;
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
 }
 
 __device__ void func0() {
   __shared__ MyStruct a;
   MyStruct *ap = &a; // composite type

test/CodeGenCUDA/device-var-init.cu

__device__ void df() {
   T_FA_T t_fa_t;
   EC_I_EC ec_i_ec;
   EC_I_EC1 ec_i_ec1;
   T_V_T t_v_t;
   T_B_NEC t_b_nec;
   T_F_NEC t_f_nec;
   T_FA_NEC t_fa_nec;
   static __shared__ UC s_uc;
+#if ERROR_CASE
+  static __device__ int ds;
+  // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  static __constant__ int dc;
+  // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  static int v;
+  // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+#endif
 }
 
 // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
 // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
 // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
 // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
 // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
 // CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)