CUDA: give correct address space to globals declared in functions

llvm-svn: 162787

clang/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());

clang/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++;
 }