#include "toy/Passes.h"
#include "mlir/Dialect/Affine/Passes.h"
+#include "mlir/Dialect/LLVMIR/Transforms/Passes.h"
#include "mlir/ExecutionEngine/ExecutionEngine.h"
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/IR/AsmState.h"
if (isLoweringToLLVM) {
// Finish lowering the toy IR to the LLVM dialect.
pm.addPass(mlir::toy::createLowerToLLVMPass());
+ // This is necessary to have line tables emitted and basic
+ // debugger working. In the future we will add proper debug information
+ // emission directly from our frontend.
+ pm.addNestedPass<mlir::LLVM::LLVMFuncOp>(
+ mlir::LLVM::createDIScopeForLLVMFuncOpPass());
}
if (mlir::failed(pm.run(*module)))
#include "toy/Passes.h"
#include "mlir/Dialect/Affine/Passes.h"
+#include "mlir/Dialect/LLVMIR/Transforms/Passes.h"
#include "mlir/ExecutionEngine/ExecutionEngine.h"
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/IR/AsmState.h"
if (isLoweringToLLVM) {
// Finish lowering the toy IR to the LLVM dialect.
pm.addPass(mlir::toy::createLowerToLLVMPass());
+ // This is necessary to have line tables emitted and basic
+ // debugger working. In the future we will add proper debug information
+ // emission directly from our frontend.
+ pm.addNestedPass<mlir::LLVM::LLVMFuncOp>(
+ mlir::LLVM::createDIScopeForLLVMFuncOpPass());
}
if (mlir::failed(pm.run(*module)))
namespace LLVM {
+/// Create a pass to add DIScope to LLVMFuncOp that are missing it.
+std::unique_ptr<Pass> createDIScopeForLLVMFuncOpPass();
+
/// Generate the code for registering conversion passes.
#define GEN_PASS_REGISTRATION
#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
let constructor = "::mlir::NVVM::createOptimizeForTargetPass()";
}
+def DIScopeForLLVMFuncOp : Pass<"ensure-debug-info-scope-on-llvm-func", "LLVM::LLVMFuncOp"> {
+ let summary = "Materialize LLVM debug info subprogram attribute on every LLVMFuncOp";
+ let description = [{
+ Having a debug info subprogram attribute on a function is required for
+ emitting line tables from MLIR FileLocCol locations.
+
+ This is not intended to be a proper replacement for frontends to emit
+ complete debug informations, however it is a convenient way to get line
+ tables for debugging purposes. This allow to step trough in a debugger
+ line-by-line or get a backtrace with line numbers.
+ }];
+
+ let constructor = "mlir::LLVM::createDIScopeForLLVMFuncOpPass()";
+}
+
#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES
add_mlir_dialect_library(MLIRLLVMIRTransforms
+ DIScopeForLLVMFuncOp.cpp
LegalizeForExport.cpp
OptimizeForNVVM.cpp
RequestCWrappers.cpp
--- /dev/null
+//===- DILineTableFromLocations.cpp - -------------------------------------===//
+//
+// 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/Dialect/LLVMIR/Transforms/Passes.h"
+
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Pass/Pass.h"
+#include "llvm/BinaryFormat/Dwarf.h"
+#include "llvm/Support/Debug.h"
+#include "llvm/Support/Path.h"
+
+namespace mlir {
+namespace LLVM {
+#define GEN_PASS_DEF_DISCOPEFORLLVMFUNCOP
+#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
+} // namespace LLVM
+} // namespace mlir
+
+using namespace mlir;
+
+/// Attempt to extract a filename for the given loc.
+static FileLineColLoc extractFileLoc(Location loc) {
+ if (auto fileLoc = loc.dyn_cast<FileLineColLoc>())
+ return fileLoc;
+ if (auto nameLoc = loc.dyn_cast<NameLoc>())
+ return extractFileLoc(nameLoc.getChildLoc());
+ if (auto opaqueLoc = loc.dyn_cast<OpaqueLoc>())
+ return extractFileLoc(opaqueLoc.getFallbackLocation());
+ return FileLineColLoc();
+}
+
+namespace {
+/// Add a debug info scope to LLVMFuncOp that are missing it.
+struct DIScopeForLLVMFuncOp
+ : public LLVM::impl::DIScopeForLLVMFuncOpBase<DIScopeForLLVMFuncOp> {
+ void runOnOperation() override {
+ LLVM::LLVMFuncOp llvmFunc = getOperation();
+ Location loc = llvmFunc.getLoc();
+ if (loc->findInstanceOf<mlir::FusedLocWith<LLVM::DISubprogramAttr>>())
+ return;
+
+ MLIRContext *context = &getContext();
+
+ // To find a DICompileUnitAttr attached to a parent (the module for
+ // example), otherwise create a default one.
+ LLVM::DICompileUnitAttr compileUnitAttr;
+ if (ModuleOp module = llvmFunc->getParentOfType<ModuleOp>()) {
+ auto fusedCompileUnitAttr =
+ module->getLoc()
+ ->findInstanceOf<mlir::FusedLocWith<LLVM::DICompileUnitAttr>>();
+ if (fusedCompileUnitAttr)
+ compileUnitAttr = fusedCompileUnitAttr.getMetadata();
+ }
+
+ // Filename, line and colmun to associate to the function.
+ LLVM::DIFileAttr fileAttr;
+ int64_t line = 1, col = 1;
+ FileLineColLoc fileLoc = extractFileLoc(loc);
+ if (!fileLoc && compileUnitAttr) {
+ fileAttr = compileUnitAttr.getFile();
+ } else if (!fileLoc) {
+ fileAttr = LLVM::DIFileAttr::get(context, "<unknown>", "");
+ } else {
+ line = fileLoc.getLine();
+ col = fileLoc.getColumn();
+ StringRef inputFilePath = fileLoc.getFilename().getValue();
+ fileAttr = LLVM::DIFileAttr::get(
+ context, llvm::sys::path::filename(inputFilePath),
+ llvm::sys::path::parent_path(inputFilePath));
+ }
+ if (!compileUnitAttr) {
+ compileUnitAttr = LLVM::DICompileUnitAttr::get(
+ context, llvm::dwarf::DW_LANG_C, fileAttr,
+ StringAttr::get(context, "MLIR"), /*isOptimized=*/true,
+ LLVM::DIEmissionKind::LineTablesOnly);
+ }
+ auto subroutineTypeAttr =
+ LLVM::DISubroutineTypeAttr::get(context, llvm::dwarf::DW_CC_normal, {});
+
+ StringAttr funcNameAttr = llvmFunc.getNameAttr();
+ auto subprogramAttr =
+ LLVM::DISubprogramAttr::get(context, compileUnitAttr, fileAttr,
+ funcNameAttr, funcNameAttr, fileAttr,
+ /*line=*/line,
+ /*scopeline=*/col,
+ LLVM::DISubprogramFlags::Definition |
+ LLVM::DISubprogramFlags::Optimized,
+ subroutineTypeAttr);
+ llvmFunc->setLoc(FusedLoc::get(context, {loc}, subprogramAttr));
+ }
+};
+
+} // end anonymous namespace
+
+std::unique_ptr<Pass> mlir::LLVM::createDIScopeForLLVMFuncOpPass() {
+ return std::make_unique<DIScopeForLLVMFuncOp>();
+}
\ No newline at end of file
--- /dev/null
+// RUN: mlir-opt %s --pass-pipeline="builtin.module(llvm.func(ensure-debug-info-scope-on-llvm-func))" --split-input-file --mlir-print-debuginfo | FileCheck %s
+
+
+
+// CHECK-LABEL: llvm.func @func_no_debug()
+// CHECK: llvm.return loc(#loc
+// CHECK: loc(#loc[[LOC:[0-9]+]])
+// CHECK: #di_file = #llvm.di_file<"<unknown>" in "">
+// CHECK: #di_subprogram = #llvm.di_subprogram<compileUnit = #di_compile_unit, scope = #di_file, name = "func_no_debug", linkageName = "func_no_debug", file = #di_file, line = 1, scopeLine = 1, subprogramFlags = "Definition|Optimized", type = #di_subroutine_type>
+// CHECK: #loc[[LOC]] = loc(fused<#di_subprogram>
+llvm.func @func_no_debug() {
+ llvm.return loc(unknown)
+} loc(unknown)
+
+
+// -----
+
+// Test that existing debug info is not overwritten.
+// CHECK-LABEL: llvm.func @func_with_debug()
+// CHECK: llvm.return loc(#loc
+// CHECK: loc(#loc[[LOC:[0-9]+]])
+// CHECK: #di_file = #llvm.di_file<"<unknown>" in "">
+// CHECK: #di_subprogram = #llvm.di_subprogram<compileUnit = #di_compile_unit, scope = #di_file, name = "func_with_debug", linkageName = "func_with_debug", file = #di_file, line = 42, scopeLine = 42, subprogramFlags = "Definition|Optimized", type = #di_subroutine_type>
+// CHECK: #loc[[LOC]] = loc(fused<#di_subprogram>
+module {
+ llvm.func @func_with_debug() {
+ llvm.return loc(#loc1)
+ } loc(#loc2)
+} loc(#loc)
+#di_file = #llvm.di_file<"<unknown>" in "">
+#di_subroutine_type = #llvm.di_subroutine_type<callingConvention = DW_CC_normal>
+#loc = loc("foo":0:0)
+#loc1 = loc(unknown)
+#di_compile_unit = #llvm.di_compile_unit<sourceLanguage = DW_LANG_C, file = #di_file, producer = "MLIR", isOptimized = true, emissionKind = LineTablesOnly>
+#di_subprogram = #llvm.di_subprogram<compileUnit = #di_compile_unit, scope = #di_file, name = "func_with_debug", linkageName = "func_with_debug", file = #di_file, line = 42, scopeLine = 42, subprogramFlags = "Definition|Optimized", type = #di_subroutine_type>
+#loc2 = loc(fused<#di_subprogram>[#loc1])
+
+// -----
+
+// Test that a compile unit present on a module op is used for the function.
+// CHECK-LABEL: llvm.func @propagate_compile_unit()
+// CHECK: llvm.return loc(#loc
+// CHECK: loc(#loc[[FUNCLOC:[0-9]+]])
+// CHECK: loc(#loc[[MODULELOC:[0-9]+]])
+// CHECK-DAG: #[[DI_FILE_MODULE:.+]] = #llvm.di_file<"bar.mlir" in "baz">
+// CHECK-DAG: #[[DI_FILE_FUNC:.+]] = #llvm.di_file<"file.mlir" in "">
+// CHECK-DAG: #loc[[FUNCFILELOC:[0-9]+]] = loc("file.mlir":9:8)
+// CHECK-DAG: #di_compile_unit = #llvm.di_compile_unit<sourceLanguage = DW_LANG_C, file = #[[DI_FILE_MODULE]], producer = "MLIR", isOptimized = true, emissionKind = LineTablesOnly>
+// CHECK-DAG: #di_subprogram = #llvm.di_subprogram<compileUnit = #di_compile_unit, scope = #[[DI_FILE_FUNC]], name = "propagate_compile_unit", linkageName = "propagate_compile_unit", file = #[[DI_FILE_FUNC]], line = 9, scopeLine = 8, subprogramFlags = "Definition|Optimized", type = #di_subroutine_type>
+// CHECK-DAG: #loc[[MODULELOC]] = loc(fused<#di_compile_unit>[#loc])
+// CHECK-DAG: #loc[[FUNCLOC]] = loc(fused<#di_subprogram>[#loc[[FUNCFILELOC]]
+module {
+ llvm.func @propagate_compile_unit() {
+ llvm.return loc(unknown)
+ } loc("file.mlir":9:8)
+} loc(#loc)
+#di_file = #llvm.di_file<"bar.mlir" in "baz">
+#di_compile_unit = #llvm.di_compile_unit<sourceLanguage = DW_LANG_C, file = #di_file, producer = "MLIR", isOptimized = true, emissionKind = LineTablesOnly>
+#loc = loc(fused<#di_compile_unit>["foo.mlir":2:1])
projects / platform / upstream / llvm.git / commitdiff