[OpenMP] Change debugging symbol to weak_odr linkage

[jhuber6]

Summary

The new device runtime uses an internal variable to set debugging. This
variable was originally privately linked because every module will have
a copy of it. This caused problems with merging the device bitcode
library because it would get renamed and there was not a way to refer to
an external, private symbol. This changes the symbol to weak_odr so it
can be defined multiply, but will not be renamed.

[jdoerfert]

You checked that this works with the usage in the RT, right?

[jhuber6]

In D109997#3007216, @jdoerfert wrote:

You checked that this works with the usage in the RT, right?

Worked when I tested it. Want me to make a patch for it off of this?

[jdoerfert]

LG, and yes.

clang/test/OpenMP/target_debug_codegen.cpp

-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex "(__omp_rtl_debug_kind|llvm\.used)"
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex "__omp_rtl_debug_kind"
 // Test target codegen - host bc file has to be created first.
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug=111 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-EQ
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-DEFAULT
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 //.
-// CHECK: @__omp_rtl_debug_kind = private constant i32 1
-// CHECK: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32* @__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK: @__omp_rtl_debug_kind = weak_odr constant i32 1
 //.
-// CHECK-EQ: @__omp_rtl_debug_kind = private constant i32 111
-// CHECK-EQ: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32* @__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK-EQ: @__omp_rtl_debug_kind = weak_odr constant i32 111
 //.
-// CHECK-DEFAULT: @__omp_rtl_debug_kind = private constant i32 0
-// CHECK-DEFAULT: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32* @__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK-DEFAULT: @__omp_rtl_debug_kind = weak_odr constant i32 0
 //.
 void foo() {
 #pragma omp target
   { }
 }
 
 #endif

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

 OpenMPIRBuilder::~OpenMPIRBuilder() {
   assert(OutlineInfos.empty() && "There must be no outstanding outlinings");
 }
 
 GlobalValue *OpenMPIRBuilder::createDebugKind(unsigned DebugKind) {
   IntegerType *I32Ty = Type::getInt32Ty(M.getContext());
   auto *GV = new GlobalVariable(
       M, I32Ty,
-      /* isConstant = */ true, GlobalValue::PrivateLinkage,
+      /* isConstant = */ true, GlobalValue::WeakODRLinkage,
       ConstantInt::get(I32Ty, DebugKind), "__omp_rtl_debug_kind");
 
-  llvm::appendToUsed(M, {GV});
-
   return GV;
 }
 
 Value *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
                                          IdentFlag LocFlags,
                                          unsigned Reserve2Flags) {
   // Enable "C-mode".
   LocFlags |= OMP_IDENT_FLAG_KMPC;