[llvm] dcb83b2 - [OpenMP] Mark device RTL variables as hidden

[Joseph Huber via llvm-commits]

Author: Joseph Huber
Date: 2022-01-18T12:53:17-05:00
New Revision: dcb83b236421d5a4d6e767ee43ae1cdab5fce0f2

URL: https://github.com/llvm/llvm-project/commit/dcb83b236421d5a4d6e767ee43ae1cdab5fce0f2
DIFF: https://github.com/llvm/llvm-project/commit/dcb83b236421d5a4d6e767ee43ae1cdab5fce0f2.diff

LOG: [OpenMP] Mark device RTL variables as hidden

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.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D117320

Added: 
    

Modified: 
    clang/test/OpenMP/target_globals_codegen.cpp
    llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
    llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

Removed: 
    


################################################################################
diff  --git a/clang/test/OpenMP/target_globals_codegen.cpp b/clang/test/OpenMP/target_globals_codegen.cpp
index 1264266340729..62e1a9ac8f9d7 100644
--- a/clang/test/OpenMP/target_globals_codegen.cpp
+++ b/clang/test/OpenMP/target_globals_codegen.cpp
@@ -12,25 +12,25 @@
 #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: @__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 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-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 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-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 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-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 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-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

diff  --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index fa92446492068..d80c521a5acfe 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -689,7 +689,8 @@ class OpenMPIRBuilder {
                              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.

diff  --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index f6e454b60d242..61194da4b589d 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -254,6 +254,7 @@ GlobalValue *OpenMPIRBuilder::createGlobalFlag(unsigned Value, StringRef Name) {
       new GlobalVariable(M, I32Ty,
                          /* isConstant = */ true, GlobalValue::WeakODRLinkage,
                          ConstantInt::get(I32Ty, Value), Name);
+  GV->setVisibility(GlobalValue::HiddenVisibility);
 
   return GV;
 }