r284144 - [CUDA] Disallow __shared__ variables in host functions.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 13 11:45:13 PDT 2016
Author: jlebar
Date: Thu Oct 13 13:45:13 2016
New Revision: 284144
URL: http://llvm.org/viewvc/llvm-project?rev=284144&view=rev
Log:
[CUDA] Disallow __shared__ variables in host functions.
Reviewers: tra, rnk
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D25143
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaDeclAttr.cpp
cfe/trunk/test/SemaCUDA/bad-attributes.cu
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=284144&r1=284143&r2=284144&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Oct 13 13:45:13 2016
@@ -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<
Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=284144&r1=284143&r2=284144&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Thu Oct 13 13:45:13 2016
@@ -3724,6 +3724,10 @@ static void handleSharedAttr(Sema &S, De
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()));
}
Modified: cfe/trunk/test/SemaCUDA/bad-attributes.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-attributes.cu?rev=284144&r1=284143&r2=284144&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu (original)
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu Thu Oct 13 13:45:13 2016
@@ -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}}
More information about the cfe-commits
mailing list