[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

GitOrigin-RevId: 807466ef28125cf7268c860b09d5563c9c93602a
diff --git a/libomptarget/include/omptarget.h b/libomptarget/include/omptarget.h
index 76180fc..7b317c4 100644
--- a/libomptarget/include/omptarget.h
+++ b/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
 }
diff --git a/libomptarget/src/exports b/libomptarget/src/exports
index 5e09a08..b7fc1c8 100644
--- a/libomptarget/src/exports
+++ b/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:
     *;
 };
diff --git a/libomptarget/src/interface.cpp b/libomptarget/src/interface.cpp
index 104f1f9..b1e9342 100644
--- a/libomptarget/src/interface.cpp
+++ b/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);
diff --git a/libomptarget/src/omptarget.cpp b/libomptarget/src/omptarget.cpp
index 9470352..64a5292 100644
--- a/libomptarget/src/omptarget.cpp
+++ b/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;
diff --git a/libomptarget/test/offloading/host_as_target.c b/libomptarget/test/offloading/host_as_target.c
index 56772a3..e64993b 100644
--- a/libomptarget/test/offloading/host_as_target.c
+++ b/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)