[OpenMP][AMDGPU] More detail in AMDGPU kernel launch info

[jplehr]

Summary

Makes the info that is printed for kernel launches configurable for
different plugins. Adds all machinery to print the detailed launch
info that the current AMD plugin provides and includes e.g. register
spill counts.

The files msgpack.cpp, msgpack.def, and msgpack.h are copied from the old plugin
and are untouched. The contents of UtilitiesHSA.cpp and .h are copied together from
various files from the old plugin. The code was originally written by
Jon Chesterfield. I updated the function and type names visible to the outside, i.e.
in headers, to respect the LLVM conventions.

[jplehr]

Rewrite using what is available upstream to read the ELF metdata for AMDGPU.
Removes the copied-over code from the old plugin.

[jhuber6]

Some initial nits. I wonder if we could move this code somewhere common, since both llvm-readobj and here wants to use it now. I don't know much about the actual msgpack format, maybe @JonChesterfield can chime in.

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

2230–2232

nit: remove braces.

2609–2610

nit. prefer early exit.

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

153–154

Probably not needed. We only support the existing llvm::object::ELF64LE alias, which is exactly what you have here. The Note alias is only used a single time so it's probably fineto just use ELF64LE::Note.

168–170

nit. no braces around single line block here and below

272

272–274

I'd guess that this should be an error. If we fail to parse the code object something is probably wrong.

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp

254–257

[jplehr]

Thanks Joseph.
My impression about the llvm-readobj was that they have quite some machinery around the printing. However, I believe I understood that the actual printing is mainly a conversion of the MsgPackDocument into a YAML and then print that. So I think that there is not too much in common.

[jplehr]

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

272–274

I'm not sure if it should be an error, but if it is not a map, we should probably skip it (and maybe warn or so). I'm not familiar enough with the msgpack format.

[jhuber6]

Also, is it possible to write a test for this?

[jplehr]

Am I right to assume you mainly asking about the code that obtains the metadata for the different kernels in the image?
Without executing a OpenMP target offload program on an AMDGPU, I see mostly one option: Construct a msgpack from a textual representation (e.g., YAML) and query that for the relevant information. This should be somewhat straight forward to do in Google Test. But, it would not automatically react to changes in the metadata produced in the backend. For that, we probably would need to compile and run something on an AMDGPU to test the output.

[jhuber6]

In D144521#4147647, @jplehr wrote:

Am I right to assume you mainly asking about the code that obtains the metadata for the different kernels in the image?
Without executing a OpenMP target offload program on an AMDGPU, I see mostly one option: Construct a msgpack from a textual representation (e.g., YAML) and query that for the relevant information. This should be somewhat straight forward to do in Google Test. But, it would not automatically react to changes in the metadata produced in the backend. For that, we probably would need to compile and run something on an AMDGPU to test the output.

Shouldn't this get populated automatically if you compile an OpenMP program? We don't need to check. We have an existing info.c test that you could copy and make AMD only.

[jhuber6]

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

377

Doesn't this crash if the identifier is not present? Should it be a std::optional or something?

[jplehr]

In D144521#4147666, @jhuber6 wrote:

In D144521#4147647, @jplehr wrote:

Am I right to assume you mainly asking about the code that obtains the metadata for the different kernels in the image?
Without executing a OpenMP target offload program on an AMDGPU, I see mostly one option: Construct a msgpack from a textual representation (e.g., YAML) and query that for the relevant information. This should be somewhat straight forward to do in Google Test. But, it would not automatically react to changes in the metadata produced in the backend. For that, we probably would need to compile and run something on an AMDGPU to test the output.

Shouldn't this get populated automatically if you compile an OpenMP program? We don't need to check. We have an existing info.c test that you could copy and make AMD only.

Yes. Apologies for me not yet being aware of the different test possibilities.
So, if we can compile / run a simple offload program to check if things get populated that would probably be even easier and account for any changes in the metadata.

[jplehr]

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

377

Good catch. Yeah, I think this should be an Expected<T> or so, as, should we not have an entry for the identifier in the map, then we did not find info for it in the image.
This would more or less mean that there is something weird in the image AFAICT as there was no metadata for the kernel with that identifier.
So, we would return an Error in the plugin in that case.

[jdoerfert]

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

224

???

317

pass it in instead of copying it out.

319

no llvm:: needed, right?

[jplehr]

• Removed: curlies for single statement blocks, unncessary llvm:: qualification, leftover empty if, unneccessary using ELFT statements
• Added check if kernel name has associated metadata and changed return type to Expected<T> to allow for error propagation

[jdoerfert]

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

444

I doubt this should be a fatal error. We should recover instead.

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

202

[jplehr]

Addressed more reviewer comments and added test

[jhuber6]

LG overall, some final nits.

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

444–445

Since we're ignoring this, it would be best to use std::optional. But just for reference, you would need to use consumeError here to actually ignore it. You could go ahead and make the KernelInfo variable an optional as well, so we don't need to bother with a zero fill or anything.

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

162

Nit, but these probably don't need to be their own variables since the getters are self-explanatory and used only once.

237

Default constructor should be fine here.

openmp/libomptarget/test/offloading/info.c

1

This should stay generic.

[jplehr]

Addressed nits and rebased. Please see if the changes to the test case are acceptable.

[jhuber6]

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

382

nit.

490

I still think this should be an optional. Then it's a direct assignment from the table lookup. The only thing you need to change is an if when you print it. I'd go ahead and just emit a secondary line for this extra info anyway. That way the first line is still generic between the plugins and this is more info.

[jplehr]

Reworked implementation to always print standard info.
Introduced optional and refactored test.

[jhuber6]

LG, few final nits about formatting that should be addressed.

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

2635

Comment above line.

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

139

these should all use the LLVM naming style.

[jplehr]

Addresssed final nits

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

struct AMDGPUDeviceImageTy : public DeviceImageTy {
 
   /// Get the executable.
   hsa_executable_t getExecutable() const { return Executable; }
 
   /// Find an HSA device symbol by its name on the executable.
   Expected<hsa_executable_symbol_t>
   findDeviceSymbol(GenericDeviceTy &Device, StringRef SymbolName) const;
 
+  /// Get additional info for kernel, e.g., register spill counts
+  std::optional<utils::KernelMetaDataTy>
+  getKernelInfo(StringRef Identifier) const {
+    auto It = KernelInfoMap.find(Identifier);
+
+    if (It == KernelInfoMap.end())
+      return {};
+
+    return It->second;
+  }
+
 private:
   /// The exectuable loaded on the agent.
   hsa_executable_t Executable;
   hsa_code_object_t CodeObject;
+  StringMap<utils::KernelMetaDataTy> KernelInfoMap;
 };
 
 /// Class implementing the AMDGPU kernel functionalities which derives from the
 /// generic kernel class.
 struct AMDGPUKernelTy : public GenericKernelTy {
   /// Create an AMDGPU kernel with a name and an execution mode.
   AMDGPUKernelTy(const char *Name, OMPTgtExecModeFlags ExecutionMode)
       : GenericKernelTy(Name, ExecutionMode),
Error initImpl(GenericDeviceTy &Device, DeviceImageTy &Image) override {
 
     // Make sure it is a kernel symbol.
     if (SymbolType != HSA_SYMBOL_KIND_KERNEL)
       return Plugin::error("Symbol %s is not a kernel function");
 
     // TODO: Read the kernel descriptor for the max threads per block. May be
     // read from the image.
 
+    // Get additional kernel info read from image
+    KernelInfo = AMDImage.getKernelInfo(getName());
+    if (!KernelInfo.has_value())
+      INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device.getDeviceId(),
+           "Could not read extra information for kernel %s.", getName());
+
     return Plugin::success();
   }
 
   /// Launch the AMDGPU kernel function.
   Error launchImpl(GenericDeviceTy &GenericDevice, uint32_t NumThreads,
                    uint64_t NumBlocks, KernelArgsTy &KernelArgs, void *Args,
                    AsyncInfoWrapperTy &AsyncInfoWrapper) const override;
 
+  /// Print more elaborate kernel launch info for AMDGPU
+  Error printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
+                               KernelArgsTy &KernelArgs, uint32_t NumThreads,
+                               uint64_t NumBlocks) const override;
+
   /// The default number of blocks is common to the whole device.
   uint32_t getDefaultNumBlocks(GenericDeviceTy &GenericDevice) const override {
     return GenericDevice.getDefaultNumBlocks();
   }
 
   /// The default number of threads is common to the whole device.
   uint32_t getDefaultNumThreads(GenericDeviceTy &GenericDevice) const override {
     return GenericDevice.getDefaultNumThreads();
private:
 
   /// The args, group and private segments sizes required by a kernel instance.
   uint32_t ArgsSize;
   uint32_t GroupSize;
   uint32_t PrivateSize;
 
   /// The size of implicit kernel arguments.
   const uint32_t ImplicitArgsSize;
+
+  /// Additional Info for the AMD GPU Kernel
+  std::optional<utils::KernelMetaDataTy> KernelInfo;
 };
 
 /// Class representing an HSA signal. Signals are used to define dependencies
 /// between asynchronous operations: kernel launches and memory transfers.
 struct AMDGPUSignalTy {
   /// Create an empty signal.
   AMDGPUSignalTy() : Signal({0}), UseCount() {}
   AMDGPUSignalTy(AMDGPUDeviceTy &Device) : Signal({0}), UseCount() {}
Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) {
   uint32_t Result;
   Status = hsa_executable_validate(Executable, &Result);
   if (auto Err = Plugin::check(Status, "Error in hsa_executable_validate: %s"))
     return Err;
 
   if (Result)
     return Plugin::error("Loaded HSA executable does not validate");
 
+  if (auto Err =
+          utils::readAMDGPUMetaDataFromImage(getMemoryBuffer(), KernelInfoMap))
+    return Err;
+
   return Plugin::success();
 }
 
 Expected<hsa_executable_symbol_t>
 AMDGPUDeviceImageTy::findDeviceSymbol(GenericDeviceTy &Device,
                                       StringRef SymbolName) const {
 
   AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(Device);
Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
   AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(GenericDevice);
   AMDGPUStreamTy &Stream = AMDGPUDevice.getStream(AsyncInfoWrapper);
 
   // Push the kernel launch into the stream.
   return Stream.pushKernelLaunch(*this, AllArgs, NumThreads, NumBlocks,
                                  GroupSize, ArgsMemoryManager);
 }
 
+Error AMDGPUKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
+                                             KernelArgsTy &KernelArgs,
+                                             uint32_t NumThreads,
+                                             uint64_t NumBlocks) const {
+  // Only do all this when the output is requested
+  if (!(getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))
+    return Plugin::success();
+
+  // 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
+
+  // Kernel Arguments Info
+  auto ArgNum = KernelArgs.NumArgs;
+  auto LoopTripCount = KernelArgs.Tripcount;
+
+  // Details for AMDGPU kernels
+  auto GroupSegmentSize = (*KernelInfo).GroupSegmentList;
+  auto SGPRCount = (*KernelInfo).SGPRCount;
+  auto VGPRCount = (*KernelInfo).VGPRCount;
+  auto SGPRSpillCount = (*KernelInfo).SGPRSpillCount;
+  auto VGPRSpillCount = (*KernelInfo).VGPRSpillCount;
+
+  // TODO set correctly once host services available
+  auto HostCallRequired = false;
+  INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),
+       "SGN:%s ConstWGSize:%d args:%d 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 rpc:%d n:%s\n",
+       getExecutionModeName(), ConstWGSize, ArgNum, NumGroups, ThreadsPerGroup,
+       NumTeams, ThreadLimit, GroupSegmentSize, SGPRCount, VGPRCount,
+       SGPRSpillCount, VGPRSpillCount, LoopTripCount, HostCallRequired,
+       getName());
+
+  return Plugin::success();
+}
+
 GenericPluginTy *Plugin::createPlugin() { return new AMDGPUPluginTy(); }
 
 GenericDeviceTy *Plugin::createDevice(int32_t DeviceId, int32_t NumDevices) {
   AMDGPUPluginTy &Plugin = get<AMDGPUPluginTy &>();
   return new AMDGPUDeviceTy(DeviceId, NumDevices, Plugin.getHostDevice(),
                             Plugin.getKernelAgent(DeviceId));
 }
 

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

 
 #include <cstdint>
 
 #include "Debug.h"
 #include "omptarget.h"
 
 #include "llvm/ADT/StringMap.h"
 #include "llvm/ADT/StringRef.h"
+#include "llvm/Support/Error.h"
+
+#include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
+#include "llvm/BinaryFormat/ELF.h"
+#include "llvm/BinaryFormat/MsgPackDocument.h"
+#include "llvm/Support/MemoryBufferRef.h"
+
+#include "llvm/Support/YAMLTraits.h"
 
 namespace llvm {
 namespace omp {
 namespace target {
 namespace plugin {
 namespace utils {
 
 // The implicit arguments of AMDGPU kernels.
bool isImageCompatibleWithEnv(const __tgt_image_info *Info,
   //   - or, the feature is missing from Image's features map i.e. it is
   //   set to ANY
   DP("Compatible: Target IDs are compatible \t[Image: %s]\t:\t[Env: %s]\n",
      ImageTargetID.data(), EnvTargetID.data());
 
   return true;
 }
 
+struct KernelMetaDataTy {
+  uint64_t KernelObject;
+  uint32_t GroupSegmentList;
+  uint32_t PrivateSegmentSize;
+  uint32_t SGPRCount;
+  uint32_t VGPRCount;
+  uint32_t SGPRSpillCount;
+  uint32_t VGPRSpillCount;
+  uint32_t KernelSegmentSize;
+  uint32_t ExplicitArgumentCount;
+  uint32_t ImplicitArgumentCount;
+};
+namespace {
+
+/// Reads the AMDGPU specific per-kernel-metadata from an image.
+class KernelInfoReader {
+public:
+  KernelInfoReader(StringMap<KernelMetaDataTy> &KIM) : KernelInfoMap(KIM) {}
+
+  /// Process ELF note to read AMDGPU metadata from respective information
+  /// fields.
+  Error processNote(const object::ELF64LE::Note &Note) {
+    if (Note.getName() != "AMDGPU")
+      return Error::success(); // We are not interested in other things
+
+    assert(Note.getType() == ELF::NT_AMDGPU_METADATA &&
+           "Parse AMDGPU MetaData");
+    auto Desc = Note.getDesc();
+    StringRef MsgPackString =
+        StringRef(reinterpret_cast<const char *>(Desc.data()), Desc.size());
+    msgpack::Document MsgPackDoc;
+    if (!MsgPackDoc.readFromBlob(MsgPackString, /*Multi=*/false))
+      return Error::success();
+
+    AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
+    if (!Verifier.verify(MsgPackDoc.getRoot()))
+      return Error::success();
+
+    auto RootMap = MsgPackDoc.getRoot().getMap(true);
+
+    if (auto Err = iterateAMDKernels(RootMap))
+      return Err;
+
+    return Error::success();
+  }
+
+private:
+  /// Extracts the relevant information via simple string look-up in the msgpack
+  /// document elements.
+  Error extractKernelData(msgpack::MapDocNode::MapTy::value_type V,
+                          std::string &KernelName,
+                          KernelMetaDataTy &KernelData) {
+    if (!V.first.isString())
+      return Error::success();
+
+    const auto isKey = [](const msgpack::DocNode &DK, StringRef SK) {
+      return DK.getString() == SK;
+    };
+
+    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();
+    }
+
+    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");
+    if (Res == MDN.end())
+      return createStringError(inconvertibleErrorCode(),
+                               "Could not find amdhsa.kernels key");
+
+    auto Pair = *Res;
+    assert(Pair.second.isArray() &&
+           "AMDGPU kernel entries are arrays of entries");
+
+    return Pair.second.getArray();
+  }
+
+  /// Iterate all entries for one "amdhsa.kernels" entry. Each entry is a
+  /// MapDocNode that either maps a string to a single value (most of them) or
+  /// to another array of things. Currently, we only handle the case that maps
+  /// to scalar value.
+  Error generateKernelInfo(msgpack::ArrayDocNode::ArrayTy::iterator It) {
+    KernelMetaDataTy KernelData;
+    std::string KernelName;
+    auto Entry = (*It).getMap();
+    for (auto MI = Entry.begin(), E = Entry.end(); MI != E; ++MI)
+      if (auto Err = extractKernelData(*MI, KernelName, KernelData))
+        return Err;
+
+    KernelInfoMap.insert({KernelName, KernelData});
+    return Error::success();
+  }
+
+  /// Go over the list of AMD kernels in the "amdhsa.kernels" entry
+  Error iterateAMDKernels(msgpack::MapDocNode &MDN) {
+    auto KernelsOrErr = getAMDKernelsArray(MDN);
+    if (auto Err = KernelsOrErr.takeError())
+      return Err;
+
+    auto KernelsArr = *KernelsOrErr;
+    for (auto It = KernelsArr.begin(), E = KernelsArr.end(); It != E; ++It) {
+      if (!It->isMap())
+        continue; // we expect <key,value> pairs
+
+      // Obtain the value for the different entries. Each array entry is a
+      // MapDocNode
+      if (auto Err = generateKernelInfo(It))
+        return Err;
+    }
+    return Error::success();
+  }
+
+  // Kernel names are the keys
+  StringMap<KernelMetaDataTy> &KernelInfoMap;
+};
+} // namespace
+
+/// Reads the AMDGPU specific metadata from the ELF file and propagates the
+/// KernelInfoMap
+Error readAMDGPUMetaDataFromImage(MemoryBufferRef MemBuffer,
+                                  StringMap<KernelMetaDataTy> &KernelInfoMap) {
+  Error Err = Error::success(); // Used later as out-parameter
+
+  auto ELFOrError = object::ELF64LEFile::create(MemBuffer.getBuffer());
+  if (auto Err = ELFOrError.takeError())
+    return Err;
+
+  const object::ELF64LEFile ELFObj = ELFOrError.get();
+  ArrayRef<object::ELF64LE::Shdr> Sections = cantFail(ELFObj.sections());
+  KernelInfoReader Reader(KernelInfoMap);
+  for (const auto &S : Sections) {
+    if (S.sh_type != ELF::SHT_NOTE)
+      continue;
+
+    for (const auto N : ELFObj.notes(S, Err)) {
+      if (Err)
+        return 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

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h

static bool isValidExecutionMode(OMPTgtExecModeFlags ExecutionMode) {
     case OMP_TGT_EXEC_MODE_SPMD:
     case OMP_TGT_EXEC_MODE_GENERIC:
     case OMP_TGT_EXEC_MODE_GENERIC_SPMD:
       return true;
     }
     return false;
   }
 
+protected:
+  /// Get the execution mode name of the kernel.
+  const char *getExecutionModeName() const {
+    switch (ExecutionMode) {
+    case OMP_TGT_EXEC_MODE_SPMD:
+      return "SPMD";
+    case OMP_TGT_EXEC_MODE_GENERIC:
+      return "Generic";
+    case OMP_TGT_EXEC_MODE_GENERIC_SPMD:
+      return "Generic-SPMD";
+    }
+    llvm_unreachable("Unknown execution mode!");
+  }
+
+  /// Prints generic kernel launch information.
+  Error printLaunchInfo(GenericDeviceTy &GenericDevice,
+                        KernelArgsTy &KernelArgs, uint32_t NumThreads,
+                        uint64_t NumBlocks) const;
+
+  /// Prints plugin-specific kernel launch information after generic kernel
+  /// launch information
+  virtual Error printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
+                                       KernelArgsTy &KernelArgs,
+                                       uint32_t NumThreads,
+                                       uint64_t NumBlocks) const;
+
 private:
   /// Prepare the arguments before launching the kernel.
   void *prepareArgs(GenericDeviceTy &GenericDevice, void **ArgPtrs,
                     ptrdiff_t *ArgOffsets, int32_t NumArgs,
                     llvm::SmallVectorImpl<void *> &Args,
                     llvm::SmallVectorImpl<void *> &Ptrs,
                     AsyncInfoWrapperTy &AsyncInfoWrapper) const;
 
private:
   bool isGenericSPMDMode() const {
     return ExecutionMode == OMP_TGT_EXEC_MODE_GENERIC_SPMD;
   }
   bool isGenericMode() const {
     return ExecutionMode == OMP_TGT_EXEC_MODE_GENERIC;
   }
   bool isSPMDMode() const { return ExecutionMode == OMP_TGT_EXEC_MODE_SPMD; }
 
-  /// Get the execution mode name of the kernel.
-  const char *getExecutionModeName() const {
-    switch (ExecutionMode) {
-    case OMP_TGT_EXEC_MODE_SPMD:
-      return "SPMD";
-    case OMP_TGT_EXEC_MODE_GENERIC:
-      return "Generic";
-    case OMP_TGT_EXEC_MODE_GENERIC_SPMD:
-      return "Generic-SPMD";
-    }
-    llvm_unreachable("Unknown execution mode!");
-  }
-
   /// The kernel name.
   const char *Name;
 
   /// The execution flags of the kernel.
   OMPTgtExecModeFlags ExecutionMode;
 
 protected:
   /// The preferred number of threads to run the kernel.

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp

Error GenericKernelTy::init(GenericDeviceTy &GenericDevice,
                             DeviceImageTy &Image) {
   PreferredNumThreads = getDefaultNumThreads(GenericDevice);
 
   MaxNumThreads = GenericDevice.getThreadLimit();
 
   return initImpl(GenericDevice, Image);
 }
 
+Error GenericKernelTy::printLaunchInfo(GenericDeviceTy &GenericDevice,
+                                       KernelArgsTy &KernelArgs,
+                                       uint32_t NumThreads,
+                                       uint64_t NumBlocks) const {
+  INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),
+       "Launching kernel %s with %" PRIu64
+       " blocks and %d threads in %s mode\n",
+       getName(), NumBlocks, NumThreads, getExecutionModeName());
+  return printLaunchInfoDetails(GenericDevice, KernelArgs, NumThreads,
+                                NumBlocks);
+}
+
+Error GenericKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
+                                              KernelArgsTy &KernelArgs,
+                                              uint32_t NumThreads,
+                                              uint64_t NumBlocks) const {
+  return Plugin::success();
+}
+
 Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
                               ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,
                               AsyncInfoWrapperTy &AsyncInfoWrapper) const {
   llvm::SmallVector<void *, 16> Args;
   llvm::SmallVector<void *, 16> Ptrs;
 
   void *KernelArgsPtr =
       prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs, Args,
                   Ptrs, AsyncInfoWrapper);
 
   uint32_t NumThreads = getNumThreads(GenericDevice, KernelArgs.ThreadLimit);
   uint64_t NumBlocks = getNumBlocks(GenericDevice, KernelArgs.NumTeams,
                                     KernelArgs.Tripcount, NumThreads);
 
-  INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),
-       "Launching kernel %s with %" PRIu64
-       " blocks and %d threads in %s mode\n",
-       getName(), NumBlocks, NumThreads, getExecutionModeName());
+  if (auto Err =
+          printLaunchInfo(GenericDevice, KernelArgs, NumThreads, NumBlocks))
+    return Err;
 
   return launchImpl(GenericDevice, NumThreads, NumBlocks, KernelArgs,
                     KernelArgsPtr, AsyncInfoWrapper);
 }
 
 void *GenericKernelTy::prepareArgs(GenericDeviceTy &GenericDevice,
                                    void **ArgPtrs, ptrdiff_t *ArgOffsets,
                                    int32_t NumArgs,
                                    llvm::SmallVectorImpl<void *> &Args,

openmp/libomptarget/test/offloading/info.c

 // RUN: %libomptarget-compile-generic \
 // RUN:     -gline-tables-only -fopenmp-extensions
 // RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | \
-// RUN:   %fcheck-generic -allow-empty -check-prefix=INFO
+// RUN:   %fcheck-generic -allow-empty -check-prefixes=INFO
+// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-amdgcn-amd-amdhsa 2>&1 | \
+// RUN:   %fcheck-amdgcn-amd-amdhsa -allow-empty -check-prefixes=INFO,AMDGPU
 
 #include <omp.h>
 #include <stdio.h>
 
 #define N 64
 
 #pragma omp declare target
 int global;

 // INFO: info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
 // INFO: info: Host Ptr           Target Ptr         Size (B) DynRefCount HoldRefCount Declaration
 // INFO: info: {{.*}}             {{.*}}             256      1           0            C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: info: {{.*}}             {{.*}}             256      0           1            B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: info: {{.*}}             {{.*}}             256      1           0            A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: info: Entering OpenMP kernel at info.c:{{[0-9]+}}:{{[0-9]+}} with 1 arguments:
 // INFO: info: firstprivate(val)[4]
 // INFO: info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
+// AMDGPU: AMDGPU device {{[0-9]}} info: SGN:Generic ConstWGSize:{{[0-9]+}} args:{{[0-9]}} teamsXthrds:({{   [0-9]+}}X {{[0-9]+}}) reqd:(   {{[0-9]+}}X   {{[0-9]+}}) lds_usage:{{[0-9]+}}B sgpr_count:{{[0-9]+}} vgpr_count:{{[0-9]+}} sgpr_spill_count:{{[0-9]+}} vgpr_spill_count:{{[0-9]+}} tripcount:{{[0-9]+}} rpc:0 n:__omp_offloading_{{.*}}main{{.*}}
 // INFO: info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
 // INFO: info: Host Ptr           Target Ptr         Size (B) DynRefCount HoldRefCount Declaration
 // INFO: info: {{.*}}             {{.*}}             256      1           0            C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: info: {{.*}}             {{.*}}             256      0           1            B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: info: {{.*}}             {{.*}}             256      1           0            A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: info: Exiting OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
 // INFO: info: alloc(A[0:64])[256]
 // INFO: info: tofrom(B[0:64])[256]