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.kala...@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.kala...@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.kala...@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() { _______________________________________________ lldb-commits mailing list lldb-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/lldb-commits
