CUDA: give static storage class to __shared__ and __constant__
variables without a storage class within a function, to implement
CUDA B.2.5: "__shared__ and __constant__ variables have implied static
storage [duration]."

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@162788 91177308-0d34-0410-b5e6-96231b3b80d8
diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index d310abf..2fe4e63 100644
--- a/lib/Sema/SemaDecl.cpp
+++ b/lib/Sema/SemaDecl.cpp
@@ -4320,6 +4320,14 @@
   // Handle attributes prior to checking for duplicates in MergeVarDecl
   ProcessDeclAttributes(S, NewVD, D);
 
+  if (getLangOpts().CUDA) {
+    // CUDA B.2.5: "__shared__ and __constant__ variables have implied static
+    // storage [duration]."
+    if (SC == SC_None && S->getFnParent() != 0 &&
+       (NewVD->hasAttr<CUDASharedAttr>() || NewVD->hasAttr<CUDAConstantAttr>()))
+      NewVD->setStorageClass(SC_Static);
+  }
+
   // In auto-retain/release, infer strong retension for variables of
   // retainable type.
   if (getLangOpts().ObjCAutoRefCount && inferObjCARCLifetime(NewVD))
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 15e4920..9df7e3f 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -24,5 +24,13 @@
   static int li;
   // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li
   li++;
+
+  __constant__ int lj;
+  // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj
+  lj++;
+
+  __shared__ int lk;
+  // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk
+  lk++;
 }