diff options
| author | Peter Collingbourne <peter@pcc.me.uk> | 2012-08-28 20:37:10 +0000 |
|---|---|---|
| committer | Peter Collingbourne <peter@pcc.me.uk> | 2012-08-28 20:37:10 +0000 |
| commit | 1aba7783ad8fe326f10f5aa494721db23bdd593e (patch) | |
| tree | 17e4c422e22a99710150ed78761b56a6e3f14fcf | |
| parent | 4112a4c001c0ffe83962e5c6755c8b3e0d74b74c (diff) | |
| download | clang-1aba7783ad8fe326f10f5aa494721db23bdd593e.tar.gz clang-1aba7783ad8fe326f10f5aa494721db23bdd593e.tar.bz2 clang-1aba7783ad8fe326f10f5aa494721db23bdd593e.tar.xz | |
CUDA: give correct address space to globals declared in functions
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@162787 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | lib/CodeGen/CGDecl.cpp | 4 | ||||
| -rw-r--r-- | test/CodeGenCUDA/address-spaces.cu | 4 |
2 files changed, 7 insertions, 1 deletions
diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp index 35d1a623a8..b9489e3f04 100644 --- a/lib/CodeGen/CGDecl.cpp +++ b/lib/CodeGen/CGDecl.cpp @@ -184,12 +184,14 @@ CodeGenFunction::CreateStaticVarDecl(const VarDecl &D, Name = GetStaticDeclName(*this, D, Separator); llvm::Type *LTy = CGM.getTypes().ConvertTypeForMem(Ty); + unsigned AddrSpace = + CGM.GetGlobalVarAddressSpace(&D, CGM.getContext().getTargetAddressSpace(Ty)); llvm::GlobalVariable *GV = new llvm::GlobalVariable(CGM.getModule(), LTy, Ty.isConstant(getContext()), Linkage, CGM.EmitNullConstant(D.getType()), Name, 0, llvm::GlobalVariable::NotThreadLocal, - CGM.getContext().getTargetAddressSpace(Ty)); + AddrSpace); GV->setAlignment(getContext().getDeclAlign(&D).getQuantity()); if (Linkage != llvm::GlobalValue::InternalLinkage) GV->setVisibility(CurFn->getVisibility()); diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index 61d4d6b6ba..15e49205b6 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -20,5 +20,9 @@ __device__ void foo() { // CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*) k++; + + static int li; + // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li + li++; } |