r284145 - [CUDA] Allow static variables in __host__ __device__ functions, so long as they're never codegen'ed for device.

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Thu Oct 13 13:45:17 2016
New Revision: 284145

URL: http://llvm.org/viewvc/llvm-project?rev=284145&view=rev
Log:
[CUDA] Allow static variables in __host__ __device__ functions, so long as they're never codegen'ed for device.

Reviewers: tra, rnk

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D25150

Added:
    cfe/trunk/test/SemaCUDA/static-vars-hd.cu
Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/lib/Sema/SemaDecl.cpp
    cfe/trunk/test/SemaCUDA/device-var-init.cu

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=284145&r1=284144&r2=284145&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Oct 13 13:45:17 2016
@@ -6741,8 +6741,8 @@ def err_dynamic_var_init : Error<
 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\"">;
+    "within a %select{__device__|__global__|__host__|__host__ __device__}0 "
+    "function, only __shared__ variables may be marked 'static'">;
 def err_cuda_vla : Error<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=284145&r1=284144&r2=284145&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu Oct 13 13:45:17 2016
@@ -10677,12 +10677,11 @@ Sema::FinalizeDeclaration(Decl *ThisDecl
       // 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);
+      if (getLangOpts().CUDA && !VD->hasAttr<CUDASharedAttr>() &&
+          CUDADiagIfDeviceCode(VD->getLocation(),
+                               diag::err_device_static_local_var)
+              << CurrentCUDATarget())
         VD->setInvalidDecl();
-      }
     }
   }
 
@@ -10696,7 +10695,7 @@ Sema::FinalizeDeclaration(Decl *ThisDecl
     if (Init && VD->hasGlobalStorage()) {
       if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
           VD->hasAttr<CUDASharedAttr>()) {
-        assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
+        assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
         bool AllowedInit = false;
         if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
           AllowedInit =

Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=284145&r1=284144&r2=284145&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu Thu Oct 13 13:45:17 2016
@@ -207,9 +207,9 @@ __device__ void df_sema() {
   // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
 
   static __device__ int ds;
-  // expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
   static __constant__ int dc;
-  // expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
   static int v;
-  // expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+  // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
 }

Added: cfe/trunk/test/SemaCUDA/static-vars-hd.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/static-vars-hd.cu?rev=284145&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/static-vars-hd.cu (added)
+++ cfe/trunk/test/SemaCUDA/static-vars-hd.cu Thu Oct 13 13:45:17 2016
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -S -o /dev/null -verify %s
+// RUN: %clang_cc1 -fcxx-exceptions -S -o /dev/null -D HOST -verify %s
+
+#include "Inputs/cuda.h"
+
+#ifdef HOST
+// expected-no-diagnostics
+#endif
+
+__host__ __device__ void f() {
+  static int x = 42;
+#ifndef HOST
+  // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+#endif
+}
+
+inline __host__ __device__ void g() {
+  static int x = 42; // no error on device because this is never codegen'ed there.
+}
+void call_g() { g(); }