diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10734,6 +10734,8 @@ "expected a reference to an integer-typed parameter">; def err_omp_at_least_one_motion_clause_required : Error< "expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'">; +def err_omp_cannot_update_with_internal_linkage : Error< + "the host cannot update a declare target variable that is not externally visible.">; def err_omp_usedeviceptr_not_a_pointer : Error< "expected pointer or reference to pointer in 'use_device_ptr' clause">; def err_omp_argument_type_isdeviceptr : Error < diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -12673,6 +12673,26 @@ return hasClauses(Clauses, K) || hasClauses(Clauses, ClauseTypes...); } +/// Check if the variables in the mapping clause are externally visible. +static bool isClauseMappable(ArrayRef Clauses) { + for (const OMPClause *C : Clauses) { + if (auto *TC = dyn_cast(C)) + return llvm::all_of(TC->all_decls(), [](ValueDecl *VD) { + return !VD || !VD->hasAttr() || + (VD->isExternallyVisible() && + VD->getVisibility() != HiddenVisibility); + }); + else if (auto *FC = dyn_cast(C)) + return llvm::all_of(FC->all_decls(), [](ValueDecl *VD) { + return !VD || !VD->hasAttr() || + (VD->isExternallyVisible() && + VD->getVisibility() != HiddenVisibility); + }); + } + + return true; +} + StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, @@ -12806,6 +12826,12 @@ Diag(StartLoc, diag::err_omp_at_least_one_motion_clause_required); return StmtError(); } + + if (!isClauseMappable(Clauses)) { + Diag(StartLoc, diag::err_omp_cannot_update_with_internal_linkage); + return StmtError(); + } + return OMPTargetUpdateDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt); } diff --git a/clang/test/OpenMP/target_update_messages.cpp b/clang/test/OpenMP/target_update_messages.cpp --- a/clang/test/OpenMP/target_update_messages.cpp +++ b/clang/test/OpenMP/target_update_messages.cpp @@ -14,6 +14,20 @@ argc = x; // expected-warning {{variable 'x' is uninitialized when used here}} } +static int y; +#pragma omp declare target(y) + +void yyy() { +#pragma omp target update to(y) // expected-error {{the host cannot update a declare target variable that is not externally visible.}} +} + +int __attribute__((visibility("hidden"))) z; +#pragma omp declare target(z) + +void zzz() { +#pragma omp target update from(z) // expected-error {{the host cannot update a declare target variable that is not externally visible.}} +} + void foo() { }