From 807466ef28125cf7268c860b09d5563c9c93602a Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Wed, 10 Mar 2021 13:25:33 -0500 Subject: [PATCH] [OpenMP] Restore backwards compatibility for libomptarget Summary: The changes introduced in D87946 changed the API for libomptarget functions. `__kmpc_push_target_tripcount` was a function in Clang 11.x but was not given a backward-compatible interface. This change will require people using Clang 13.x or 12.x to recompile their offloading programs. Reviewed By: jdoerfert cchen Differential Revision: https://reviews.llvm.org/D98358 --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 2 +- .../OpenMP/target_teams_distribute_parallel_for_codegen.cpp | 4 ++-- .../OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp | 10 +++++----- .../target_teams_distribute_parallel_for_order_codegen.cpp | 2 +- .../target_teams_distribute_parallel_for_simd_codegen.cpp | 2 +- .../target_teams_distribute_parallel_for_simd_if_codegen.cpp | 10 +++++----- clang/test/OpenMP/teams_distribute_codegen.cpp | 2 +- clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp | 2 +- .../test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp | 2 +- clang/test/OpenMP/teams_distribute_simd_codegen.cpp | 2 +- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def | 4 ++-- llvm/test/Transforms/OpenMP/add_attributes.ll | 6 +++--- openmp/libomptarget/include/omptarget.h | 6 ++++-- openmp/libomptarget/src/exports | 3 ++- openmp/libomptarget/src/interface.cpp | 7 ++++++- openmp/libomptarget/src/omptarget.cpp | 4 ++-- openmp/libomptarget/test/offloading/host_as_target.c | 8 ++++---- 17 files changed, 42 insertions(+), 34 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index f9c79e6..e41eeef 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -9948,7 +9948,7 @@ void CGOpenMPRuntime::emitTargetNumIterationsCall( llvm::Value *Args[] = {RTLoc, DeviceID, NumIterations}; CGF.EmitRuntimeCall( OMPBuilder.getOrCreateRuntimeFunction( - CGM.getModule(), OMPRTL___kmpc_push_target_tripcount), + CGM.getModule(), OMPRTL___kmpc_push_target_tripcount_mapper), Args); } }; diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp index 0229ace..c0f5323 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp @@ -39,7 +39,7 @@ #ifdef CK1 -// HCK_NO_TGT-NOT: @__kmpc_push_target_tripcount +// HCK_NO_TGT-NOT: @__kmpc_push_target_tripcount_mapper // HCK1: define{{.*}} i32 @{{.+}}target_teams_fun{{.*}}( int target_teams_fun(int *g){ @@ -60,7 +60,7 @@ int target_teams_fun(int *g){ // HCK1: [[N_PAR:%.+]] = load{{.+}}, {{.+}} [[N_CAST]], // HCK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // HCK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // HCK1: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) + // HCK1: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) // HCK1: call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, // HCK1: call void @[[OFFL1:.+]](i{{32|64}} [[N_PAR]], {{.+}}, i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]]) diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp index 6650e05..efe7df8 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp @@ -49,10 +49,10 @@ int Arg; // CHECK-LABEL: define {{.*}}void @{{.+}}gtid_test void gtid_test() { -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( #pragma omp target teams distribute parallel for @@ -107,12 +107,12 @@ int tmain(T Arg) { // CHECK-LABEL: define {{.*}}i{{[0-9]+}} @main() int main() { -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK-NOT: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK-NOT: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_2:@.+]]( // CHECK: = call {{.*}}i{{.+}} @{{.+}}tmain diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp index b2ab37f..b99ba9d 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp @@ -14,7 +14,7 @@ // CHECK-LABEL: define {{.*}}void @{{.+}}gtid_test void gtid_test() { -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: %0 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null, i8** null, i32 0, i32 0) // CHECK: call void [[TARGET_OUTLINE:@.+]]() // CHECK: ret void diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp index e604914..39ccb87 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp @@ -60,7 +60,7 @@ int target_teams_fun(int *g){ // HCK1: [[N_PAR:%.+]] = load{{.+}}, {{.+}} [[N_CAST]], // HCK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // HCK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], -// HCK1: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) +// HCK1: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) // HCK1: call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, // HCK1: call void @[[OFFL1:.+]](i{{32|64}} [[I_PAR]], i{{32|64}} [[N_PAR]], {{.+}}, i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]]) diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp index 8b0eaba..19dc15b 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp @@ -43,10 +43,10 @@ int Arg; // CHECK-LABEL: define {{.*}}void @{{.+}}gtid_test void gtid_test() { -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( #ifdef OMP5 @@ -110,12 +110,12 @@ int tmain(T Arg) { // CHECK-LABEL: define {{.*}}i{{[0-9]+}} @main() int main() { -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK-NOT: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK-NOT: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( -// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 100) +// CHECK: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 100) // CHECK: call i{{[0-9]+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, // CHECK: call void [[OFFLOADING_FUN_2:@.+]]( // CHECK: = call {{.*}}i{{.+}} @{{.+}}tmain diff --git a/clang/test/OpenMP/teams_distribute_codegen.cpp b/clang/test/OpenMP/teams_distribute_codegen.cpp index 5bbb100..aab5cce 100644 --- a/clang/test/OpenMP/teams_distribute_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_codegen.cpp @@ -33,7 +33,7 @@ int teams_argument_global(int n) { // CK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // CK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // CK1: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) + // CK1: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) // CK1: call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i8** null, i32 {{.+}}, i32 {{.+}}) // CK1: call void @[[OFFL1:.+]](i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]], diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp index b63e5ae..8fa73e7 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp @@ -32,7 +32,7 @@ int teams_argument_global(int n){ // CK1: [[TH_CAST:%.+]] = alloca i{{32|64}}, // CK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // CK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // CK1: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) + // CK1: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) // CK1: call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i8** null, i32 {{.+}}, i32 {{.+}}) // CK1: call void @[[OFFL1:.+]](i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]], diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp index 3d479c4..9b3855c 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp @@ -33,7 +33,7 @@ int teams_argument_global(int n){ // CK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // CK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // CK1: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) + // CK1: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) // CK1: call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null // CK1: call void @[[OFFL1:.+]](i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]], diff --git a/clang/test/OpenMP/teams_distribute_simd_codegen.cpp b/clang/test/OpenMP/teams_distribute_simd_codegen.cpp index fd1214d..6e5d06b 100644 --- a/clang/test/OpenMP/teams_distribute_simd_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_simd_codegen.cpp @@ -35,7 +35,7 @@ int teams_argument_global(int n) { // CK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // CK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // CK1: call void @__kmpc_push_target_tripcount(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) + // CK1: call void @__kmpc_push_target_tripcount_mapper(%struct.ident_t* @{{.+}}, i64 -1, i64 %{{.+}}) // CK1: call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i8** null, i32 {{.+}}, i32 1) // CK1: call void @[[OFFL1:.+]](i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]], diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def index 8440461..75d360b 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -375,7 +375,7 @@ __OMP_RTL(__kmpc_init_allocator, false, /* omp_allocator_handle_t */ VoidPtr, __OMP_RTL(__kmpc_destroy_allocator, false, Void, /* Int */ Int32, /* omp_allocator_handle_t */ VoidPtr) -__OMP_RTL(__kmpc_push_target_tripcount, false, Void, IdentPtr, Int64, Int64) +__OMP_RTL(__kmpc_push_target_tripcount_mapper, false, Void, IdentPtr, Int64, Int64) __OMP_RTL(__tgt_target_mapper, false, Int32, IdentPtr, Int64, VoidPtr, Int32, VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr) __OMP_RTL(__tgt_target_nowait_mapper, false, Int32, IdentPtr, Int64, VoidPtr, Int32, @@ -844,7 +844,7 @@ __OMP_RTL_ATTRS(__kmpc_free, AllocAttrs, AttributeSet(), {}) __OMP_RTL_ATTRS(__kmpc_init_allocator, DefaultAttrs, ReturnPtrAttrs, {}) __OMP_RTL_ATTRS(__kmpc_destroy_allocator, AllocAttrs, AttributeSet(), {}) -__OMP_RTL_ATTRS(__kmpc_push_target_tripcount, SetterAttrs, AttributeSet(), {}) +__OMP_RTL_ATTRS(__kmpc_push_target_tripcount_mapper, SetterAttrs, AttributeSet(), {}) __OMP_RTL_ATTRS(__tgt_target_mapper, ForkAttrs, AttributeSet(), {}) __OMP_RTL_ATTRS(__tgt_target_nowait_mapper, ForkAttrs, AttributeSet(), {}) __OMP_RTL_ATTRS(__tgt_target_teams_mapper, ForkAttrs, AttributeSet(), {}) diff --git a/llvm/test/Transforms/OpenMP/add_attributes.ll b/llvm/test/Transforms/OpenMP/add_attributes.ll index b294542..8476f42 100644 --- a/llvm/test/Transforms/OpenMP/add_attributes.ll +++ b/llvm/test/Transforms/OpenMP/add_attributes.ll @@ -627,7 +627,7 @@ declare i8* @__kmpc_init_allocator(i32, i8*, i32, i8*) declare void @__kmpc_destroy_allocator(i32, i8*) -declare void @__kmpc_push_target_tripcount(%struct.ident_t*, i64, i64) +declare void @__kmpc_push_target_tripcount_mapper(%struct.ident_t*, i64, i64) declare i32 @__kmpc_warp_active_thread_mask() @@ -1144,7 +1144,7 @@ declare void @__kmpc_proxy_task_completed_ooo(i8*) ; CHECK-NEXT: declare void @__kmpc_destroy_allocator(i32, i8*) ; CHECK: ; Function Attrs: nounwind -; CHECK-NEXT: declare void @__kmpc_push_target_tripcount(%struct.ident_t*, i64, i64) +; CHECK-NEXT: declare void @__kmpc_push_target_tripcount_mapper(%struct.ident_t*, i64, i64) ; CHECK: ; Function Attrs: convergent nounwind ; CHECK-NEXT: declare i32 @__kmpc_warp_active_thread_mask() @@ -1669,7 +1669,7 @@ declare void @__kmpc_proxy_task_completed_ooo(i8*) ; OPTIMISTIC-NEXT: declare void @__kmpc_destroy_allocator(i32, i8*) ; OPTIMISTIC: ; Function Attrs: inaccessiblememonly nofree nosync nounwind willreturn writeonly -; OPTIMISTIC-NEXT: declare void @__kmpc_push_target_tripcount(%struct.ident_t*, i64, i64) +; OPTIMISTIC-NEXT: declare void @__kmpc_push_target_tripcount_mapper(%struct.ident_t*, i64, i64) ; OPTIMISTIC: ; Function Attrs: convergent nounwind ; OPTIMISTIC-NEXT: declare i32 @__kmpc_warp_active_thread_mask() diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h index 76180fc..7b317c4 100644 --- a/openmp/libomptarget/include/omptarget.h +++ b/openmp/libomptarget/include/omptarget.h @@ -312,8 +312,10 @@ int __tgt_target_teams_nowait_mapper( int32_t thread_limit, int32_t depNum, void *depList, int32_t noAliasDepNum, void *noAliasDepList); -void __kmpc_push_target_tripcount(ident_t *loc, int64_t device_id, - uint64_t loop_tripcount); +void __kmpc_push_target_tripcount(int64_t device_id, uint64_t loop_tripcount); + +void __kmpc_push_target_tripcount_mapper(ident_t *loc, int64_t device_id, + uint64_t loop_tripcount); #ifdef __cplusplus } diff --git a/openmp/libomptarget/src/exports b/openmp/libomptarget/src/exports index 5e09a08..b7fc1c8 100644 --- a/openmp/libomptarget/src/exports +++ b/openmp/libomptarget/src/exports @@ -25,6 +25,8 @@ VERS1.0 { __tgt_target_teams_nowait_mapper; __tgt_mapper_num_components; __tgt_push_mapper_component; + __kmpc_push_target_tripcount; + __kmpc_push_target_tripcount_mapper; omp_get_num_devices; omp_get_initial_device; omp_target_alloc; @@ -34,7 +36,6 @@ VERS1.0 { omp_target_memcpy_rect; omp_target_associate_ptr; omp_target_disassociate_ptr; - __kmpc_push_target_tripcount; local: *; }; diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp index 104f1f9..b1e9342 100644 --- a/openmp/libomptarget/src/interface.cpp +++ b/openmp/libomptarget/src/interface.cpp @@ -437,8 +437,13 @@ EXTERN void __tgt_push_mapper_component(void *rt_mapper_handle, void *base, MapComponentInfoTy(base, begin, size, type, name)); } -EXTERN void __kmpc_push_target_tripcount(ident_t *loc, int64_t device_id, +EXTERN void __kmpc_push_target_tripcount(int64_t device_id, uint64_t loop_tripcount) { + __kmpc_push_target_tripcount_mapper(nullptr, device_id, loop_tripcount); +} + +EXTERN void __kmpc_push_target_tripcount_mapper(ident_t *loc, int64_t device_id, + uint64_t loop_tripcount) { TIMESCOPE_WITH_IDENT(loc); if (checkDeviceAndCtors(device_id, loc) != OFFLOAD_SUCCESS) { DP("Not offloading to device %" PRId64 "\n", device_id); diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 9470352..64a5292 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -1023,8 +1023,8 @@ TableMap *getTableMap(void *HostPtr) { /// Get loop trip count /// FIXME: This function will not work right if calling -/// __kmpc_push_target_tripcount in one thread but doing offloading in another -/// thread, which might occur when we call task yield. +/// __kmpc_push_target_tripcount_mapper in one thread but doing offloading in +/// another thread, which might occur when we call task yield. uint64_t getLoopTripCount(int64_t DeviceId) { DeviceTy &Device = PM->Devices[DeviceId]; uint64_t LoopTripCount = 0; diff --git a/openmp/libomptarget/test/offloading/host_as_target.c b/openmp/libomptarget/test/offloading/host_as_target.c index 56772a3..e64993b 100644 --- a/openmp/libomptarget/test/offloading/host_as_target.c +++ b/openmp/libomptarget/test/offloading/host_as_target.c @@ -55,8 +55,8 @@ int main(void) { printf("omp_is_initial_device() = %d\n", omp_is_initial_device()); CHECK_DATA(); - // Check that __kmpc_push_target_tripcount doesn't fail. I'm not sure how to - // check that it actually pushes to the initial device. + // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure + // how to check that it actually pushes to the initial device. #pragma omp target teams device(DevInit) num_teams(1) #pragma omp distribute for (int i = 0; i < 2; ++i) @@ -112,8 +112,8 @@ int main(void) { printf("omp_is_initial_device() = %d\n", omp_is_initial_device()); CHECK_DATA(); - // Check that __kmpc_push_target_tripcount doesn't fail. I'm not sure how to - // check that it actually pushes to the initial device. + // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure + // how to check that it actually pushes to the initial device. #pragma omp target teams num_teams(1) #pragma omp distribute for (int i = 0; i < 2; ++i) -- 2.7.4