[compiler-rt] [libc] [lld] [flang] [mlir] [clang-tools-extra] [lldb] [llvm] [libcxx] [clang] Don't emit relax relocs like R_X86_64_REX_GOTPCRELX on X86 target for OPENMP internal vars. (PR #75564)

via llvm-commits llvm-commits at lists.llvm.org
Wed Jan 10 18:59:58 PST 2024


https://github.com/UmeshKalappa0 updated https://github.com/llvm/llvm-project/pull/75564

>From 4125e4a709c594562fa6c52f045ba7442e3cb523 Mon Sep 17 00:00:00 2001
From: Umesh Kalappa <Umesh.Kalappa at amd.com>
Date: Fri, 15 Dec 2023 11:52:52 +0530
Subject: [PATCH 1/4] Problem :For Kernel Modules ,emitting the relocs like
 R_X86_64_REX_GOTPCRELX  for the OPENMP internal vars like
 https://godbolt.org/z/hhh7ozojz.

Solution : Mark the OpenMP internal variables with dso_local
conditionally for no-pic and no-pie ,then
a)reset the dso_local for thread_local and weak linkage vars.
---
 .../test/OpenMP/gomp_critical_dso_local_var.c | 23 +++++++++++++++++++
 1 file changed, 23 insertions(+)
 create mode 100644 clang/test/OpenMP/gomp_critical_dso_local_var.c

diff --git a/clang/test/OpenMP/gomp_critical_dso_local_var.c b/clang/test/OpenMP/gomp_critical_dso_local_var.c
new file mode 100644
index 00000000000000..915f6773bf67bf
--- /dev/null
+++ b/clang/test/OpenMP/gomp_critical_dso_local_var.c
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -fopenmp -x c -emit-llvm %s -o - | FileCheck %s --check-prefix=DSO_LOCAL
+
+// DSO_LOCAL-DAG: @.gomp_critical_user_.var = common dso_local global [8 x i32] zeroinitializer, align 8
+int omp_critical_test()
+{
+  int sum;
+  int known_sum;
+
+  sum=0;
+#pragma omp parallel
+  {
+    int mysum=0;
+    int i;
+#pragma omp for
+    for (i = 0; i < 1000; i++)
+      mysum = mysum + i;
+#pragma omp critical
+    sum = mysum +sum;
+  }
+  known_sum = 999 * 1000 / 2;
+  return (known_sum == sum);
+}
+

>From 842245de490ab15f8a901b94576ae4539c760e1e Mon Sep 17 00:00:00 2001
From: Umesh Kalappa <Umesh.Kalappa at amd.com>
Date: Fri, 15 Dec 2023 12:49:48 +0530
Subject: [PATCH 2/4] testcases are changed accordignly.

---
 clang/lib/CodeGen/CGOpenMPRuntime.cpp           | 2 ++
 clang/test/OpenMP/critical_codegen.cpp          | 6 +++---
 clang/test/OpenMP/critical_codegen_attr.cpp     | 6 +++---
 clang/test/OpenMP/for_reduction_codegen.cpp     | 8 ++++----
 clang/test/OpenMP/gomp_critical_dso_local_var.c | 1 -
 clang/test/OpenMP/simd_codegen.cpp              | 4 ++--
 llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp       | 8 ++++++++
 7 files changed, 22 insertions(+), 13 deletions(-)

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 7f7e6f53066644..183c757d72b8a7 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1793,6 +1793,8 @@ Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
   if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPUseTLS &&
       CGM.getTarget().isTLSSupported()) {
     GAddr->setThreadLocal(/*Val=*/true);
+    /// reset the dso_local for thread_local.
+    GAddr->setDSOLocal(/*Val=*/false);
     return Address(GAddr, GAddr->getValueType(),
                    CGM.getContext().getTypeAlignInChars(VarType));
   }
diff --git a/clang/test/OpenMP/critical_codegen.cpp b/clang/test/OpenMP/critical_codegen.cpp
index 24145d44d962e5..9a613161ac294a 100644
--- a/clang/test/OpenMP/critical_codegen.cpp
+++ b/clang/test/OpenMP/critical_codegen.cpp
@@ -16,9 +16,9 @@
 #define HEADER
 
 // ALL:       [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, ptr }
-// ALL:       [[UNNAMED_LOCK:@.+]] = common global [8 x i32] zeroinitializer
-// ALL:       [[THE_NAME_LOCK:@.+]] = common global [8 x i32] zeroinitializer
-// ALL:       [[THE_NAME_LOCK1:@.+]] = common global [8 x i32] zeroinitializer
+// ALL:       [[UNNAMED_LOCK:@.+]] = common dso_local global [8 x i32] zeroinitializer
+// ALL:       [[THE_NAME_LOCK:@.+]] = common dso_local global [8 x i32] zeroinitializer
+// ALL:       [[THE_NAME_LOCK1:@.+]] = common dso_local global [8 x i32] zeroinitializer
 
 // ALL:       define {{.*}}void [[FOO:@.+]]()
 
diff --git a/clang/test/OpenMP/critical_codegen_attr.cpp b/clang/test/OpenMP/critical_codegen_attr.cpp
index 34d90a9e3a6e48..5f1a76e2ad0f1f 100644
--- a/clang/test/OpenMP/critical_codegen_attr.cpp
+++ b/clang/test/OpenMP/critical_codegen_attr.cpp
@@ -16,9 +16,9 @@
 #define HEADER
 
 // ALL:       [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, ptr }
-// ALL:       [[UNNAMED_LOCK:@.+]] = common global [8 x i32] zeroinitializer
-// ALL:       [[THE_NAME_LOCK:@.+]] = common global [8 x i32] zeroinitializer
-// ALL:       [[THE_NAME_LOCK1:@.+]] = common global [8 x i32] zeroinitializer
+// ALL:       [[UNNAMED_LOCK:@.+]] = common dso_local global [8 x i32] zeroinitializer
+// ALL:       [[THE_NAME_LOCK:@.+]] = common dso_local global [8 x i32] zeroinitializer
+// ALL:       [[THE_NAME_LOCK1:@.+]] = common dso_local global [8 x i32] zeroinitializer
 
 // ALL:       define {{.*}}void [[FOO:@.+]]()
 
diff --git a/clang/test/OpenMP/for_reduction_codegen.cpp b/clang/test/OpenMP/for_reduction_codegen.cpp
index 893c606f8d7b9f..b128bd5d79c251 100644
--- a/clang/test/OpenMP/for_reduction_codegen.cpp
+++ b/clang/test/OpenMP/for_reduction_codegen.cpp
@@ -528,12 +528,12 @@ int main() {
 
 #endif
 //.
-// CHECK1: @.gomp_critical_user_.reduction.var = common global [8 x i32] zeroinitializer, align 8
-// CHECK1: @.gomp_critical_user_.atomic_reduction.var = common global [8 x i32] zeroinitializer, align 8
+// CHECK1: @.gomp_critical_user_.reduction.var = common dso_local global [8 x i32] zeroinitializer, align 8
+// CHECK1: @.gomp_critical_user_.atomic_reduction.var = common dso_local global [8 x i32] zeroinitializer, align 8
 //.
-// CHECK3: @.gomp_critical_user_.reduction.var = common global [8 x i32] zeroinitializer, align 8
+// CHECK3: @.gomp_critical_user_.reduction.var = common dso_local global [8 x i32] zeroinitializer, align 8
 //.
-// CHECK4: @.gomp_critical_user_.reduction.var = common global [8 x i32] zeroinitializer, align 8
+// CHECK4: @.gomp_critical_user_.reduction.var = common dso_local global [8 x i32] zeroinitializer, align 8
 //.
 // CHECK1-LABEL: define {{[^@]+}}@main
 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] {
diff --git a/clang/test/OpenMP/gomp_critical_dso_local_var.c b/clang/test/OpenMP/gomp_critical_dso_local_var.c
index 915f6773bf67bf..331c8cbad27eb7 100644
--- a/clang/test/OpenMP/gomp_critical_dso_local_var.c
+++ b/clang/test/OpenMP/gomp_critical_dso_local_var.c
@@ -20,4 +20,3 @@ int omp_critical_test()
   known_sum = 999 * 1000 / 2;
   return (known_sum == sum);
 }
-
diff --git a/clang/test/OpenMP/simd_codegen.cpp b/clang/test/OpenMP/simd_codegen.cpp
index b96e4213e8e0e1..e85aea8b77a0e1 100644
--- a/clang/test/OpenMP/simd_codegen.cpp
+++ b/clang/test/OpenMP/simd_codegen.cpp
@@ -23,8 +23,8 @@
 #define CONDITIONAL
 #endif //OMP5
 // CHECK: [[SS_TY:%.+]] = type { i32 }
-// OMP5-DAG: [[LAST_IV:@.+]] = {{.*}}common global i64 0
-// OMP5-DAG: [[LAST_A:@.+]] = {{.*}}common global i32 0
+// OMP5-DAG: [[LAST_IV:@.+]] = {{.*}}common dso_local global i64 0
+// OMP5-DAG: [[LAST_A:@.+]] = {{.*}}common dso_local global i32 0
 
 long long get_val() { extern void mayThrow(); mayThrow(); return 0; }
 double *g_ptr;
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index ce428f78dc843e..e1aa6efc82eaf3 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -5224,6 +5224,12 @@ OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name,
     const llvm::Align TypeAlign = DL.getABITypeAlign(Ty);
     const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpace);
     GV->setAlignment(std::max(TypeAlign, PtrAlign));
+
+    if (!GV->isDSOLocal() && !GV->isThreadLocal()) {
+      bool IsPIE = GV->getParent()->getPIELevel() != llvm::PIELevel::Default;
+      bool IsPIC = GV->getParent()->getPICLevel() != llvm::PICLevel::NotPIC;
+      GV->setDSOLocal(!IsPIC || IsPIE);
+    }
     Elem.second = GV;
   }
 
@@ -6684,6 +6690,8 @@ Constant *OpenMPIRBuilder::getAddrOfDeclareTargetVar(
 
       auto *GV = cast<GlobalVariable>(Ptr);
       GV->setLinkage(GlobalValue::WeakAnyLinkage);
+      /// reset dso_local for weak linkage.
+      GV->setDSOLocal(false);
 
       if (!Config.isTargetDevice()) {
         if (GlobalInitializer)

>From 1bc9b01b28343b46faf72ab8e047cb4349dc1bd7 Mon Sep 17 00:00:00 2001
From: Umesh Kalappa <Umesh.Kalappa at amd.com>
Date: Tue, 19 Dec 2023 10:17:53 +0530
Subject: [PATCH 3/4] Fix : Updated the testcases and added the dso_local
 attribute for OPENMP internal global vars.

---
 clang/lib/CodeGen/CGOpenMPRuntime.cpp              |  2 --
 clang/test/OpenMP/declare_target_codegen.cpp       |  8 ++++----
 clang/test/OpenMP/declare_target_link_codegen.cpp  |  6 +++---
 clang/test/OpenMP/taskloop_reduction_codegen.cpp   |  4 ++--
 llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp          |  4 +---
 .../omptarget-declare-target-llvm-device.mlir      |  2 +-
 .../LLVMIR/omptarget-declare-target-llvm-host.mlir | 14 +++++++-------
 mlir/test/Target/LLVMIR/openmp-llvm.mlir           |  2 +-
 8 files changed, 19 insertions(+), 23 deletions(-)

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 183c757d72b8a7..7f7e6f53066644 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1793,8 +1793,6 @@ Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
   if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPUseTLS &&
       CGM.getTarget().isTLSSupported()) {
     GAddr->setThreadLocal(/*Val=*/true);
-    /// reset the dso_local for thread_local.
-    GAddr->setDSOLocal(/*Val=*/false);
     return Address(GAddr, GAddr->getValueType(),
                    CGM.getContext().getTypeAlignInChars(VarType));
   }
diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp
index a5a9b790b4689f..243bf565ee5ca7 100644
--- a/clang/test/OpenMP/declare_target_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_codegen.cpp
@@ -33,10 +33,10 @@
 // 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
+// CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak dso_local global ptr null
+// CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak dso_local global ptr null
+// CHECK-DAG: @fff_decl_tgt_ref_ptr = weak dso_local global ptr null
+// CHECK-DAG: @eee_decl_tgt_ref_ptr = weak dso_local global ptr null
 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
 // CHECK-DAG: @pair = {{.*}}addrspace(3) global %struct.PAIR undef
 // CHECK-DAG: @_ZN2SS3SSSE ={{ protected | }}global i32 1,
diff --git a/clang/test/OpenMP/declare_target_link_codegen.cpp b/clang/test/OpenMP/declare_target_link_codegen.cpp
index 2372b2738b5bea..d93358b4516355 100644
--- a/clang/test/OpenMP/declare_target_link_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_link_codegen.cpp
@@ -19,11 +19,11 @@
 #define HEADER
 
 // HOST-DAG: @c = external global i32,
-// HOST-DAG: @c_decl_tgt_ref_ptr = weak global ptr @c
+// HOST-DAG: @c_decl_tgt_ref_ptr = weak dso_local global ptr @c
 // HOST-DAG: @[[D:.+]] = internal global i32 2
-// HOST-DAG: @[[D_PTR:.+]] = weak global ptr @[[D]]
+// HOST-DAG: @[[D_PTR:.+]] = weak dso_local global ptr @[[D]]
 // DEVICE-NOT: @c =
-// DEVICE: @c_decl_tgt_ref_ptr = weak global ptr null
+// DEVICE: @c_decl_tgt_ref_ptr = weak dso_local global ptr null
 // HOST: [[SIZES:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 4, i64 4]
 // HOST: [[MAPTYPES:@.+]] = private unnamed_addr constant [3 x i64] [i64 35, i64 531, i64 531]
 // HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_decl_tgt_ref_ptr\00"
diff --git a/clang/test/OpenMP/taskloop_reduction_codegen.cpp b/clang/test/OpenMP/taskloop_reduction_codegen.cpp
index 762f2801312618..583f52be144cab 100644
--- a/clang/test/OpenMP/taskloop_reduction_codegen.cpp
+++ b/clang/test/OpenMP/taskloop_reduction_codegen.cpp
@@ -4,8 +4,8 @@
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 // expected-no-diagnostics
 
-// CHECK: [[RED_SIZE1:@reduction_size[.].+]] = common thread_local global i64 0
-// CHECK: [[RED_SIZE2:@reduction_size[.].+]] = common thread_local global i64 0
+// CHECK: [[RED_SIZE1:@reduction_size[.].+]] = common dso_local thread_local global i64 0
+// CHECK: [[RED_SIZE2:@reduction_size[.].+]] = common dso_local thread_local global i64 0
 
 struct S {
   float a;
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index e1aa6efc82eaf3..910d7c6e76e561 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -5225,7 +5225,7 @@ OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name,
     const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpace);
     GV->setAlignment(std::max(TypeAlign, PtrAlign));
 
-    if (!GV->isDSOLocal() && !GV->isThreadLocal()) {
+    if (!GV->isDSOLocal()) {
       bool IsPIE = GV->getParent()->getPIELevel() != llvm::PIELevel::Default;
       bool IsPIC = GV->getParent()->getPICLevel() != llvm::PICLevel::NotPIC;
       GV->setDSOLocal(!IsPIC || IsPIE);
@@ -6690,8 +6690,6 @@ Constant *OpenMPIRBuilder::getAddrOfDeclareTargetVar(
 
       auto *GV = cast<GlobalVariable>(Ptr);
       GV->setLinkage(GlobalValue::WeakAnyLinkage);
-      /// reset dso_local for weak linkage.
-      GV->setDSOLocal(false);
 
       if (!Config.isTargetDevice()) {
         if (GlobalInitializer)
diff --git a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device.mlir b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device.mlir
index cf08761981fb3a..1c932b3d342fcf 100644
--- a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-device.mlir
@@ -8,7 +8,7 @@
 // file created by the host and appended as an attribute to the module.
 
 module attributes {omp.is_target_device = true} {
-  // CHECK-DAG: @_QMtest_0Esp_decl_tgt_ref_ptr = weak global ptr null, align 8
+  // CHECK-DAG: @_QMtest_0Esp_decl_tgt_ref_ptr = weak dso_local global ptr null, align 8
   llvm.mlir.global external @_QMtest_0Esp() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : i32 {
     %0 = llvm.mlir.constant(0 : i32) : i32
     llvm.return %0 : i32
diff --git a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir
index 2baa20010d0558..763f12de8aaef3 100644
--- a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir
@@ -5,14 +5,14 @@
 module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_target_device = false} {
 
   // CHECK-DAG: @_QMtest_0Earray_1d = global [3 x i32] [i32 1, i32 2, i32 3]
-  // CHECK-DAG: @_QMtest_0Earray_1d_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Earray_1d
+  // CHECK-DAG: @_QMtest_0Earray_1d_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Earray_1d
   // CHECK-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [36 x i8] c"_QMtest_0Earray_1d_decl_tgt_ref_ptr\00"
   // CHECK-DAG: @.omp_offloading.entry._QMtest_0Earray_1d_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Earray_1d_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Earray_1d_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
   llvm.mlir.global external @_QMtest_0Earray_1d(dense<[1, 2, 3]> : tensor<3xi32>) {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : !llvm.array<3 x i32>
 
   // CHECK-DAG: @_QMtest_0Earray_2d = global [2 x [2 x i32]] {{.*}}
-  // CHECK-DAG: @_QMtest_0Earray_2d_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Earray_2d
+  // CHECK-DAG: @_QMtest_0Earray_2d_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Earray_2d
   // CHECK-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [36 x i8] c"_QMtest_0Earray_2d_decl_tgt_ref_ptr\00"
   // CHECK-DAG: @.omp_offloading.entry._QMtest_0Earray_2d_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Earray_2d_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Earray_2d_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
@@ -32,7 +32,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
   }
 
   // CHECK-DAG:  @_QMtest_0Edata_extended_link_1 = global float 2.000000e+00
-  // CHECK-DAG:  @_QMtest_0Edata_extended_link_1_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Edata_extended_link_1
+  // CHECK-DAG:  @_QMtest_0Edata_extended_link_1_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Edata_extended_link_1
   // CHECK-DAG:  @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [48 x i8] c"_QMtest_0Edata_extended_link_1_decl_tgt_ref_ptr\00"
   // CHECK-DAG:  @.omp_offloading.entry._QMtest_0Edata_extended_link_1_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Edata_extended_link_1_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG:  !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_extended_link_1_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
@@ -42,7 +42,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
   }
 
   // CHECK-DAG:  @_QMtest_0Edata_extended_link_2 = global float 3.000000e+00
-  // CHECK-DAG:  @_QMtest_0Edata_extended_link_2_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Edata_extended_link_2
+  // CHECK-DAG:  @_QMtest_0Edata_extended_link_2_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Edata_extended_link_2
   // CHECK-DAG:  @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [48 x i8] c"_QMtest_0Edata_extended_link_2_decl_tgt_ref_ptr\00"
   // CHECK-DAG:  @.omp_offloading.entry._QMtest_0Edata_extended_link_2_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Edata_extended_link_2_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG:  !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_extended_link_2_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
@@ -88,7 +88,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
   }
 
   // CHECK-DAG: @_QMtest_0Edata_int = global i32 1
-  // CHECK-DAG: @_QMtest_0Edata_int_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Edata_int
+  // CHECK-DAG: @_QMtest_0Edata_int_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Edata_int
   // CHECK-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [36 x i8] c"_QMtest_0Edata_int_decl_tgt_ref_ptr\00"
   // CHECK-DAG: @.omp_offloading.entry._QMtest_0Edata_int_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Edata_int_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_int_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
@@ -134,7 +134,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
   }
 
   // CHECK-DAG: @_QMtest_0Ept1 = global { ptr, i64, i32, i8, i8, i8, i8 } { ptr null, i64 ptrtoint (ptr getelementptr (i32, ptr null, i32 1) to i64), i32 20180515, i8 0, i8 9, i8 1, i8 0 }
-  // CHECK-DAG: @_QMtest_0Ept1_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Ept1
+  // CHECK-DAG: @_QMtest_0Ept1_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Ept1
   // CHECK-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [31 x i8] c"_QMtest_0Ept1_decl_tgt_ref_ptr\00"
   // CHECK-DAG: @.omp_offloading.entry._QMtest_0Ept1_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Ept1_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Ept1_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
@@ -165,7 +165,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
   }
 
   // CHECK-DAG: @_QMtest_0Ept2_tar = global i32 5
-  // CHECK-DAG: @_QMtest_0Ept2_tar_decl_tgt_ref_ptr = weak global ptr @_QMtest_0Ept2_tar
+  // CHECK-DAG: @_QMtest_0Ept2_tar_decl_tgt_ref_ptr = weak dso_local global ptr @_QMtest_0Ept2_tar
   // CHECK-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [35 x i8] c"_QMtest_0Ept2_tar_decl_tgt_ref_ptr\00"
   // CHECK-DAG: @.omp_offloading.entry._QMtest_0Ept2_tar_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { ptr @_QMtest_0Ept2_tar_decl_tgt_ref_ptr, ptr @.omp_offloading.entry_name{{.*}}, i64 8, i32 1, i32 0 }, section "omp_offloading_entries", align 1
   // CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Ept2_tar_decl_tgt_ref_ptr", i32 {{.*}}, i32 {{.*}}}
diff --git a/mlir/test/Target/LLVMIR/openmp-llvm.mlir b/mlir/test/Target/LLVMIR/openmp-llvm.mlir
index 1c02c0265462c2..a8a90bb3cb1fb8 100644
--- a/mlir/test/Target/LLVMIR/openmp-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/openmp-llvm.mlir
@@ -2166,7 +2166,7 @@ llvm.func @single_nowait(%x: i32, %y: i32, %zaddr: !llvm.ptr) {
 // -----
 
 // CHECK: @_QFsubEx = internal global i32 undef
-// CHECK: @_QFsubEx.cache = common global ptr null
+// CHECK: @_QFsubEx.cache = common dso_local global ptr null
 
 // CHECK-LABEL: @omp_threadprivate
 llvm.func @omp_threadprivate() {

>From f39d169abcc04f70f03397e8101e787fb00a0ec0 Mon Sep 17 00:00:00 2001
From: Umesh Kalappa <Umesh.Kalappa at amd.com>
Date: Wed, 20 Dec 2023 15:33:36 +0530
Subject: [PATCH 4/4] Updated with relevant comment.

---
 llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 910d7c6e76e561..4588be4a703f05 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -5224,7 +5224,10 @@ OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name,
     const llvm::Align TypeAlign = DL.getABITypeAlign(Ty);
     const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpace);
     GV->setAlignment(std::max(TypeAlign, PtrAlign));
-
+    // Mark the GV with dso_local attribute to prevent the GOT relocations
+    // for non-pic/pie code.
+    // TODO :later we refactor these changes if you think otherwise like
+    // GV is preemptible(for no-pic object,not sure case exist).
     if (!GV->isDSOLocal()) {
       bool IsPIE = GV->getParent()->getPIELevel() != llvm::PIELevel::Default;
       bool IsPIC = GV->getParent()->getPICLevel() != llvm::PICLevel::NotPIC;



More information about the llvm-commits mailing list