diff options
| -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++; } |