Author: jingyue Date: Wed Mar 25 15:06:28 2015 New Revision: 233208 URL: http://llvm.org/viewvc/llvm-project?rev=233208&view=rev Log: Fix addrspace when emitting constructors of static local variables
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. Test Plan: Clang used to crash on the added test. Reviewers: nlewycky, pcc, eliben, rsmith Reviewed By: eliben, rsmith Subscribers: llvm-commits Differential Revision: http://reviews.llvm.org/D8575 Modified: cfe/trunk/lib/CodeGen/CGDeclCXX.cpp cfe/trunk/test/CodeGenCUDA/address-spaces.cu Modified: cfe/trunk/lib/CodeGen/CGDeclCXX.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDeclCXX.cpp?rev=233208&r1=233207&r2=233208&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGDeclCXX.cpp (original) +++ cfe/trunk/lib/CodeGen/CGDeclCXX.cpp Wed Mar 25 15:06:28 2015 @@ -139,6 +139,29 @@ void CodeGenFunction::EmitCXXGlobalVarDe 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( Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=233208&r1=233207&r2=233208&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original) +++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Wed Mar 25 15:06:28 2015 @@ -100,3 +100,20 @@ __device__ float *func5() { } // 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*)) +} _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
