}
/*! \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 {
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.
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.
*/
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 {
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;
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());
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";
}
// 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);
}
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.