[PATCH] D156368: [OpenMP] Do not always emit unused extern variables

[Joseph Huber via Phabricator via cfe-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG141c4e7a9403: [OpenMP] Do not always emit unused extern variables (authored by jhuber6).
Herald added a project: OpenMP.
Herald added a subscriber: openmp-commits.

Changed prior to commit:
  https://reviews.llvm.org/D156368?vs=545168&id=545211#toc

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D156368/new/

https://reviews.llvm.org/D156368

Files:
  clang/lib/CodeGen/CGOpenMPRuntime.cpp
  clang/lib/CodeGen/CodeGenModule.cpp
  clang/test/OpenMP/declare_target_codegen.cpp
  openmp/libomptarget/test/offloading/extern.c


Index: openmp/libomptarget/test/offloading/extern.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/offloading/extern.c
@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compile-generic -DVAR -c -o %t.o
+// RUN: %libomptarget-compile-generic %t.o && %libomptarget-run-generic | %fcheck-generic
+
+#ifdef VAR
+int x = 1;
+#else
+#include <stdio.h>
+#include <assert.h>
+extern int x;
+
+int main() {
+  int value = 0;
+#pragma omp target map(from : value)
+  value = x;
+  assert(value == 1);
+
+  x = 999;
+#pragma omp target update to(x)
+
+#pragma omp target map(from : value)
+  value = x;
+  assert(value == 999);
+
+  // CHECK: PASS
+  printf ("PASS\n");
+}
+#endif
Index: clang/test/OpenMP/declare_target_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_target_codegen.cpp
+++ clang/test/OpenMP/declare_target_codegen.cpp
@@ -29,7 +29,6 @@
 // CHECK-DAG: @flag = protected global i8 undef,
 // CHECK-DAG: @dx = {{protected | }}global i32 0,
 // CHECK-DAG: @dy = {{protected | }}global i32 0,
-// CHECK-DAG: @aaa = external global i32,
 // CHECK-DAG: @bbb = {{protected | }}global i32 0,
 // CHECK-DAG: weak constant %struct.__tgt_offload_entry { ptr @bbb,
 // CHECK-DAG: @ccc = external global i32,
@@ -80,7 +79,7 @@
 extern int aaa;
 int bbb = 0;
 extern int ccc;
-int ddd = 0;
+int ddd = ccc;
 #pragma omp end declare target
 
 #pragma omp declare target
@@ -260,8 +259,6 @@
 
 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
 
-// CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}}
-// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
 // CHECK-DAG: !{{{.+}}virtual_foo
 
 #ifdef OMP5
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3605,6 +3605,13 @@
         // Emit declaration of the must-be-emitted declare target variable.
         if (std::optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
                 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
+
+          // If this variable has external storage and doesn't require special
+          // link handling we defer to its canonical definition.
+          if (VD->hasExternalStorage() &&
+              Res != OMPDeclareTargetDeclAttr::MT_Link)
+            return;
+
           bool UnifiedMemoryEnabled =
               getOpenMPRuntime().hasRequiresUnifiedSharedMemory();
           if ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10166,6 +10166,13 @@
 
   std::optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
+
+  // If this is an 'extern' declaration we defer to the canonical definition and
+  // do not emit an offloading entry.
+  if (Res && *Res != OMPDeclareTargetDeclAttr::MT_Link &&
+      VD->hasExternalStorage())
+    return;
+
   if (!Res) {
     if (CGM.getLangOpts().OpenMPIsTargetDevice) {
       // Register non-target variables being emitted in device code (debug info


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D156368.545211.patch
Type: text/x-patch
Size: 3325 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230728/4dc05e54/attachment.bin>