[clang] [flang] [lld] [llvm] [lldb] [libc] [mlir] [libcxx] [compiler-rt] [clang-tools-extra] Don't emit relax relocs like R_X86_64_REX_GOTPCRELX on X86 target for OPENMP internal vars. (PR #75564)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Dec 18 20:50:31 PST 2023
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/3] 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/3] 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/3] 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() {
More information about the cfe-commits
mailing list