Add a pass that builds a debug info scope for LLVMFuncOp (adding a DISubprogramAttr)

[mehdi_amini]

Summary

This may be seen as a hack, but it allows for any piece of MLIR to be able
to end up with DWARF debug info through LLVM.
Assuming the operations in the function have location such as FileLineCol,
this provides backtraces with line tables and allows to step in a debugger.
That makes this pass a perfect companion to -snapshot-op-locations

It was also the default behavior of MLIR to LLVM IR translation until MLIR
got support for proper debug info attributes.

[mehdi_amini]

Add a test

[mehdi_amini]

add end line to test file

[rriddle]

Nice

mlir/lib/Dialect/LLVMIR/Transforms/DIScopeForLLVMFuncOp.cpp

44

Can you spell out some of these autos?

[mehdi_amini]

replace auto with qualified type

mlir/examples/toy/Ch6/toyc.cpp

 //===----------------------------------------------------------------------===//
 
 #include "toy/Dialect.h"
 #include "toy/MLIRGen.h"
 #include "toy/Parser.h"
 #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"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser/Parser.h"
if (enableOpt) {
       optPM.addPass(mlir::createLoopFusionPass());
       optPM.addPass(mlir::createAffineScalarReplacementPass());
     }
   }
 
   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)))
     return 4;
   return 0;
 }
 
 int dumpAST() {

mlir/examples/toy/Ch7/toyc.cpp

 //===----------------------------------------------------------------------===//
 
 #include "toy/Dialect.h"
 #include "toy/MLIRGen.h"
 #include "toy/Parser.h"
 #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"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser/Parser.h"
if (enableOpt) {
       optPM.addPass(mlir::createLoopFusionPass());
       optPM.addPass(mlir::createAffineScalarReplacementPass());
     }
   }
 
   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)))
     return 4;
   return 0;
 }
 
 int dumpAST() {

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h

 #include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"
 #include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
 #include "mlir/Pass/Pass.h"
 
 namespace mlir {
 
 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"
 
 } // namespace LLVM
 } // namespace mlir
 
 #endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_PASSES_H

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td

def LLVMRequestCWrappers
   let constructor = "::mlir::LLVM::createRequestCWrappersPass()";
 }
 
 def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> {
   let summary = "Optimize NVVM IR";
   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

mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt

 add_mlir_dialect_library(MLIRLLVMIRTransforms
+  DIScopeForLLVMFuncOp.cpp
   LegalizeForExport.cpp
   OptimizeForNVVM.cpp
   RequestCWrappers.cpp
 
   DEPENDS
   MLIRLLVMPassIncGen
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRFuncDialect
   MLIRLLVMDialect
   MLIRPass
   MLIRTransforms
   MLIRNVVMDialect
   )

mlir/lib/Dialect/LLVMIR/Transforms/DIScopeForLLVMFuncOp.cpp

+//===- 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();
+    auto loc = llvmFunc.getLoc();
+    if (loc->findInstanceOf<mlir::FusedLocWith<LLVM::DISubprogramAttr>>())
+      return;
+
+    auto *context = &getContext();
+
+    // To find a DICompileUnitAttr attached to a parent (the module for
+    // example), otherwise create a default one.
+    LLVM::DICompileUnitAttr compileUnitAttr;
+    if (auto 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;
+    auto 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

mlir/test/Dialect/LLVMIR/add-debuginfo-func-scope.mlir

+// 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])
 No newline at end of file