maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

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

llvm-svn: 233208
This commit is contained in:
Jingyue Wu 2015-03-25 20:06:28 +00:00
parent f9dc7036d3
commit 4f7b9eb217
2 changed files with 40 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
clang/lib/CodeGen/CGDeclCXX.cpp
View File

@ -139,6 +139,29 @@ void CodeGenFunction::EmitCXXGlobalVarDeclInit(const VarDecl &D,
const Expr *Init = D.getInit(); const Expr *Init = D.getInit();
QualType T = D.getType(); 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 (!T->isReferenceType()) {
if (getLangOpts().OpenMP && D.hasAttr<OMPThreadPrivateDeclAttr>()) if (getLangOpts().OpenMP && D.hasAttr<OMPThreadPrivateDeclAttr>())
(void)CGM.getOpenMPRuntime().emitThreadPrivateVarDefinition( (void)CGM.getOpenMPRuntime().emitThreadPrivateVarDefinition(

17
clang/test/CodeGenCUDA/address-spaces.cu
View File

@ -100,3 +100,20 @@ __device__ float *func5() {
} }
// CHECK: define float* @_Z5func5v() // CHECK: define float* @_Z5func5v()
// CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*) // 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*))
}