dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[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
This commit is contained in:
Justin Lebar 2016-10-13 18:45:13 +00:00
parent fff8e6c7b6
commit a1f58c8582
3 changed files with 8 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -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<

4
lib/Sema/SemaDeclAttr.cpp
View File

@ -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()));
}

1
test/SemaCUDA/bad-attributes.cu
View File

@ -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}}