Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -6723,6 +6723,9 @@ "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< Index: clang/lib/Sema/SemaDeclAttr.cpp =================================================================== --- clang/lib/Sema/SemaDeclAttr.cpp +++ clang/lib/Sema/SemaDeclAttr.cpp @@ -3718,6 +3718,10 @@ 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())); } Index: clang/test/SemaCUDA/bad-attributes.cu =================================================================== --- clang/test/SemaCUDA/bad-attributes.cu +++ clang/test/SemaCUDA/bad-attributes.cu @@ -65,6 +65,7 @@ __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}}