Fix addrspace when emitting constructors of static local variables

[jingyue]

Summary

Due to CUDA's implicit address space casting, the type of a static local
variable may be more specific (i.e. with address space qualifiers) than
the type expected by the constructor. Emit an addrspacecast in that
case.

[eliben]

lib/CodeGen/CGDeclCXX.cpp

145

Why is DeclaredAddrSpace passed to GetGlobalVarAddressSpace here?

[eliben]

In general this looks good; there are quite a few layers of calls to emit the global var initialization, and I'm not familiar enough to say if this is inserted at the right layer. Hopefully one of the other reviewers will be able to clarify this.

lib/CodeGen/CGDeclCXX.cpp

143

IMHO this could use a bit more clarification - when can this happen and what does "more generic" mean. Perhaps a small input code example in the comment + showing what VarDecl and DeclPtr end up being in that case?

[jingyue]

compute ActualAddrSpace in a different way

[jingyue]

Speaking of which layer we should add the conversion, I actually worked on another patch (http://reviews.llvm.org/differential/diff/22588/) that has EmitStaticVarDecl pass the addrspacecasted global variable to AddInitializerToStaticVarDecl. They both fix the issue. The patch here is to me simpler, because the other patch requires several interface-level changes.

lib/CodeGen/CGDeclCXX.cpp

143

Done.

145

DeclaredAddrSpace is the type GetGlobalVarAddressSpace should return when D is not address space qualified. I agree it's a little confusing, so I now compute ActualAddrSpace in a different way. It's equivalent, but more understandable.

[eliben]

LGTM - the comment w/ the example really helps. Thanks

[eliben]

Oh another something - do we have tests that other special members work in this case? E.g. if StructWithCtor also had a copy ctor or dtor, for example

[jingyue]

Test the copy constructor and regular member functions

[jingyue]

Interesting, clang crashes if StructWithCtor has a dtor. However, as far as I can see, it is a separate issue unrelated to address space casting and should be fixed in another patch. In this version, I only tested the default contructor, the copy constructor and a regular member function.

FYI, the call stack if I add the dtor:

#0 0x43b0124 llvm::sys::PrintStackTrace(llvm::raw_ostream&) /home/jingyue/Work/llvm/llvm-git/lib/Support/Unix/Signals.inc:425:0            
#1 0x43b0439 PrintStackTraceSignalHandler(void*) /home/jingyue/Work/llvm/llvm-git/lib/Support/Unix/Signals.inc:483:0
#2 0x43aef39 SignalHandler(int) /home/jingyue/Work/llvm/llvm-git/lib/Support/Unix/Signals.inc:199:0
#3 0x7fc8111cf340 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10340)   
#4 0x135f330 clang::Decl::hasAttrs() const /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/Frontend/../../include/clang/AST/DeclBase.h:422:0
#5 0x16cf9f2 bool clang::Decl::hasAttr<clang::CUDAConstantAttr>() const /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/../../include/clang/AST/DeclBase.h:483:0
#6 0x16c49fd clang::CodeGen::CodeGenModule::GetGlobalVarAddressSpace(clang::VarDecl const*, unsigned int) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:1903:0
#7 0x16c4151 clang::CodeGen::CodeGenModule::GetOrCreateLLVMGlobal(llvm::StringRef, llvm::PointerType*, clang::VarDecl const*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:1757:0
#8 0x16c486d clang::CodeGen::CodeGenModule::CreateRuntimeVariable(llvm::Type*, llvm::StringRef) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:1875:0
#9 0x1766e3a emitGlobalDtorWithCXAAtExit(clang::CodeGen::CodeGenFunction&, llvm::Constant*, llvm::Constant*, bool) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/ItaniumCXXABI.cpp:1930:0
#10 0x1766f42 (anonymous namespace)::ItaniumCXXABI::registerGlobalDtor(clang::CodeGen::CodeGenFunction&, clang::VarDecl const&, llvm::Constant*, llvm::Constant*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/ItaniumCXXABI.cpp:1947:0
#11 0x180600a EmitDeclDestroy(clang::CodeGen::CodeGenFunction&, clang::VarDecl const&, llvm::Constant*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDeclCXX.cpp:112:0
#12 0x18063af clang::CodeGen::CodeGenFunction::EmitCXXGlobalVarDeclInit(clang::VarDecl const&, llvm::Constant*, bool) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDeclCXX.cpp:176:0
#13 0x1766b33 (anonymous namespace)::ItaniumCXXABI::EmitGuardedInit(clang::CodeGen::CodeGenFunction&, clang::VarDecl const&, llvm::GlobalVariable*, bool) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/ItaniumCXXABI.cpp:1888:0
#14 0x18069d5 clang::CodeGen::CodeGenFunction::EmitCXXGuardedInit(clang::VarDecl const&, llvm::GlobalVariable*, bool) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDeclCXX.cpp:249:0
#15 0x17fba61 clang::CodeGen::CodeGenFunction::AddInitializerToStaticVarDecl(clang::VarDecl const&, llvm::GlobalVariable*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDecl.cpp:293:0
#16 0x17fbdfb clang::CodeGen::CodeGenFunction::EmitStaticVarDecl(clang::VarDecl const&, llvm::GlobalValue::LinkageTypes) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDecl.cpp:363:0
#17 0x17faf27 clang::CodeGen::CodeGenFunction::EmitVarDecl(clang::VarDecl const&) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDecl.cpp:136:0
#18 0x17fae7f clang::CodeGen::CodeGenFunction::EmitDecl(clang::Decl const&) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGDecl.cpp:111:0
#19 0x167d3b5 clang::CodeGen::CodeGenFunction::EmitDeclStmt(clang::DeclStmt const&) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGStmt.cpp:1098:0
#20 0x1679d44 clang::CodeGen::CodeGenFunction::EmitSimpleStmt(clang::Stmt const*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGStmt.cpp:248:0
#21 0x1679495 clang::CodeGen::CodeGenFunction::EmitStmt(clang::Stmt const*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGStmt.cpp:50:0
#22 0x1679ff5 clang::CodeGen::CodeGenFunction::EmitCompoundStmtWithoutScope(clang::CompoundStmt const&, bool, clang::CodeGen::AggValueSlot) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CGStmt.cpp:282:0                                                                 
#23 0x16b1035 clang::CodeGen::CodeGenFunction::EmitFunctionBody(clang::CodeGen::FunctionArgList&, clang::Stmt const*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenFunction.cpp:760:0
#24 0x16b19a2 clang::CodeGen::CodeGenFunction::GenerateCode(clang::GlobalDecl, llvm::Function*, clang::CodeGen::CGFunctionInfo const&) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenFunction.cpp:879:0
#25 0x16c68ac clang::CodeGen::CodeGenModule::EmitGlobalFunctionDefinition(clang::GlobalDecl, llvm::GlobalValue*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:2476:0
#26 0x16c34e3 clang::CodeGen::CodeGenModule::EmitGlobalDefinition(clang::GlobalDecl, llvm::GlobalValue*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:1533:0
#27 0x16c2d6a clang::CodeGen::CodeGenModule::EmitGlobal(clang::GlobalDecl) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:1390:0
#28 0x16ca527 clang::CodeGen::CodeGenModule::EmitTopLevelDecl(clang::Decl*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenModule.cpp:3247:0
#29 0x1625ad9 (anonymous namespace)::CodeGeneratorImpl::HandleTopLevelDecl(clang::DeclGroupRef) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/ModuleBuilder.cpp:121:0
#30 0x160a83a clang::BackendConsumer::HandleTopLevelDecl(clang::DeclGroupRef) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenAction.cpp:105:0
#31 0x18e84ed clang::ParseAST(clang::Sema&, bool, bool) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/Parse/ParseAST.cpp:142:0
#32 0x1319840 clang::ASTFrontendAction::ExecuteAction() /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/Frontend/FrontendAction.cpp:539:0
#33 0x160d228 clang::CodeGenAction::ExecuteAction() /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/CodeGen/CodeGenAction.cpp:734:0
#34 0x131931b clang::FrontendAction::Execute() /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/Frontend/FrontendAction.cpp:443:0
#35 0x12deb8e clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/Frontend/CompilerInstance.cpp:807:0
#36 0x12a0c81 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) /home/jingyue/Work/llvm/llvm-git/tools/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp:222:0
#37 0x128cd9c cc1_main(llvm::ArrayRef<char const*>, char const*, void*) /home/jingyue/Work/llvm/llvm-git/tools/clang/tools/driver/cc1_main.cpp:110:0
#38 0x129a54e ExecuteCC1Tool(llvm::ArrayRef<char const*>, llvm::StringRef) /home/jingyue/Work/llvm/llvm-git/tools/clang/tools/driver/driver.cpp:369:0                                                                                                                                 
#39 0x129ab30 main /home/jingyue/Work/llvm/llvm-git/tools/clang/tools/driver/driver.cpp:415:0
#40 0x7fc8101ceec5 __libc_start_main /build/buildd/eglibc-2.19/csu/libc-start.c:321:0
#41 0x128b549 _start (/home/jingyue/Work/llvm/install-git/bin/clang+0x128b549)

[eliben]

I think it's alright to handle the dtor crash in a separate change (if it looks like a separate issue). Just open a bug about it

[rsmith]

LGTM

lib/CodeGen/CGDeclCXX.cpp

 
 void CodeGenFunction::EmitCXXGlobalVarDeclInit(const VarDecl &D,
                                                llvm::Constant *DeclPtr,
                                                bool PerformInit) {
 
   const Expr *Init = D.getInit();
   QualType T = D.getType();
 
+  // The address space of a static local variable (DeclPtr) may be different
+  // from the address space of the "this" argument of the constructor. In that
+  // case, we need an addrspacecast before calling the constructor.
+  //
+  // struct StructWithCtor {
+  //   __device__ StructWithCtor() {...}
+  // };
+  // __device__ void foo() {
+  //   __shared__ StructWithCtor s;
+  //   ...
+  // }
+  //
+  // For example, in the above CUDA code, the static local variable s has a
+  // "shared" address space qualifier, but the constructor of StructWithCtor
+  // expects "this" in the "generic" address space.
+  unsigned ExpectedAddrSpace = getContext().getTargetAddressSpace(T);
+  unsigned ActualAddrSpace = DeclPtr->getType()->getPointerAddressSpace();
+  if (ActualAddrSpace != ExpectedAddrSpace) {
+    llvm::Type *LTy = CGM.getTypes().ConvertTypeForMem(T);
+    llvm::PointerType *PTy = llvm::PointerType::get(LTy, ExpectedAddrSpace);
+    DeclPtr = llvm::ConstantExpr::getAddrSpaceCast(DeclPtr, PTy);
+  }
+
   if (!T->isReferenceType()) {
     if (getLangOpts().OpenMP && D.hasAttr<OMPThreadPrivateDeclAttr>())
       (void)CGM.getOpenMPRuntime().emitThreadPrivateVarDefinition(
           &D, DeclPtr, D.getAttr<OMPThreadPrivateDeclAttr>()->getLocation(),
           PerformInit, this);
     if (PerformInit)
       EmitDeclInit(*this, D, DeclPtr);
     if (CGM.isTypeConstant(D.getType(), true))

test/CodeGenCUDA/address-spaces.cu

 
 __shared__ float b;
 
 __device__ float *func5() {
   return &b; // implicit cast from a return value
 }
 // CHECK: define float* @_Z5func5v()
 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
+
+struct StructWithCtor {
+  __device__ StructWithCtor(): data(1) {}
+  __device__ StructWithCtor(const StructWithCtor &second): data(second.data) {}
+  __device__ int getData() { return data; }
+  int data;
+};
+
+__device__ int construct_shared_struct() {
+// 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*))
+}