From: Artem Belevich <tra@google.com>
Date: Mon, 9 May 2016 19:36:08 +0000 (+0000)
Subject: [CUDA] Only __shared__ variables can be static local on device side.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=a93fdcf11017c5e79d5e130d7f981763b218d8bf;p=clang

[CUDA] Only __shared__ variables can be static local on device side.

According to CUDA programming guide (v7.5):
> E.2.9.4: Within the body of a device or global function, only
> shared variables may be declared with static storage class.

Differential Revision: http://reviews.llvm.org/D20034

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@268962 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index ea0d52a05f..13eb56d31b 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6554,7 +6554,9 @@ def err_dynamic_var_init : Error<
     "__device__, __constant__, and __shared__ variables.">;
 def err_shared_var_init : Error<
     "initialization is not supported for __shared__ variables.">;
-
+def err_device_static_local_var : Error<
+    "Within a __device__/__global__ function, "
+    "only __shared__ variables may be marked \"static\"">;
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
   "%select{function|block|method|constructor}2; expected type from format "
diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index 5360cbbc06..a2c28de5f5 100644
--- a/lib/Sema/SemaDecl.cpp
+++ b/lib/Sema/SemaDecl.cpp
@@ -10391,15 +10391,24 @@ Sema::FinalizeDeclaration(Decl *ThisDecl) {
     }
   }
 
-  // Static locals inherit dll attributes from their function.
   if (VD->isStaticLocal()) {
     if (FunctionDecl *FD =
             dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
+      // Static locals inherit dll attributes from their function.
       if (Attr *A = getDLLAttr(FD)) {
         auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext()));
         NewAttr->setInherited(true);
         VD->addAttr(NewAttr);
       }
+      // CUDA E.2.9.4: Within the body of a __device__ or __global__
+      // function, only __shared__ variables may be declared with
+      // static storage class.
+      if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+          (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>()) &&
+          !VD->hasAttr<CUDASharedAttr>()) {
+        Diag(VD->getLocation(), diag::err_device_static_local_var);
+        VD->setInvalidDecl();
+      }
     }
   }
 
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 31cba958e1..4670f46db9 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -38,14 +38,6 @@ __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
   k++;
 
-  static int li;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
-  li++;
-
-  __constant__ int lj;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
-  lj++;
-
   __shared__ int lk;
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu
index 864cc6daee..23c9fe1376 100644
--- a/test/CodeGenCUDA/device-var-init.cu
+++ b/test/CodeGenCUDA/device-var-init.cu
@@ -368,6 +368,14 @@ __device__ void df() {
   T_F_NEC t_f_nec;
   T_FA_NEC t_fa_nec;
   static __shared__ UC s_uc;
+#if ERROR_CASE
+  static __device__ int ds;
+  // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  static __constant__ int dc;
+  // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  static int v;
+  // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+#endif
 }
 
 // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)