[LLVM] Expand tvm::Type to DWARF conversion

src/target/llvm/codegen_cpu.cc

@@ -55,6 +55,7 @@
 
 #include <algorithm>
 #include <memory>
+#include <tuple>
 #include <unordered_map>
 #include <unordered_set>
 
@@ -285,12 +286,7 @@ llvm::DIType* CodeGenCPU::GetDebugType(const Type& ty_tir) {
 llvm::DIType* CodeGenCPU::GetDebugType(const Type& ty_tir, llvm::Type* ty_llvm) {
   if (ty_llvm == t_void_) {
     return nullptr;
-  } else if (ty_llvm == llvm::Type::getFloatTy(*llvm_target_->GetContext())) {
-    return dbg_info_->di_builder_->createBasicType("float", 32, llvm::dwarf::DW_ATE_float);
-  } else if (ty_llvm == t_int8_) {
-    return dbg_info_->di_builder_->createBasicType("int8", 8, llvm::dwarf::DW_ATE_signed);
-  } else if (ty_llvm == t_int32_) {
-    return dbg_info_->di_builder_->createBasicType("int32", 32, llvm::dwarf::DW_ATE_signed);
+
   } else if (ty_llvm->isPointerTy()) {
     auto* ptr_type = ty_tir.as<PointerTypeNode>();
     ICHECK(ptr_type != nullptr || GetRuntimeDataType(ty_tir).is_handle())
@@ -300,6 +296,26 @@ llvm::DIType* CodeGenCPU::GetDebugType(const Type& ty_tir, llvm::Type* ty_llvm)
                                              : nullptr;
     return dbg_info_->di_builder_->createPointerType(pointee_type,
                                                      ty_llvm->getPrimitiveSizeInBits());
+
+  } else if (auto* prim_type = ty_tir.as<PrimTypeNode>()) {
+    DataType dtype = prim_type->dtype;
+    auto dwarf_type = [&]() -> llvm::dwarf::TypeKind {
+      if (dtype.is_bool()) {
+        return llvm::dwarf::DW_ATE_boolean;
+      } else if (dtype.is_float()) {
+        return llvm::dwarf::DW_ATE_float;
+      } else if (dtype.is_int()) {
+        return llvm::dwarf::DW_ATE_signed;
+      } else if (dtype.is_uint()) {
+        return llvm::dwarf::DW_ATE_unsigned;
+      } else {
+        LOG(FATAL) << "No DWARF representation for TIR type " << dtype;
+      }
+    }();
+
+    return dbg_info_->di_builder_->createBasicType(DLDataType2String(dtype),
+                                                   dtype.bits() * dtype.lanes(), dwarf_type);
+
   } else {
     std::string type_str;
     llvm::raw_string_ostream rso(type_str);

tests/python/unittest/test_target_codegen_llvm.py

@@ -1002,5 +1002,25 @@ def tir_assume_func(A: T.Buffer((4, 4), "int32"), B: T.Buffer((14,), "int32")):
     m = tvm.build(mod, [inp, out], target="llvm")
 
 
+@tvm.testing.requires_llvm
+def test_debug_symbol_for_float64():
+    """Check that LLVM can define DWARF debug type for float64
+
+    In previous versions, only specific data types could exist in the
+    function signature.  In this test, the "calling_conv" attribute
+    prevents lowering to the PackedFunc API.
+    """
+
+    @T.prim_func
+    def func(a: T.handle("float64"), b: T.handle("float64"), n: T.int64):
+        T.func_attr({"calling_conv": 2})
+        A = T.Buffer(16, "float64", data=a)
+        B = T.Buffer(16, "float64", data=b)
+        for i in range(n):
+            B[i] = A[i]
+
+    tvm.build(func, target="llvm")
+
+
 if __name__ == "__main__":
     tvm.testing.main()