[OpenMP][libomptarget][AMDGPU] Update print launch info

[jplehr]

Summary

Clean up for the AMD-specific kernel launch info in the NextGen Plugins.

• Fixes a mistake introduced with the initial commit that added printing of an AMD-only property.
• Removes another AMD-only property (not clear on upstream status)
• Adds some more comment to what info is printed.

[jplehr]

Updated test and removed two unused variables.

[jdoerfert]

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

204

Can we assert that there are at most three? Also, add a message to each assert.

[jplehr]

Adds assert and assert messages.
It felt safer to assert exactly three elements to be consistent with the name of the lambda.

[jdoerfert]

LG, one more suggestion

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

2648

I would try to shorten this, like # for count, maybe SPGR/VPGR and then <N>/<M>, etc. But this is already better than what we had.

[jplehr]

Followed Johannes suggestion and rebased

[jplehr]

Actually followed Johannes' suggestion.

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

Error AMDGPUKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
 
   // We don't have data to print additional info, but no hard error
   if (!KernelInfo.has_value())
     return Plugin::success();
 
   // General Info
   auto ConstWGSize = getDefaultNumThreads(GenericDevice);
   auto NumGroups = NumBlocks;
-  auto ThreadsPerGroup = getDefaultNumThreads(GenericDevice);
-  auto NumTeams = KernelArgs.NumTeams[0];       // Only first dimension
-  auto ThreadLimit = KernelArgs.ThreadLimit[0]; // Only first dimension
+  auto ThreadsPerGroup = NumThreads;
 
   // Kernel Arguments Info
   auto ArgNum = KernelArgs.NumArgs;
   auto LoopTripCount = KernelArgs.Tripcount;
+  auto ThreadLimitX = KernelArgs.ThreadLimit[0];
 
-  // Details for AMDGPU kernels
+  // Details for AMDGPU kernels (read from image)
+  // https://www.llvm.org/docs/AMDGPUUsage.html#code-object-v4-metadata
   auto GroupSegmentSize = (*KernelInfo).GroupSegmentList;
   auto SGPRCount = (*KernelInfo).SGPRCount;
   auto VGPRCount = (*KernelInfo).VGPRCount;
   auto SGPRSpillCount = (*KernelInfo).SGPRSpillCount;
   auto VGPRSpillCount = (*KernelInfo).VGPRSpillCount;
+  auto MaxFlatWorkgroupSize = (*KernelInfo).MaxFlatWorkgroupSize;
 
-  // TODO set correctly once host services available
-  auto HostCallRequired = false;
+  // Prints additional launch info that contains the following.
+  // Num Args: The number of kernel arguments
+  // Teams x Thrds: The number of teams and the number of threads actually
+  // running.
+  // MaxFlatWorkgroupSize: Maximum flat work-group size supported by the
+  // kernel in work-items
+  // LDS Usage: Amount of bytes used in LDS storage
+  // S/VGPR Count: the number of S/V GPRs occupied by the kernel
+  // S/VGPR Spill Count: how many S/VGPRs are spilled by the kernel
+  // Tripcount: loop tripcount for the kernel
   INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),
-       "SGN:%s ConstWGSize:%d args:%d teamsXthrds:(%4luX%4d) "
-       "reqd:(%4dX%4d) lds_usage:%uB sgpr_count:%u vgpr_count:%u "
-       "sgpr_spill_count:%u vgpr_spill_count:%u tripcount:%lu rpc:%d n:%s\n",
-       getExecutionModeName(), ConstWGSize, ArgNum, NumGroups, ThreadsPerGroup,
-       NumTeams, ThreadLimit, GroupSegmentSize, SGPRCount, VGPRCount,
-       SGPRSpillCount, VGPRSpillCount, LoopTripCount, HostCallRequired,
-       getName());
+       "Num Args: %d Teams x Thrds: %4lux%4u (MaxFlatWorkGroupSize: %u) LDS "
+       "Usage: %uB SGPR Count: %u VGPR Count:%u "
+       "SGPR Spill Count: %u VGPR Spill Count: %u Tripcount: %lu\n",
+       ArgNum, NumGroups, ThreadsPerGroup, MaxFlatWorkgroupSize,
+       GroupSegmentSize, SGPRCount, VGPRCount, SGPRSpillCount, VGPRSpillCount,
+       LoopTripCount);
 
   return Plugin::success();
 }
 
 GenericPluginTy *Plugin::createPlugin() { return new AMDGPUPluginTy(); }
 
 GenericDeviceTy *Plugin::createDevice(int32_t DeviceId, int32_t NumDevices) {
   AMDGPUPluginTy &Plugin = get<AMDGPUPluginTy &>();

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

struct KernelMetaDataTy {
   uint32_t PrivateSegmentSize;
   uint32_t SGPRCount;
   uint32_t VGPRCount;
   uint32_t SGPRSpillCount;
   uint32_t VGPRSpillCount;
   uint32_t KernelSegmentSize;
   uint32_t ExplicitArgumentCount;
   uint32_t ImplicitArgumentCount;
+  uint32_t RequestedWorkgroupSize[3];
+  uint32_t WorkgroupSizeHint[3];
+  uint32_t WavefronSize;
+  uint32_t MaxFlatWorkgroupSize;
 };
 namespace {
 
 /// Reads the AMDGPU specific per-kernel-metadata from an image.
 class KernelInfoReader {
 public:
   KernelInfoReader(StringMap<KernelMetaDataTy> &KIM) : KernelInfoMap(KIM) {}
 
Error extractKernelData(msgpack::MapDocNode::MapTy::value_type V,
                           KernelMetaDataTy &KernelData) {
     if (!V.first.isString())
       return Error::success();
 
     const auto isKey = [](const msgpack::DocNode &DK, StringRef SK) {
       return DK.getString() == SK;
     };
 
+    const auto getSequenceOfThreeInts = [](msgpack::DocNode &DN,
+                                           uint32_t *Vals) {
+      assert(DN.isArray());
+      auto DNA = DN.getArray();
+
+      int i = 0;
+      for (auto DNABegin = DNA.begin(), DNAEnd = DNA.end(); DNABegin != DNAEnd;
+           ++DNABegin) {
+        Vals[i++] = DNABegin->getUInt();
+      }
+    };
+
     if (isKey(V.first, ".name")) {
       KernelName = V.second.toString();
     } else if (isKey(V.first, ".sgpr_count")) {
       KernelData.SGPRCount = V.second.getUInt();
     } else if (isKey(V.first, ".sgpr_spill_count")) {
       KernelData.SGPRSpillCount = V.second.getUInt();
     } else if (isKey(V.first, ".vgpr_count")) {
       KernelData.VGPRCount = V.second.getUInt();
     } else if (isKey(V.first, ".vgpr_spill_count")) {
       KernelData.VGPRSpillCount = V.second.getUInt();
     } else if (isKey(V.first, ".private_segment_fixed_size")) {
       KernelData.PrivateSegmentSize = V.second.getUInt();
     } else if (isKey(V.first, ".group_segement_fixed_size")) {
       KernelData.GroupSegmentList = V.second.getUInt();
+    } else if (isKey(V.first, ".reqd_workgroup_size")) {
+      getSequenceOfThreeInts(V.second, KernelData.RequestedWorkgroupSize);
+    } else if (isKey(V.first, ".workgroup_size_hint")) {
+      getSequenceOfThreeInts(V.second, KernelData.WorkgroupSizeHint);
+    } else if (isKey(V.first, ".wavefront_size")) {
+      KernelData.WavefronSize = V.second.getUInt();
+    } else if (isKey(V.first, ".max_flat_workgroup_size")) {
+      KernelData.MaxFlatWorkgroupSize = V.second.getUInt();
     }
 
     return Error::success();
   }
 
   /// Get the "amdhsa.kernels" element from the msgpack Document
   Expected<msgpack::ArrayDocNode> getAMDKernelsArray(msgpack::MapDocNode &MDN) {
     auto Res = MDN.find("amdhsa.kernels");
for (const auto N : ELFObj.notes(S, Err)) {
       // Fills the KernelInfoTabel entries in the reader
       if ((Err = Reader.processNote(N)))
         return Err;
     }
   }
 
   return Error::success();
 }
+
 } // namespace utils
 } // namespace plugin
 } // namespace target
 } // namespace omp
 } // namespace llvm