[Mlir-commits] [mlir] 9a5a9a8 - [MLIR][Remark] Import LLVM optimization remarks into the remark engine (#222300)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Sep 9 05:27:07 PDT 2026
Author: Guray Ozen
Date: 2026-09-09T14:27:01+02:00
New Revision: 9a5a9a8a59bbfd0253ed005e231509de4436f2e8
URL: https://github.com/llvm/llvm-project/commit/9a5a9a8a59bbfd0253ed005e231509de4436f2e8
DIFF: https://github.com/llvm/llvm-project/commit/9a5a9a8a59bbfd0253ed005e231509de4436f2e8.diff
LOG: [MLIR][Remark] Import LLVM optimization remarks into the remark engine (#222300)
Collect LLVM side of remarks with MLIR's remark infra.
LLVM passes emit their own [optimization
remarks](https://llvm.org/docs/Remarks.html) and MLIR has own. This PR
bridges MLIR and LLVM remarks.
`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.
Assisted by Claude Fable 5.1
---------
Co-authored-by: Claude Fable 5.1 <noreply at anthropic.com>
Added:
mlir/include/mlir/Remark/LLVMRemarkImport.h
mlir/lib/Remark/LLVMRemarkImport.cpp
mlir/test/Dialect/GPU/module-to-binary-nvvm-remarks.mlir
mlir/unittests/Remark/CMakeLists.txt
mlir/unittests/Remark/LLVMRemarkImportTest.cpp
Modified:
mlir/docs/Remarks.md
mlir/include/mlir/Target/LLVM/ModuleToObject.h
mlir/lib/Remark/CMakeLists.txt
mlir/lib/Target/LLVM/CMakeLists.txt
mlir/lib/Target/LLVM/ModuleToObject.cpp
mlir/unittests/CMakeLists.txt
mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
Removed:
################################################################################
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/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);
More information about the Mlir-commits
mailing list