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