[NFC][OpenMP] Prepare amdgpu plugin for asynchronous implementation of target region launch

[carlo.bertolli]

Summary

Part 2 of
At present, amdgpu plugin merges both asynchronous and synchronous kernel launch implementations into a single synchronous version.
This patch prepares the plugin for asynchronous implementation by:

• Privatizing actual kernel launch code (valid in both cases) into an anonymous namespace base function (submitted at D115267)

• - Separating the control flow path of asynchronous and synchronous kernel launch functions** (this diff)

[carlo.bertolli]

format

[JonChesterfield]

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

2249

Naming seems wrong -the NowaitLocked function does wait, and if it didn't, we'd be missing a wait somewhere.

Diff has been created without context making it difficult to say anything further. 'Context not available.'

[carlo.bertolli]

[OpenMP][NFC] Apply requested change: remove "nowait" from function name.
Add context (sorry about that!)

[JonChesterfield]

LG, new entry point is a precondition on doing an async launch, else NFC.

Doesn't apply to main successfully though, arc patch D115273 fails.

[carlo.bertolli]

rebase onto main

[carlo.bertolli]

Trying rebase again.

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

static uint64_t acquire_available_packet_id(hsa_queue_t *queue) {
   bool full = true;
   while (full) {
     full =
         packet_id >= (queue->size + hsa_queue_load_read_index_scacquire(queue));
   }
   return packet_id;
 }
 
-int32_t __tgt_rtl_run_target_team_region_locked(
-    int32_t device_id, void *tgt_entry_ptr, void **tgt_args,
-    ptrdiff_t *tgt_offsets, int32_t arg_num, int32_t num_teams,
-    int32_t thread_limit, uint64_t loop_tripcount) {
+int32_t runRegionLocked(int32_t device_id, void *tgt_entry_ptr, void **tgt_args,
+                        ptrdiff_t *tgt_offsets, int32_t arg_num,
+                        int32_t num_teams, int32_t thread_limit,
+                        uint64_t loop_tripcount) {
   // Set the context we are using
   // update thread limit content in gpu memory if un-initialized or specified
   // from host
 
   DP("Run target team region thread_limit %d\n", thread_limit);
 
   // All args are references.
   std::vector<void *> args(arg_num);

 int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
                                          void **tgt_args,
                                          ptrdiff_t *tgt_offsets,
                                          int32_t arg_num, int32_t num_teams,
                                          int32_t thread_limit,
                                          uint64_t loop_tripcount) {
 
   DeviceInfo.load_run_lock.lock_shared();
-  int32_t res = __tgt_rtl_run_target_team_region_locked(
-      device_id, tgt_entry_ptr, tgt_args, tgt_offsets, arg_num, num_teams,
-      thread_limit, loop_tripcount);
+  int32_t res =
+      runRegionLocked(device_id, tgt_entry_ptr, tgt_args, tgt_offsets, arg_num,
+                      num_teams, thread_limit, loop_tripcount);
 
   DeviceInfo.load_run_lock.unlock_shared();
   return res;
 }
 
 int32_t __tgt_rtl_run_target_region(int32_t device_id, void *tgt_entry_ptr,
                                     void **tgt_args, ptrdiff_t *tgt_offsets,
                                     int32_t arg_num) {
   // use one team and one thread
   // fix thread num
   int32_t team_num = 1;
   int32_t thread_limit = 0; // use default
   return __tgt_rtl_run_target_team_region(device_id, tgt_entry_ptr, tgt_args,
                                           tgt_offsets, arg_num, team_num,
                                           thread_limit, 0);
 }
 
+int32_t __tgt_rtl_run_target_team_region_async(
+    int32_t device_id, void *tgt_entry_ptr, void **tgt_args,
+    ptrdiff_t *tgt_offsets, int32_t arg_num, int32_t num_teams,
+    int32_t thread_limit, uint64_t loop_tripcount) {
+
+  DeviceInfo.load_run_lock.lock_shared();
+  int32_t res =
+      runRegionLocked(device_id, tgt_entry_ptr, tgt_args, tgt_offsets, arg_num,
+                      num_teams, thread_limit, loop_tripcount);
+
+  DeviceInfo.load_run_lock.unlock_shared();
+  return res;
+}
+
 int32_t __tgt_rtl_run_target_region_async(int32_t device_id,
                                           void *tgt_entry_ptr, void **tgt_args,
                                           ptrdiff_t *tgt_offsets,
                                           int32_t arg_num,
                                           __tgt_async_info *AsyncInfo) {
   assert(AsyncInfo && "AsyncInfo is nullptr");
   initAsyncInfo(AsyncInfo);
 
   // use one team and one thread
   // fix thread num
   int32_t team_num = 1;
   int32_t thread_limit = 0; // use default
-  return __tgt_rtl_run_target_team_region(device_id, tgt_entry_ptr, tgt_args,
-                                          tgt_offsets, arg_num, team_num,
-                                          thread_limit, 0);
+  return __tgt_rtl_run_target_team_region_async(device_id, tgt_entry_ptr,
+                                                tgt_args, tgt_offsets, arg_num,
+                                                team_num, thread_limit, 0);
 }
 
 int32_t __tgt_rtl_synchronize(int32_t device_id, __tgt_async_info *AsyncInfo) {
   assert(AsyncInfo && "AsyncInfo is nullptr");
 
   // Cuda asserts that AsyncInfo->Queue is non-null, but this invariant
   // is not ensured by devices.cpp for amdgcn
   // assert(AsyncInfo->Queue && "AsyncInfo->Queue is nullptr");