[PATCH] Fix addrspace when emitting constructors of static local variables

[Jingyue Wu]

Hi eliben, nlewycky, pcc, rsmith,

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.

http://reviews.llvm.org/D8575

Files:
  lib/CodeGen/CGDeclCXX.cpp
  test/CodeGenCUDA/address-spaces.cu

Index: lib/CodeGen/CGDeclCXX.cpp
===================================================================
--- lib/CodeGen/CGDeclCXX.cpp
+++ lib/CodeGen/CGDeclCXX.cpp
@@ -139,6 +139,16 @@
   const Expr *Init = D.getInit();
   QualType T = D.getType();
 
+  // The address space of D may be more generic than the address space of
+  // DeclPtr. In that case, we perform an addrspacecast.
+  unsigned DeclaredAddrSpace = getContext().getTargetAddressSpace(T);
+  unsigned ActualAddrSpace = CGM.GetGlobalVarAddressSpace(&D, DeclaredAddrSpace);
+  if (ActualAddrSpace != DeclaredAddrSpace) {
+    llvm::Type *LTy = CGM.getTypes().ConvertTypeForMem(T);
+    llvm::PointerType *PTy = llvm::PointerType::get(LTy, DeclaredAddrSpace);
+    DeclPtr = llvm::ConstantExpr::getAddrSpaceCast(DeclPtr, PTy);
+  }
+
   if (!T->isReferenceType()) {
     if (getLangOpts().OpenMP && D.hasAttr<OMPThreadPrivateDeclAttr>())
       (void)CGM.getOpenMPRuntime().emitThreadPrivateVarDefinition(
Index: test/CodeGenCUDA/address-spaces.cu
===================================================================
--- test/CodeGenCUDA/address-spaces.cu
+++ test/CodeGenCUDA/address-spaces.cu
@@ -100,3 +100,14 @@
 }
 // CHECK: define float* @_Z5func5v()
 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
+
+struct StructWithCtor {
+  __device__ StructWithCtor(): data(1) {}
+  int data;
+};
+
+__device__ void construct_shared_struct() {
+// CHECK-LABEL: define void @_Z23construct_shared_structv()
+  __shared__ StructWithCtor s;
+// CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
+}

EMAIL PREFERENCES
  http://reviews.llvm.org/settings/panel/emailpreferences/
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D8575.22539.patch
Type: text/x-patch
Size: 1702 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150324/48800ac0/attachment.bin>