This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch master
in repository https://gitbox.apache.org/repos/asf/incubator-tvm.git
The following commit(s) were added to refs/heads/master by this push:
new 12e737f Make "none" DataType explicit (#5491)
12e737f is described below
commit 12e737f5289acb1e6ef2ad0aa590bf7b12c679b5
Author: Krzysztof Parzyszek <[email protected]>
AuthorDate: Fri May 1 00:06:06 2020 -0500
Make "none" DataType explicit (#5491)
* Make "none" DataType explicit
The None data type is created when converting an empty string to DataType.
Add functions to create it and recognize it. Convert it to the "void" LLVM
type in LLVM codegen.
* Rename "none" to "void"
* Map VoidType:Type -> Void:DataType in GetRuntimeDataType
* Map Void:DataType -> VoidType:Type in GetType
---
include/tvm/runtime/data_type.h | 20 +++++++++++++++++---
src/target/llvm/codegen_llvm.cc | 3 +++
src/tir/ir/op.cc | 7 ++++---
tests/python/unittest/test_target_codegen_llvm.py | 12 ++++++++++++
4 files changed, 36 insertions(+), 6 deletions(-)
diff --git a/include/tvm/runtime/data_type.h b/include/tvm/runtime/data_type.h
index 44385d6..940818a 100644
--- a/include/tvm/runtime/data_type.h
+++ b/include/tvm/runtime/data_type.h
@@ -107,7 +107,7 @@ class DataType {
}
/*! \return whether type is a handle type. */
bool is_handle() const {
- return code() == DataType::kHandle;
+ return code() == DataType::kHandle && !is_void();
}
/*! \return whether type is a vector type. */
bool is_vector() const {
@@ -117,6 +117,10 @@ class DataType {
bool is_vector_bool() const {
return is_vector() && bits() == 1;
}
+ /*! \return whether type is a Void type. */
+ bool is_void() const {
+ return code() == DataType::kHandle && bits() == 0 && lanes() == 0;
+ }
/*!
* \brief Create a new data type by change lanes to a specified value.
* \param lanes The target number of lanes.
@@ -212,6 +216,13 @@ class DataType {
return DataType(kHandle, bits, lanes);
}
/*!
+ * \brief Construct a Void type.
+ * \return The constructed data type.
+ */
+ static DataType Void() {
+ return DataType(kHandle, 0, 0);
+ }
+ /*!
* \brief Get the corresponding type of TVMShapeIndex.
* \return The type of TVM shape index.
*/
@@ -335,6 +346,9 @@ inline std::ostream& operator<<(std::ostream& os,
DLDataType t) { // NOLINT(*)
if (t.bits == 1 && t.lanes == 1 && t.code == kDLUInt) {
os << "bool"; return os;
}
+ if (DataType(t).is_void()) {
+ return os << "void";
+ }
if (t.code < kTVMCustomBegin) {
os << TypeCode2Str(t.code);
} else {
@@ -361,9 +375,9 @@ inline std::string DLDataType2String(DLDataType t) {
inline DLDataType String2DLDataType(std::string s) {
DLDataType t;
- // handle None type
+ // handle void type
if (s.length() == 0) {
- t.bits = 0; t.lanes = 0; t.code = kTVMOpaqueHandle;
+ t = DataType::Void();
return t;
}
t.bits = 32; t.lanes = 1;
diff --git a/src/target/llvm/codegen_llvm.cc b/src/target/llvm/codegen_llvm.cc
index 86cd5a3..74bda71 100644
--- a/src/target/llvm/codegen_llvm.cc
+++ b/src/target/llvm/codegen_llvm.cc
@@ -309,6 +309,9 @@ llvm::Type* CodeGenLLVM::DTypeToLLVMType(const DataType&
dtype) const {
CHECK_EQ(dtype.lanes(), 1);
return t_void_p_;
}
+ if (dtype.is_void()) {
+ return t_void_;
+ }
llvm::Type* etype = nullptr;
if (dtype.is_int() || dtype.is_uint()) {
etype = llvm::Type::getIntNTy(*ctx_, dtype.bits());
diff --git a/src/tir/ir/op.cc b/src/tir/ir/op.cc
index 4ad244f..6224321 100644
--- a/src/tir/ir/op.cc
+++ b/src/tir/ir/op.cc
@@ -38,6 +38,8 @@ runtime::DataType GetRuntimeDataType(const Type& type) {
return n->dtype;
} else if (type.as<PointerTypeNode>()) {
return DataType::Handle();
+ } else if (IsVoidType(type)) {
+ return DataType::Void();
} else {
LOG(FATAL) << "Type " << type
<< " does not have a corresponding runtime::DataType";
@@ -57,9 +59,8 @@ Type GetType(const PrimExpr& expr) {
}
// Default: return the type indicated by the dtype.
runtime::DataType dtype = expr.dtype();
- // These types already implies the specific type.
- if (dtype.is_int() || dtype.is_uint() || dtype.is_float()) {
- return PrimType(dtype);
+ if (dtype.is_void()) {
+ return VoidType();
}
return PrimType(dtype);
}
diff --git a/tests/python/unittest/test_target_codegen_llvm.py
b/tests/python/unittest/test_target_codegen_llvm.py
index a7e1e57..c659172 100644
--- a/tests/python/unittest/test_target_codegen_llvm.py
+++ b/tests/python/unittest/test_target_codegen_llvm.py
@@ -43,6 +43,18 @@ def test_llvm_intrin():
fcode = tvm.build(mod, None, "llvm")
+def test_llvm_void_intrin():
+ ib = tvm.tir.ir_builder.create()
+ A = ib.pointer("uint8", name="A")
+ # Create an intrinsic that returns void.
+ x = tvm.tir.call_llvm_intrin('', 'llvm.va_start', tvm.tir.const(1,
'uint32'), A)
+ ib.emit(x)
+ body = ib.get()
+ mod = tvm.IRModule.from_expr(
+ tvm.tir.PrimFunc([A], body).with_attr("global_symbol", "main"))
+ fcode = tvm.build(mod, None, "llvm")
+
+
def test_llvm_overloaded_intrin():
# Name lookup for overloaded intrinsics in LLVM 4- requires a name
# that includes the overloaded types.
