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

2 files changed, 16 insertions, 0 deletions

diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index d310abfe4d..2fe4e63dbb 100644
--- a/lib/Sema/SemaDecl.cpp
+++ b/lib/Sema/SemaDecl.cpp
@@ -4320,6 +4320,14 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC,
   // 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 15e49205b6..9df7e3f4d2 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -24,5 +24,13 @@ __device__ void foo() {
   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++;
 }