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