[PATCH] D122403: [OpenMP] Add a sematnic check for updating hidden or internal values

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 created this revision.
jhuber6 added reviewers: jdoerfert, tianshilei1992, JonChesterfield, ABataev.
Herald added subscribers: guansong, yaxunl.
Herald added a project: All.
jhuber6 requested review of this revision.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.

A previous patch removed the compiler from generting 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 to check if the user is
attempting the update a variable on the device from the host that is not
externally visible.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D122403

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


Index: clang/test/OpenMP/target_update_messages.cpp
===================================================================
--- clang/test/OpenMP/target_update_messages.cpp
+++ 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() {
 }
 
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ 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);
 }
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ 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 <


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D122403.417932.patch
Type: text/x-patch
Size: 3539 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220324/3e3f7122/attachment.bin>