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 @@ -12619,6 +12619,26 @@ return hasClauses(Clauses, K) || hasClauses(Clauses, ClauseTypes...); } +/// Check if the variables in the mapping clause are externally visible. +static bool isClauseMappable(ArrayRef<OMPClause *> Clauses) { + for (const OMPClause *C : Clauses) { + if (auto *TC = dyn_cast<OMPToClause>(C)) + return llvm::all_of(TC->all_decls(), [](ValueDecl *VD) { + return !VD || !VD->hasAttr<OMPDeclareTargetDeclAttr>() || + (VD->isExternallyVisible() && + VD->getVisibility() != HiddenVisibility); + }); + else if (auto *FC = dyn_cast<OMPFromClause>(C)) + return llvm::all_of(FC->all_decls(), [](ValueDecl *VD) { + return !VD || !VD->hasAttr<OMPDeclareTargetDeclAttr>() || + (VD->isExternallyVisible() && + VD->getVisibility() != HiddenVisibility); + }); + } + + return true; +} + StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc, @@ -12752,6 +12772,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 @@ -8,12 +8,27 @@ // RUN: %clang_cc1 -verify=expected,ge50,ge51,cxx2b -fopenmp -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++2b %s -Wuninitialized + void xxx(int argc) { int x; // expected-note {{initialize the variable 'x' to silence this warning}} #pragma omp target update to(x) 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() { }