[CUDA] Restrict init of local __shared__ variables to empty constructors only.

[tra]

Summary

While shared variables look like any other variable with a static storage class to compiler, they behave differently on device side.

• one instance is created per block of GPUS, so standard "initialize once using guard variable" model does not quite work.
• lifetime of the variables ends when the global function exits. Again, it does not fit current assumption about static local vars as we will need to init them again if that function is called again.
• with that in mind, deinitialization on app exit does not work either as the variable no longer exists past its kernel's exit.

nvcc takes a rather dangerous shortcut and allows non-empty constructors for local static variables. It calls initializer on every entry into the scope and produces a warning that there's going to be a data race as there will be many kernels doing init on many instances of that shared variable. It also calls destructors on exit from the scope. Now, imagine recursive call of a function with a local static variable...

Until we figure out better way to deal with this, clang will only allow empty constructors for local shared variables in a way identical to restrictions imposed on dynamic initializers for global variables.

[jlebar]

While I think this is 100% the right thing to do, I am worried about breaking existing targets. Maybe we need an escape valve, at least until we get that sorted out? Unless you're pretty confident this isn't happening / will be easy enough to fix.

lib/CodeGen/CGDecl.cpp

376

Please set off "which is ensured by Sema" somehow. I'd probably say

initializers. (This is ensured by Sema.)

lib/Sema/SemaDecl.cpp

10416

s/have the same checks applied/apply the same checks

[tra]

Reworded comments.

Removed tests that no longer apply as we don't generate constructors for static local variables on device side.
Empty constructor cases are already covered by test/CodeGenCUDA/device-var-init.cu.

[jlebar]

Art makes the good point that any code which is broken by this change is itself currently racy at the very best. And we have a way to fix any code which hits this new error (just manually initialize your shared variable using placement new from within thread 0). So, lgtm.

[tra]

In D20039#424067, @jlebar wrote:

While I think this is 100% the right thing to do, I am worried about breaking existing targets. Maybe we need an escape valve, at least until we get that sorted out? Unless you're pretty confident this isn't happening / will be easy enough to fix.

While empty constructors for shared variables are not unusual, I haven't seen any non-empty constructors in practice.
Considering that such constructor would crash clang until this patch, it would be hard to miss such cases.

Escape hatch would require to make clang to do something reasonable.
Currently llvm crashes because we attempt to generate a load from a guard variable using atomic load instruction that's not supported by NVPTX backend. Even if it was supported, we'd also need to implement _cxa_guard_acquire/release.

We can disable thread-safe guard variants. This will be broken, too, because there will be many per-block instances of the variable, but only one global guard.

We can spend an effort to do the same thing as nvcc -- no guard of any kind, ctor/dtor on function entry/exit and compiler warning about guaranteed data race, but I don't see much point doing *that*.

IMO fixing the source code to initialize static variable explicitly would be the right thing to do.

lib/CodeGen/CGDecl.cpp

void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D,
   if (D.getType()->isVariablyModifiedType())
     EmitVariablyModifiedType(D.getType());
 
   // Save the type in case adding the initializer forces a type change.
   llvm::Type *expectedType = addr->getType();
 
   llvm::GlobalVariable *var =
     cast<llvm::GlobalVariable>(addr->stripPointerCasts());
+
+  // CUDA's local and local static __shared__ variables should not
+  // have any non-empty initializers which is ensured by Sema.
+  // Whatever initializer such variable may have when it gets here is
+  // a no-op and should not be emitted.
+  bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+                         D.hasAttr<CUDASharedAttr>();
   // If this value has an initializer, emit it.
-  if (D.getInit())
+  if (D.getInit() && !isCudaSharedVar)
     var = AddInitializerToStaticVarDecl(D, var);
 
   var->setAlignment(alignment.getQuantity());
 
   if (D.hasAttr<AnnotateAttr>())
     CGM.AddGlobalAnnotations(&D, var);
 
   if (const SectionAttr *SA = D.getAttr<SectionAttr>())

 }
 
 void CodeGenModule::EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D,
                                             CodeGenFunction *CGF) {
   if (!LangOpts.OpenMP || (!LangOpts.EmitAllDecls && !D->isUsed()))
     return;
   getOpenMPRuntime().emitUserDefinedReduction(CGF, D);
 }
-

lib/Sema/SemaDecl.cpp

if (FunctionDecl *FD =
         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.
+  // __device__ variables.  We also must have the same checks applied
+  // to all __shared__ variables whether they are local or
+  // not.
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     const Expr *Init = VD->getInit();
-    const bool IsGlobal = VD->hasGlobalStorage() && !VD->isStaticLocal();
-    if (Init && IsGlobal &&
+    if (Init && VD->hasGlobalStorage() &&
         (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
          VD->hasAttr<CUDASharedAttr>())) {
+      assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
       bool AllowedInit = false;
       if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
         AllowedInit =
             isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
       // We'll allow constant initializers even if it's a non-empty
       // constructor according to CUDA rules. This deviates from NVCC,
       // but allows us to handle things like constexpr constructors.
       if (!AllowedInit &&

test/CodeGenCUDA/device-var-init.cu

 // dynamic in-class field initializer -- not allowed
 __device__ int f();
 struct NCF {
   int ncf = f();
 };
 
 // static in-class field initializer.  NVCC does not allow it, but
 // clang generates static initializer for this, so we'll accept it.
+// We still can't use it on __shared__ vars as they don't allow *any*
+// initializers.
 struct NCFS {
   int ncfs = 3;
 };
 
 // undefined templated constructor -- not allowed
 struct UTC {
   template <typename... T> __device__ UTC(T...);
 };
__device__ void df() {
   T_F_T t_f_t;
   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;
+  static __shared__ EC s_ec;
+  static __shared__ ETC s_etc;
 #if ERROR_CASE
+  static __shared__ NCFS s_ncfs;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __shared__ UC s_uc;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   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 @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
 // CHECK:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
 // CHECK:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
 // CHECK:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
 // CHECK:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
 // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
 // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
 // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
-// CHECK:   call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*))
+// CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+// CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
 // CHECK: ret void
 
 // We should not emit global init function.
 // CHECK-NOT: @_GLOBAL__sub_I