[OpenMP] Add a semantic check for updating hidden or internal values

A previous patch removed the compiler generating offloading entries
for variables that were declared on the device but were internal or
hidden. This allowed us to compile programs but turns any attempt to run
'#pragma omp target update' on one of those variables a silent failure.
This patch adds a check in the semantic analysis for if the user is
attempting the update a variable on the device from the host that is not
externally visible.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D122403

Author: jhuber6

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10734,6 +10734,8 @@ def err_omp_expected_int_param : Error<
   "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 <

clang/lib/Sema/SemaOpenMP.cpp

@@ -12673,6 +12673,26 @@ static bool hasClauses(ArrayRef<OMPClause *> Clauses, const OpenMPClauseKind K,
   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,
@@ -12806,6 +12826,12 @@ StmtResult Sema::ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
     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);
 }

clang/test/OpenMP/target_update_messages.cpp

@@ -14,6 +14,20 @@ void xxx(int argc) {
   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() {
 }