[OpenMP] Do not always emit unused extern variables

[jhuber6]

Summary

Currently, the precense of the OpenMP target declare metadata requires
that we always codegen a global declaration. This is undesirable in the
case that we could defer or omit this declaration as is common with
unused extern variables. This is important as it allows us, in the
runtime, to rely on static linking semantics to omit unused symbols so
they are not included when the user links it in.

This patch changes the check for always emitting these variables.
Because of this we also need to extend this logic to the generation of
the offloading entries. This has the result of derring the offload entry
generation to the canonical definitoin. So we are effectively assuming
whoever owns the storage for this variable will perform that operation.
This makes an exception for link attributes as those require their own
special handling.

Let me know if this is sound in the implementation, I do not have the
largest view of the standards here.

Fixes: https://github.com/llvm/llvm-project/issues/64133

[jhuber6]

Add OpenMP runtime test

[tianshilei1992]

I think this is reasonable.

clang/test/OpenMP/declare_target_codegen.cpp

264

Since ccc is used here, it is not supposed to be removed right?

[jhuber6]

clang/test/OpenMP/declare_target_codegen.cpp

264

This is the declare target metadta, which this patch no longer emits for external storage variables. Since we assume that whoever defined ccc will handle it.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

 void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
                                                    llvm::Constant *Addr) {
   if (CGM.getLangOpts().OMPTargetTriples.empty() &&
       !CGM.getLangOpts().OpenMPIsTargetDevice)
     return;
 
   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
       // may cause this).
       StringRef VarName = CGM.getMangledName(VD);
       EmittedNonTargetVariables.try_emplace(VarName, Addr);
     }
     return;

clang/lib/CodeGen/CodeGenModule.cpp

if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
     const auto *VD = cast<VarDecl>(Global);
     assert(VD->isFileVarDecl() && "Cannot emit local var decl as global.");
     if (VD->isThisDeclarationADefinition() != VarDecl::Definition &&
         !Context.isMSStaticDataMemberInlineDefinition(VD)) {
       if (LangOpts.OpenMP) {
         // 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 ||
                *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
               !UnifiedMemoryEnabled) {
             (void)GetAddrOfGlobalVar(VD);
           } else {
             assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||

clang/test/OpenMP/declare_target_codegen.cpp

 // KMPC-ONLY-NOT: __tgt
 
 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
 // CHECK-DAG: Bake
 // CHECK-NOT: @{{hhh|ggg|fff|eee}} =
 // 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,
 // CHECK-DAG: @ddd = {{protected | }}global i32 0,
 // CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global ptr null
 // CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global ptr null
 // CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global ptr null
 // CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global ptr null

 #pragma omp declare target
 extern int bbb;
 #pragma omp end declare target
 
 #pragma omp declare target
 extern int aaa;
 int bbb = 0;
 extern int ccc;
-int ddd = 0;
+int ddd = ccc;
 #pragma omp end declare target
 
 #pragma omp declare target
 extern int bbb;
 #pragma omp end declare target
 
 extern int eee;
 int fff = 0;

 };
 
 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-5]]()
 
 // CHECK-DAG: declare extern_weak noundef signext i32 @__create()
 
 // 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
 void host_fun() {}
 #pragma omp declare target to(host_fun) device_type(host)
 void device_fun() {}
 #pragma omp declare target to(device_fun) device_type(nohost)
 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}

openmp/libomptarget/test/offloading/extern.c

+// 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