[clang] 8c1449a - [OpenMP] Make kernels have protected visibility
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 18 14:37:36 PDT 2022
Author: Joseph Huber
Date: 2022-10-18T16:37:28-05:00
New Revision: 8c1449a84d61f6caa4b6e63f262b2c8976949548
URL: https://github.com/llvm/llvm-project/commit/8c1449a84d61f6caa4b6e63f262b2c8976949548
DIFF: https://github.com/llvm/llvm-project/commit/8c1449a84d61f6caa4b6e63f262b2c8976949548.diff
LOG: [OpenMP] Make kernels have protected visibility
This patch changes the kernels generated by OpenMP to have protected
visibility. This is unlikely to change anything functionally. However,
protected visibility better matches the behaviour of these GPU kernels.
We do not expect any pending shared library load to preempt these
kernels so we can specify a more restrictive visibility.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D136198
Added:
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/amdgcn_target_codegen.cpp
clang/test/OpenMP/assumes_include_nvptx.cpp
clang/test/OpenMP/declare_target_codegen.cpp
clang/test/OpenMP/declare_target_link_codegen.cpp
clang/test/OpenMP/metadirective_device_arch_codegen.cpp
clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
clang/test/OpenMP/target_firstprivate_codegen.cpp
clang/test/OpenMP/target_private_codegen.cpp
clang/test/OpenMP/target_reduction_codegen.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 1c211b2d7ad3d..84536363d6053 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1883,6 +1883,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
llvm::Function *Fn = CGM.CreateGlobalInitOrCleanUpFunction(
FTy, Twine(Buffer, "_ctor"), FI, Loc, false,
llvm::GlobalValue::WeakODRLinkage);
+ Fn->setVisibility(llvm::GlobalValue::ProtectedVisibility);
if (CGM.getTriple().isAMDGCN())
Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
auto NL = ApplyDebugLocation::CreateEmpty(CtorCGF);
@@ -1930,6 +1931,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
llvm::Function *Fn = CGM.CreateGlobalInitOrCleanUpFunction(
FTy, Twine(Buffer, "_dtor"), FI, Loc, false,
llvm::GlobalValue::WeakODRLinkage);
+ Fn->setVisibility(llvm::GlobalValue::ProtectedVisibility);
if (CGM.getTriple().isAMDGCN())
Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
auto NL = ApplyDebugLocation::CreateEmpty(DtorCGF);
@@ -6332,6 +6334,7 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
OutlinedFnID = llvm::ConstantExpr::getBitCast(OutlinedFn, CGM.Int8PtrTy);
OutlinedFn->setLinkage(llvm::GlobalValue::WeakODRLinkage);
OutlinedFn->setDSOLocal(false);
+ OutlinedFn->setVisibility(llvm::GlobalValue::ProtectedVisibility);
if (CGM.getTriple().isAMDGCN())
OutlinedFn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
} else {
diff --git a/clang/test/OpenMP/amdgcn_target_codegen.cpp b/clang/test/OpenMP/amdgcn_target_codegen.cpp
index 2edb706a591ed..99486703af416 100644
--- a/clang/test/OpenMP/amdgcn_target_codegen.cpp
+++ b/clang/test/OpenMP/amdgcn_target_codegen.cpp
@@ -9,7 +9,7 @@
#define N 1000
int test_amdgcn_target_tid_threads() {
-// CHECK-LABEL: define weak_odr amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads
+// CHECK-LABEL: define weak_odr protected amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads
int arr[N];
@@ -23,7 +23,7 @@ int test_amdgcn_target_tid_threads() {
}
int test_amdgcn_target_tid_threads_simd() {
-// CHECK-LABEL: define weak_odr amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads_simd
+// CHECK-LABEL: define weak_odr protected amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads_simd
int arr[N];
diff --git a/clang/test/OpenMP/assumes_include_nvptx.cpp b/clang/test/OpenMP/assumes_include_nvptx.cpp
index 8b720f3d334b8..c12e9833b9291 100644
--- a/clang/test/OpenMP/assumes_include_nvptx.cpp
+++ b/clang/test/OpenMP/assumes_include_nvptx.cpp
@@ -11,11 +11,11 @@
// TODO: Think about teaching the OMPIRBuilder about default attributes as well so the __kmpc* declarations are annotated.
-// CHECK: define weak_odr void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}() [[attr0:#[0-9]]]
+// CHECK: define weak_odr protected void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}() [[attr0:#[0-9]]]
// CHECK: call i32 @__kmpc_target_init(
// CHECK: declare noundef float @_Z3sinf(float noundef) [[attr1:#[0-9]*]]
// CHECK: declare void @__kmpc_target_deinit(
-// CHECK: define weak_odr void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}() [[attr0]]
+// CHECK: define weak_odr protected void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}() [[attr0]]
// CHECK: %call = call noundef double @_Z3sind(double noundef 0.000000e+00) [[attr2:#[0-9]]]
// CHECK: declare noundef double @_Z3sind(double noundef) [[attr1]]
diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp
index 43a79a98f12e8..e43289e37dfcc 100644
--- a/clang/test/OpenMP/declare_target_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_codegen.cpp
@@ -139,7 +139,7 @@ int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
int maini1() {
int a;
static long aa = 32 + bbb + ccc + fff + ggg;
-// CHECK-DAG: define weak_odr void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](ptr noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}, i64 {{.*}}, i64 {{.*}})
+// CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](ptr noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}, i64 {{.*}}, i64 {{.*}})
#pragma omp target map(tofrom \
: a, b)
{
@@ -152,7 +152,7 @@ int maini1() {
int baz3() { return 2 + baz2(); }
int baz2() {
-// CHECK-DAG: define weak_odr void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
+// CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
#pragma omp target parallel
++c;
return 2 + baz3();
@@ -164,7 +164,7 @@ static __typeof(create) __t_create __attribute__((__weakref__("__create")));
int baz5() {
bool a;
-// CHECK-DAG: define weak_odr void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}})
+// CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}})
#pragma omp target
a = __extension__(void *) & __t_create != 0;
return a;
diff --git a/clang/test/OpenMP/declare_target_link_codegen.cpp b/clang/test/OpenMP/declare_target_link_codegen.cpp
index 3e6ca6a542b8b..bfb2e6726edbf 100644
--- a/clang/test/OpenMP/declare_target_link_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_link_codegen.cpp
@@ -50,7 +50,7 @@ int maini1() {
return 0;
}
-// DEVICE: define weak_odr void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* noundef nonnull align {{[0-9]+}} dereferenceable{{[^,]*}}
+// DEVICE: define weak_odr protected void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* noundef nonnull align {{[0-9]+}} dereferenceable{{[^,]*}}
// DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_ref_ptr,
// DEVICE: [[C:%.+]] = load i32, i32* [[C_REF]],
// DEVICE: store i32 [[C]], i32* %
diff --git a/clang/test/OpenMP/metadirective_device_arch_codegen.cpp b/clang/test/OpenMP/metadirective_device_arch_codegen.cpp
index eac71d0e5b5a5..8289f1e2d8db4 100644
--- a/clang/test/OpenMP/metadirective_device_arch_codegen.cpp
+++ b/clang/test/OpenMP/metadirective_device_arch_codegen.cpp
@@ -38,7 +38,7 @@ int metadirective1() {
return errors;
}
-// CHECK-LABEL: define weak_odr amdgpu_kernel void {{.+}}metadirective1
+// CHECK-LABEL: define weak_odr protected amdgpu_kernel void {{.+}}metadirective1
// CHECK: entry:
// CHECK: %{{[0-9]}} = call i32 @__kmpc_target_init
// CHECK: user_code.entry:
diff --git a/clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp b/clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
index a80e6278d4d74..27632a707ba53 100644
--- a/clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
+++ b/clang/test/OpenMP/metadirective_device_isa_codegen_amdgcn.cpp
@@ -22,7 +22,7 @@ int amdgcn_device_isa_selected() {
return threadCount;
}
-// CHECK: define weak_odr amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected
+// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_selected
// CHECK: user_code.entry:
// CHECK: call void @__kmpc_parallel_51
// CHECK-NOT: call i32 @__kmpc_single
@@ -44,7 +44,7 @@ int amdgcn_device_isa_not_selected() {
return threadCount;
}
-// CHECK: define weak_odr amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_not_selected
+// CHECK: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}amdgcn_device_isa_not_selected
// CHECK: user_code.entry:
// CHECK: call i32 @__kmpc_single
// CHECK-NOT: call void @__kmpc_parallel_51
diff --git a/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp b/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
index c38e958f74f36..cb857a6a5d430 100644
--- a/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
@@ -44,7 +44,7 @@ int caz() { return 0; }
static int c = foo() + bar() + baz();
#pragma omp declare target (c)
// HOST-DAG: @[[C_CTOR:__omp_offloading__.+_c_l44_ctor]] = private constant i8 0
-// DEVICE-DAG: define weak_odr void [[C_CTOR:@__omp_offloading__.+_c_l44_ctor]]()
+// DEVICE-DAG: define weak_odr protected void [[C_CTOR:@__omp_offloading__.+_c_l44_ctor]]()
// DEVICE-DAG: call noundef i32 [[FOO]]()
// DEVICE-DAG: call noundef i32 [[BAR]]()
// DEVICE-DAG: call noundef i32 [[BAZ]]()
@@ -61,14 +61,14 @@ struct S {
S cd = doo() + car() + caz() + baz();
#pragma omp end declare target
// HOST-DAG: @[[CD_CTOR:__omp_offloading__.+_cd_l61_ctor]] = private constant i8 0
-// DEVICE-DAG: define weak_odr void [[CD_CTOR:@__omp_offloading__.+_cd_l61_ctor]]()
+// DEVICE-DAG: define weak_odr protected void [[CD_CTOR:@__omp_offloading__.+_cd_l61_ctor]]()
// DEVICE-DAG: call noundef i32 [[DOO]]()
// DEVICE-DAG: call noundef i32 [[CAR]]()
// DEVICE-DAG: call noundef i32 [[CAZ]]()
// DEVICE-DAG: ret void
// HOST-DAG: @[[CD_DTOR:__omp_offloading__.+_cd_l61_dtor]] = private constant i8 0
-// DEVICE-DAG: define weak_odr void [[CD_DTOR:@__omp_offloading__.+_cd_l61_dtor]]()
+// DEVICE-DAG: define weak_odr protected void [[CD_DTOR:@__omp_offloading__.+_cd_l61_dtor]]()
// DEVICE-DAG: call void
// DEVICE-DAG: ret void
diff --git a/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp b/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
index b1b5ea80ffc23..875cc47544b0e 100644
--- a/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
@@ -95,7 +95,7 @@ int foo(int n, double *ptr) {
ptr[0]++;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(double* noundef [[PTR_IN:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(double* noundef [[PTR_IN:%.+]])
// TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
// TCHECK-NOT: alloca double*,
// TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]],
diff --git a/clang/test/OpenMP/target_firstprivate_codegen.cpp b/clang/test/OpenMP/target_firstprivate_codegen.cpp
index 8fb6c7b3c1ef3..4f8f5ae37472c 100644
--- a/clang/test/OpenMP/target_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/target_firstprivate_codegen.cpp
@@ -143,7 +143,7 @@ int foo(int n, double *ptr) {
// CHECK: [[PTR_GEP_ARG:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
// CHECK: {{.+}} = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[A_IN:%.+]], i32** noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[P_IN:%.+]], i{{[0-9]+}} noundef [[GA_IN:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[A_IN:%.+]], i32** noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[P_IN:%.+]], i{{[0-9]+}} noundef [[GA_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[P_ADDR:%.+]] = alloca i32**,
// TCHECK: [[GA_ADDR:%.+]] = alloca i{{64|32}},
@@ -352,7 +352,7 @@ int foo(int n, double *ptr) {
// CHECK: [[PTR_GEP_ARG3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
// CHECK: {{.+}} = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(double* noundef [[PTR_IN:%.+]], [[TTII]]* noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[E:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(double* noundef [[PTR_IN:%.+]], [[TTII]]* noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[E:%.+]])
// TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
// TCHECK-NOT: alloca [[TTII]],
// TCHECK-NOT: alloca double*,
@@ -391,7 +391,7 @@ static int fstatic(int n) {
return a;
}
-// TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[A_IN:%.+]], i{{[0-9]+}} noundef [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[A_IN:%.+]], i{{[0-9]+}} noundef [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
@@ -479,7 +479,7 @@ struct S1 {
// only check that we use the map types stored in the global variable
// CHECK: {{.+}} = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i{{[0-9]+}} noundef [[B_IN:%.+]], i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i{{[0-9]+}}{{.+}} [[C_IN:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i{{[0-9]+}} noundef [[B_IN:%.+]], i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i{{[0-9]+}}{{.+}} [[C_IN:%.+]])
// TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
// TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
@@ -587,7 +587,7 @@ int bar(int n, double *ptr) {
// CHECK: {{.+}} = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]])
-// TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK-NOT: alloca i{{[0-9]+}},
diff --git a/clang/test/OpenMP/target_private_codegen.cpp b/clang/test/OpenMP/target_private_codegen.cpp
index d34ae99c18a02..de060df17af94 100644
--- a/clang/test/OpenMP/target_private_codegen.cpp
+++ b/clang/test/OpenMP/target_private_codegen.cpp
@@ -45,7 +45,7 @@ int foo(int n) {
{
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}()
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}()
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
// TCHECK-NOT: store {{.+}}, {{.+}} [[A]],
// TCHECK: ret void
@@ -55,7 +55,7 @@ int foo(int n) {
a = 1;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}()
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}()
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
// TCHECK: ret void
@@ -66,7 +66,7 @@ int foo(int n) {
aa = 1;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}()
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}()
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
@@ -85,7 +85,7 @@ int foo(int n) {
}
// make sure that private variables are generated in all cases and that we use those instances for operations inside the
// target region
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i{{[0-9]+}} noundef [[VLA3:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i{{[0-9]+}} noundef [[VLA3:%.+]])
// TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[VLA_ADDR4:%.+]] = alloca i{{[0-9]+}},
@@ -179,7 +179,7 @@ int fstatic(int n) {
return a;
}
-// TCHECK: define weak_odr void @__omp_offloading_{{.+}}()
+// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}()
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}},
@@ -207,7 +207,7 @@ struct S1 {
return c[1][1] + (int)b;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]])
// TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
// TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
@@ -261,7 +261,7 @@ int bar(int n){
}
// template
-// TCHECK: define weak_odr void @__omp_offloading_{{.+}}()
+// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}()
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
diff --git a/clang/test/OpenMP/target_reduction_codegen.cpp b/clang/test/OpenMP/target_reduction_codegen.cpp
index e61a409129ff8..ad20ac8dd68f3 100644
--- a/clang/test/OpenMP/target_reduction_codegen.cpp
+++ b/clang/test/OpenMP/target_reduction_codegen.cpp
@@ -45,7 +45,7 @@ int foo(int n) {
{
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: load i32*, i32** [[A]],
@@ -56,7 +56,7 @@ int foo(int n) {
a = 1;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: [[REF:%.+]] = load i32*, i32** [[A]],
@@ -69,7 +69,7 @@ int foo(int n) {
aa = 1;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[AA:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
@@ -118,7 +118,7 @@ int fstatic(int n) {
return a;
}
-// TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
+// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}}*,
@@ -154,7 +154,7 @@ struct S1 {
return c[1][1] + (int)b;
}
- // TCHECK: define weak_odr void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i16*{{.+}})
+ // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}([[S1]]* noundef [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], i16*{{.+}})
// TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
// TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
@@ -206,7 +206,7 @@ int bar(int n){
}
// template
-// TCHECK: define weak_odr void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
+// TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
More information about the cfe-commits
mailing list