[Mlir-commits] [mlir] [MLIR][Remark] Import LLVM optimization remarks into the remark engine (PR #222300)
Guray Ozen
llvmlistbot at llvm.org
Wed Sep 9 04:56:51 PDT 2026
https://github.com/grypp updated https://github.com/llvm/llvm-project/pull/222300
>From 85d56060cd808a58768052cf71e3078ac71f6de3 Mon Sep 17 00:00:00 2001
From: Guray Ozen <gozen at nvidia.com>
Date: Wed, 9 Sep 2026 11:56:14 +0200
Subject: [PATCH 1/2] [MLIR][Remark] Import LLVM optimization remarks into the
remark engine
Collect LLVM side of remarks with MLIR's remark infra.
`ModuleToObject::run` installs `remark::LLVMToMLIRDiagnosticHandler` on
its LLVM context through the new `setupLLVMContext` hook. The handler
answers LLVM's remark-enable queries from the `RemarkEngine` filters, so
no remark is built when none is requested, and imports enabled remarks
as MLIR remarks: the category is the LLVM pass name prefixed with
`llvm-` (`--remarks-filter=llvm-.*` selects all LLVM remarks,
`llvm-inline` a single pass), the remark name is kept, the message is
stored under `Remark` and structured arguments become metrics. Remarks
are attached to their debug location if present, otherwise to the MLIR
symbol named after the LLVM function. Remarks not requested by MLIR are
left to LLVM's default handling, so `-pass-remarks` keeps working, and
other LLVM diagnostics become MLIR diagnostics.
`remark::importLLVMRemarks` replays a serialized LLVM remark file (YAML
or bitstream) through the same conversion.
Co-Authored-By: Claude Fable 5.1 <noreply at anthropic.com>
---
mlir/docs/Remarks.md | 35 ++
mlir/include/mlir/Remark/LLVMRemarkImport.h | 70 ++++
.../include/mlir/Target/LLVM/ModuleToObject.h | 6 +
mlir/lib/Remark/CMakeLists.txt | 1 +
mlir/lib/Remark/LLVMRemarkImport.cpp | 315 ++++++++++++++
mlir/lib/Target/LLVM/CMakeLists.txt | 1 +
mlir/lib/Target/LLVM/ModuleToObject.cpp | 8 +
.../GPU/module-to-binary-nvvm-remarks.mlir | 30 ++
.../TestOneShotModuleBufferize.cpp | 1 +
mlir/test/lib/Dialect/Test/TestOpDefs.cpp | 1 +
mlir/unittests/CMakeLists.txt | 1 +
mlir/unittests/Remark/CMakeLists.txt | 9 +
.../unittests/Remark/LLVMRemarkImportTest.cpp | 395 ++++++++++++++++++
.../Target/LLVM/SerializeNVVMTarget.cpp | 81 ++++
14 files changed, 954 insertions(+)
create mode 100644 mlir/include/mlir/Remark/LLVMRemarkImport.h
create mode 100644 mlir/lib/Remark/LLVMRemarkImport.cpp
create mode 100644 mlir/test/Dialect/GPU/module-to-binary-nvvm-remarks.mlir
create mode 100644 mlir/unittests/Remark/CMakeLists.txt
create mode 100644 mlir/unittests/Remark/LLVMRemarkImportTest.cpp
diff --git a/mlir/docs/Remarks.md b/mlir/docs/Remarks.md
index 43f5d4da12247..c70e20cd44ec4 100644
--- a/mlir/docs/Remarks.md
+++ b/mlir/docs/Remarks.md
@@ -326,3 +326,38 @@ public:
auto streamer = std::make_unique<MyStreamer>();
remark::enableOptimizationRemarks(context, std::move(streamer), cats);
```
+
+***
+
+## Importing LLVM Remarks
+
+LLVM passes emit their own
+[optimization remarks](https://llvm.org/docs/Remarks.html) while MLIR compiles
+a module with LLVM, for example in `gpu-module-to-binary`. `ModuleToObject`
+installs a `remark::LLVMToMLIRDiagnosticHandler` on the `llvm::LLVMContext` it
+compiles with, which imports these remarks into the remark engine: the category
+is the LLVM pass name prefixed with `llvm-`, the LLVM remark name becomes the
+name, the message is stored under the `Remark` argument and the structured LLVM
+arguments are copied as metrics. The remark filters therefore select LLVM
+remarks by pass name, and `llvm-.*` selects all of them:
+
+```
+mlir-opt kernel.mlir --gpu-module-to-binary="format=isa" --remarks-filter-passed=llvm-inline
+```
+
+```
+[Passed] Inlined | Category:llvm-inline | Function=kernel | Callee=helper, Caller=kernel, Cost=-30, Remark="'helper' inlined into 'kernel' with (cost=-30, threshold=337)", Threshold=337
+```
+
+Remarks that are not enabled in the engine are left to LLVM, so LLVM's
+`-pass-remarks` flags keep working. Other LLVM diagnostics, such as errors
+reported by a backend, become MLIR diagnostics.
+
+LLVM remarks carry a source location only when the LLVM IR has debug
+information; run `ensure-debug-info-scope-on-llvm-func` before serializing to
+get `file:line:column` locations. Otherwise a remark is attached to the MLIR
+symbol named after its LLVM function, or to the serialized operation.
+
+Remarks saved by an LLVM-based compiler to a YAML or bitstream file can be
+replayed into the engine with `remark::importLLVMRemarks`, which applies the
+same conversion.
diff --git a/mlir/include/mlir/Remark/LLVMRemarkImport.h b/mlir/include/mlir/Remark/LLVMRemarkImport.h
new file mode 100644
index 0000000000000..ccacb70e4c1b0
--- /dev/null
+++ b/mlir/include/mlir/Remark/LLVMRemarkImport.h
@@ -0,0 +1,70 @@
+//===- LLVMRemarkImport.h - Import LLVM remarks into MLIR -------*- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Utilities to import LLVM optimization remarks into the MLIR remark engine,
+// either as they are emitted by LLVM passes or from a serialized remark file.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_REMARK_LLVMREMARKIMPORT_H
+#define MLIR_REMARK_LLVMREMARKIMPORT_H
+
+#include "mlir/IR/Location.h"
+#include "mlir/IR/Remarks.h"
+#include "llvm/ADT/StringMap.h"
+#include "llvm/IR/DiagnosticHandler.h"
+#include "llvm/Remarks/RemarkFormat.h"
+
+namespace mlir {
+class Operation;
+
+namespace remark {
+
+/// Prefix of the category of remarks imported from LLVM. The category is the
+/// prefix followed by the LLVM pass name, e.g. `llvm-inline`, so that remark
+/// filters can select LLVM remarks (`llvm-.*`) or a single LLVM pass.
+constexpr llvm::StringLiteral llvmRemarkCategoryPrefix = "llvm-";
+
+/// Parses `buffer`, a serialized LLVM remark file, and reports its remarks into
+/// the remark engine of the context of `anchor`, the operation the remarks
+/// refer to. The category is the prefixed LLVM pass name, the LLVM remark name
+/// becomes the remark name and the message is stored under the "Remark"
+/// argument. A remark is attached to its debug location if it has one,
+/// otherwise to the symbol named after its LLVM function inside `anchor`,
+/// otherwise to `anchor`. Fails if the buffer cannot be parsed.
+LogicalResult importLLVMRemarks(Operation *anchor, StringRef buffer,
+ llvm::remarks::Format format);
+
+/// Diagnostic handler forwarding LLVM diagnostics to MLIR. Optimization remarks
+/// enabled in the remark engine of the context of `anchor` are imported as
+/// described for `importLLVMRemarks`, so the remark filters of the engine
+/// select LLVM remarks by their prefixed pass name. Other remarks are left to
+/// LLVM, which prints those enabled with its own `-pass-remarks` flags. All
+/// other diagnostics are reported as MLIR diagnostics. Install the handler with
+/// `RespectFilters` set, so that LLVM only forwards enabled remarks.
+class LLVMToMLIRDiagnosticHandler : public llvm::DiagnosticHandler {
+public:
+ explicit LLVMToMLIRDiagnosticHandler(Operation *anchor);
+
+ bool handleDiagnostics(const llvm::DiagnosticInfo &diag) override;
+ bool isAnalysisRemarkEnabled(StringRef passName) const override;
+ bool isMissedOptRemarkEnabled(StringRef passName) const override;
+ bool isPassedOptRemarkEnabled(StringRef passName) const override;
+ using llvm::DiagnosticHandler::isAnyRemarkEnabled;
+ bool isAnyRemarkEnabled() const override;
+
+private:
+ Operation *anchor;
+ detail::RemarkEngine *engine;
+ llvm::StringMap<Location> functionLocations;
+};
+
+} // namespace remark
+} // namespace mlir
+
+#endif // MLIR_REMARK_LLVMREMARKIMPORT_H
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index 7d232dcf9ab90..4ae634a372a8e 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -53,6 +53,12 @@ class ModuleToObject {
protected:
// Hooks to be implemented by derived classes.
+ /// Hook for configuring the LLVM context, called by `run` before translating
+ /// the operation. By default it installs a
+ /// `remark::LLVMToMLIRDiagnosticHandler` to report LLVM diagnostics and
+ /// remarks through MLIR.
+ virtual void setupLLVMContext(llvm::LLVMContext &llvmContext);
+
/// Hook for computing the Datalayout
virtual void setDataLayoutAndTriple(llvm::Module &module);
diff --git a/mlir/lib/Remark/CMakeLists.txt b/mlir/lib/Remark/CMakeLists.txt
index 920a95db4f07b..a75c17288d5b6 100644
--- a/mlir/lib/Remark/CMakeLists.txt
+++ b/mlir/lib/Remark/CMakeLists.txt
@@ -1,4 +1,5 @@
add_mlir_library(MLIRRemarkStreamer
+ LLVMRemarkImport.cpp
RemarkStreamer.cpp
ADDITIONAL_HEADER_DIRS
diff --git a/mlir/lib/Remark/LLVMRemarkImport.cpp b/mlir/lib/Remark/LLVMRemarkImport.cpp
new file mode 100644
index 0000000000000..5b89e5fda6aa1
--- /dev/null
+++ b/mlir/lib/Remark/LLVMRemarkImport.cpp
@@ -0,0 +1,315 @@
+//===- LLVMRemarkImport.cpp - Import LLVM remarks into MLIR ---------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Remark/LLVMRemarkImport.h"
+
+#include "mlir/IR/Diagnostics.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/SymbolTable.h"
+
+#include "llvm/ADT/Twine.h"
+#include "llvm/IR/DiagnosticInfo.h"
+#include "llvm/IR/DiagnosticPrinter.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/GlobalValue.h"
+#include "llvm/Remarks/Remark.h"
+#include "llvm/Remarks/RemarkParser.h"
+#include "llvm/Support/Error.h"
+#include "llvm/Support/raw_ostream.h"
+
+using namespace mlir;
+
+//===----------------------------------------------------------------------===//
+// Locations
+//===----------------------------------------------------------------------===//
+
+/// Returns the location of the symbol `functionName` inside `anchor`, or the
+/// location of `anchor` if there is no such symbol.
+static Location getFunctionLocation(Operation *anchor,
+ llvm::StringMap<Location> &cache,
+ StringRef functionName) {
+ if (functionName.empty() || !anchor->hasTrait<OpTrait::SymbolTable>())
+ return anchor->getLoc();
+ auto it = cache.find(functionName);
+ if (it != cache.end())
+ return it->second;
+ Location loc = anchor->getLoc();
+ if (Operation *symbol = SymbolTable::lookupSymbolIn(anchor, functionName))
+ loc = symbol->getLoc();
+ cache.try_emplace(functionName, loc);
+ return loc;
+}
+
+static Location resolveLocation(Operation *anchor,
+ llvm::StringMap<Location> &cache,
+ StringRef file, unsigned line, unsigned column,
+ StringRef functionName) {
+ if (!file.empty() && line != 0)
+ return FileLineColLoc::get(anchor->getContext(), file, line, column);
+ return getFunctionLocation(anchor, cache, functionName);
+}
+
+static Location
+resolveLocation(Operation *anchor, llvm::StringMap<Location> &cache,
+ const llvm::DiagnosticInfoWithLocationBase &diag) {
+ StringRef functionName = diag.getFunction().getName();
+ if (!diag.isLocationAvailable())
+ return resolveLocation(anchor, cache, "", 0, 0, functionName);
+ llvm::DiagnosticLocation loc = diag.getLocation();
+ return resolveLocation(anchor, cache, loc.getAbsolutePath(), loc.getLine(),
+ loc.getColumn(), functionName);
+}
+
+//===----------------------------------------------------------------------===//
+// Remark conversion
+//===----------------------------------------------------------------------===//
+
+namespace {
+/// The fields of an LLVM remark that are imported, shared by the live and the
+/// serialized paths.
+struct ImportedRemark {
+ remark::RemarkKind kind = remark::RemarkKind::RemarkUnknown;
+ StringRef passName;
+ StringRef remarkName;
+ StringRef functionName;
+ std::string message;
+ std::optional<uint64_t> hotness;
+ SmallVector<std::pair<StringRef, StringRef>> args;
+};
+} // namespace
+
+static remark::RemarkKind getRemarkKind(llvm::DiagnosticKind kind) {
+ switch (kind) {
+ case llvm::DK_OptimizationRemark:
+ case llvm::DK_MachineOptimizationRemark:
+ return remark::RemarkKind::RemarkPassed;
+ case llvm::DK_OptimizationRemarkMissed:
+ case llvm::DK_MachineOptimizationRemarkMissed:
+ return remark::RemarkKind::RemarkMissed;
+ case llvm::DK_OptimizationRemarkAnalysis:
+ case llvm::DK_OptimizationRemarkAnalysisFPCommute:
+ case llvm::DK_OptimizationRemarkAnalysisAliasing:
+ case llvm::DK_MachineOptimizationRemarkAnalysis:
+ return remark::RemarkKind::RemarkAnalysis;
+ case llvm::DK_OptimizationFailure:
+ return remark::RemarkKind::RemarkFailure;
+ default:
+ return remark::RemarkKind::RemarkUnknown;
+ }
+}
+
+static remark::RemarkKind getRemarkKind(llvm::remarks::Type type) {
+ switch (type) {
+ case llvm::remarks::Type::Passed:
+ return remark::RemarkKind::RemarkPassed;
+ case llvm::remarks::Type::Missed:
+ return remark::RemarkKind::RemarkMissed;
+ case llvm::remarks::Type::Analysis:
+ case llvm::remarks::Type::AnalysisFPCommute:
+ case llvm::remarks::Type::AnalysisAliasing:
+ return remark::RemarkKind::RemarkAnalysis;
+ case llvm::remarks::Type::Failure:
+ return remark::RemarkKind::RemarkFailure;
+ case llvm::remarks::Type::Unknown:
+ return remark::RemarkKind::RemarkUnknown;
+ }
+ llvm_unreachable("unknown remark type");
+}
+
+/// Returns the MLIR remark category of the LLVM pass `passName`.
+static std::string getCategory(StringRef passName) {
+ return (llvm::Twine(remark::llvmRemarkCategoryPrefix) + passName).str();
+}
+
+/// Renames the argument keys that the remark engine uses itself.
+static std::string getArgKey(StringRef key) {
+ if (key == "Remark" || key == "RemarkId" || key == "RelatedTo")
+ return ("LLVM" + key).str();
+ return key.str();
+}
+
+static void emitImportedRemark(remark::detail::RemarkEngine &engine,
+ Location loc, const ImportedRemark &imported) {
+ std::string category = getCategory(imported.passName);
+ remark::RemarkOpts opts = remark::RemarkOpts::name(imported.remarkName)
+ .category(category)
+ .function(imported.functionName);
+ remark::detail::InFlightRemark inFlight;
+ switch (imported.kind) {
+ case remark::RemarkKind::RemarkPassed:
+ inFlight = engine.emitOptimizationRemark(loc, opts);
+ break;
+ case remark::RemarkKind::RemarkMissed:
+ inFlight = engine.emitOptimizationRemarkMiss(loc, opts);
+ break;
+ case remark::RemarkKind::RemarkFailure:
+ inFlight = engine.emitOptimizationRemarkFailure(loc, opts);
+ break;
+ case remark::RemarkKind::RemarkAnalysis:
+ inFlight = engine.emitOptimizationRemarkAnalysis(loc, opts);
+ break;
+ case remark::RemarkKind::RemarkUnknown:
+ return;
+ }
+ if (!inFlight)
+ return;
+
+ inFlight << StringRef(imported.message);
+ for (const auto &[key, value] : imported.args) {
+ // Plain strings are already part of the message.
+ if (key == "String")
+ continue;
+ inFlight << remark::detail::Remark::Arg(getArgKey(key), value);
+ }
+ if (imported.hotness)
+ inFlight << remark::detail::Remark::Arg("Hotness", *imported.hotness);
+}
+
+static void importLLVMRemark(remark::detail::RemarkEngine &engine, Location loc,
+ const llvm::DiagnosticInfoOptimizationBase &diag) {
+ ImportedRemark imported;
+ imported.kind =
+ getRemarkKind(static_cast<llvm::DiagnosticKind>(diag.getKind()));
+ imported.passName = diag.getPassName();
+ imported.remarkName = diag.getRemarkName();
+ imported.functionName =
+ llvm::GlobalValue::dropLLVMManglingEscape(diag.getFunction().getName());
+ imported.message = diag.getMsg();
+ imported.hotness = diag.getHotness();
+ for (const llvm::DiagnosticInfoOptimizationBase::Argument &arg :
+ diag.getArgs())
+ imported.args.emplace_back(arg.Key, arg.Val);
+ emitImportedRemark(engine, loc, imported);
+}
+
+static void importLLVMRemark(remark::detail::RemarkEngine &engine, Location loc,
+ const llvm::remarks::Remark &remark) {
+ ImportedRemark imported;
+ imported.kind = getRemarkKind(remark.RemarkType);
+ imported.passName = remark.PassName;
+ imported.remarkName = remark.RemarkName;
+ imported.functionName = remark.FunctionName;
+ imported.message = remark.getArgsAsMsg();
+ imported.hotness = remark.Hotness;
+ for (const llvm::remarks::Argument &arg : remark.Args)
+ imported.args.emplace_back(arg.Key, arg.Val);
+ emitImportedRemark(engine, loc, imported);
+}
+
+LogicalResult mlir::remark::importLLVMRemarks(Operation *anchor,
+ StringRef buffer,
+ llvm::remarks::Format format) {
+ remark::detail::RemarkEngine *engine =
+ anchor->getContext()->getRemarkEngine();
+ if (!engine)
+ return success();
+
+ llvm::Expected<std::unique_ptr<llvm::remarks::RemarkParser>> parser =
+ llvm::remarks::createRemarkParser(format, buffer);
+ if (!parser) {
+ llvm::consumeError(parser.takeError());
+ return failure();
+ }
+
+ llvm::StringMap<Location> functionLocations;
+ while (true) {
+ llvm::Expected<std::unique_ptr<llvm::remarks::Remark>> next =
+ (*parser)->next();
+ if (!next) {
+ llvm::Error error = next.takeError();
+ bool endOfFile = error.isA<llvm::remarks::EndOfFileError>();
+ llvm::consumeError(std::move(error));
+ return success(endOfFile);
+ }
+ const llvm::remarks::Remark &remark = **next;
+ Location loc = remark.Loc ? resolveLocation(anchor, functionLocations,
+ remark.Loc->SourceFilePath,
+ remark.Loc->SourceLine,
+ remark.Loc->SourceColumn,
+ remark.FunctionName)
+ : resolveLocation(anchor, functionLocations, "",
+ 0, 0, remark.FunctionName);
+ importLLVMRemark(*engine, loc, remark);
+ }
+}
+
+//===----------------------------------------------------------------------===//
+// LLVMToMLIRDiagnosticHandler
+//===----------------------------------------------------------------------===//
+
+remark::LLVMToMLIRDiagnosticHandler::LLVMToMLIRDiagnosticHandler(
+ Operation *anchor)
+ : anchor(anchor), engine(anchor->getContext()->getRemarkEngine()) {}
+
+bool remark::LLVMToMLIRDiagnosticHandler::isAnalysisRemarkEnabled(
+ StringRef passName) const {
+ return (engine &&
+ engine->isAnalysisOptRemarkEnabled(getCategory(passName))) ||
+ DiagnosticHandler::isAnalysisRemarkEnabled(passName);
+}
+
+bool remark::LLVMToMLIRDiagnosticHandler::isMissedOptRemarkEnabled(
+ StringRef passName) const {
+ return (engine && engine->isMissedOptRemarkEnabled(getCategory(passName))) ||
+ DiagnosticHandler::isMissedOptRemarkEnabled(passName);
+}
+
+bool remark::LLVMToMLIRDiagnosticHandler::isPassedOptRemarkEnabled(
+ StringRef passName) const {
+ return (engine && engine->isPassedOptRemarkEnabled(getCategory(passName))) ||
+ DiagnosticHandler::isPassedOptRemarkEnabled(passName);
+}
+
+bool remark::LLVMToMLIRDiagnosticHandler::isAnyRemarkEnabled() const {
+ return (engine && engine->isAnyRemarkEnabled()) ||
+ DiagnosticHandler::isAnyRemarkEnabled();
+}
+
+bool remark::LLVMToMLIRDiagnosticHandler::handleDiagnostics(
+ const llvm::DiagnosticInfo &diag) {
+ if (const auto *optRemark =
+ dyn_cast<llvm::DiagnosticInfoOptimizationBase>(&diag)) {
+ remark::RemarkKind kind =
+ getRemarkKind(static_cast<llvm::DiagnosticKind>(optRemark->getKind()));
+ if (!engine ||
+ !engine->isRemarkEnabled(kind, getCategory(optRemark->getPassName())))
+ return false;
+ importLLVMRemark(*engine,
+ resolveLocation(anchor, functionLocations, *optRemark),
+ *optRemark);
+ return true;
+ }
+
+ std::string message;
+ llvm::raw_string_ostream os(message);
+ llvm::DiagnosticPrinterRawOStream printer(os);
+ diag.print(printer);
+ StringRef text = StringRef(message).rtrim();
+
+ Location loc = anchor->getLoc();
+ if (const auto *withLoc = dyn_cast<llvm::DiagnosticInfoUnsupported>(&diag)) {
+ loc = resolveLocation(anchor, functionLocations, *withLoc);
+ // The location is carried by `loc` already.
+ text.consume_front(withLoc->getLocationStr() + ": ");
+ }
+
+ switch (diag.getSeverity()) {
+ case llvm::DS_Error:
+ emitError(loc) << text;
+ break;
+ case llvm::DS_Warning:
+ emitWarning(loc) << text;
+ break;
+ case llvm::DS_Remark:
+ case llvm::DS_Note:
+ emitRemark(loc) << text;
+ break;
+ }
+ return true;
+}
diff --git a/mlir/lib/Target/LLVM/CMakeLists.txt b/mlir/lib/Target/LLVM/CMakeLists.txt
index bd1d66d1e43af..d59138cd9db23 100644
--- a/mlir/lib/Target/LLVM/CMakeLists.txt
+++ b/mlir/lib/Target/LLVM/CMakeLists.txt
@@ -19,6 +19,7 @@ add_mlir_library(MLIRTargetLLVM
Target
LINK_LIBS PUBLIC
MLIRExecutionEngineUtils
+ MLIRRemarkStreamer
MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index 8eaf8d9cadcd2..0f43e9c384048 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -16,6 +16,7 @@
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/IR/BuiltinAttributeInterfaces.h"
#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/Remark/LLVMRemarkImport.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Target/LLVMIR/ModuleTranslation.h"
@@ -242,9 +243,16 @@ ModuleToObject::moduleToObject(llvm::Module &llvmModule) {
return binaryData;
}
+void ModuleToObject::setupLLVMContext(llvm::LLVMContext &llvmContext) {
+ llvmContext.setDiagnosticHandler(
+ std::make_unique<remark::LLVMToMLIRDiagnosticHandler>(&getOperation()),
+ /*RespectFilters=*/true);
+}
+
std::optional<SmallVector<char, 0>> ModuleToObject::run() {
// Translate the module to LLVM IR.
llvm::LLVMContext llvmContext;
+ setupLLVMContext(llvmContext);
std::unique_ptr<llvm::Module> llvmModule = translateToLLVMIR(llvmContext);
if (!llvmModule) {
getOperation().emitError() << "Failed creating the llvm::Module.";
diff --git a/mlir/test/Dialect/GPU/module-to-binary-nvvm-remarks.mlir b/mlir/test/Dialect/GPU/module-to-binary-nvvm-remarks.mlir
new file mode 100644
index 0000000000000..f89ca441748ab
--- /dev/null
+++ b/mlir/test/Dialect/GPU/module-to-binary-nvvm-remarks.mlir
@@ -0,0 +1,30 @@
+// REQUIRES: host-supports-nvptx
+// RUN: mlir-opt %s --pass-pipeline='builtin.module(gpu-module-to-binary{format=isa})' --remarks-filter-passed=llvm-inline 2>&1 | FileCheck %s
+// RUN: mlir-opt %s --pass-pipeline='builtin.module(ensure-debug-info-scope-on-llvm-func,gpu-module-to-binary{format=isa})' --remarks-filter-passed=llvm-inline 2>&1 | FileCheck %s --check-prefix=CHECK-LOC
+// RUN: mlir-opt %s --pass-pipeline='builtin.module(gpu-module-to-binary{format=isa})' --remarks-filter-passed=llvm-loop-unroll 2>&1 | FileCheck %s --check-prefix=CHECK-NONE
+// RUN: mlir-opt %s --pass-pipeline='builtin.module(gpu-module-to-binary{format=isa})' --remarks-filter-missed=llvm-inline 2>&1 | FileCheck %s --check-prefix=CHECK-NONE
+
+// Remarks emitted by LLVM while serializing the GPU module are reported through
+// the MLIR remark engine, with the LLVM pass name, prefixed with `llvm-`, as the category.
+
+// CHECK-NONE-NOT: remark:
+// CHECK-NONE: gpu.binary @kernel_module
+// CHECK-NONE-NOT: remark:
+
+module attributes {gpu.container_module} {
+ gpu.module @kernel_module [#nvvm.target<chip = "sm_70">] {
+ llvm.func @helper(%a: f32, %b: f32) -> f32 {
+ %0 = llvm.fadd %a, %b : f32
+ llvm.return %0 : f32
+ }
+ // Without debug info, the remark is attached to the function.
+ // CHECK: remarks.mlir:[[#@LINE+1]]:{{[0-9]+}}: remark: [Passed] Inlined | Category:llvm-inline | Function=kernel | Callee=helper, Caller=kernel, Cost={{.*}}, Remark="'helper' inlined into 'kernel'
+ llvm.func @kernel(%arg0: f32, %arg1: !llvm.ptr) attributes {gpu.kernel} {
+ // With debug info, the remark is attached to the call.
+ // CHECK-LOC: remarks.mlir:[[#@LINE+1]]:{{[0-9]+}}: remark: [Passed] Inlined | Category:llvm-inline | Function=kernel
+ %0 = llvm.call @helper(%arg0, %arg0) : (f32, f32) -> f32
+ llvm.store %0, %arg1 : f32, !llvm.ptr
+ llvm.return
+ }
+ }
+}
diff --git a/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp b/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp
index 036e05bfc2e02..cdd0a90837b57 100644
--- a/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp
+++ b/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp
@@ -13,6 +13,7 @@
#include "mlir/Dialect/Bufferization/Transforms/Transforms.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Pass/Pass.h"
+#include "llvm/ADT/TypeSwitch.h"
#include "TestAttributes.h" // TestTensorEncodingAttr, TestMemRefLayoutAttr
#include "TestDialect.h"
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 4c683358e71a5..52c8d5ecd195b 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -14,6 +14,7 @@
#include "mlir/Interfaces/FunctionImplementation.h"
#include "mlir/Interfaces/MemorySlotInterfaces.h"
#include "llvm/ADT/SmallVectorExtras.h"
+#include "llvm/ADT/TypeSwitch.h"
using namespace mlir;
using namespace test;
diff --git a/mlir/unittests/CMakeLists.txt b/mlir/unittests/CMakeLists.txt
index 654ec44d90b04..25514840c0f58 100644
--- a/mlir/unittests/CMakeLists.txt
+++ b/mlir/unittests/CMakeLists.txt
@@ -15,6 +15,7 @@ add_subdirectory(Interfaces)
add_subdirectory(IR)
add_subdirectory(Parser)
add_subdirectory(Pass)
+add_subdirectory(Remark)
add_subdirectory(Support)
add_subdirectory(Rewrite)
add_subdirectory(TableGen)
diff --git a/mlir/unittests/Remark/CMakeLists.txt b/mlir/unittests/Remark/CMakeLists.txt
new file mode 100644
index 0000000000000..af687a89537a6
--- /dev/null
+++ b/mlir/unittests/Remark/CMakeLists.txt
@@ -0,0 +1,9 @@
+add_mlir_unittest(MLIRRemarkTests
+ LLVMRemarkImportTest.cpp
+)
+mlir_target_link_libraries(MLIRRemarkTests
+ PRIVATE
+ MLIRIR
+ MLIRParser
+)
+target_link_libraries(MLIRRemarkTests PRIVATE MLIRRemarkStreamer)
diff --git a/mlir/unittests/Remark/LLVMRemarkImportTest.cpp b/mlir/unittests/Remark/LLVMRemarkImportTest.cpp
new file mode 100644
index 0000000000000..39a4b116ad844
--- /dev/null
+++ b/mlir/unittests/Remark/LLVMRemarkImportTest.cpp
@@ -0,0 +1,395 @@
+//===- LLVMRemarkImportTest.cpp - LLVM remark import unit tests -----------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Remark/LLVMRemarkImport.h"
+
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/Diagnostics.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/Remarks.h"
+#include "mlir/IR/SymbolTable.h"
+#include "mlir/Parser/Parser.h"
+
+#include "llvm/BinaryFormat/Dwarf.h"
+#include "llvm/IR/DIBuilder.h"
+#include "llvm/IR/DebugInfoMetadata.h"
+#include "llvm/IR/DiagnosticInfo.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/LLVMContext.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Remarks/RemarkFormat.h"
+#include "llvm/Support/Path.h"
+
+#include "gmock/gmock.h"
+#include "gtest/gtest.h"
+
+using namespace mlir;
+using namespace testing;
+
+namespace {
+
+class CollectingStreamer : public remark::detail::MLIRRemarkStreamerBase {
+public:
+ CollectingStreamer(std::vector<remark::detail::Remark> &remarks)
+ : remarks(remarks) {}
+
+ void streamOptimizationRemark(const remark::detail::Remark &remark) override {
+ remarks.push_back(remark);
+ }
+
+private:
+ std::vector<remark::detail::Remark> &remarks;
+};
+
+std::optional<std::string> getArg(const remark::detail::Remark &remark,
+ StringRef key) {
+ for (const remark::detail::Remark::Arg &arg : remark.getArgs())
+ if (arg.key == key)
+ return arg.val;
+ return std::nullopt;
+}
+
+class LLVMRemarkImportTest : public ::testing::Test {
+protected:
+ void SetUp() override {
+ // The nested module is the symbol named after the LLVM function.
+ module = parseSourceString<ModuleOp>(R"mlir(
+ module @anchor {
+ module @kernel {
+ }
+ %0 = builtin.unrealized_conversion_cast to i32
+ }
+ )mlir",
+ &context);
+ ASSERT_TRUE(module);
+ anchor = module->getOperation();
+ kernel = SymbolTable::lookupSymbolIn(anchor, "kernel");
+ ASSERT_TRUE(kernel);
+ nonSymbolTable = &module->getBody()->back();
+ }
+
+ void enableRemarks(const remark::RemarkCategories &cats) {
+ ASSERT_TRUE(succeeded(remark::enableOptimizationRemarks(
+ context, std::make_unique<CollectingStreamer>(remarks),
+ std::make_unique<remark::RemarkEmittingPolicyAll>(), cats)));
+ }
+
+ void installHandler(llvm::LLVMContext &llvmContext) {
+ llvmContext.setDiagnosticHandler(
+ std::make_unique<remark::LLVMToMLIRDiagnosticHandler>(anchor),
+ /*RespectFilters=*/true);
+ }
+
+ llvm::Function *createKernel(llvm::Module &llvmModule) {
+ llvm::FunctionType *type = llvm::FunctionType::get(
+ llvm::Type::getVoidTy(llvmModule.getContext()), /*isVarArg=*/false);
+ return llvm::Function::Create(type, llvm::GlobalValue::ExternalLinkage,
+ "kernel", llvmModule);
+ }
+
+ MLIRContext context;
+ OwningOpRef<ModuleOp> module;
+ Operation *anchor = nullptr;
+ Operation *kernel = nullptr;
+ Operation *nonSymbolTable = nullptr;
+ std::vector<remark::detail::Remark> remarks;
+};
+
+TEST_F(LLVMRemarkImportTest, HandlerAnswersRemarkQueriesFromEngine) {
+ enableRemarks({/*all=*/std::nullopt, /*passed=*/"llvm-loop-unroll",
+ /*missed=*/"llvm-inline", /*analysis=*/std::nullopt,
+ /*failed=*/std::nullopt});
+
+ llvm::LLVMContext llvmContext;
+ installHandler(llvmContext);
+ const llvm::DiagnosticHandler *handler = llvmContext.getDiagHandlerPtr();
+ ASSERT_TRUE(handler);
+
+ EXPECT_TRUE(handler->isAnyRemarkEnabled());
+ EXPECT_TRUE(handler->isPassedOptRemarkEnabled("loop-unroll"));
+ EXPECT_FALSE(handler->isPassedOptRemarkEnabled("inline"));
+ EXPECT_TRUE(handler->isMissedOptRemarkEnabled("inline"));
+ EXPECT_FALSE(handler->isMissedOptRemarkEnabled("loop-unroll"));
+ EXPECT_FALSE(handler->isAnalysisRemarkEnabled("loop-unroll"));
+ EXPECT_TRUE(handler->isAnyRemarkEnabled("loop-unroll"));
+ EXPECT_FALSE(handler->isAnyRemarkEnabled("licm"));
+}
+
+TEST_F(LLVMRemarkImportTest, HandlerWithoutEngineDisablesRemarks) {
+ llvm::LLVMContext llvmContext;
+ installHandler(llvmContext);
+ const llvm::DiagnosticHandler *handler = llvmContext.getDiagHandlerPtr();
+ ASSERT_TRUE(handler);
+ EXPECT_FALSE(handler->isAnyRemarkEnabled());
+ EXPECT_FALSE(handler->isPassedOptRemarkEnabled("loop-unroll"));
+ EXPECT_FALSE(handler->isMissedOptRemarkEnabled("loop-unroll"));
+ EXPECT_FALSE(handler->isAnalysisRemarkEnabled("loop-unroll"));
+}
+
+TEST_F(LLVMRemarkImportTest, ImportsLiveRemarkWithFunctionLocation) {
+ enableRemarks({/*all=*/std::nullopt, /*passed=*/"llvm-loop-unroll",
+ /*missed=*/std::nullopt, /*analysis=*/std::nullopt,
+ /*failed=*/std::nullopt});
+
+ llvm::LLVMContext llvmContext;
+ installHandler(llvmContext);
+ llvm::Module llvmModule("m", llvmContext);
+ llvm::Function *fn = createKernel(llvmModule);
+
+ llvmContext.diagnose(
+ llvm::OptimizationRemark("loop-unroll", "FullyUnrolled", fn)
+ << "unrolled loop by a factor of "
+ << llvm::DiagnosticInfoOptimizationBase::Argument("UnrollCount", 4));
+ // Not enabled: missed remarks and the `inline` category.
+ llvmContext.diagnose(
+ llvm::OptimizationRemarkMissed("loop-unroll", "NotUnrolled", fn)
+ << "not unrolled");
+ llvmContext.diagnose(llvm::OptimizationRemark("inline", "Inlined", fn)
+ << "inlined");
+
+ ASSERT_EQ(remarks.size(), 1u);
+ const remark::detail::Remark &remark = remarks[0];
+ EXPECT_EQ(remark.getRemarkKind(), remark::RemarkKind::RemarkPassed);
+ EXPECT_EQ(remark.getRemarkName(), "FullyUnrolled");
+ EXPECT_EQ(remark.getCategoryName(), "llvm-loop-unroll");
+ EXPECT_EQ(remark.getFunction(), "kernel");
+ EXPECT_EQ(remark.getLocation(), kernel->getLoc());
+ EXPECT_EQ(getArg(remark, "Remark"), "unrolled loop by a factor of 4");
+ EXPECT_EQ(getArg(remark, "UnrollCount"), "4");
+ EXPECT_FALSE(getArg(remark, "String"));
+}
+
+TEST_F(LLVMRemarkImportTest, ImportsLiveRemarkWithDebugLocation) {
+ enableRemarks({/*all=*/std::nullopt, /*passed=*/std::nullopt,
+ /*missed=*/"llvm-licm", /*analysis=*/std::nullopt,
+ /*failed=*/std::nullopt});
+
+ llvm::LLVMContext llvmContext;
+ installHandler(llvmContext);
+ llvm::Module llvmModule("m", llvmContext);
+ llvm::Function *fn = createKernel(llvmModule);
+
+ llvm::DIBuilder dib(llvmModule);
+ llvm::DIFile *file = dib.createFile("kernel.mlir", "/src");
+ llvm::DICompileUnit *cu = dib.createCompileUnit(
+ llvm::dwarf::DW_LANG_C, file, "mlir", /*isOptimized=*/false, "", 0);
+ llvm::DISubroutineType *fnType =
+ dib.createSubroutineType(dib.getOrCreateTypeArray({}));
+ llvm::DISubprogram *sp = dib.createFunction(
+ cu, "kernel", "kernel", file, /*LineNo=*/1, fnType, /*ScopeLine=*/1,
+ llvm::DINode::FlagZero, llvm::DISubprogram::SPFlagDefinition);
+ fn->setSubprogram(sp);
+ llvm::BasicBlock *bb = llvm::BasicBlock::Create(llvmContext, "entry", fn);
+ llvm::IRBuilder<> builder(bb);
+ llvm::Instruction *inst = builder.CreateRetVoid();
+ inst->setDebugLoc(llvm::DILocation::get(llvmContext, 7, 3, sp));
+ dib.finalize();
+
+ llvmContext.diagnose(
+ llvm::OptimizationRemarkMissed("licm", "LoadNotHoisted", inst)
+ << "failed to hoist load");
+
+ ASSERT_EQ(remarks.size(), 1u);
+ const remark::detail::Remark &remark = remarks[0];
+ EXPECT_EQ(remark.getRemarkKind(), remark::RemarkKind::RemarkMissed);
+ EXPECT_EQ(remark.getRemarkName(), "LoadNotHoisted");
+ EXPECT_EQ(remark.getCategoryName(), "llvm-licm");
+ auto loc = dyn_cast<FileLineColLoc>(remark.getLocation());
+ ASSERT_TRUE(loc);
+ SmallString<32> expectedFile;
+ llvm::sys::path::append(expectedFile, "/src", "kernel.mlir");
+ EXPECT_EQ(loc.getFilename().getValue(), expectedFile);
+ EXPECT_EQ(loc.getLine(), 7u);
+ EXPECT_EQ(loc.getColumn(), 3u);
+}
+
+TEST_F(LLVMRemarkImportTest, ImportsOptimizationFailureAsFailedRemark) {
+ enableRemarks({/*all=*/std::nullopt, /*passed=*/std::nullopt,
+ /*missed=*/std::nullopt, /*analysis=*/std::nullopt,
+ /*failed=*/"llvm-transform-warning"});
+
+ llvm::LLVMContext llvmContext;
+ installHandler(llvmContext);
+ llvm::Module llvmModule("m", llvmContext);
+ llvm::Function *fn = createKernel(llvmModule);
+ llvm::BasicBlock *bb = llvm::BasicBlock::Create(llvmContext, "entry", fn);
+
+ llvmContext.diagnose(llvm::DiagnosticInfoOptimizationFailure(
+ "transform-warning", "FailedRequestedUnrolling",
+ llvm::DiagnosticLocation(), bb)
+ << "loop not unrolled");
+
+ ASSERT_EQ(remarks.size(), 1u);
+ EXPECT_EQ(remarks[0].getRemarkKind(), remark::RemarkKind::RemarkFailure);
+ EXPECT_EQ(remarks[0].getRemarkName(), "FailedRequestedUnrolling");
+ EXPECT_EQ(remarks[0].getCategoryName(), "llvm-transform-warning");
+ EXPECT_EQ(getArg(remarks[0], "Remark"), "loop not unrolled");
+}
+
+TEST_F(LLVMRemarkImportTest, LeavesRemarksNotEnabledInEngineToLLVM) {
+ enableRemarks({/*all=*/std::nullopt, /*passed=*/"llvm-loop-unroll",
+ /*missed=*/std::nullopt, /*analysis=*/std::nullopt,
+ /*failed=*/std::nullopt});
+
+ llvm::LLVMContext llvmContext;
+ llvm::Module llvmModule("m", llvmContext);
+ llvm::Function *fn = createKernel(llvmModule);
+ llvm::BasicBlock *bb = llvm::BasicBlock::Create(llvmContext, "entry", fn);
+ remark::LLVMToMLIRDiagnosticHandler handler(anchor);
+
+ EXPECT_TRUE(handler.handleDiagnostics(
+ llvm::OptimizationRemark("loop-unroll", "FullyUnrolled", fn)
+ << "unrolled"));
+ EXPECT_FALSE(handler.handleDiagnostics(
+ llvm::OptimizationRemarkMissed("loop-unroll", "NotUnrolled", fn)
+ << "not unrolled"));
+ EXPECT_FALSE(handler.handleDiagnostics(
+ llvm::DiagnosticInfoOptimizationFailure("transform-warning",
+ "FailedRequestedUnrolling",
+ llvm::DiagnosticLocation(), bb)
+ << "loop not unrolled"));
+ EXPECT_EQ(remarks.size(), 1u);
+}
+
+TEST_F(LLVMRemarkImportTest, ReportsLLVMDiagnosticsAsMLIRDiagnostics) {
+ std::vector<std::pair<DiagnosticSeverity, std::string>> diagnostics;
+ std::vector<Location> locations;
+ ScopedDiagnosticHandler diagHandler(&context, [&](Diagnostic &diag) {
+ diagnostics.emplace_back(diag.getSeverity(), diag.str());
+ locations.push_back(diag.getLocation());
+ return success();
+ });
+
+ llvm::LLVMContext llvmContext;
+ installHandler(llvmContext);
+ llvm::Module llvmModule("m", llvmContext);
+ llvm::Function *fn = createKernel(llvmModule);
+
+ llvmContext.diagnose(
+ llvm::DiagnosticInfoUnsupported(*fn, "unsupported operation"));
+ llvmContext.diagnose(llvm::DiagnosticInfoUnsupported(
+ *fn, "suspicious operation", llvm::DiagnosticLocation(),
+ llvm::DS_Warning));
+
+ ASSERT_EQ(diagnostics.size(), 2u);
+ EXPECT_EQ(diagnostics[0].first, DiagnosticSeverity::Error);
+ EXPECT_THAT(diagnostics[0].second, HasSubstr("unsupported operation"));
+ EXPECT_THAT(diagnostics[0].second, Not(HasSubstr("<unknown>:0:0")));
+ EXPECT_EQ(locations[0], kernel->getLoc());
+ EXPECT_EQ(diagnostics[1].first, DiagnosticSeverity::Warning);
+ EXPECT_THAT(diagnostics[1].second, HasSubstr("suspicious operation"));
+ EXPECT_TRUE(llvmContext.getDiagHandlerPtr()->HasErrors);
+}
+
+TEST_F(LLVMRemarkImportTest, ImportsSerializedRemarks) {
+ enableRemarks({/*all=*/"llvm-inline|llvm-loop-unroll", /*passed=*/"",
+ /*missed=*/"",
+ /*analysis=*/"", /*failed=*/""});
+
+ const char *yaml = R"yaml(--- !Passed
+Pass: inline
+Name: Inlined
+DebugLoc: { File: kernel.mlir, Line: 12, Column: 5 }
+Function: kernel
+Hotness: 300
+Args:
+ - String: 'inlined '
+ - Callee: helper
+ - String: ' into '
+ - Caller: kernel
+ - Remark: collides
+...
+--- !Missed
+Pass: loop-unroll
+Name: NotUnrolled
+Function: kernel
+Args:
+ - String: 'not unrolled: '
+ - Reason: trip count unknown
+...
+--- !Missed
+Pass: loop-unroll
+Name: NotUnrolled
+Function: not_a_symbol
+...
+--- !Analysis
+Pass: asm-printer
+Name: InstructionCount
+Function: kernel
+Args:
+ - NumInstructions: '42'
+...
+)yaml";
+
+ ASSERT_TRUE(succeeded(
+ remark::importLLVMRemarks(anchor, yaml, llvm::remarks::Format::YAML)));
+
+ // The `asm-printer` remark is not enabled.
+ ASSERT_EQ(remarks.size(), 3u);
+
+ const remark::detail::Remark &passed = remarks[0];
+ EXPECT_EQ(passed.getRemarkKind(), remark::RemarkKind::RemarkPassed);
+ EXPECT_EQ(passed.getRemarkName(), "Inlined");
+ EXPECT_EQ(passed.getCategoryName(), "llvm-inline");
+ EXPECT_EQ(passed.getFunction(), "kernel");
+ auto loc = dyn_cast<FileLineColLoc>(passed.getLocation());
+ ASSERT_TRUE(loc);
+ EXPECT_EQ(loc.getFilename().getValue(), "kernel.mlir");
+ EXPECT_EQ(loc.getLine(), 12u);
+ EXPECT_EQ(loc.getColumn(), 5u);
+ EXPECT_EQ(getArg(passed, "Remark"), "inlined helper into kernelcollides");
+ EXPECT_EQ(getArg(passed, "Callee"), "helper");
+ EXPECT_EQ(getArg(passed, "Caller"), "kernel");
+ EXPECT_EQ(getArg(passed, "LLVMRemark"), "collides");
+ EXPECT_EQ(getArg(passed, "Hotness"), "300");
+ EXPECT_FALSE(getArg(passed, "String"));
+
+ const remark::detail::Remark &missed = remarks[1];
+ EXPECT_EQ(missed.getRemarkKind(), remark::RemarkKind::RemarkMissed);
+ EXPECT_EQ(missed.getRemarkName(), "NotUnrolled");
+ EXPECT_EQ(missed.getCategoryName(), "llvm-loop-unroll");
+ EXPECT_EQ(missed.getLocation(), kernel->getLoc());
+ EXPECT_EQ(getArg(missed, "Remark"), "not unrolled: trip count unknown");
+ EXPECT_EQ(getArg(missed, "Reason"), "trip count unknown");
+
+ // Unknown function: attached to the anchor.
+ EXPECT_EQ(remarks[2].getLocation(), anchor->getLoc());
+}
+
+TEST_F(LLVMRemarkImportTest, ImportsSerializedRemarksWithoutSymbolTable) {
+ enableRemarks({/*all=*/"llvm-inline", /*passed=*/"", /*missed=*/"",
+ /*analysis=*/"", /*failed=*/""});
+ const char *yaml = R"yaml(--- !Passed
+Pass: inline
+Name: Inlined
+Function: kernel
+...
+)yaml";
+ ASSERT_TRUE(succeeded(remark::importLLVMRemarks(
+ nonSymbolTable, yaml, llvm::remarks::Format::YAML)));
+ ASSERT_EQ(remarks.size(), 1u);
+ EXPECT_EQ(remarks[0].getLocation(), nonSymbolTable->getLoc());
+}
+
+TEST_F(LLVMRemarkImportTest, ImportSerializedRemarksReportsParseErrors) {
+ enableRemarks({/*all=*/".*", /*passed=*/"", /*missed=*/"", /*analysis=*/"",
+ /*failed=*/""});
+ EXPECT_TRUE(failed(remark::importLLVMRemarks(anchor, "--- !Passed\nFoo: [",
+ llvm::remarks::Format::YAML)));
+ EXPECT_TRUE(failed(remark::importLLVMRemarks(
+ anchor, "not a bitstream", llvm::remarks::Format::Bitstream)));
+ EXPECT_TRUE(remarks.empty());
+}
+
+TEST_F(LLVMRemarkImportTest, ImportSerializedRemarksWithoutEngine) {
+ EXPECT_TRUE(succeeded(remark::importLLVMRemarks(
+ anchor, "not a remark file", llvm::remarks::Format::YAML)));
+}
+
+} // namespace
diff --git a/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp b/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
index 31c5aa3195db7..fa795f3fcf8b6 100644
--- a/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
+++ b/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
@@ -10,6 +10,8 @@
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/Remarks.h"
+#include "mlir/IR/SymbolTable.h"
#include "mlir/InitAllDialects.h"
#include "mlir/Parser/Parser.h"
#include "mlir/Target/LLVM/NVVM/Target.h"
@@ -224,6 +226,85 @@ TEST_F(MLIRTargetLLVMNVVM,
}
}
+// Test that LLVM optimization remarks are imported into the remark engine.
+TEST_F(MLIRTargetLLVMNVVM, SKIP_WITHOUT_NVPTX(LLVMRemarksAreImported)) {
+ MLIRContext context(registry);
+ std::string moduleStr = R"mlir(
+ gpu.module @nvvm_test {
+ llvm.func @helper(%a: f32, %b: f32) -> f32 {
+ %0 = llvm.fadd %a, %b : f32
+ llvm.return %0 : f32
+ }
+ llvm.func @nvvm_kernel(%arg0: f32, %arg1: !llvm.ptr)
+ attributes {gpu.kernel, nvvm.kernel} {
+ %0 = llvm.call @helper(%arg0, %arg0) : (f32, f32) -> f32
+ llvm.store %0, %arg1 : f32, !llvm.ptr
+ llvm.return
+ }
+ }
+ )mlir";
+ OwningOpRef<ModuleOp> module =
+ parseSourceString<ModuleOp>(moduleStr, &context);
+ ASSERT_TRUE(!!module);
+
+ struct CollectingStreamer : public remark::detail::MLIRRemarkStreamerBase {
+ CollectingStreamer(std::vector<remark::detail::Remark> &remarks)
+ : remarks(remarks) {}
+ void
+ streamOptimizationRemark(const remark::detail::Remark &remark) override {
+ remarks.push_back(remark);
+ }
+ std::vector<remark::detail::Remark> &remarks;
+ };
+ std::vector<remark::detail::Remark> remarks;
+ remark::RemarkCategories cats{/*all=*/std::nullopt, /*passed=*/"llvm-inline",
+ /*missed=*/std::nullopt,
+ /*analysis=*/"llvm-asm-printer",
+ /*failed=*/std::nullopt};
+ ASSERT_TRUE(succeeded(remark::enableOptimizationRemarks(
+ context, std::make_unique<CollectingStreamer>(remarks),
+ std::make_unique<remark::RemarkEmittingPolicyAll>(), cats)));
+
+ NVVM::NVVMTargetAttr target = NVVM::NVVMTargetAttr::get(&context);
+ auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
+ ASSERT_TRUE(!!serializer);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Assembly);
+ for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
+ std::optional<mlir::gpu::SerializedObject> object =
+ serializer.serializeToObject(gpuModule, options);
+ ASSERT_TRUE(object != std::nullopt);
+ ASSERT_TRUE(!object->getObject().empty());
+
+ Operation *kernel =
+ SymbolTable::lookupSymbolIn(gpuModule.getOperation(), "nvvm_kernel");
+ ASSERT_TRUE(kernel);
+
+ // Middle-end remark.
+ auto inlined =
+ llvm::find_if(remarks, [](const remark::detail::Remark &remark) {
+ return remark.getCategoryName() == "llvm-inline" &&
+ remark.getRemarkName() == "Inlined";
+ });
+ ASSERT_TRUE(inlined != remarks.end());
+ EXPECT_EQ(inlined->getRemarkKind(), remark::RemarkKind::RemarkPassed);
+ EXPECT_EQ(inlined->getFunction(), "nvvm_kernel");
+ EXPECT_EQ(inlined->getLocation(), kernel->getLoc());
+
+ // Codegen remark.
+ auto instructionCount =
+ llvm::find_if(remarks, [](const remark::detail::Remark &remark) {
+ return remark.getCategoryName() == "llvm-asm-printer" &&
+ remark.getRemarkName() == "InstructionCount" &&
+ remark.getFunction() == "nvvm_kernel";
+ });
+ ASSERT_TRUE(instructionCount != remarks.end());
+ EXPECT_EQ(instructionCount->getRemarkKind(),
+ remark::RemarkKind::RemarkAnalysis);
+ EXPECT_EQ(instructionCount->getLocation(), kernel->getLoc());
+ remarks.clear();
+ }
+}
+
// Test linking LLVM IR from a resource attribute.
TEST_F(MLIRTargetLLVMNVVM, SKIP_WITHOUT_NVPTX(LinkedLLVMIRResource)) {
MLIRContext context(registry);
>From 4ce5233989602bd84e07aa26a89f5491e6515340 Mon Sep 17 00:00:00 2001
From: Guray Ozen <gozen at nvidia.com>
Date: Wed, 9 Sep 2026 13:56:37 +0200
Subject: [PATCH 2/2] remove unnecessary files
---
.../lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp | 1 -
mlir/test/lib/Dialect/Test/TestOpDefs.cpp | 1 -
2 files changed, 2 deletions(-)
diff --git a/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp b/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp
index cdd0a90837b57..036e05bfc2e02 100644
--- a/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp
+++ b/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp
@@ -13,7 +13,6 @@
#include "mlir/Dialect/Bufferization/Transforms/Transforms.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Pass/Pass.h"
-#include "llvm/ADT/TypeSwitch.h"
#include "TestAttributes.h" // TestTensorEncodingAttr, TestMemRefLayoutAttr
#include "TestDialect.h"
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 52c8d5ecd195b..4c683358e71a5 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -14,7 +14,6 @@
#include "mlir/Interfaces/FunctionImplementation.h"
#include "mlir/Interfaces/MemorySlotInterfaces.h"
#include "llvm/ADT/SmallVectorExtras.h"
-#include "llvm/ADT/TypeSwitch.h"
using namespace mlir;
using namespace test;
More information about the Mlir-commits
mailing list