[OpenMP] Mark device RTL variables as hidden

[jhuber6]

Summary

This patch changes the visibility of the __omp_rtl_debug_kind variable
to be hidden. These variables are only used by the plugin so they do not
need to be read externally. Previously the default visibility prevented
these variables from being completely eliminated in the module.

[tianshilei1992]

LG. However, we might want to document OpenMPIRBuilder::createGlobalFlag saying that function is only for creating flags for what kind of variables in case it is misused.

[jhuber6]

Fix test and add the fact that its hidden to the comment.

clang/test/OpenMP/target_globals_codegen.cpp

 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex "__omp_rtl_"
 // 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
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-threads-oversubscription -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-THREADS
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-teams-oversubscription -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-TEAMS
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 //.
-// CHECK: @__omp_rtl_debug_kind = weak_odr constant i32 1
-// CHECK: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0
-// CHECK: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0
-//.
-// CHECK-EQ: @__omp_rtl_debug_kind = weak_odr constant i32 111
-// CHECK-EQ: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0
-// CHECK-EQ: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0
-//.
-// CHECK-DEFAULT: @__omp_rtl_debug_kind = weak_odr constant i32 0
-// CHECK-DEFAULT: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0
-// CHECK-DEFAULT: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0
-//.
-// CHECK-THREADS: @__omp_rtl_debug_kind = weak_odr constant i32 0
-// CHECK-THREADS: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0
-// CHECK-THREADS: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 1
-//.
-// CHECK-TEAMS: @__omp_rtl_debug_kind = weak_odr constant i32 0
-// CHECK-TEAMS: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 1
-// CHECK-TEAMS: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0
+// CHECK: @__omp_rtl_debug_kind = weak_odr hidden constant i32 1
+// CHECK: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 0
+// CHECK: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 0
+//.
+// CHECK-EQ: @__omp_rtl_debug_kind = weak_odr hidden constant i32 111
+// CHECK-EQ: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 0
+// CHECK-EQ: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 0
+//.
+// CHECK-DEFAULT: @__omp_rtl_debug_kind = weak_odr hidden constant i32 0
+// CHECK-DEFAULT: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 0
+// CHECK-DEFAULT: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 0
+//.
+// CHECK-THREADS: @__omp_rtl_debug_kind = weak_odr hidden constant i32 0
+// CHECK-THREADS: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 0
+// CHECK-THREADS: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 1
+//.
+// CHECK-TEAMS: @__omp_rtl_debug_kind = weak_odr hidden constant i32 0
+// CHECK-TEAMS: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 1
+// CHECK-TEAMS: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 0
 //.
 void foo() {
 #pragma omp target
   { }
 }
 
 #endif

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

Constant *getOrCreateSrcLocStr(const LocationDescription &Loc,
                                  uint32_t &SrcLocStrSize);
 
   /// Return an ident_t* encoding the source location \p SrcLocStr and \p Flags.
   /// TODO: Create a enum class for the Reserve2Flags
   Constant *getOrCreateIdent(Constant *SrcLocStr, uint32_t SrcLocStrSize,
                              omp::IdentFlag Flags = omp::IdentFlag(0),
                              unsigned Reserve2Flags = 0);
 
-  /// Create a global flag \p Namein the module with initial value \p Value.
+  /// Create a hidden global flag \p Name in the module with initial value \p
+  /// Value.
   GlobalValue *createGlobalFlag(unsigned Value, StringRef Name);
 
   /// Generate control flow and cleanup for cancellation.
   ///
   /// \param CancelFlag Flag indicating if the cancellation is performed.
   /// \param CanceledDirective The kind of directive that is cancled.
   /// \param ExitCB Extra code to be generated in the exit block.
   void emitCancelationCheckImpl(Value *CancelFlag,

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

 }
 
 GlobalValue *OpenMPIRBuilder::createGlobalFlag(unsigned Value, StringRef Name) {
   IntegerType *I32Ty = Type::getInt32Ty(M.getContext());
   auto *GV =
       new GlobalVariable(M, I32Ty,
                          /* isConstant = */ true, GlobalValue::WeakODRLinkage,
                          ConstantInt::get(I32Ty, Value), Name);
+  GV->setVisibility(GlobalValue::HiddenVisibility);
 
   return GV;
 }
 
 Constant *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
                                             uint32_t SrcLocStrSize,
                                             IdentFlag LocFlags,
                                             unsigned Reserve2Flags) {