[CUDA] Allow trivial constructors as initializer for __shared__ variables.

[tra]

Summary

r245786 (D12241) disabled all initializers for shared variables. It also prevents placement of records in shared space as they have constructor as their initializer.

[jingyue]

Thanks for fixing this.

lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D) {
   const VarDecl *InitDecl;
   const Expr *InitExpr = D->getAnyInitializer(InitDecl);
 
   // CUDA E.2.4.1 "__shared__ variables cannot have an initialization as part
   // of their declaration."
   if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice
       && D->hasAttr<CUDASharedAttr>()) {
     if (InitExpr) {
+      const auto *C = dyn_cast<CXXConstructExpr>(InitExpr);
+      if (C == nullptr || !C->getConstructor()->hasTrivialBody())
         Error(D->getLocation(),
               "__shared__ variable cannot have an initialization.");
     }
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   } else if (!InitExpr) {
     // This is a tentative definition; tentative definitions are
     // implicitly initialized with { 0 }.
     //
     // Note that tentative definitions are only emitted at the end of
     // a translation unit, so they should never have incomplete

test/CodeGenCUDA/address-spaces.cu

 };
 
 // CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer
 // CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @b = addrspace(3) global float undef
+// CHECK: @c = addrspace(3) global %struct.c undef
+// CHECK  @d = addrspace(3) global %struct.d undef
 
 __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
   i++;
 
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
   j++;
 

 // CHECK-LABEL: define i32 @_Z23construct_shared_structv()
   __shared__ StructWithCtor s;
 // CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
   __shared__ StructWithCtor t(s);
 // CHECK: call void @_ZN14StructWithCtorC1ERKS_(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*), %struct.StructWithCtor* dereferenceable(4) addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
   return t.getData();
 // CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
 }
+
+// Make sure we allow __shared__ structures with default or empty constructors.
+struct c {
+  int i;
+};
+__shared__ struct c c;
+
+struct d {
+  int i;
+  d() {}
+};
+__shared__ struct d d;