This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG807466ef2812: [OpenMP] Restore backwards compatibility for
libomptarget (authored by jhuber6).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D98358/new/
https://reviews.llvm.org/D98358
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp
clang/test/OpenMP/teams_distribute_codegen.cpp
clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp
clang/test/OpenMP/teams_distribute_simd_codegen.cpp
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
llvm/test/Transforms/OpenMP/add_attributes.ll
openmp/libomptarget/include/omptarget.h
openmp/libomptarget/src/exports
openmp/libomptarget/src/interface.cpp
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/test/offloading/host_as_target.c
Index: openmp/libomptarget/test/offloading/host_as_target.c
===================================================================
--- openmp/libomptarget/test/offloading/host_as_target.c
+++ openmp/libomptarget/test/offloading/host_as_target.c
@@ -55,8 +55,8 @@
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 @@
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)
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -1023,8 +1023,8 @@
/// 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;
Index: openmp/libomptarget/src/interface.cpp
===================================================================
--- openmp/libomptarget/src/interface.cpp
+++ openmp/libomptarget/src/interface.cpp
@@ -437,8 +437,13 @@
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);
Index: openmp/libomptarget/src/exports
===================================================================
--- openmp/libomptarget/src/exports
+++ openmp/libomptarget/src/exports
@@ -25,6 +25,8 @@
__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 @@
omp_target_memcpy_rect;
omp_target_associate_ptr;
omp_target_disassociate_ptr;
- __kmpc_push_target_tripcount;
local:
*;
};
Index: openmp/libomptarget/include/omptarget.h
===================================================================
--- openmp/libomptarget/include/omptarget.h
+++ openmp/libomptarget/include/omptarget.h
@@ -312,8 +312,10 @@
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
}
Index: llvm/test/Transforms/OpenMP/add_attributes.ll
===================================================================
--- llvm/test/Transforms/OpenMP/add_attributes.ll
+++ llvm/test/Transforms/OpenMP/add_attributes.ll
@@ -627,7 +627,7 @@
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 @@
; 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 @@
; 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()
Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -375,7 +375,7 @@
__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_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(), {})
Index: clang/test/OpenMP/teams_distribute_simd_codegen.cpp
===================================================================
--- clang/test/OpenMP/teams_distribute_simd_codegen.cpp
+++ clang/test/OpenMP/teams_distribute_simd_codegen.cpp
@@ -35,7 +35,7 @@
// 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]],
Index: clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp
===================================================================
--- clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp
+++ clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp
@@ -33,7 +33,7 @@
// 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]],
Index: clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
===================================================================
--- clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
+++ clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
@@ -32,7 +32,7 @@
// 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]],
Index: clang/test/OpenMP/teams_distribute_codegen.cpp
===================================================================
--- clang/test/OpenMP/teams_distribute_codegen.cpp
+++ clang/test/OpenMP/teams_distribute_codegen.cpp
@@ -33,7 +33,7 @@
// 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]],
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_if_codegen.cpp
@@ -43,10 +43,10 @@
// 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 @@
// 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
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_codegen.cpp
@@ -60,7 +60,7 @@
// 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]])
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_order_codegen.cpp
+++ 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
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_if_codegen.cpp
@@ -49,10 +49,10 @@
// 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 @@
// 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
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_codegen.cpp
+++ 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 @@
// 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]])
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9948,7 +9948,7 @@
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);
}
};
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
