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

diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp
index 35d1a623a833b5084ee87aae998306281290d71c..b9489e3f0409b02cb0ac29e655ce88e026420f54 100644 (file)
--- 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 61d4d6b6ba48d89cb5ce7bbd3d2bf07479396bce..15e49205b6aba15d299fd9b54fceba7f65cc769c 100644 (file)
--- 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++;
 }