[tvm] branch main updated: [LLVM] Expand tvm::Type to DWARF conversion (#14568)

[kparzysz]

This is an automated email from the ASF dual-hosted git repository.

kparzysz pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git


The following commit(s) were added to refs/heads/main by this push:
     new ca7c3d8a14 [LLVM] Expand tvm::Type to DWARF conversion (#14568)
ca7c3d8a14 is described below

commit ca7c3d8a14dca982635627f03ebbf6f1195e6586
Author: Eric Lunderberg <lunderb...@users.noreply.github.com>
AuthorDate: Tue Apr 11 18:51:01 2023 -0500

    [LLVM] Expand tvm::Type to DWARF conversion (#14568)
    
    * [LLVM] Expand tvm::Type to DWARF conversion
    
    Prior to this commit, DWARF debug symbols for `float32`, `int8`, and
    `int32` could be generated, with other datatypes resulting in an
    error.  This commit expands the range of types with debug symbols to
    allow for any `DLDataTypeCode::kDLInt`, `kDLUint`, or `kDLFloat`,
    regardless of the bitsize.
    
    * Use existing DLDataType2String
    
    * Added unit test
---
 src/target/llvm/codegen_cpu.cc                    | 28 ++++++++++++++++++-----
 tests/python/unittest/test_target_codegen_llvm.py | 20 ++++++++++++++++
 2 files changed, 42 insertions(+), 6 deletions(-)

diff --git a/src/target/llvm/codegen_cpu.cc b/src/target/llvm/codegen_cpu.cc
index 21d2c6ebe0..59575c370f 100644
--- a/src/target/llvm/codegen_cpu.cc
+++ b/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);
diff --git a/tests/python/unittest/test_target_codegen_llvm.py 
b/tests/python/unittest/test_target_codegen_llvm.py
index 3190115aa6..856de716fc 100644
--- a/tests/python/unittest/test_target_codegen_llvm.py
+++ b/tests/python/unittest/test_target_codegen_llvm.py
@@ -1002,5 +1002,25 @@ def test_llvm_assume():
     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()