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 <[email protected]>
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")
[email protected]_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()
