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; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits