[cfe-commits] r162788 - in /cfe/trunk: lib/Sema/SemaDecl.cpp test/CodeGenCUDA/address-spaces.cu

[Peter Collingbourne]

Author: pcc
Date: Tue Aug 28 15:37:50 2012
New Revision: 162788

URL: http://llvm.org/viewvc/llvm-project?rev=162788&view=rev
Log:
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]."

Modified:
    cfe/trunk/lib/Sema/SemaDecl.cpp
    cfe/trunk/test/CodeGenCUDA/address-spaces.cu

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=162788&r1=162787&r2=162788&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Aug 28 15:37:50 2012
@@ -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))

Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=162788&r1=162787&r2=162788&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Tue Aug 28 15:37:50 2012
@@ -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++;
 }