jhuber6 created this revision.
jhuber6 added a reviewer: jdoerfert.
Herald added subscribers: guansong, hiraditya, yaxunl.
jhuber6 requested review of this revision.
Herald added subscribers: llvm-commits, cfe-commits, sstefan1.
Herald added projects: clang, LLVM.
The new device runtime uses an internal variable to set debugging. This
variable was originally privately linked because every module will have
a copy of it. This caused problems with merging the device bitcode
library because it would get renamed and there was not a way to refer to
an external, private symbol. This changes the symbol to weak_odr so it
can be defined multiply, but will not be renamed.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D109997
Files:
clang/test/OpenMP/target_debug_codegen.cpp
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
Index: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
===================================================================
--- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -249,11 +249,9 @@
IntegerType *I32Ty = Type::getInt32Ty(M.getContext());
auto *GV = new GlobalVariable(
M, I32Ty,
- /* isConstant = */ true, GlobalValue::PrivateLinkage,
+ /* isConstant = */ true, GlobalValue::WeakODRLinkage,
ConstantInt::get(I32Ty, DebugKind), "__omp_rtl_debug_kind");
- llvm::appendToUsed(M, {GV});
-
return GV;
}
Index: clang/test/OpenMP/target_debug_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_debug_codegen.cpp
+++ clang/test/OpenMP/target_debug_codegen.cpp
@@ -1,4 +1,4 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --check-globals --global-value-regex
"(__omp_rtl_debug_kind|llvm\.used)"
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --check-globals --global-value-regex "__omp_rtl_debug_kind"
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown
-fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown
-fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime
-fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path
%t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK
@@ -10,14 +10,11 @@
#define HEADER
//.
-// CHECK: @__omp_rtl_debug_kind = private constant i32 1
-// CHECK: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32*
@__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK: @__omp_rtl_debug_kind = weak_odr constant i32 1
//.
-// CHECK-EQ: @__omp_rtl_debug_kind = private constant i32 111
-// CHECK-EQ: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32*
@__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK-EQ: @__omp_rtl_debug_kind = weak_odr constant i32 111
//.
-// CHECK-DEFAULT: @__omp_rtl_debug_kind = private constant i32 0
-// CHECK-DEFAULT: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32*
@__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK-DEFAULT: @__omp_rtl_debug_kind = weak_odr constant i32 0
//.
void foo() {
#pragma omp target
Index: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
===================================================================
--- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -249,11 +249,9 @@
IntegerType *I32Ty = Type::getInt32Ty(M.getContext());
auto *GV = new GlobalVariable(
M, I32Ty,
- /* isConstant = */ true, GlobalValue::PrivateLinkage,
+ /* isConstant = */ true, GlobalValue::WeakODRLinkage,
ConstantInt::get(I32Ty, DebugKind), "__omp_rtl_debug_kind");
- llvm::appendToUsed(M, {GV});
-
return GV;
}
Index: clang/test/OpenMP/target_debug_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_debug_codegen.cpp
+++ clang/test/OpenMP/target_debug_codegen.cpp
@@ -1,4 +1,4 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex "(__omp_rtl_debug_kind|llvm\.used)"
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex "__omp_rtl_debug_kind"
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK
@@ -10,14 +10,11 @@
#define HEADER
//.
-// CHECK: @__omp_rtl_debug_kind = private constant i32 1
-// CHECK: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32* @__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK: @__omp_rtl_debug_kind = weak_odr constant i32 1
//.
-// CHECK-EQ: @__omp_rtl_debug_kind = private constant i32 111
-// CHECK-EQ: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32* @__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK-EQ: @__omp_rtl_debug_kind = weak_odr constant i32 111
//.
-// CHECK-DEFAULT: @__omp_rtl_debug_kind = private constant i32 0
-// CHECK-DEFAULT: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32* @__omp_rtl_debug_kind to i8*)], section "llvm.metadata"
+// CHECK-DEFAULT: @__omp_rtl_debug_kind = weak_odr constant i32 0
//.
void foo() {
#pragma omp target
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
