address reviewer comments

Author: fabianmcg

mlir/include/mlir/Dialect/GPU/IR/CompilationAttrs.td

@@ -120,10 +120,10 @@ def GPU_KernelTableAttr : GPU_Attr<"KernelTable", "kernel_table"> {
 
     Examples:
     ```mlir
-      // Empty table
+      // Empty table.
       #gpu.kernel_table<>
 
-      // Table with a single kernel
+      // Table with a single kernel.
       #gpu.kernel_table<[#gpu.kernel<kernel0, () -> () >]>
 
       // Table with multiple kernels.

mlir/lib/Dialect/GPU/IR/GPUDialect.cpp

@@ -2274,15 +2274,13 @@ KernelTableAttr::verify(function_ref<InFlightDiagnostic()> emitError,
 }
 
 KernelAttr KernelTableAttr::lookup(StringRef key) const {
-  std::pair<ArrayRef<KernelAttr>::iterator, bool> it =
-      impl::findAttrSorted(begin(), end(), key);
-  return it.second ? *it.first : KernelAttr();
+  auto [iterator, found] = impl::findAttrSorted(begin(), end(), key);
+  return found ? *iterator : KernelAttr();
 }
 
 KernelAttr KernelTableAttr::lookup(StringAttr key) const {
-  std::pair<ArrayRef<KernelAttr>::iterator, bool> it =
-      impl::findAttrSorted(begin(), end(), key);
-  return it.second ? *it.first : KernelAttr();
+  auto [iterator, found] = impl::findAttrSorted(begin(), end(), key);
+  return found ? *iterator : KernelAttr();
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Target/LLVM/ROCDL/Target.cpp

@@ -506,12 +506,15 @@ ROCDLTargetAttrImpl::createObject(Attribute attribute, Operation *module,
   gpu::CompilationTarget format = options.getCompilationTarget();
   // If format is `fatbin` transform it to binary as `fatbin` is not yet
   // supported.
-  if (format > gpu::CompilationTarget::Binary)
+  gpu::KernelTableAttr kernels;
+  if (format > gpu::CompilationTarget::Binary) {
     format = gpu::CompilationTarget::Binary;
+    kernels = ROCDL::getKernelMetadata(module, object);
+  }
   DictionaryAttr properties{};
   Builder builder(attribute.getContext());
   StringAttr objectStr =
       builder.getStringAttr(StringRef(object.data(), object.size()));
   return builder.getAttr<gpu::ObjectAttr>(attribute, format, objectStr,
-                                          properties, nullptr);
+                                          properties, kernels);
 }

mlir/lib/Target/LLVM/ROCDL/Utils.cpp

@@ -31,7 +31,7 @@ mlir::ROCDL::getAMDHSAKernelsELFMetadata(Builder &builder,
   llvm::Error error = llvm::offloading::amdgpu::getAMDGPUMetaDataFromImage(
       buffer, kernels, elfABIVersion);
   // Return `nullopt` if the metadata couldn't be retrieved.
-  if (!error) {
+  if (error) {
     llvm::consumeError(std::move(error));
     return std::nullopt;
   }

mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp

@@ -116,7 +116,7 @@ TargetAttrImpl::createObject(Attribute attribute, Operation *module,
       module->getContext(), attribute, gpu::CompilationTarget::Offload,
       StringAttr::get(module->getContext(),
                       StringRef(object.data(), object.size())),
-      module->getAttrDictionary());
+      module->getAttrDictionary(), nullptr);
 }
 
 TEST_F(MLIRTargetLLVM, SKIP_WITHOUT_NATIVE(TargetAttrAPI)) {