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

[Joseph Huber via cfe-commits]

Author: Joseph Huber
Date: 2022-03-24T19:38:30-04:00
New Revision: bfda79341bb5e5d8eb9a8fa63958ecfe75f57d78

URL: https://github.com/llvm/llvm-project/commit/bfda79341bb5e5d8eb9a8fa63958ecfe75f57d78
DIFF: https://github.com/llvm/llvm-project/commit/bfda79341bb5e5d8eb9a8fa63958ecfe75f57d78.diff

LOG: [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

Added: 
    

Modified: 
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/lib/Sema/SemaOpenMP.cpp
    clang/test/OpenMP/target_update_messages.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 14b276cf82295..4b35aaec82295 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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 <

diff  --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 3424393f7e5e6..a7b58707c8ecc 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/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);
 }

diff  --git a/clang/test/OpenMP/target_update_messages.cpp b/clang/test/OpenMP/target_update_messages.cpp
index fd9d5455c9f1e..f936a075e1b48 100644
--- a/clang/test/OpenMP/target_update_messages.cpp
+++ b/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() {
 }