[libomptarget][nfc] Move grid size computation

[JonChesterfield]

Summary

Change getLaunchVals to return the integers used for launch

[JonChesterfield]

There are some review comments for this in D105239

[pdhaliwal]

Looks good! Thanks.

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

 }
 
 // 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.
 struct launchVals {
-  int threadsPerGroup;
-  int num_groups;
+  int WorkgroupSize;
+  int GridSize;
 };
 
 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;
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;
+  res.WorkgroupSize = threadsPerGroup;
+  res.GridSize = threadsPerGroup * 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 =
int32_t __tgt_rtl_run_target_team_region_locked(
    */
   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
                     DeviceInfo.NumTeams[KernelInfo->device_id]);
-  int num_groups = LV.num_groups;
-  int threadsPerGroup = LV.threadsPerGroup;
+  const int GridSize = LV.GridSize;
+  const int WorkgroupSize = LV.WorkgroupSize;
 
   if (print_kernel_trace >= LAUNCH) {
+    int num_groups = GridSize / WorkgroupSize;
     // 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 "
             "sgpr_spill_count:%u vgpr_spill_count:%u tripcount:%lu n:%s\n",
             device_id, KernelInfo->ExecutionMode, KernelInfo->ConstWGSize,
-            arg_num, num_groups, threadsPerGroup, num_teams, thread_limit,
+            arg_num, num_groups, WorkgroupSize, num_teams, thread_limit,
             group_segment_size, sgpr_count, vgpr_count, sgpr_spill_count,
             vgpr_spill_count, loop_tripcount, KernelInfo->Name);
   }
 
   // Run on the device.
   {
     hsa_queue_t *queue = DeviceInfo.HSAQueues[device_id];
     if (!queue) {
       return OFFLOAD_FAIL;
     }
     uint64_t packet_id = acquire_available_packet_id(queue);
 
     const uint32_t mask = queue->size - 1; // size is a power of 2
     hsa_kernel_dispatch_packet_t *packet =
         (hsa_kernel_dispatch_packet_t *)queue->base_address +
         (packet_id & mask);
 
     // packet->header is written last
     packet->setup = UINT16_C(1) << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
-    packet->workgroup_size_x = threadsPerGroup;
+    packet->workgroup_size_x = WorkgroupSize;
     packet->workgroup_size_y = 1;
     packet->workgroup_size_z = 1;
     packet->reserved0 = 0;
-    packet->grid_size_x = num_groups * threadsPerGroup;
+    packet->grid_size_x = GridSize;
     packet->grid_size_y = 1;
     packet->grid_size_z = 1;
     packet->private_segment_size = KernelInfoEntry.private_segment_size;
     packet->group_segment_size = KernelInfoEntry.group_segment_size;
     packet->kernel_object = KernelInfoEntry.kernel_object;
     packet->kernarg_address = 0;     // use the block allocator
     packet->reserved2 = 0;           // atmi writes id_ here
     packet->completion_signal = {0}; // may want a pool of signals