[libomptarget][nfc] Replace out arguments with struct return

[JonChesterfield]

Summary

A step towards making this function adequately self contained that it
can be tested easily. No functional change intended here, left variable
names unchanged.

[JonChesterfield]

• hoist access to device id

[JonChesterfield]

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

1891

replaces passing device_id with passing the one integer it was used to look up. Rolling into this change as still NFC and it spares me a merge conflict fixup

[ronlieb]

lgtm

[dhruvachak]

Anything that is accessible directly from getLaunchVals (such as DeviceInfo.*), can we not pass them and access them directly? e.g. DeviceInfo.EnvTeamLimit, etc. That would minimize the number of arguments to getLaunchVals.

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

int32_t __tgt_rtl_data_delete(int device_id, void *tgt_ptr) {
   return OFFLOAD_SUCCESS;
 }
 
 // Determine launch values for threadsPerGroup and num_groups.
 // Outputs: treadsPerGroup, num_groups
 // Inputs: Max_Teams, Max_WG_Size, Warp_Size, ExecutionMode,
 //         EnvTeamLimit, EnvNumTeams, num_teams, thread_limit,
 //         loop_tripcount.
-void getLaunchVals(int &threadsPerGroup, int &num_groups, int ConstWGSize,
-                   int ExecutionMode, int EnvTeamLimit, int EnvNumTeams,
-                   int num_teams, int thread_limit, uint64_t loop_tripcount,
-                   int32_t device_id) {
+struct launchVals {
+  int threadsPerGroup;
+  int num_groups;
+};
+
+launchVals getLaunchVals(int ConstWGSize, int ExecutionMode, int EnvTeamLimit,
+                         int EnvNumTeams, int num_teams, int thread_limit,
+                         uint64_t loop_tripcount, int DeviceNumTeams) {
+
+  int threadsPerGroup = RTLDeviceInfoTy::Default_WG_Size;
+  int num_groups = 0;
 
   int Max_Teams = DeviceInfo.EnvMaxTeamsDefault > 0
                       ? DeviceInfo.EnvMaxTeamsDefault
-                      : DeviceInfo.NumTeams[device_id];
+                      : DeviceNumTeams;
   if (Max_Teams > DeviceInfo.HardTeamLimit)
     Max_Teams = DeviceInfo.HardTeamLimit;
 
   if (print_kernel_trace & STARTUP_DETAILS) {
     fprintf(stderr, "RTLDeviceInfoTy::Max_Teams: %d\n",
             RTLDeviceInfoTy::Max_Teams);
     fprintf(stderr, "Max_Teams: %d\n", Max_Teams);
     fprintf(stderr, "RTLDeviceInfoTy::Warp_Size: %d\n",
launchVals getLaunchVals(int ConstWGSize, int ExecutionMode, int EnvTeamLimit,
   }
   if (print_kernel_trace & STARTUP_DETAILS) {
     fprintf(stderr, "threadsPerGroup: %d\n", threadsPerGroup);
     fprintf(stderr, "num_groups: %d\n", num_groups);
     fprintf(stderr, "loop_tripcount: %ld\n", loop_tripcount);
   }
   DP("Final %d num_groups and %d threadsPerGroup\n", num_groups,
      threadsPerGroup);
+
+  launchVals res;
+  res.threadsPerGroup = threadsPerGroup;
+  res.num_groups = num_groups;
+  return res;
 }
 
 static uint64_t acquire_available_packet_id(hsa_queue_t *queue) {
   uint64_t packet_id = hsa_queue_add_write_index_relaxed(queue, 1);
   bool full = true;
   while (full) {
     full =
         packet_id >= (queue->size + hsa_queue_load_read_index_scacquire(queue));
int32_t __tgt_rtl_run_target_team_region_locked(
   const uint32_t sgpr_spill_count = KernelInfoEntry.sgpr_spill_count;
   const uint32_t vgpr_spill_count = KernelInfoEntry.vgpr_spill_count;
 
   assert(arg_num == (int)KernelInfoEntry.num_args);
 
   /*
    * Set limit based on ThreadsPerGroup and GroupsPerDevice
    */
-  int num_groups = 0;
-
-  int threadsPerGroup = RTLDeviceInfoTy::Default_WG_Size;
-
-  getLaunchVals(threadsPerGroup, num_groups, KernelInfo->ConstWGSize,
-                KernelInfo->ExecutionMode, DeviceInfo.EnvTeamLimit,
-                DeviceInfo.EnvNumTeams,
+  launchVals LV =
+      getLaunchVals(KernelInfo->ConstWGSize, KernelInfo->ExecutionMode,
+                    DeviceInfo.EnvTeamLimit, DeviceInfo.EnvNumTeams,
                     num_teams,      // From run_region arg
                     thread_limit,   // From run_region arg
                     loop_tripcount, // From run_region arg
-                KernelInfo->device_id);
+                    DeviceInfo.NumTeams[KernelInfo->device_id]);
+  int num_groups = LV.num_groups;
+  int threadsPerGroup = LV.threadsPerGroup;
 
   if (print_kernel_trace >= LAUNCH) {
     // enum modes are SPMD, GENERIC, NONE 0,1,2
     // if doing rtl timing, print to stderr, unless stdout requested.
     bool traceToStdout = print_kernel_trace & (RTL_TO_STDOUT | RTL_TIMING);
     fprintf(traceToStdout ? stdout : stderr,
             "DEVID:%2d SGN:%1d ConstWGSize:%-4d args:%2d teamsXthrds:(%4dX%4d) "
             "reqd:(%4dX%4d) lds_usage:%uB sgpr_count:%u vgpr_count:%u "