From 1aba7783ad8fe326f10f5aa494721db23bdd593e Mon Sep 17 00:00:00 2001 From: Peter Collingbourne Date: Tue, 28 Aug 2012 20:37:10 +0000 Subject: [PATCH] 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 --- lib/CodeGen/CGDecl.cpp | 4 +++- test/CodeGenCUDA/address-spaces.cu | 4 ++++ 2 files changed, 7 insertions(+), 1 deletion(-) 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++; } -- 2.50.1