From: Justin Lebar <jlebar@google.com>
Date: Thu, 13 Oct 2016 18:45:13 +0000 (+0000)
Subject: [CUDA] Disallow __shared__ variables in host functions.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=a1f58c85829f2c491e9cf32445f8c8287f11d80f;p=clang

[CUDA] Disallow __shared__ variables in host functions.

Reviewers: tra, rnk

Subscribers: cfe-commits

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

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

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index ff7b3422a9..58029b0303 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6747,6 +6747,9 @@ def err_cuda_vla : Error<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
 def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
+def err_cuda_host_shared : Error<
+    "__shared__ local variables not allowed in "
+    "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
 def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
 
 def warn_non_pod_vararg_with_format_string : Warning<
diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 3e8a39a578..aafb2cbb4c 100644
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -3724,6 +3724,10 @@ static void handleSharedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
     S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
     return;
   }
+  if (S.getLangOpts().CUDA && VD->hasLocalStorage() &&
+      S.CUDADiagIfHostCode(Attr.getLoc(), diag::err_cuda_host_shared)
+          << S.CurrentCUDATarget())
+    return;
   D->addAttr(::new (S.Context) CUDASharedAttr(
       Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex()));
 }
diff --git a/test/SemaCUDA/bad-attributes.cu b/test/SemaCUDA/bad-attributes.cu
index 4fb584aec4..5aaff97596 100644
--- a/test/SemaCUDA/bad-attributes.cu
+++ b/test/SemaCUDA/bad-attributes.cu
@@ -65,6 +65,7 @@ __global__ static inline void foobar() {};
 __constant__ int global_constant;
 void host_fn() {
   __constant__ int c; // expected-error {{__constant__ variables must be global}}
+  __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
 }
 __device__ void device_fn() {
   __constant__ int c; // expected-error {{__constant__ variables must be global}}