Fix location emission

.gitignore

@@ -271,3 +271,6 @@ gallery/how_to/work_with_microtvm/micro_tvmc.py
 
 # Printed TIR code on disk
 *.tir
+
+# GDB history file
+.gdb_history

gallery/tutorial/debug_tir.py

@@ -0,0 +1,74 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements.  See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership.  The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License.  You may obtain a copy of the License at
+#
+#   http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied.  See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""
+.. _tutorial-topi:
+
+Debugging TIR
+=============
+
+"""
+
+# sphinx_gallery_start_ignore
+from tvm import testing
+
+testing.utils.install_request_hook(depth=3)
+# sphinx_gallery_end_ignore
+
+import tvm
+import tvm.testing
+import numpy as np
+from tvm.script import tir as T
+
+# Installing dependencies
+#
+#     .. code-block:: bash
+#
+#       pip install -q tensorflow
+#       apt-get -qq install curl
+
+
+@tvm.script.ir_module
+class MyModule:
+    @T.prim_func
+    def main(a: T.handle, b: T.handle):
+        # We exchange data between function by handles, which are similar to pointer.
+        T.func_attr({"global_symbol": "main", "tir.noalias": True})
+        # Create buffer from handles.
+        A = T.match_buffer(a, (8,), dtype="float32")
+        B = T.match_buffer(b, (8,), dtype="float32")
+        for i in range(8):
+            # A block is an abstraction for computation.
+            with T.block("B"):
+                # Define a spatial block iterator and bind it to value i.
+                vi = T.axis.spatial(8, i)
+                assert 1 == 0, "Some numbers"
+                B[vi] = A[vi] + 1.0
+
+
+print("Actually starting ------")
+with tvm.transform.PassContext(opt_level=3, config={"tir.enable_debug": True}):
+    runtime_module = tvm.build(MyModule, target="llvm")
+
+# print(runtime_module.get_source())
+print(type(runtime_module))
+
+a = tvm.nd.array(np.arange(8).astype("float32"))
+b = tvm.nd.array(np.zeros((8,)).astype("float32"))
+print("EXECUTING ------")
+runtime_module(a, b)
+print(a)
+print(b)

src/printer/tir_text_printer_debug.cc

@@ -25,25 +25,64 @@
 
 #include "tir_text_printer_debug.h"
 
+#include <optional>
 #include <string>
 
 #include "text_printer.h"
 
 namespace tvm {
 namespace tir {
 
-std::string span_text(const Span& span) {
+std::optional<std::string> span_text(const Span& span) {
   if (!span.defined()) {
-    return "missing";
+    return std::nullopt;
+  }
+
+  std::string source("main.tir");
+  if (span->source_name.defined() && span->source_name->name.get()) {
+    source = span->source_name->name;
   }
-  std::string source("file");
   return source + ":" + std::to_string(span->line) + ":" + std::to_string(span->column);
 }
 
+template <typename ObjectPtr>
+void add_all_relevant_lines(const std::vector<std::tuple<const ObjectPtr*, size_t>>& data,
+                            size_t current_line, Doc* output) {
+  ICHECK(output) << "output must be a valid Doc";
+  for (const auto& item : data) {
+    if (std::get<1>(item) != current_line - 1) {
+      // Item is not relevant for this line, skip it
+      continue;
+    }
+
+    // Print out the item's span info if present
+    auto text = span_text(std::get<0>(item)->span);
+    if (text.has_value()) {
+      *output << *text;
+    } else {
+      *output << "missing";
+    }
+    *output << ", ";
+  }
+}
+
 Doc TIRTextPrinterDebug::NewLine() {
   current_line_ += 1;
 
-  return TIRTextPrinter::NewLine();
+  if (!show_spans_) {
+    return TIRTextPrinter::NewLine();
+  }
+
+  Doc output;
+
+  output << " [";
+
+  add_all_relevant_lines(exprs_by_line_, current_line_, &output);
+  add_all_relevant_lines(stmts_by_line_, current_line_, &output);
+
+  output << "]" << TIRTextPrinter::NewLine();
+
+  return output;
 }
 
 #define X(TypeName)                                               \

src/printer/tir_text_printer_debug.h

@@ -37,7 +37,8 @@ namespace tir {
 
 class TIRTextPrinterDebug : public TIRTextPrinter {
  public:
-  TIRTextPrinterDebug() : TIRTextPrinter(false, &meta_), current_line_(1) {}
+  explicit TIRTextPrinterDebug(bool show_spans)
+      : TIRTextPrinter(false, &meta_), current_line_(1), show_spans_(show_spans) {}
 
   std::vector<std::tuple<const PrimExprNode*, size_t>> GetExprsByLine() const {
     return exprs_by_line_;
@@ -61,6 +62,9 @@ class TIRTextPrinterDebug : public TIRTextPrinter {
   // Line that the printer is currently printing
   size_t current_line_;
 
+  // Whether to include spans relevant to each line before a newline or not
+  bool show_spans_;
+
   // Record of all stmts and exprs and their corresponding line
   std::vector<std::tuple<const StmtNode*, size_t>> stmts_by_line_;
   std::vector<std::tuple<const PrimExprNode*, size_t>> exprs_by_line_;

src/target/llvm/codegen_cpu.cc

@@ -207,6 +207,8 @@ llvm::DISubprogram* CodeGenCPU::CreateDebugFunction(const PrimFunc& f) {
       /*Flags=*/llvm::DINode::FlagPrototyped, /*isOptimized=*/true);
 #endif
   return DIFunction;
+#else
+  return nullptr;
 #endif
 }
 
@@ -952,7 +954,6 @@ llvm::Value* CodeGenCPU::CreateCallPacked(const CallNode* op, bool use_string_lo
 }
 
 llvm::Value* CodeGenCPU::CreateCallTracePacked(const CallNode* op) {
-  EmitDebugLocation(op);
   ICHECK_EQ(op->args.size(), 6U);
   PackedCall pc = MakeCallPackedLowered(op->args, op->dtype, op->args[3].as<IntImmNode>()->value,
                                         op->args[4].as<IntImmNode>()->value, true);
@@ -1388,7 +1389,6 @@ void CodeGenCPU::AddStartupFunction() {
 }
 
 llvm::Value* CodeGenCPU::CreateIntrinsic(const CallNode* op) {
-  EmitDebugLocation(op);
   if (op->op.same_as(builtin::tvm_call_packed_lowered())) {
     return CreateCallPacked(op, true /* use_string_lookup */);
   } else if (op->op.same_as(builtin::tvm_call_trace_packed_lowered())) {