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*))
+}