diff options
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 3 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaDeclAttr.cpp | 4 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/bad-attributes.cu | 1 |
3 files changed, 8 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index ff7b3422a94..58029b03039 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/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/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 3e8a39a5786..aafb2cbb4cc 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/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/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu index 4fb584aec4c..5aaff975967 100644 --- a/clang/test/SemaCUDA/bad-attributes.cu +++ b/clang/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}} |
