[Mlir-commits] [mlir] [mlir][gpu] Add metadata attributes for storing kernel metadata in GPU objects (PR #95292)

[Oleksandr Alex Zinenko]

================
@@ -0,0 +1,87 @@
+//===- Utils.cpp - MLIR ROCDL target utils ----------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This files defines ROCDL target related utility classes and functions.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Target/LLVM/ROCDL/Utils.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+
+#include "llvm/ADT/StringMap.h"
+#include "llvm/Frontend/Offloading/Utility.h"
+
+using namespace mlir;
+using namespace mlir::ROCDL;
+
+std::optional<DenseMap<StringAttr, NamedAttrList>>
+mlir::ROCDL::getAMDHSAKernelsELFMetadata(Builder &builder,
+                                         ArrayRef<char> elfData) {
+  uint16_t elfABIVersion;
+  llvm::StringMap<llvm::offloading::amdgpu::AMDGPUKernelMetaData> kernels;
+  llvm::MemoryBufferRef buffer(StringRef(elfData.data(), elfData.size()),
+                               "buffer");
+  // Get the metadata.
+  llvm::Error error = llvm::offloading::amdgpu::getAMDGPUMetaDataFromImage(
+      buffer, kernels, elfABIVersion);
+  // Return `nullopt` if the metadata couldn't be retrieved.
+  if (!error) {
+    llvm::consumeError(std::move(error));
----------------
ftynse wrote:

I can't parse this. Error converts to `false` when there is no error, making this condition true. It will then consume a non-error (which shouldn't be necessary) and return nullopt. 


https://github.com/llvm/llvm-project/pull/95292