diff --git a/include/tvm/ir/node_functor.h b/include/tvm/ir/node_functor.h index c27dc9ec4b87..c7be2188d314 100644 --- a/include/tvm/ir/node_functor.h +++ b/include/tvm/ir/node_functor.h @@ -23,7 +23,7 @@ #ifndef TVM_IR_NODE_FUNCTOR_H_ #define TVM_IR_NODE_FUNCTOR_H_ -#include +#include #include #include diff --git a/include/tvm/ir/op.h b/include/tvm/ir/op.h index d8b2cf07d11f..dc8f99cd4789 100644 --- a/include/tvm/ir/op.h +++ b/include/tvm/ir/op.h @@ -25,6 +25,7 @@ #ifndef TVM_IR_OP_H_ #define TVM_IR_OP_H_ +#include #include #include #include @@ -32,7 +33,6 @@ #include #include #include -#include #include #include diff --git a/include/tvm/relax/op_attr_types.h b/include/tvm/relax/op_attr_types.h index 2e686035b20c..1fd9b45c323c 100644 --- a/include/tvm/relax/op_attr_types.h +++ b/include/tvm/relax/op_attr_types.h @@ -53,6 +53,12 @@ enum OpPatternKind { kOpaque = 8 }; +/*! + * \brief Packed function implementation for operators. The relax operator will be lowered to + * this packed function call during codegen. + */ +using FCallPacked = ffi::String; + /*! * \brief Infer output struct info given the call * @@ -61,12 +67,6 @@ enum OpPatternKind { */ using FInferStructInfo = ffi::TypedFunction; -/*! - * \brief Packed function implementation for operators. The relax operator will be lowered to - * this packed function call during codegen. - */ -using FCallPacked = ffi::String; - /*! * \brief The function type of a normalization function. * diff --git a/include/tvm/relax/utils.h b/include/tvm/relax/utils.h index 792f7dd11f90..bfbcaa069818 100644 --- a/include/tvm/relax/utils.h +++ b/include/tvm/relax/utils.h @@ -25,9 +25,9 @@ #define TVM_RELAX_UTILS_H_ #include +#include #include #include -#include namespace tvm { namespace relax { diff --git a/include/tvm/runtime/data_type.h b/include/tvm/runtime/data_type.h index 67fe50350d2f..9f230cac824e 100644 --- a/include/tvm/runtime/data_type.h +++ b/include/tvm/runtime/data_type.h @@ -26,8 +26,8 @@ #include #include +#include #include -#include #include #include @@ -36,8 +36,6 @@ namespace tvm { namespace runtime { -using tvm_index_t = ffi::Shape::index_type; - /*! * \brief Runtime primitive data type. * @@ -404,10 +402,10 @@ class DataType { * \return The type of TVM shape index. */ static DataType ShapeIndex() { - if (std::is_signed::value) { - return DataType::Int(sizeof(tvm_index_t) * 8); + if (std::is_signed::value) { + return DataType::Int(sizeof(ffi::Shape::index_type) * 8); } else { - return DataType::UInt(sizeof(tvm_index_t) * 8); + return DataType::UInt(sizeof(ffi::Shape::index_type) * 8); } } @@ -451,9 +449,6 @@ inline bool TypeEqual(DLDataType lhs, DLDataType rhs) { return lhs.code == rhs.code && lhs.bits == rhs.bits && lhs.lanes == rhs.lanes; } -using ffi::DLDataTypeToString; -using ffi::StringToDLDataType; - inline std::ostream& operator<<(std::ostream& os, const DataType& dtype) { // NOLINT(*) return os << dtype.operator DLDataType(); } diff --git a/include/tvm/runtime/device_api.h b/include/tvm/runtime/device_api.h index 47607c5b8875..be5d4e89005b 100644 --- a/include/tvm/runtime/device_api.h +++ b/include/tvm/runtime/device_api.h @@ -25,10 +25,10 @@ #define TVM_RUNTIME_DEVICE_API_H_ #include +#include #include #include #include -#include #include /*! diff --git a/include/tvm/runtime/logging.h b/include/tvm/runtime/logging.h index d051a01da4c4..68718acc7cb9 100644 --- a/include/tvm/runtime/logging.h +++ b/include/tvm/runtime/logging.h @@ -60,8 +60,6 @@ namespace tvm { namespace runtime { -using ffi::EnvErrorAlreadySet; - /*! \brief Internal implementation */ namespace detail { // Provide support for customized logging. diff --git a/include/tvm/runtime/tensor.h b/include/tvm/runtime/tensor.h index f0ea3508bc91..33a78a48d6ae 100644 --- a/include/tvm/runtime/tensor.h +++ b/include/tvm/runtime/tensor.h @@ -42,17 +42,12 @@ namespace tvm { namespace runtime { -using ffi::GetDataSize; -using ffi::IsAligned; -using ffi::IsContiguous; - /*! * \brief Managed Tensor. * The array is backed by reference counted blocks. */ class Tensor : public tvm::ffi::Tensor { public: - using Container = ffi::TensorObj; Tensor() = default; /*! * \brief constructor. diff --git a/include/tvm/runtime/vm/bytecode.h b/include/tvm/runtime/vm/bytecode.h index 5a60febf8443..0f1927e0cbcb 100644 --- a/include/tvm/runtime/vm/bytecode.h +++ b/include/tvm/runtime/vm/bytecode.h @@ -24,8 +24,8 @@ #ifndef TVM_RUNTIME_VM_BYTECODE_H_ #define TVM_RUNTIME_VM_BYTECODE_H_ +#include #include -#include #include #include diff --git a/include/tvm/s_tir/random_engine.h b/include/tvm/s_tir/random_engine.h index d594e1ba0c35..0acfd50fbed2 100644 --- a/include/tvm/s_tir/random_engine.h +++ b/include/tvm/s_tir/random_engine.h @@ -23,7 +23,7 @@ */ #ifndef TVM_S_TIR_RANDOM_ENGINE_H_ #define TVM_S_TIR_RANDOM_ENGINE_H_ -#include +#include #include #include diff --git a/include/tvm/script/printer/doc.h b/include/tvm/script/printer/doc.h index 0430c8d8f172..8803e846c08f 100644 --- a/include/tvm/script/printer/doc.h +++ b/include/tvm/script/printer/doc.h @@ -63,7 +63,7 @@ class DocNode : public ffi::Object { * this Doc is generated, in order to position the diagnostic * message. */ - mutable ffi::Array source_paths; + mutable ffi::Array source_paths; static void RegisterReflection() { namespace refl = tvm::ffi::reflection; @@ -308,7 +308,7 @@ class LiteralDoc : public ExprDoc { * \param p The object path */ static LiteralDoc DataType(const runtime::DataType& v, const ffi::Optional& p) { - std::string dtype = v.is_void() ? "void" : runtime::DLDataTypeToString(v); + std::string dtype = v.is_void() ? "void" : ffi::DLDataTypeToString(v); return LiteralDoc::Str(dtype, p); } /*! diff --git a/include/tvm/target/codegen.h b/include/tvm/target/codegen.h index 32baf16d4a3a..0274592d6604 100644 --- a/include/tvm/target/codegen.h +++ b/include/tvm/target/codegen.h @@ -34,10 +34,6 @@ namespace tvm { /*! \brief namespace for target translation and codegen. */ namespace codegen { -// use packed function from runtime. -using ffi::Any; -using ffi::Function; -using ffi::PackedArgs; /*! * \brief Build a module from array of lowered function. diff --git a/include/tvm/target/virtual_device.h b/include/tvm/target/virtual_device.h index 7829bb61d4ad..5ff282adb68b 100644 --- a/include/tvm/target/virtual_device.h +++ b/include/tvm/target/virtual_device.h @@ -36,7 +36,7 @@ namespace tvm { /*! - * Abstract label for an area of memory. + * \brief Abstract label for an area of memory. * * Currently uninterpreted and arbitrary. Likely to be replaced by a structured representation * of a memory pool in the future. Please try to use this alias instead of ffi::String to aid future diff --git a/include/tvm/topi/detail/constant_utils.h b/include/tvm/topi/detail/constant_utils.h index a77177984734..07df5c470bf4 100644 --- a/include/tvm/topi/detail/constant_utils.h +++ b/include/tvm/topi/detail/constant_utils.h @@ -25,6 +25,7 @@ #define TVM_TOPI_DETAIL_CONSTANT_UTILS_H_ #include +#include #include #include #include diff --git a/src/arith/analyzer.cc b/src/arith/analyzer.cc index 4650cfb43b1c..8bce80f4ef8f 100644 --- a/src/arith/analyzer.cc +++ b/src/arith/analyzer.cc @@ -24,6 +24,7 @@ #include #include #include +#include #include #include diff --git a/src/arith/const_fold.h b/src/arith/const_fold.h index 8464443118f9..91db540f2e82 100644 --- a/src/arith/const_fold.h +++ b/src/arith/const_fold.h @@ -25,6 +25,7 @@ #define TVM_ARITH_CONST_FOLD_H_ #include +#include #include #include diff --git a/src/arith/domain_touched.cc b/src/arith/domain_touched.cc index b218eb9e57e9..977ea779f450 100644 --- a/src/arith/domain_touched.cc +++ b/src/arith/domain_touched.cc @@ -23,6 +23,7 @@ */ #include #include +#include #include #include #include diff --git a/src/arith/int_set.cc b/src/arith/int_set.cc index 66c148d47857..6b3e2b953270 100644 --- a/src/arith/int_set.cc +++ b/src/arith/int_set.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include diff --git a/src/arith/solve_linear_equation.cc b/src/arith/solve_linear_equation.cc index 8d6b58351359..4b6ac036e8bb 100644 --- a/src/arith/solve_linear_equation.cc +++ b/src/arith/solve_linear_equation.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include diff --git a/src/ir/diagnostic.cc b/src/ir/diagnostic.cc index 18fa77f62658..f234963c9261 100644 --- a/src/ir/diagnostic.cc +++ b/src/ir/diagnostic.cc @@ -25,6 +25,7 @@ #include #include #include +#include namespace tvm { diff --git a/src/ir/instrument.cc b/src/ir/instrument.cc index ad47ccf2ed44..e88713a50632 100644 --- a/src/ir/instrument.cc +++ b/src/ir/instrument.cc @@ -26,8 +26,10 @@ #include #include #include +#include #include +#include #include namespace tvm { diff --git a/src/ir/repr.cc b/src/ir/repr.cc index cf15ecbbf685..addbd33209f3 100644 --- a/src/ir/repr.cc +++ b/src/ir/repr.cc @@ -24,7 +24,7 @@ * The legacy ReprPrinter has been replaced by ffi::ReprPrint. This file: * - Implements the Dump() debug helpers (they call ffi::ReprPrint). * - Registers node.AsRepr (for backward Python compatibility) via ffi::ReprPrint. - * - Registers __ffi_repr__ hooks for AccessPath and AccessStep. + * - Registers __ffi_repr__ hooks for ffi::reflection::AccessPath and AccessStep. */ #include #include @@ -48,7 +48,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { // Python's tvm.runtime._ffi_node_api sets __object_repr__ = AsRepr via init_ffi_api. refl::GlobalDef().def("node.AsRepr", [](ffi::Any obj) -> ffi::String { return ffi::ReprPrint(obj); }); - // Register __ffi_repr__ for AccessPath/AccessStep so that ffi.ReprPrint + // Register __ffi_repr__ for ffi::reflection::AccessPath/AccessStep so that ffi.ReprPrint // uses the concise ".field[idx]" format. // // AccessStep: format one step fragment (e.g. ".field", "[0]", "[key]?"). @@ -79,7 +79,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { } return os.str(); }); - // AccessPath: recurse through parent via fn_repr rather than walking the + // ffi::reflection::AccessPath: recurse through parent via fn_repr rather than walking the // linked list manually. Root (no step) emits ""; each non-root node // prepends its parent's repr and appends the current step's repr. refl::TypeAttrDef().def( diff --git a/src/ir/script_printer.cc b/src/ir/script_printer.cc index ea0c3d031eae..dc1f035f5cb3 100644 --- a/src/ir/script_printer.cc +++ b/src/ir/script_printer.cc @@ -28,8 +28,6 @@ namespace tvm { -using AccessPath = ffi::reflection::AccessPath; - TVM_FFI_STATIC_INIT_BLOCK() { PrinterConfigNode::RegisterReflection(); } TVMScriptPrinter::FType& TVMScriptPrinter::vtable() { @@ -94,11 +92,13 @@ PrinterConfig::PrinterConfig(ffi::Map config_dict) { } if (auto v = config_dict.Get("path_to_underline")) { n->path_to_underline = - Downcast>>(v).value_or(ffi::Array()); + Downcast>>(v).value_or( + ffi::Array()); } if (auto v = config_dict.Get("path_to_annotate")) { - n->path_to_annotate = Downcast>>(v).value_or( - ffi::Map()); + n->path_to_annotate = + Downcast>>(v).value_or( + ffi::Map()); } if (auto v = config_dict.Get("obj_to_underline")) { n->obj_to_underline = Downcast>>(v).value_or( diff --git a/src/ir/source_map.cc b/src/ir/source_map.cc index e61cd8db753d..96f4b8fda973 100644 --- a/src/ir/source_map.cc +++ b/src/ir/source_map.cc @@ -25,6 +25,7 @@ #include #include #include +#include #include diff --git a/src/ir/structural_hash.cc b/src/ir/structural_hash.cc index e1903871e175..9f33c2f50a03 100644 --- a/src/ir/structural_hash.cc +++ b/src/ir/structural_hash.cc @@ -60,9 +60,9 @@ TVM_FFI_STATIC_INIT_BLOCK() { return rtmod; }); - refl::TypeAttrDef() + refl::TypeAttrDef() .def("__data_to_json__", - [](const runtime::Tensor::Container* node) { + [](const ffi::TensorObj* node) { std::string result; support::BytesOutStream mstrm(&result); support::Base64OutStream b64strm(&mstrm); diff --git a/src/ir/transform.cc b/src/ir/transform.cc index c301037732d2..82c3f13c5618 100644 --- a/src/ir/transform.cc +++ b/src/ir/transform.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include diff --git a/src/relax/analysis/graph_partitioner.h b/src/relax/analysis/graph_partitioner.h index 7084139e299b..1ae994842b1f 100644 --- a/src/relax/analysis/graph_partitioner.h +++ b/src/relax/analysis/graph_partitioner.h @@ -28,6 +28,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/analysis/layout_transformation.cc b/src/relax/analysis/layout_transformation.cc index 4fa4ed48534e..dcee90c9a7ec 100644 --- a/src/relax/analysis/layout_transformation.cc +++ b/src/relax/analysis/layout_transformation.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include diff --git a/src/relax/analysis/well_formed.cc b/src/relax/analysis/well_formed.cc index 875489d43815..ec654c6eb0ef 100644 --- a/src/relax/analysis/well_formed.cc +++ b/src/relax/analysis/well_formed.cc @@ -72,6 +72,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/backend/contrib/clml/codegen.cc b/src/relax/backend/contrib/clml/codegen.cc index 24122ddf5241..eaa57f8315e4 100644 --- a/src/relax/backend/contrib/clml/codegen.cc +++ b/src/relax/backend/contrib/clml/codegen.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/backend/contrib/codegen_json/codegen_json.h b/src/relax/backend/contrib/codegen_json/codegen_json.h index bb0f80b82fcd..d7874eb84679 100644 --- a/src/relax/backend/contrib/codegen_json/codegen_json.h +++ b/src/relax/backend/contrib/codegen_json/codegen_json.h @@ -91,7 +91,7 @@ class OpAttrExtractor { void Visit(const char* key, DataType* value) { if (!value->is_void()) { - SetNodeAttr(key, ffi::String(runtime::DLDataTypeToString(*value))); + SetNodeAttr(key, ffi::String(ffi::DLDataTypeToString(*value))); } else { SetNodeAttr(key, ffi::String("")); } diff --git a/src/relax/backend/contrib/cutlass/codegen.cc b/src/relax/backend/contrib/cutlass/codegen.cc index 4a91dcab9cff..91840f6936e5 100644 --- a/src/relax/backend/contrib/cutlass/codegen.cc +++ b/src/relax/backend/contrib/cutlass/codegen.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/backend/contrib/nnapi/codegen.cc b/src/relax/backend/contrib/nnapi/codegen.cc index 9d85d5ef82d1..757570e69ad5 100644 --- a/src/relax/backend/contrib/nnapi/codegen.cc +++ b/src/relax/backend/contrib/nnapi/codegen.cc @@ -67,7 +67,7 @@ class CollectFromCompositeFunctionBody : public ExprVisitor { void SetAstypeAttribute(const CallNode* call_node) { const auto* astype_attrs = call_node->attrs.as(); TVM_FFI_ICHECK(astype_attrs); - node_->SetAttr("astype_dtype", ffi::String(runtime::DLDataTypeToString(astype_attrs->dtype))); + node_->SetAttr("astype_dtype", ffi::String(ffi::DLDataTypeToString(astype_attrs->dtype))); } void SetMeanAttribute(const CallNode* call_node) { diff --git a/src/relax/backend/contrib/tensorrt/codegen.cc b/src/relax/backend/contrib/tensorrt/codegen.cc index 011b8138a595..2be214ed941c 100644 --- a/src/relax/backend/contrib/tensorrt/codegen.cc +++ b/src/relax/backend/contrib/tensorrt/codegen.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/backend/vm/codegen_vm_tir.cc b/src/relax/backend/vm/codegen_vm_tir.cc index b143b1473d05..10da7d983619 100644 --- a/src/relax/backend/vm/codegen_vm_tir.cc +++ b/src/relax/backend/vm/codegen_vm_tir.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include diff --git a/src/relax/ir/block_builder.cc b/src/relax/ir/block_builder.cc index b2a70a39c266..1061c02eb1f8 100644 --- a/src/relax/ir/block_builder.cc +++ b/src/relax/ir/block_builder.cc @@ -33,6 +33,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/ir/dataflow_matcher.cc b/src/relax/ir/dataflow_matcher.cc index 37e90c530930..22e3a7bbc31a 100644 --- a/src/relax/ir/dataflow_matcher.cc +++ b/src/relax/ir/dataflow_matcher.cc @@ -33,6 +33,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/ir/transform.cc b/src/relax/ir/transform.cc index a7d7047c3095..4b4c7077c64d 100644 --- a/src/relax/ir/transform.cc +++ b/src/relax/ir/transform.cc @@ -30,6 +30,7 @@ #include #include #include +#include namespace tvm { namespace relax { diff --git a/src/relax/op/tensor/index.cc b/src/relax/op/tensor/index.cc index f954e901a8ea..efa221fd64f3 100644 --- a/src/relax/op/tensor/index.cc +++ b/src/relax/op/tensor/index.cc @@ -26,6 +26,7 @@ #include #include +#include #include #include diff --git a/src/relax/op/tensor/manipulate.cc b/src/relax/op/tensor/manipulate.cc index b8045a12c2cf..c0b82a760d13 100644 --- a/src/relax/op/tensor/manipulate.cc +++ b/src/relax/op/tensor/manipulate.cc @@ -26,6 +26,7 @@ #include #include +#include #include #include diff --git a/src/relax/script/printer/binding.cc b/src/relax/script/printer/binding.cc index ec158a0b6773..d756a82a0e18 100644 --- a/src/relax/script/printer/binding.cc +++ b/src/relax/script/printer/binding.cc @@ -24,7 +24,8 @@ namespace tvm { namespace script { namespace printer { -IfDoc PrintIfExpr(const relax::If& n, const AccessPath& n_p, const IRDocsifier& d, // +IfDoc PrintIfExpr(const relax::If& n, const AccessPath& n_p, + const IRDocsifier& d, // const ffi::Optional& var, const ffi::Optional& ann) { using relax::SeqExpr; ExprDoc cond = d->AsDoc(n->cond, n_p->Attr("cond")); diff --git a/src/relax/transform/bundle_model_params.cc b/src/relax/transform/bundle_model_params.cc index b8b4825e35ba..b4e4f186d19d 100644 --- a/src/relax/transform/bundle_model_params.cc +++ b/src/relax/transform/bundle_model_params.cc @@ -23,12 +23,12 @@ */ #include +#include #include #include #include #include #include -#include #include "utils.h" diff --git a/src/relax/transform/eliminate_common_subexpr.cc b/src/relax/transform/eliminate_common_subexpr.cc index 7be779984ce9..0d0b8de82a1d 100644 --- a/src/relax/transform/eliminate_common_subexpr.cc +++ b/src/relax/transform/eliminate_common_subexpr.cc @@ -32,6 +32,7 @@ #include #include #include +#include #include "../../support/utils.h" diff --git a/src/relax/transform/fold_constant.cc b/src/relax/transform/fold_constant.cc index 934d93edf494..ed28e5dbc8da 100644 --- a/src/relax/transform/fold_constant.cc +++ b/src/relax/transform/fold_constant.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include #include diff --git a/src/relax/transform/fuse_ops.cc b/src/relax/transform/fuse_ops.cc index 5803dad48514..7af1bb0c8a6a 100644 --- a/src/relax/transform/fuse_ops.cc +++ b/src/relax/transform/fuse_ops.cc @@ -35,6 +35,7 @@ #include #include #include +#include #include #include #include diff --git a/src/relax/transform/fuse_tir.cc b/src/relax/transform/fuse_tir.cc index 7e8ca65b0d7c..bb29a798dc4c 100644 --- a/src/relax/transform/fuse_tir.cc +++ b/src/relax/transform/fuse_tir.cc @@ -237,7 +237,7 @@ class FuseTIRBufferSubstitutor : private StmtExprMutator { auto f_mutate_match_buffers = [this](const MatchBufferRegion& match_buffer) { const Buffer& src_buffer = SubstituteBuffer(match_buffer->source->buffer); const Buffer& tgt_buffer = SubstituteAllocatedBuffer(match_buffer->buffer); - Region region = MutateRegion(match_buffer->source->region); + ffi::Array region = MutateRegion(match_buffer->source->region); if (src_buffer.same_as(match_buffer->source->buffer) && tgt_buffer.same_as(match_buffer->buffer) && region.same_as(match_buffer->source->region)) { @@ -252,7 +252,7 @@ class FuseTIRBufferSubstitutor : private StmtExprMutator { auto f_mutate_read_write_region = [this](const BufferRegion& buffer_region) { const Buffer& buffer = SubstituteBuffer(buffer_region->buffer); - const Region& region = MutateRegion(buffer_region->region); + const ffi::Array& region = MutateRegion(buffer_region->region); if (buffer.same_as(buffer_region->buffer) && region.same_as(buffer_region->region)) { return buffer_region; } else { @@ -302,7 +302,7 @@ class FuseTIRBufferSubstitutor : private StmtExprMutator { // However, `A[vi, vj], A[vi, vj + 1]` is not allow for now. // Note: the order of return region should remain the same as the first occurrence of the region ffi::Array ret; - std::unordered_map buffer_region_set; + std::unordered_map> buffer_region_set; for (const BufferRegion& region : regions) { auto it = buffer_region_set.find(region->buffer.get()); @@ -328,7 +328,7 @@ class FuseTIRBufferSubstitutor : private StmtExprMutator { } } - inline Region MutateRegion(const Region& region) { + inline ffi::Array MutateRegion(const ffi::Array& region) { return MutateArray(region, [this](const Range& range) { const PrimExpr& min = this->VisitExpr(range->min); const PrimExpr& extent = this->VisitExpr(range->extent); diff --git a/src/relax/transform/lambda_lift.cc b/src/relax/transform/lambda_lift.cc index 387a22b7b385..6fb78bfb1422 100644 --- a/src/relax/transform/lambda_lift.cc +++ b/src/relax/transform/lambda_lift.cc @@ -23,12 +23,12 @@ */ #include +#include #include #include #include #include #include -#include #include #include diff --git a/src/relax/transform/legalize_ops.cc b/src/relax/transform/legalize_ops.cc index 31e1625526c0..a6d74d91721d 100644 --- a/src/relax/transform/legalize_ops.cc +++ b/src/relax/transform/legalize_ops.cc @@ -30,6 +30,7 @@ #include #include #include +#include #include #include @@ -328,7 +329,7 @@ class LegalizeMutator : public ExprMutator { // Second choice, use a default legalization legalization_func = legalize_map[op]; } else if (call_packed_map.count(op)) { - // Third choice, use an explicit FCallPacked replacement. This does not require the shape + // Third choice, use an explicit ffi::String replacement. This does not require the shape ffi::String packed_func_name = call_packed_map[op]; legalization_func = [packed_func_name](const BlockBuilder& bb, const Call& call) -> Expr { return Call(ExternFunc(packed_func_name), call->args, Attrs(), {GetStructInfo(call)}); diff --git a/src/relax/transform/lift_transform_params.cc b/src/relax/transform/lift_transform_params.cc index 9be91ace1e01..5430e181bca4 100644 --- a/src/relax/transform/lift_transform_params.cc +++ b/src/relax/transform/lift_transform_params.cc @@ -23,13 +23,13 @@ */ #include +#include #include #include #include #include #include #include -#include #include #include diff --git a/src/relax/transform/meta_schedule.cc b/src/relax/transform/meta_schedule.cc index 3d1df9773f5d..a9dd126a3e61 100644 --- a/src/relax/transform/meta_schedule.cc +++ b/src/relax/transform/meta_schedule.cc @@ -24,6 +24,7 @@ #include #include #include +#include #include #include diff --git a/src/runtime/const_loader_module.cc b/src/runtime/const_loader_module.cc index b7ce95dd2dbb..8592d228d0bd 100644 --- a/src/runtime/const_loader_module.cc +++ b/src/runtime/const_loader_module.cc @@ -36,6 +36,7 @@ #include #include #include +#include #include #include diff --git a/src/runtime/contrib/cblas/cblas.cc b/src/runtime/contrib/cblas/cblas.cc index a91db72e5dab..926ce0195245 100644 --- a/src/runtime/contrib/cblas/cblas.cc +++ b/src/runtime/contrib/cblas/cblas.cc @@ -20,10 +20,10 @@ /*! * \file Use external cblas library call. */ +#include #include #include #include -#include extern "C" { #include diff --git a/src/runtime/contrib/cblas/dnnl_blas.cc b/src/runtime/contrib/cblas/dnnl_blas.cc index d6a9baa21bc8..420e244301b2 100644 --- a/src/runtime/contrib/cblas/dnnl_blas.cc +++ b/src/runtime/contrib/cblas/dnnl_blas.cc @@ -20,10 +20,10 @@ /*! * \file Use external cblas library call. */ +#include #include #include #include -#include extern "C" { #include diff --git a/src/runtime/contrib/cblas/mkl.cc b/src/runtime/contrib/cblas/mkl.cc index 59783134157c..60fecc11bd66 100644 --- a/src/runtime/contrib/cblas/mkl.cc +++ b/src/runtime/contrib/cblas/mkl.cc @@ -20,10 +20,10 @@ /*! * \file Use external mkl library call. */ +#include #include #include #include -#include extern "C" { #include diff --git a/src/runtime/contrib/clml/clml_runtime.cc b/src/runtime/contrib/clml/clml_runtime.cc index 5ea6c1398eeb..dd66987bdd11 100644 --- a/src/runtime/contrib/clml/clml_runtime.cc +++ b/src/runtime/contrib/clml/clml_runtime.cc @@ -36,6 +36,7 @@ #include "clml_utils.h" #endif +#include #include namespace tvm { diff --git a/src/runtime/contrib/cublas/cublas.cc b/src/runtime/contrib/cublas/cublas.cc index dcaf93d2da2e..e58ffdeee0ba 100644 --- a/src/runtime/contrib/cublas/cublas.cc +++ b/src/runtime/contrib/cublas/cublas.cc @@ -20,11 +20,11 @@ /*! * \file Use external cblas library call. */ +#include #include #include #include #include -#include #include "../../3rdparty/compiler-rt/builtin_fp16.h" #include "../cblas/gemm_common.h" diff --git a/src/runtime/contrib/cublas/cublas_utils.h b/src/runtime/contrib/cublas/cublas_utils.h index 429e9831146b..ad67eb1ee9e8 100644 --- a/src/runtime/contrib/cublas/cublas_utils.h +++ b/src/runtime/contrib/cublas/cublas_utils.h @@ -28,7 +28,7 @@ #include #include #include -#include +#include #include #if CUDART_VERSION >= 10010 diff --git a/src/runtime/contrib/cudnn/conv_backward.cc b/src/runtime/contrib/cudnn/conv_backward.cc index d26f82645eaf..bfc65baaff93 100644 --- a/src/runtime/contrib/cudnn/conv_backward.cc +++ b/src/runtime/contrib/cudnn/conv_backward.cc @@ -24,6 +24,7 @@ #include #include #include +#include #include "cudnn_utils.h" @@ -78,7 +79,7 @@ void BackwardDataFindAlgo(int format, int dims, int groups, const int pad[], con dx_dim_int64[i] = dx_dim[i]; } SetConvDescriptors(entry_ptr, format, dims, groups, pad, stride, dilation, dx_dim_int64.data(), - w_dim_int64.data(), dy_dim_int64.data(), StringToDLDataType(data_dtype), + w_dim_int64.data(), dy_dim_int64.data(), ffi::StringToDLDataType(data_dtype), conv_dtype); int returned_algo_count = 0; @@ -157,7 +158,7 @@ void BackwardFilterFindAlgo(int format, int dims, int groups, const int pad[], c dw_dim_int64[i] = dw_dim[i]; } SetConvDescriptors(entry_ptr, format, dims, groups, pad, stride, dilation, x_dim_int64.data(), - dw_dim_int64.data(), dy_dim_int64.data(), StringToDLDataType(data_dtype), + dw_dim_int64.data(), dy_dim_int64.data(), ffi::StringToDLDataType(data_dtype), conv_dtype); int returned_algo_count = 0; diff --git a/src/runtime/contrib/cudnn/conv_forward.cc b/src/runtime/contrib/cudnn/conv_forward.cc index 6a5737c183b0..6c6fd7eb4036 100644 --- a/src/runtime/contrib/cudnn/conv_forward.cc +++ b/src/runtime/contrib/cudnn/conv_forward.cc @@ -24,6 +24,7 @@ #include #include #include +#include #include "cudnn_utils.h" @@ -123,7 +124,7 @@ void FindAlgo(int format, int dims, int groups, const int pad[], const int strid y_dim_int64[i] = y_dim[i]; } SetConvDescriptors(entry_ptr, format, dims, groups, pad, stride, dilation, x_dim_int64.data(), - w_dim_int64.data(), y_dim_int64.data(), StringToDLDataType(data_dtype), + w_dim_int64.data(), y_dim_int64.data(), ffi::StringToDLDataType(data_dtype), conv_dtype); int returned_algo_count = 0; diff --git a/src/runtime/contrib/cudnn/cudnn_utils.h b/src/runtime/contrib/cudnn/cudnn_utils.h index 58eac57c679d..91f50dfc1c92 100644 --- a/src/runtime/contrib/cudnn/cudnn_utils.h +++ b/src/runtime/contrib/cudnn/cudnn_utils.h @@ -25,8 +25,8 @@ #define TVM_RUNTIME_CONTRIB_CUDNN_CUDNN_UTILS_H_ #include +#include #include -#include #include diff --git a/src/runtime/contrib/dnnl/dnnl_kernel.h b/src/runtime/contrib/dnnl/dnnl_kernel.h index a407f5589c61..b7a0ec0f7314 100644 --- a/src/runtime/contrib/dnnl/dnnl_kernel.h +++ b/src/runtime/contrib/dnnl/dnnl_kernel.h @@ -25,9 +25,9 @@ #ifndef TVM_RUNTIME_CONTRIB_DNNL_DNNL_KERNEL_H_ #define TVM_RUNTIME_CONTRIB_DNNL_DNNL_KERNEL_H_ +#include #include #include -#include #include diff --git a/src/runtime/contrib/example_npu/example_npu_runtime.cc b/src/runtime/contrib/example_npu/example_npu_runtime.cc index 440a5d9715ec..0408a3fe9acd 100644 --- a/src/runtime/contrib/example_npu/example_npu_runtime.cc +++ b/src/runtime/contrib/example_npu/example_npu_runtime.cc @@ -31,6 +31,7 @@ #include #include +#include #include #include diff --git a/src/runtime/contrib/hipblas/hipblas.cc b/src/runtime/contrib/hipblas/hipblas.cc index b2cc7331117a..eca971e06606 100644 --- a/src/runtime/contrib/hipblas/hipblas.cc +++ b/src/runtime/contrib/hipblas/hipblas.cc @@ -20,10 +20,10 @@ /*! * \file Use external hipblas library call. */ +#include #include #include #include -#include #include "../../3rdparty/compiler-rt/builtin_fp16.h" #include "../cblas/gemm_common.h" diff --git a/src/runtime/contrib/hipblas/hipblas_utils.h b/src/runtime/contrib/hipblas/hipblas_utils.h index a44c984d9a3f..90c8c489d370 100644 --- a/src/runtime/contrib/hipblas/hipblas_utils.h +++ b/src/runtime/contrib/hipblas/hipblas_utils.h @@ -28,7 +28,7 @@ #include #include #include -#include +#include #include diff --git a/src/runtime/contrib/nnapi/nnapi_builder.cc b/src/runtime/contrib/nnapi/nnapi_builder.cc index 044ff1ccd4a8..8491e1a75939 100644 --- a/src/runtime/contrib/nnapi/nnapi_builder.cc +++ b/src/runtime/contrib/nnapi/nnapi_builder.cc @@ -22,7 +22,7 @@ #include "nnapi_builder.h" #include -#include +#include #include #include @@ -138,7 +138,7 @@ NNAPIModelBuilder::~NNAPIModelBuilder() { ANeuralNetworksModel_free(model_); } NNAPIOperand NNAPIModelBuilder::CreateOperandWithValue(const DLTensor& tensor) { NNAPIOperand operand(next_operand_index_++, &tensor); - const size_t operand_data_size = GetDataSize(tensor); + const size_t operand_data_size = ffi::GetDataSize(tensor); TVM_FFI_ICHECK_EQ(ANeuralNetworksModel_addOperand(model_, operand.GetOperandType().Get()), ANEURALNETWORKS_NO_ERROR); diff --git a/src/runtime/contrib/nnapi/nnapi_ops.cc b/src/runtime/contrib/nnapi/nnapi_ops.cc index a6b5a9c221a7..4a8bf5ba97aa 100644 --- a/src/runtime/contrib/nnapi/nnapi_ops.cc +++ b/src/runtime/contrib/nnapi/nnapi_ops.cc @@ -273,7 +273,7 @@ void CastOpConverter::Convert(NNAPIModelBuilder& builder, const JSONGraphNode& n // Extract the dtype attribute and check that the output operand type matches the dtype specified. const auto dtype_str = node.GetAttr("astype_dtype"); - const DLDataType dtype = StringToDLDataType(std::string(dtype_str)); + const DLDataType dtype = ffi::StringToDLDataType(std::string(dtype_str)); TVM_FFI_ICHECK(outputs.size() == 1); const auto output_tensor_type = outputs[0].GetTensorType(); TVM_FFI_ICHECK(TensorTypeFromDLDataType(dtype) == output_tensor_type) diff --git a/src/runtime/contrib/nnapi/nnapi_runtime.cc b/src/runtime/contrib/nnapi/nnapi_runtime.cc index 1939e90992e7..46329f201ac2 100644 --- a/src/runtime/contrib/nnapi/nnapi_runtime.cc +++ b/src/runtime/contrib/nnapi/nnapi_runtime.cc @@ -145,7 +145,7 @@ class NNAPIRuntime : public JSONRuntimeBase { const uint32_t eid = EntryID(nid, j); const auto entry = data_entry_[eid]; - const auto operand_data_size = GetDataSize(*entry); + const auto operand_data_size = ffi::GetDataSize(*entry); TVM_FFI_ICHECK_EQ( ANeuralNetworksExecution_setInput(execution, i, operand.GetOperandType().Get(), entry->data, operand_data_size), @@ -161,7 +161,7 @@ class NNAPIRuntime : public JSONRuntimeBase { const auto eid = EntryID(node); const auto entry = data_entry_[eid]; - const auto operand_data_size = GetDataSize(*entry); + const auto operand_data_size = ffi::GetDataSize(*entry); TVM_FFI_ICHECK_EQ( ANeuralNetworksExecution_setOutput(execution, i, operand.GetOperandType().Get(), entry->data, operand_data_size), diff --git a/src/runtime/contrib/random/mt_random_engine.cc b/src/runtime/contrib/random/mt_random_engine.cc index 64c3ff66a7cf..c01fe9267326 100644 --- a/src/runtime/contrib/random/mt_random_engine.cc +++ b/src/runtime/contrib/random/mt_random_engine.cc @@ -22,9 +22,9 @@ * \brief mt19937 random engine */ #include +#include #include #include -#include #include #include diff --git a/src/runtime/contrib/random/random.cc b/src/runtime/contrib/random/random.cc index 1f3fdf869e99..af94f97ef16f 100644 --- a/src/runtime/contrib/random/random.cc +++ b/src/runtime/contrib/random/random.cc @@ -20,10 +20,10 @@ /*! * \file External random functions for tensor. */ +#include #include #include #include -#include #include diff --git a/src/runtime/contrib/sort/sort.cc b/src/runtime/contrib/sort/sort.cc index 541548d18250..0d072c963846 100644 --- a/src/runtime/contrib/sort/sort.cc +++ b/src/runtime/contrib/sort/sort.cc @@ -22,10 +22,10 @@ */ #include +#include #include #include #include -#include #include #include @@ -334,8 +334,8 @@ void RegisterSort() { "input ndim " << input->ndim; - auto data_dtype = DLDataTypeToString(input->dtype); - auto out_dtype = DLDataTypeToString(output->dtype); + auto data_dtype = ffi::DLDataTypeToString(input->dtype); + auto out_dtype = ffi::DLDataTypeToString(output->dtype); TVM_FFI_ICHECK_EQ(data_dtype, out_dtype); diff --git a/src/runtime/contrib/tensorrt/tensorrt_builder.cc b/src/runtime/contrib/tensorrt/tensorrt_builder.cc index 63d886e520a7..4caa8e383e15 100644 --- a/src/runtime/contrib/tensorrt/tensorrt_builder.cc +++ b/src/runtime/contrib/tensorrt/tensorrt_builder.cc @@ -24,6 +24,7 @@ #include "tensorrt_builder.h" +#include #include #include @@ -227,10 +228,10 @@ nvinfer1::Weights TensorRTBuilder::GetDLTensorAsWeights(const DLTensor* dptr, const auto trt_dtype = (static_cast(dptr->dtype.bits) == 16) ? nvinfer1::DataType::kHALF : nvinfer1::DataType::kFLOAT; - const size_t weight_bytes = GetDataSize(*dptr); + const size_t weight_bytes = ffi::GetDataSize(*dptr); nvinfer1::Weights weight{trt_dtype, nullptr, 0}; size_t count = 1; - for (tvm_index_t i = 0; i < dptr->ndim; ++i) { + for (ffi::Shape::index_type i = 0; i < dptr->ndim; ++i) { count *= dptr->shape[i]; } weight.count = count; diff --git a/src/runtime/contrib/tensorrt/tensorrt_runtime.cc b/src/runtime/contrib/tensorrt/tensorrt_runtime.cc index df8443dd590a..d4fcffd541bb 100644 --- a/src/runtime/contrib/tensorrt/tensorrt_runtime.cc +++ b/src/runtime/contrib/tensorrt/tensorrt_runtime.cc @@ -25,6 +25,7 @@ #include #include #include +#include #include #include diff --git a/src/runtime/cpu_device_api.cc b/src/runtime/cpu_device_api.cc index 9762d2c3b46f..b549182dab31 100644 --- a/src/runtime/cpu_device_api.cc +++ b/src/runtime/cpu_device_api.cc @@ -20,10 +20,10 @@ /*! * \file cpu_device_api.cc */ +#include #include #include #include -#include #include #include diff --git a/src/runtime/cuda/cuda_device_api.cc b/src/runtime/cuda/cuda_device_api.cc index a01d223ff6f3..5de47bd3e431 100644 --- a/src/runtime/cuda/cuda_device_api.cc +++ b/src/runtime/cuda/cuda_device_api.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include #include @@ -454,7 +455,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { break; default: TVM_FFI_THROW(InternalError) - << "Unsupported data type " << runtime::DLDataTypeToString(tensor_dtype); + << "Unsupported data type " << ffi::DLDataTypeToString(tensor_dtype); } break; case DataType::kUInt: @@ -474,7 +475,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { break; default: TVM_FFI_THROW(InternalError) - << "Unsupported data type " << runtime::DLDataTypeToString(tensor_dtype); + << "Unsupported data type " << ffi::DLDataTypeToString(tensor_dtype); } break; case DataType::kFloat: @@ -491,7 +492,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { break; default: TVM_FFI_THROW(InternalError) - << "Unsupported data type " << runtime::DLDataTypeToString(tensor_dtype); + << "Unsupported data type " << ffi::DLDataTypeToString(tensor_dtype); } break; case DataType::kBFloat: @@ -502,7 +503,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { break; default: TVM_FFI_THROW(InternalError) - << "Unsupported data type " << runtime::DLDataTypeToString(tensor_dtype); + << "Unsupported data type " << ffi::DLDataTypeToString(tensor_dtype); } break; case DataType::kFloat8_e4m3fn: @@ -515,7 +516,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { break; default: TVM_FFI_THROW(InternalError) - << "Unsupported data type " << runtime::DLDataTypeToString(tensor_dtype); + << "Unsupported data type " << ffi::DLDataTypeToString(tensor_dtype); } // sanity checks per cuTensorMapEncodeTiled requirements diff --git a/src/runtime/device_api.cc b/src/runtime/device_api.cc index d8aff594ea95..959cd619abbc 100644 --- a/src/runtime/device_api.cc +++ b/src/runtime/device_api.cc @@ -132,7 +132,7 @@ void* DeviceAPI::AllocDataSpace(Device dev, int ndim, const int64_t* shape, DLDa temp.shape = const_cast(shape); temp.strides = nullptr; temp.byte_offset = 0; - size_t size = GetDataSize(temp); + size_t size = ffi::GetDataSize(temp); size_t alignment = GetDataAlignment(temp.dtype); return AllocDataSpace(dev, size, alignment, dtype); } @@ -143,8 +143,8 @@ void* DeviceAPI::AllocDataSpace(Device dev, int ndim, const int64_t* shape, DLDa void DeviceAPI::CopyDataFromTo(DLTensor* from, DLTensor* to, TVMStreamHandle stream) { // by default, we can always redirect to the flat memory copy operation. - size_t nbytes = GetDataSize(*from); - TVM_FFI_ICHECK_EQ(nbytes, GetDataSize(*to)); + size_t nbytes = ffi::GetDataSize(*from); + TVM_FFI_ICHECK_EQ(nbytes, ffi::GetDataSize(*to)); TVM_FFI_ICHECK(ffi::IsContiguous(*from) && ffi::IsContiguous(*to)) << "CopyDataFromTo only support contiguous array for now"; diff --git a/src/runtime/disco/distributed/socket_session.cc b/src/runtime/disco/distributed/socket_session.cc index d0b8d7df1640..a9d9d912aa82 100644 --- a/src/runtime/disco/distributed/socket_session.cc +++ b/src/runtime/disco/distributed/socket_session.cc @@ -18,6 +18,7 @@ */ #include #include +#include #include diff --git a/src/runtime/disco/nccl/nccl.cc b/src/runtime/disco/nccl/nccl.cc index 1230cf15f8a7..3167ab243ca7 100644 --- a/src/runtime/disco/nccl/nccl.cc +++ b/src/runtime/disco/nccl/nccl.cc @@ -18,6 +18,7 @@ */ #include +#include #include #include diff --git a/src/runtime/file_utils.cc b/src/runtime/file_utils.cc index 095b0288bad8..180f04da7dd7 100644 --- a/src/runtime/file_utils.cc +++ b/src/runtime/file_utils.cc @@ -22,10 +22,10 @@ */ #include "file_utils.h" +#include #include #include #include -#include #include #include diff --git a/src/runtime/hexagon/hexagon_buffer.cc b/src/runtime/hexagon/hexagon_buffer.cc index 84297c25e66b..b77638e45176 100644 --- a/src/runtime/hexagon/hexagon_buffer.cc +++ b/src/runtime/hexagon/hexagon_buffer.cc @@ -18,6 +18,8 @@ */ #include "hexagon_buffer.h" +#include + #include #include #include diff --git a/src/runtime/hexagon/hexagon_buffer.h b/src/runtime/hexagon/hexagon_buffer.h index 2dd7c127e3ed..0e578fccf477 100644 --- a/src/runtime/hexagon/hexagon_buffer.h +++ b/src/runtime/hexagon/hexagon_buffer.h @@ -20,10 +20,10 @@ #ifndef TVM_RUNTIME_HEXAGON_HEXAGON_BUFFER_H_ #define TVM_RUNTIME_HEXAGON_HEXAGON_BUFFER_H_ +#include #include #include #include -#include #include #include diff --git a/src/runtime/hexagon/hexagon_common.h b/src/runtime/hexagon/hexagon_common.h index 335e611d603e..7ffc4457192a 100644 --- a/src/runtime/hexagon/hexagon_common.h +++ b/src/runtime/hexagon/hexagon_common.h @@ -26,7 +26,7 @@ #include #include #include -#include +#include #if defined(__hexagon__) #include diff --git a/src/runtime/hexagon/hexagon_device_api.cc b/src/runtime/hexagon/hexagon_device_api.cc index 0d1c432571c6..ae0e0862dfc2 100644 --- a/src/runtime/hexagon/hexagon_device_api.cc +++ b/src/runtime/hexagon/hexagon_device_api.cc @@ -168,7 +168,7 @@ void HexagonDeviceAPI::FreeWorkspace(Device dev, void* data) { void HexagonDeviceAPI::CopyDataFromTo(DLTensor* from, DLTensor* to, TVMStreamHandle stream) { TVM_FFI_ICHECK_EQ(from->byte_offset, 0); TVM_FFI_ICHECK_EQ(to->byte_offset, 0); - TVM_FFI_ICHECK_EQ(GetDataSize(*from), GetDataSize(*to)); + TVM_FFI_ICHECK_EQ(ffi::GetDataSize(*from), ffi::GetDataSize(*to)); TVM_FFI_ICHECK(runtime_hexbuffs) << "Attempted to copy Hexagon data with " << "HexagonDeviceAPI::CopyDataFromTo before initializing resources. " @@ -182,11 +182,11 @@ void HexagonDeviceAPI::CopyDataFromTo(DLTensor* from, DLTensor* to, TVMStreamHan HexagonBuffer* hex_to_buf = lookup_hexagon_buffer(to->data); if (hex_from_buf && hex_to_buf) { - hex_to_buf->CopyFrom(*hex_from_buf, GetDataSize(*from)); + hex_to_buf->CopyFrom(*hex_from_buf, ffi::GetDataSize(*from)); } else if (hex_to_buf) { - hex_to_buf->CopyFrom(from->data, GetDataSize(*from)); + hex_to_buf->CopyFrom(from->data, ffi::GetDataSize(*from)); } else if (hex_from_buf) { - hex_from_buf->CopyTo(to->data, GetDataSize(*to)); + hex_from_buf->CopyTo(to->data, ffi::GetDataSize(*to)); } else { TVM_FFI_ICHECK(false) << "CopyDataFromTo requested between src and dst which are not managed by the " diff --git a/src/runtime/hexagon/hexagon_thread_manager.cc b/src/runtime/hexagon/hexagon_thread_manager.cc index c1c3eadc3126..76e57c67e8a1 100644 --- a/src/runtime/hexagon/hexagon_thread_manager.cc +++ b/src/runtime/hexagon/hexagon_thread_manager.cc @@ -18,6 +18,7 @@ */ #include "hexagon_thread_manager.h" +#include namespace tvm { namespace runtime { diff --git a/src/runtime/hexagon/hexagon_thread_manager.h b/src/runtime/hexagon/hexagon_thread_manager.h index 83c5316a7259..c02e23f29c34 100644 --- a/src/runtime/hexagon/hexagon_thread_manager.h +++ b/src/runtime/hexagon/hexagon_thread_manager.h @@ -22,7 +22,7 @@ #include #include -#include +#include #include #include diff --git a/src/runtime/hexagon/hexagon_vtcm_pool.cc b/src/runtime/hexagon/hexagon_vtcm_pool.cc index ef3dc592f003..f96ba975da0d 100644 --- a/src/runtime/hexagon/hexagon_vtcm_pool.cc +++ b/src/runtime/hexagon/hexagon_vtcm_pool.cc @@ -17,6 +17,7 @@ * under the License. */ #include "hexagon_vtcm_pool.h" +#include #include "HAP_compute_res.h" #include "hexagon_common.h" diff --git a/src/runtime/hexagon/hexagon_vtcm_pool.h b/src/runtime/hexagon/hexagon_vtcm_pool.h index 0f7153eb54f6..5159c458c8d6 100644 --- a/src/runtime/hexagon/hexagon_vtcm_pool.h +++ b/src/runtime/hexagon/hexagon_vtcm_pool.h @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/hexagon/qhl/qhl_wrapper.cc b/src/runtime/hexagon/qhl/qhl_wrapper.cc index e1515ecc7e08..a90b8eb1618f 100644 --- a/src/runtime/hexagon/qhl/qhl_wrapper.cc +++ b/src/runtime/hexagon/qhl/qhl_wrapper.cc @@ -19,7 +19,7 @@ #if defined(__hexagon__) #include #include -#include +#include #define restrict __restrict__ #define LOG2VLEN 7 diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index cd78591b4dbb..9f20a8f6d229 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -29,6 +29,7 @@ extern "C" { #include #include #include +#include #include #include diff --git a/src/runtime/hexagon/rpc/simulator/session.cc b/src/runtime/hexagon/rpc/simulator/session.cc index 0864796a9ad9..918614afcde7 100644 --- a/src/runtime/hexagon/rpc/simulator/session.cc +++ b/src/runtime/hexagon/rpc/simulator/session.cc @@ -20,6 +20,7 @@ #include #include #include +#include // POSIX includes #include #include diff --git a/src/runtime/memory/memory_manager.cc b/src/runtime/memory/memory_manager.cc index ba96c0071e0d..626222e6c87f 100644 --- a/src/runtime/memory/memory_manager.cc +++ b/src/runtime/memory/memory_manager.cc @@ -25,6 +25,7 @@ #include #include #include +#include #include #include diff --git a/src/runtime/memory/naive_allocator.h b/src/runtime/memory/naive_allocator.h index 6a968c86ef3b..134d05762286 100644 --- a/src/runtime/memory/naive_allocator.h +++ b/src/runtime/memory/naive_allocator.h @@ -24,6 +24,7 @@ #define TVM_RUNTIME_MEMORY_NAIVE_ALLOCATOR_H_ #include +#include #include #include diff --git a/src/runtime/memory/pooled_allocator.h b/src/runtime/memory/pooled_allocator.h index 620393466867..2862dde1ae6d 100644 --- a/src/runtime/memory/pooled_allocator.h +++ b/src/runtime/memory/pooled_allocator.h @@ -24,6 +24,7 @@ #define TVM_RUNTIME_MEMORY_POOLED_ALLOCATOR_H_ #include +#include #include #include diff --git a/src/runtime/metadata.h b/src/runtime/metadata.h index e85d53b07cbe..c034041ce4a4 100644 --- a/src/runtime/metadata.h +++ b/src/runtime/metadata.h @@ -74,7 +74,7 @@ class FunctionInfoObj : public ffi::Object { obj.Set("name", name); json::Array sarg_types; for (const auto& t : arg_types) { - sarg_types.push_back(ffi::String(DLDataTypeToString(t))); + sarg_types.push_back(ffi::String(ffi::DLDataTypeToString(t))); } obj.Set("arg_types", std::move(sarg_types)); { @@ -96,7 +96,7 @@ class FunctionInfoObj : public ffi::Object { auto sarg_types_arr = src.at("arg_types").cast(); arg_types = ffi::Array(); for (size_t i = 0; i < sarg_types_arr.size(); ++i) { - arg_types.push_back(StringToDLDataType(std::string(sarg_types_arr[i].cast()))); + arg_types.push_back(ffi::StringToDLDataType(std::string(sarg_types_arr[i].cast()))); } auto lt = src.find("launch_param_tags"); if (lt != src.end()) { diff --git a/src/runtime/metal/metal_common.h b/src/runtime/metal/metal_common.h index cc538f84dce0..ebbbcde071b4 100644 --- a/src/runtime/metal/metal_common.h +++ b/src/runtime/metal/metal_common.h @@ -33,7 +33,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/minrpc/minrpc_server.h b/src/runtime/minrpc/minrpc_server.h index 434c88b693e9..84cf45a6ca92 100644 --- a/src/runtime/minrpc/minrpc_server.h +++ b/src/runtime/minrpc/minrpc_server.h @@ -31,7 +31,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index a43b29d5ec59..d80a52e5e705 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -27,7 +27,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/src/runtime/opencl/opencl_device_api.cc b/src/runtime/opencl/opencl_device_api.cc index 0b63f497dbff..952a9b67141c 100644 --- a/src/runtime/opencl/opencl_device_api.cc +++ b/src/runtime/opencl/opencl_device_api.cc @@ -23,6 +23,7 @@ #include #include #include +#include #include @@ -507,9 +508,9 @@ void OpenCLWorkspace::FreeDataSpace(Device dev, void* ptr) { void OpenCLWorkspace::CopyDataFromTo(DLTensor* from, DLTensor* to, TVMStreamHandle stream) { this->Init(); - size_t nbytes = GetDataSize(*from); - TVM_FFI_ICHECK_EQ(nbytes, GetDataSize(*to)); - TVM_FFI_ICHECK(IsContiguous(*from) && IsContiguous(*to)) + size_t nbytes = ffi::GetDataSize(*from); + TVM_FFI_ICHECK_EQ(nbytes, ffi::GetDataSize(*to)); + TVM_FFI_ICHECK(ffi::IsContiguous(*from) && ffi::IsContiguous(*to)) << "CopyDataFromTo only support contiguous array for now"; if (IsOpenCLDevice(from->device) && IsOpenCLDevice(to->device)) { diff --git a/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc index e4c4a1a9af31..a10b1a81b837 100644 --- a/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc +++ b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc @@ -33,7 +33,7 @@ #include #endif -#include +#include #include #include diff --git a/src/runtime/rpc/rpc_channel.cc b/src/runtime/rpc/rpc_channel.cc index f462dac3d257..11e14a9a8dbd 100644 --- a/src/runtime/rpc/rpc_channel.cc +++ b/src/runtime/rpc/rpc_channel.cc @@ -22,7 +22,7 @@ */ #include "rpc_channel.h" -#include +#include #include diff --git a/src/runtime/rpc/rpc_device_api.cc b/src/runtime/rpc/rpc_device_api.cc index 579b45abb31c..6e0dd162b3ba 100644 --- a/src/runtime/rpc/rpc_device_api.cc +++ b/src/runtime/rpc/rpc_device_api.cc @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include @@ -98,14 +98,14 @@ class RPCDeviceAPI final : public DeviceAPI { from_tensor.device = RemoveRPCSessionMask(dev_from); from_tensor.data = static_cast(from->data)->data; void* to_bytes = static_cast(to->data) + to->byte_offset; - size_t nbytes = GetDataSize(*to); + size_t nbytes = ffi::GetDataSize(*to); GetSess(dev_from)->CopyFromRemote(&from_tensor, to_bytes, nbytes); } else if (dev_from.device_type == kDLCPU && IsRPCSessionDevice(dev_to)) { DLTensor to_tensor = *to; to_tensor.device = RemoveRPCSessionMask(dev_to); to_tensor.data = static_cast(to->data)->data; void* from_bytes = static_cast(from->data) + from->byte_offset; - size_t nbytes = GetDataSize(*from); + size_t nbytes = ffi::GetDataSize(*from); GetSess(dev_to)->CopyToRemote(from_bytes, &to_tensor, nbytes); } else { TVM_FFI_THROW(InternalError) << "expect copy from/to remote or between remote"; diff --git a/src/runtime/rpc/rpc_local_session.cc b/src/runtime/rpc/rpc_local_session.cc index 5094bc678bac..0a670ceb941c 100644 --- a/src/runtime/rpc/rpc_local_session.cc +++ b/src/runtime/rpc/rpc_local_session.cc @@ -106,7 +106,7 @@ void LocalSession::CallFunc(RPCSession::PackedFuncHandle func, ffi::PackedArgs a } void LocalSession::CopyToRemote(void* from_bytes, DLTensor* to, uint64_t nbytes) { - TVM_FFI_ICHECK_EQ(nbytes, GetDataSize(*to)); + TVM_FFI_ICHECK_EQ(nbytes, ffi::GetDataSize(*to)); DLTensor from; from.data = from_bytes; from.device = {kDLCPU, 0}; diff --git a/src/runtime/rpc/rpc_module.cc b/src/runtime/rpc/rpc_module.cc index 3bf7da474ee2..74f19ab3e3bb 100644 --- a/src/runtime/rpc/rpc_module.cc +++ b/src/runtime/rpc/rpc_module.cc @@ -399,7 +399,7 @@ inline void CPUCacheFlushImpl(const char* addr, unsigned int len) { inline void CPUCacheFlush(int begin_index, const ffi::PackedArgs& args) { for (int i = begin_index; i < args.size(); i++) { CPUCacheFlushImpl(static_cast((args[i].cast()->data)), - GetDataSize(*(args[i].cast()))); + ffi::GetDataSize(*(args[i].cast()))); } } diff --git a/src/runtime/rpc/rpc_server_env.cc b/src/runtime/rpc/rpc_server_env.cc index a51d98b17f93..b8c08b19c413 100644 --- a/src/runtime/rpc/rpc_server_env.cc +++ b/src/runtime/rpc/rpc_server_env.cc @@ -23,6 +23,7 @@ */ #include #include +#include #include "../file_utils.h" diff --git a/src/runtime/static_library.cc b/src/runtime/static_library.cc index e800ed231b24..d3ea7b345838 100644 --- a/src/runtime/static_library.cc +++ b/src/runtime/static_library.cc @@ -31,6 +31,7 @@ #include #include #include +#include #include diff --git a/src/runtime/static_library.h b/src/runtime/static_library.h index 65ee6f8808c8..0ce4d9e003c6 100644 --- a/src/runtime/static_library.h +++ b/src/runtime/static_library.h @@ -28,7 +28,7 @@ #include #include -#include +#include #include #include diff --git a/src/runtime/tensor.cc b/src/runtime/tensor.cc index d4fe1772b978..61f037caec55 100644 --- a/src/runtime/tensor.cc +++ b/src/runtime/tensor.cc @@ -25,7 +25,7 @@ #include #include #include -#include +#include #include #include "tvm/runtime/data_type.h" @@ -60,9 +60,9 @@ inline void VerifyDataType(DLDataType dtype) { } void TensorCopyFromBytes(DLTensor* handle, const void* data, size_t nbytes) { - size_t arr_size = GetDataSize(*handle); + size_t arr_size = ffi::GetDataSize(*handle); TVM_FFI_ICHECK_EQ(arr_size, nbytes) << "TensorCopyFromBytes: size mismatch"; - TVM_FFI_ICHECK(IsContiguous(*handle)) + TVM_FFI_ICHECK(ffi::IsContiguous(*handle)) << "TensorCopyFromBytes only support contiguous array for now"; DLTensor from; @@ -80,7 +80,7 @@ void TensorCopyFromBytes(DLTensor* handle, const void* data, size_t nbytes) { void Tensor::CopyToBytes(const DLTensor* handle, void* data, size_t nbytes, TVMStreamHandle stream) { - size_t arr_size = GetDataSize(*handle); + size_t arr_size = ffi::GetDataSize(*handle); TVM_FFI_ICHECK_EQ(arr_size, nbytes) << "ArrayCopyToBytes: size mismatch"; TVM_FFI_ICHECK(ffi::IsContiguous(*handle)) << "ArrayCopyToBytes only support contiguous array for now"; @@ -101,7 +101,7 @@ void Tensor::CopyToBytes(const DLTensor* handle, void* data, size_t nbytes, void Tensor::CopyFromBytes(const DLTensor* handle, void* data, size_t nbytes, TVMStreamHandle stream) { - size_t arr_size = GetDataSize(*handle); + size_t arr_size = ffi::GetDataSize(*handle); TVM_FFI_ICHECK_EQ(arr_size, nbytes) << "ArrayCopyToBytes: size mismatch"; TVM_FFI_ICHECK(ffi::IsContiguous(*handle)) << "ArrayCopyToBytes only support contiguous array for now"; @@ -160,7 +160,7 @@ Tensor Tensor::CreateView(ffi::Shape shape, DLDataType dtype, uint64_t relative_ return ss.str(); }(); const auto& curr_dl_tensor = *get_mutable(); - size_t curr_size = GetDataSize(curr_dl_tensor); + size_t curr_size = ffi::GetDataSize(curr_dl_tensor); size_t view_size = ffi::GetDataSize(shape.Product(), dtype); TVM_FFI_CHECK_LE(relative_byte_offset + view_size, curr_size, ValueError) << "View with shape " << shape << " and datatype " << dtype << " would have a size of " @@ -215,8 +215,8 @@ Tensor Tensor::CopyTo(const Device& dev, ffi::Optional mem_scope) c } void Tensor::CopyFromTo(const DLTensor* from, DLTensor* to, TVMStreamHandle stream) { - size_t from_size = GetDataSize(*from); - size_t to_size = GetDataSize(*to); + size_t from_size = ffi::GetDataSize(*from); + size_t to_size = ffi::GetDataSize(*to); TVM_FFI_ICHECK_EQ(from_size, to_size) << "TVMTensorCopyFromTo: The size in bytes must exactly match."; diff --git a/src/runtime/thread_pool.cc b/src/runtime/thread_pool.cc index 63eba5eba23f..ba2b89770bd7 100644 --- a/src/runtime/thread_pool.cc +++ b/src/runtime/thread_pool.cc @@ -26,7 +26,7 @@ #include #include #include -#include +#include #include "threading_backend.h" #if TVM_THREADPOOL_USE_OPENMP diff --git a/src/runtime/timer.cc b/src/runtime/timer.cc index 075f56337e77..f2adcd353342 100644 --- a/src/runtime/timer.cc +++ b/src/runtime/timer.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include #include diff --git a/src/runtime/vm/attn_backend.h b/src/runtime/vm/attn_backend.h index 5db83ff499e1..ae88843667c3 100644 --- a/src/runtime/vm/attn_backend.h +++ b/src/runtime/vm/attn_backend.h @@ -28,7 +28,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/vm/builtin.cc b/src/runtime/vm/builtin.cc index 706d74339097..f5485e7a3326 100644 --- a/src/runtime/vm/builtin.cc +++ b/src/runtime/vm/builtin.cc @@ -27,7 +27,7 @@ #include #include #include -#include +#include #include #include #include @@ -611,7 +611,7 @@ bool ReadIfCond(ffi::AnyView cond) { break; } default: - TVM_FFI_THROW(InternalError) << "Unknown scalar int type: " << DLDataTypeToString(arr->dtype); + TVM_FFI_THROW(InternalError) << "Unknown scalar int type: " << ffi::DLDataTypeToString(arr->dtype); throw; } return result != 0; @@ -702,7 +702,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { } default: TVM_FFI_THROW(InternalError) - << "Unknown scalar int type: " << DLDataTypeToString(arr->dtype); + << "Unknown scalar int type: " << ffi::DLDataTypeToString(arr->dtype); throw; } out_shape.push_back(result); diff --git a/src/runtime/vm/bytecode.cc b/src/runtime/vm/bytecode.cc index 4356305521a3..552f19091bc2 100644 --- a/src/runtime/vm/bytecode.cc +++ b/src/runtime/vm/bytecode.cc @@ -22,7 +22,7 @@ * \brief The bytecode for Relax virtual machine. */ -#include +#include #include #include diff --git a/src/runtime/vm/hexagon/builtin.cc b/src/runtime/vm/hexagon/builtin.cc index 54fd70b2800f..c7429975647f 100644 --- a/src/runtime/vm/hexagon/builtin.cc +++ b/src/runtime/vm/hexagon/builtin.cc @@ -45,8 +45,8 @@ TVM_FFI_STATIC_INIT_BLOCK() { void* src = sptr->data; int ret = DMA_RETRY; - TVM_FFI_ICHECK_EQ(GetDataSize(*dptr), GetDataSize(*sptr)); - auto size = GetDataSize(*dptr); + TVM_FFI_ICHECK_EQ(ffi::GetDataSize(*dptr), ffi::GetDataSize(*sptr)); + auto size = ffi::GetDataSize(*dptr); TVM_FFI_ICHECK(size > 0); if (bypass_cache) qurt_mem_cache_clean(reinterpret_cast(src), size, @@ -65,7 +65,7 @@ TVM_FFI_STATIC_INIT_BLOCK() { if (bypass_cache) { const DLTensor* dptr = dst_arr.operator->(); void* dst = dptr->data; - auto size = GetDataSize(*dptr); + auto size = ffi::GetDataSize(*dptr); qurt_mem_cache_clean(reinterpret_cast(dst), size, QURT_MEM_CACHE_FLUSH, QURT_MEM_DCACHE); } diff --git a/src/runtime/vm/kv_state.h b/src/runtime/vm/kv_state.h index 4578a8a30690..fd001f8048a2 100644 --- a/src/runtime/vm/kv_state.h +++ b/src/runtime/vm/kv_state.h @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include namespace tvm { diff --git a/src/runtime/vm/lm_support.cc b/src/runtime/vm/lm_support.cc index fc1c84cffaef..d07f84be1647 100644 --- a/src/runtime/vm/lm_support.cc +++ b/src/runtime/vm/lm_support.cc @@ -40,7 +40,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/src/runtime/vm/paged_kv_cache.cc b/src/runtime/vm/paged_kv_cache.cc index 879066d9cf30..bb3aee7e340b 100644 --- a/src/runtime/vm/paged_kv_cache.cc +++ b/src/runtime/vm/paged_kv_cache.cc @@ -24,7 +24,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/vm/tensor_cache_support.cc b/src/runtime/vm/tensor_cache_support.cc index 1804dcc622b1..ee77c5ddd8f0 100644 --- a/src/runtime/vm/tensor_cache_support.cc +++ b/src/runtime/vm/tensor_cache_support.cc @@ -137,7 +137,7 @@ void CopyTensorFromBytes(Tensor param, const void* data, size_t nbytes, // It creates a host side memory mirror, for every cl_mem that tries to copy data from host // which can cause memory issue. Her we use a large staging buffer to postpone deallocation if (staging_buffer->defined()) { - size_t curr_size = runtime::GetDataSize(*(staging_buffer->value().operator->())); + size_t curr_size = ffi::GetDataSize(*(staging_buffer->value().operator->())); if (curr_size < nbytes) { *staging_buffer = std::nullopt; } diff --git a/src/runtime/vm/vm.cc b/src/runtime/vm/vm.cc index d3fa356792ee..b7e29710aff9 100644 --- a/src/runtime/vm/vm.cc +++ b/src/runtime/vm/vm.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include @@ -779,7 +780,7 @@ void VirtualMachineImpl::RunInstrCall(VMFrame* curr_frame, Instruction instr) { for (int i = 0; i < instr.num_args; ++i) { if (call_args[i + args_begin_offset].type_index() == ffi::TypeIndex::kTVMFFIDataType) { std::string str_dtype = - DLDataTypeToString(call_args[i + args_begin_offset].cast()); + ffi::DLDataTypeToString(call_args[i + args_begin_offset].cast()); temp_dtype.emplace_back(std::make_unique(str_dtype)); call_args[i + args_begin_offset] = *temp_dtype.back(); } diff --git a/src/runtime/vulkan/spirv_shader.h b/src/runtime/vulkan/spirv_shader.h index f290d0dbd195..e9575defd110 100644 --- a/src/runtime/vulkan/spirv_shader.h +++ b/src/runtime/vulkan/spirv_shader.h @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/vulkan/vulkan_common.h b/src/runtime/vulkan/vulkan_common.h index d25817a2d787..826048d8578d 100644 --- a/src/runtime/vulkan/vulkan_common.h +++ b/src/runtime/vulkan/vulkan_common.h @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include diff --git a/src/runtime/vulkan/vulkan_device.h b/src/runtime/vulkan/vulkan_device.h index 324c3c319cc0..c327149cc2b0 100644 --- a/src/runtime/vulkan/vulkan_device.h +++ b/src/runtime/vulkan/vulkan_device.h @@ -20,7 +20,7 @@ #ifndef TVM_RUNTIME_VULKAN_VULKAN_DEVICE_H_ #define TVM_RUNTIME_VULKAN_VULKAN_DEVICE_H_ -#include +#include #include #include diff --git a/src/runtime/vulkan/vulkan_instance.cc b/src/runtime/vulkan/vulkan_instance.cc index e23e3b7f1ec2..fc88db7644cd 100644 --- a/src/runtime/vulkan/vulkan_instance.cc +++ b/src/runtime/vulkan/vulkan_instance.cc @@ -18,6 +18,7 @@ */ #include "vulkan_instance.h" +#include #include #include diff --git a/src/s_tir/analysis/is_pure_function.cc b/src/s_tir/analysis/is_pure_function.cc index 40feab7f5c80..2ca557b171d1 100644 --- a/src/s_tir/analysis/is_pure_function.cc +++ b/src/s_tir/analysis/is_pure_function.cc @@ -33,7 +33,6 @@ namespace tvm { namespace s_tir { using namespace tvm::tirx; -using AccessPath = ffi::reflection::AccessPath; namespace { class PurityChecker : TIRVisitorWithPath { @@ -47,12 +46,12 @@ class PurityChecker : TIRVisitorWithPath { private: explicit PurityChecker(bool assert_on_error) : assert_on_error_(assert_on_error) {} - void VisitStmt_(const AllocBufferNode* op, AccessPath path) override { + void VisitStmt_(const AllocBufferNode* op, ffi::reflection::AccessPath path) override { internal_allocations_.insert(op->buffer->data); TIRVisitorWithPath::VisitStmt_(op, path); } - void VisitStmt_(const BufferStoreNode* op, AccessPath path) override { + void VisitStmt_(const BufferStoreNode* op, ffi::reflection::AccessPath path) override { TIRVisitorWithPath::VisitStmt_(op, path); if (!internal_allocations_.count(op->buffer->data)) { @@ -65,7 +64,7 @@ class PurityChecker : TIRVisitorWithPath { } } - void VisitExpr_(const CallNode* call, AccessPath path) override { + void VisitExpr_(const CallNode* call, ffi::reflection::AccessPath path) override { TIRVisitorWithPath::VisitExpr_(call, path); static auto op_call_effect = Op::GetAttrMap("TCallEffectKind"); diff --git a/src/s_tir/analysis/sblock_access_region_detector.cc b/src/s_tir/analysis/sblock_access_region_detector.cc index 19fc8fd090fb..918b9d2815d8 100644 --- a/src/s_tir/analysis/sblock_access_region_detector.cc +++ b/src/s_tir/analysis/sblock_access_region_detector.cc @@ -204,7 +204,7 @@ void BlockReadWriteDetector::VisitExpr_(const CallNode* op) { if (it != buffer_var_map_.end()) { const Buffer& buffer = (*it).second; const BufferRegion buffer_region = BufferRegion::FullRegion(buffer); - const Region& region = buffer_region->region; + const ffi::Array& region = buffer_region->region; std::vector int_set; int_set.reserve(region.size()); for (const Range& range : region) { @@ -287,7 +287,7 @@ std::vector BlockReadWriteDetector::ConvertMatchedRegion( const MatchBufferRegion& match_buffer, const std::vector& int_sets) const { const Buffer& buffer = match_buffer->buffer; - Region region; + ffi::Array region; region.reserve(int_sets.size()); TVM_FFI_ICHECK_EQ(buffer->shape.size(), int_sets.size()); for (size_t i = 0; i < int_sets.size(); ++i) { @@ -363,7 +363,7 @@ void BlockReadWriteDetector::UpdateOpaque(const Var& buffer_var) { if (it != buffer_var_map_.end()) { const Buffer& buffer = (*it).second; const BufferRegion buffer_region = BufferRegion::FullRegion(buffer); - const Region& region = buffer_region->region; + const ffi::Array& region = buffer_region->region; std::vector int_set; int_set.reserve(region.size()); for (const Range& range : region) { diff --git a/src/s_tir/meta_schedule/arg_info.cc b/src/s_tir/meta_schedule/arg_info.cc index 5411d46cc20b..87c6715a9841 100644 --- a/src/s_tir/meta_schedule/arg_info.cc +++ b/src/s_tir/meta_schedule/arg_info.cc @@ -126,7 +126,7 @@ TensorInfo::TensorInfo(runtime::DataType dtype, ffi::Shape shape) { ffi::ObjectRef TensorInfoNode::AsJSON() const { static ffi::String tag = "TENSOR"; - ffi::String dtype = DLDataTypeToString(this->dtype); + ffi::String dtype = ffi::DLDataTypeToString(this->dtype); ffi::Array shape = support::AsArray(this->shape); return ffi::Array{tag, dtype, shape}; } @@ -140,7 +140,7 @@ TensorInfo TensorInfo::FromJSON(const ffi::ObjectRef& json_obj) { // Load json[1] => dtype { ffi::String dtype_str = json_array->at(1).cast(); - dtype = StringToDLDataType(dtype_str); + dtype = ffi::StringToDLDataType(dtype_str); } // Load json[2] => shape shape = AsIntArray(json_array->at(2).cast()); diff --git a/src/s_tir/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc b/src/s_tir/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc index c699ea65136b..27c3ded758ad 100644 --- a/src/s_tir/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc +++ b/src/s_tir/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc @@ -18,6 +18,7 @@ */ #include #include +#include #include "../utils.h" diff --git a/src/s_tir/meta_schedule/postproc/rewrite_tensorize.cc b/src/s_tir/meta_schedule/postproc/rewrite_tensorize.cc index 958bf5f9227f..01d619302a5a 100644 --- a/src/s_tir/meta_schedule/postproc/rewrite_tensorize.cc +++ b/src/s_tir/meta_schedule/postproc/rewrite_tensorize.cc @@ -19,6 +19,7 @@ #include #include #include +#include #include diff --git a/src/s_tir/meta_schedule/schedule_rule/apply_custom_rule.cc b/src/s_tir/meta_schedule/schedule_rule/apply_custom_rule.cc index 73cae90cd48e..dfd8f99aee8e 100644 --- a/src/s_tir/meta_schedule/schedule_rule/apply_custom_rule.cc +++ b/src/s_tir/meta_schedule/schedule_rule/apply_custom_rule.cc @@ -17,6 +17,7 @@ * under the License. */ #include +#include #include "../utils.h" diff --git a/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling.cc b/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling.cc index 69311ea3c8d5..87244c8809e4 100644 --- a/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling.cc +++ b/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling.cc @@ -21,6 +21,7 @@ #include #include #include +#include #include #include diff --git a/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc b/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc index 00e51cbb1ebf..2dc9de361e8f 100644 --- a/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc +++ b/src/s_tir/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc @@ -820,7 +820,7 @@ ffi::Optional MultiLevelTilingTensorCoreNode::TransformWithTensorIntrin( rhs_to_index_map_tgt[mapping_info->rhs_iters[i - offset]->var] = index_map->final_indices[i]; } - auto f_get_sub_index_map = [&](const tirx::Buffer& lhs_buffer, const tirx::Region& lhs_region) { + auto f_get_sub_index_map = [&](const tirx::Buffer& lhs_buffer, const ffi::Array& lhs_region) { std::vector sub_index_map_src; std::vector sub_index_map_tgt; const tirx::Buffer& rhs_buffer = mapping_info->lhs_buffer_map[lhs_buffer]; diff --git a/src/s_tir/meta_schedule/space_generator/space_generator.cc b/src/s_tir/meta_schedule/space_generator/space_generator.cc index 9bddc18d839e..da5f5f399833 100644 --- a/src/s_tir/meta_schedule/space_generator/space_generator.cc +++ b/src/s_tir/meta_schedule/space_generator/space_generator.cc @@ -17,6 +17,7 @@ * under the License. */ #include +#include #include "../../../target/canonicalizer/llvm/arm_aprofile.h" #include "../utils.h" diff --git a/src/s_tir/meta_schedule/utils.h b/src/s_tir/meta_schedule/utils.h index 664a00cf3268..5dc99d744c28 100644 --- a/src/s_tir/meta_schedule/utils.h +++ b/src/s_tir/meta_schedule/utils.h @@ -43,6 +43,7 @@ #include #include #include +#include #include #include diff --git a/src/s_tir/schedule/concrete_schedule.cc b/src/s_tir/schedule/concrete_schedule.cc index 89bebd33f833..21f5454040a6 100644 --- a/src/s_tir/schedule/concrete_schedule.cc +++ b/src/s_tir/schedule/concrete_schedule.cc @@ -19,6 +19,7 @@ #include "./concrete_schedule.h" #include +#include #include diff --git a/src/s_tir/schedule/primitive/blockize_tensorize.cc b/src/s_tir/schedule/primitive/blockize_tensorize.cc index 282e167ee55c..da4deb01bc87 100644 --- a/src/s_tir/schedule/primitive/blockize_tensorize.cc +++ b/src/s_tir/schedule/primitive/blockize_tensorize.cc @@ -18,6 +18,7 @@ */ #include +#include #include diff --git a/src/s_tir/schedule/primitive/cache_read_write.cc b/src/s_tir/schedule/primitive/cache_read_write.cc index 39d3bacfbe8b..3c754d1fa3af 100644 --- a/src/s_tir/schedule/primitive/cache_read_write.cc +++ b/src/s_tir/schedule/primitive/cache_read_write.cc @@ -1087,7 +1087,7 @@ class ReindexCacheReadRewriter : public CacheReadRewriter { ffi::Array new_reads; for (const BufferRegion& buf_region : reads) { if (buf_region->buffer.same_as(info_->read_buffer)) { - ffi::Array region; + Region region; for (const PrimExpr index : new_indices_) { region.push_back(Range::FromMinExtent(index, Integer(1))); } @@ -1103,7 +1103,7 @@ class ReindexCacheReadRewriter : public CacheReadRewriter { for (const MatchBufferRegion& match_buffer_region : match_buffers) { BufferRegion source = match_buffer_region->source; if (source->buffer.same_as(info_->read_buffer)) { - ffi::Array region; + Region region; for (const PrimExpr index : new_indices_) { region.push_back(Range::FromMinExtent(index, Integer(1))); } @@ -1376,7 +1376,7 @@ class ReindexCacheWriteRewriter : public CacheWriteRewriter { ffi::Array new_reads; for (const BufferRegion& buf_region : reads) { if (buf_region->buffer.same_as(info_->write_buffer)) { - ffi::Array region; + Region region; for (const PrimExpr index : new_indices_) { region.push_back(Range::FromMinExtent(index, Integer(1))); } @@ -1392,7 +1392,7 @@ class ReindexCacheWriteRewriter : public CacheWriteRewriter { for (const MatchBufferRegion& match_buffer_region : match_buffers) { BufferRegion source = match_buffer_region->source; if (source->buffer.same_as(info_->write_buffer)) { - ffi::Array region; + Region region; for (const PrimExpr index : new_indices_) { region.push_back(Range::FromMinExtent(index, Integer(1))); } diff --git a/src/s_tir/schedule/primitive/layout_transformation.cc b/src/s_tir/schedule/primitive/layout_transformation.cc index 4208873b4637..d9c729dd9078 100644 --- a/src/s_tir/schedule/primitive/layout_transformation.cc +++ b/src/s_tir/schedule/primitive/layout_transformation.cc @@ -20,6 +20,7 @@ #include #include #include +#include #include #include diff --git a/src/s_tir/support/parallel_for.h b/src/s_tir/support/parallel_for.h index 9374027c421a..1b2c5fa18fbb 100644 --- a/src/s_tir/support/parallel_for.h +++ b/src/s_tir/support/parallel_for.h @@ -25,7 +25,7 @@ #define TVM_S_TIR_SUPPORT_PARALLEL_FOR_H_ #include -#include +#include #include #include diff --git a/src/s_tir/support/table_printer.h b/src/s_tir/support/table_printer.h index 6ccaa23eca75..eb29b9706e74 100644 --- a/src/s_tir/support/table_printer.h +++ b/src/s_tir/support/table_printer.h @@ -19,7 +19,7 @@ #ifndef TVM_S_TIR_SUPPORT_TABLE_PRINTER_H_ #define TVM_S_TIR_SUPPORT_TABLE_PRINTER_H_ -#include +#include #include #include diff --git a/src/s_tir/transform/inject_double_buffer.cc b/src/s_tir/transform/inject_double_buffer.cc index b91a5214ae95..9c5e9bf0b8b5 100644 --- a/src/s_tir/transform/inject_double_buffer.cc +++ b/src/s_tir/transform/inject_double_buffer.cc @@ -28,6 +28,7 @@ #include #include #include +#include #include "../../tirx/transform/ir_utils.h" diff --git a/src/s_tir/transform/loop_partition.cc b/src/s_tir/transform/loop_partition.cc index e68b465dd263..d47c861873a7 100644 --- a/src/s_tir/transform/loop_partition.cc +++ b/src/s_tir/transform/loop_partition.cc @@ -31,6 +31,7 @@ #include #include #include +#include #include #include diff --git a/src/s_tir/transform/lower_async_dma.cc b/src/s_tir/transform/lower_async_dma.cc index 628e20be88bc..756461b0dd08 100644 --- a/src/s_tir/transform/lower_async_dma.cc +++ b/src/s_tir/transform/lower_async_dma.cc @@ -32,6 +32,7 @@ #include #include #include +#include #include #include diff --git a/src/s_tir/transform/lower_match_buffer.cc b/src/s_tir/transform/lower_match_buffer.cc index bd27d5189321..4caa02bc713c 100644 --- a/src/s_tir/transform/lower_match_buffer.cc +++ b/src/s_tir/transform/lower_match_buffer.cc @@ -25,6 +25,7 @@ #include #include #include +#include #include #include #include diff --git a/src/script/ir_builder/ir/ir.cc b/src/script/ir_builder/ir/ir.cc index 153c42e9d14a..683806768dc2 100644 --- a/src/script/ir_builder/ir/ir.cc +++ b/src/script/ir_builder/ir/ir.cc @@ -20,6 +20,7 @@ #include #include #include +#include #include "./utils.h" diff --git a/src/script/printer/doc.cc b/src/script/printer/doc.cc index 8ff66df53bde..5cd9edca79dc 100644 --- a/src/script/printer/doc.cc +++ b/src/script/printer/doc.cc @@ -18,9 +18,9 @@ */ #include #include +#include #include #include -#include #include namespace tvm { diff --git a/src/script/printer/doc_printer/python_doc_printer.cc b/src/script/printer/doc_printer/python_doc_printer.cc index db04e7427acd..78b9b9fa986f 100644 --- a/src/script/printer/doc_printer/python_doc_printer.cc +++ b/src/script/printer/doc_printer/python_doc_printer.cc @@ -18,7 +18,7 @@ */ #include #include -#include +#include #include #include diff --git a/src/script/printer/ir_docsifier.cc b/src/script/printer/ir_docsifier.cc index 5fb247a4882a..dd5762973b73 100644 --- a/src/script/printer/ir_docsifier.cc +++ b/src/script/printer/ir_docsifier.cc @@ -16,10 +16,10 @@ * specific language governing permissions and limitations * under the License. */ +#include #include #include #include -#include #include #include diff --git a/src/script/printer/utils.h b/src/script/printer/utils.h index 84d1854b756d..eed29f102dfc 100644 --- a/src/script/printer/utils.h +++ b/src/script/printer/utils.h @@ -25,6 +25,7 @@ #include #include #include +#include #include #include @@ -122,7 +123,7 @@ inline ExprDoc Relax(const IRDocsifier& d, const ffi::String& attr) { } inline std::string DType2Str(const runtime::DataType& dtype) { - return dtype.is_void() ? "void" : runtime::DLDataTypeToString(dtype); + return dtype.is_void() ? "void" : ffi::DLDataTypeToString(dtype); } /*! \brief Add headers as comments to doc if needed */ diff --git a/src/support/base64.h b/src/support/base64.h index af011e317538..36cb81cb5447 100644 --- a/src/support/base64.h +++ b/src/support/base64.h @@ -26,7 +26,7 @@ #ifndef TVM_SUPPORT_BASE64_H_ #define TVM_SUPPORT_BASE64_H_ -#include +#include #include #include diff --git a/src/support/pipe.h b/src/support/pipe.h index ec7a8ea14d9e..c208d223bd61 100644 --- a/src/support/pipe.h +++ b/src/support/pipe.h @@ -24,7 +24,7 @@ #ifndef TVM_SUPPORT_PIPE_H_ #define TVM_SUPPORT_PIPE_H_ -#include +#include #include #ifdef _WIN32 diff --git a/src/support/ring_buffer.h b/src/support/ring_buffer.h index 40d741a762e4..912b4c8d4b46 100644 --- a/src/support/ring_buffer.h +++ b/src/support/ring_buffer.h @@ -24,7 +24,7 @@ #ifndef TVM_SUPPORT_RING_BUFFER_H_ #define TVM_SUPPORT_RING_BUFFER_H_ -#include +#include #include #include diff --git a/src/target/canonicalizer/llvm/arm_aprofile.cc b/src/target/canonicalizer/llvm/arm_aprofile.cc index 97f0071394a9..0ad87ad66bf8 100644 --- a/src/target/canonicalizer/llvm/arm_aprofile.cc +++ b/src/target/canonicalizer/llvm/arm_aprofile.cc @@ -24,6 +24,8 @@ #include "arm_aprofile.h" +#include + #include #include diff --git a/src/target/codegen.cc b/src/target/codegen.cc index f24cb6a49497..39500a0451fa 100644 --- a/src/target/codegen.cc +++ b/src/target/codegen.cc @@ -35,6 +35,7 @@ #include #include +#include #include #include #include @@ -109,7 +110,7 @@ class ModuleSerializer { uint64_t module_index = 0; auto fpush_imports_to_stack = [&](ffi::ModuleObj* node) { - for (Any m : node->imports()) { + for (ffi::Any m : node->imports()) { ffi::ModuleObj* next = m.cast().operator->(); if (visited.count(next) == 0) { visited.insert(next); @@ -177,7 +178,7 @@ class ModuleSerializer { for (size_t parent_index = 0; parent_index < mod_group_vec_.size(); ++parent_index) { child_indices.clear(); for (const auto* m : mod_group_vec_[parent_index]) { - for (Any im : m->imports()) { + for (ffi::Any im : m->imports()) { uint64_t mod_index = mod2index_.at(im.cast().operator->()); // skip cycle when dso modules are merged together if (mod_index != parent_index) { diff --git a/src/target/cuda/codegen_cuda.cc b/src/target/cuda/codegen_cuda.cc index d76e5fbac187..ec5f014e8e0b 100644 --- a/src/target/cuda/codegen_cuda.cc +++ b/src/target/cuda/codegen_cuda.cc @@ -31,6 +31,7 @@ #include #include +#include #include #include #include diff --git a/src/target/cuda/ptx.h b/src/target/cuda/ptx.h index b82a9c6ad3f3..7bdc16e3ae0c 100644 --- a/src/target/cuda/ptx.h +++ b/src/target/cuda/ptx.h @@ -24,7 +24,7 @@ #ifndef TVM_TARGET_SOURCE_PTX_H_ #define TVM_TARGET_SOURCE_PTX_H_ -#include +#include #include #include diff --git a/src/target/hexagon/llvm/codegen_hexagon.cc b/src/target/hexagon/llvm/codegen_hexagon.cc index a7dbda398d8f..c83af58c4ce7 100644 --- a/src/target/hexagon/llvm/codegen_hexagon.cc +++ b/src/target/hexagon/llvm/codegen_hexagon.cc @@ -45,6 +45,7 @@ #include #include #include +#include #include #include diff --git a/src/target/intrin_rule.cc b/src/target/intrin_rule.cc index 31e8b6a83290..9e1a8ce068cc 100644 --- a/src/target/intrin_rule.cc +++ b/src/target/intrin_rule.cc @@ -26,6 +26,7 @@ #include #include #include +#include namespace tvm { namespace codegen { diff --git a/src/target/llvm/codegen_aarch64.cc b/src/target/llvm/codegen_aarch64.cc index 7c328f18ab12..18da2e66d7a8 100644 --- a/src/target/llvm/codegen_aarch64.cc +++ b/src/target/llvm/codegen_aarch64.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include "../../arith/scalable_expression.h" #include "codegen_cpu.h" diff --git a/src/target/llvm/codegen_cpu.cc b/src/target/llvm/codegen_cpu.cc index 974cbb1e8a5c..09308a6ebbfd 100644 --- a/src/target/llvm/codegen_cpu.cc +++ b/src/target/llvm/codegen_cpu.cc @@ -51,6 +51,7 @@ #include #include #include +#include #include #include diff --git a/src/target/llvm/codegen_llvm.cc b/src/target/llvm/codegen_llvm.cc index 06c521a85502..a0e237500c19 100644 --- a/src/target/llvm/codegen_llvm.cc +++ b/src/target/llvm/codegen_llvm.cc @@ -81,6 +81,7 @@ #include #include #include +#include #include #include @@ -2265,7 +2266,7 @@ llvm::DIType* CodeGenLLVM::GetDebugType(const Type& ty_tir, llvm::Type* ty_llvm) if (dtype.is_scalable_vector()) return nullptr; - return dbg_info_->di_builder_->createBasicType(DLDataTypeToString(dtype).operator std::string(), + return dbg_info_->di_builder_->createBasicType(ffi::DLDataTypeToString(dtype).operator std::string(), dtype.bits() * dtype.lanes(), dwarf_type); } else { diff --git a/src/target/metal/codegen_metal.cc b/src/target/metal/codegen_metal.cc index e0cc10fe3e5c..c84df824a14f 100644 --- a/src/target/metal/codegen_metal.cc +++ b/src/target/metal/codegen_metal.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include #include diff --git a/src/target/rocm/llvm/codegen_amdgpu.cc b/src/target/rocm/llvm/codegen_amdgpu.cc index ef31805be7d2..2da399231e31 100644 --- a/src/target/rocm/llvm/codegen_amdgpu.cc +++ b/src/target/rocm/llvm/codegen_amdgpu.cc @@ -47,6 +47,7 @@ #include #include #include +#include #include "../../../runtime/metadata.h" #include "../../build_common.h" diff --git a/src/target/target_kind.cc b/src/target/target_kind.cc index 8d4cd50a9a6c..e593852e43ad 100644 --- a/src/target/target_kind.cc +++ b/src/target/target_kind.cc @@ -27,6 +27,7 @@ #include #include #include +#include #include diff --git a/src/tirx/analysis/verify_memory.cc b/src/tirx/analysis/verify_memory.cc index d97be6c5daae..6c4ba1193400 100644 --- a/src/tirx/analysis/verify_memory.cc +++ b/src/tirx/analysis/verify_memory.cc @@ -29,6 +29,7 @@ #include #include #include +#include namespace tvm { namespace tirx { diff --git a/src/tirx/analysis/verify_well_formed.cc b/src/tirx/analysis/verify_well_formed.cc index 6f9858dd074e..cc33e59d5690 100644 --- a/src/tirx/analysis/verify_well_formed.cc +++ b/src/tirx/analysis/verify_well_formed.cc @@ -40,7 +40,6 @@ namespace tvm { namespace tirx { -using AccessPath = ffi::reflection::AccessPath; namespace { @@ -234,19 +233,19 @@ class UndefinedVarVerifier : public Verifier { private: using Verifier::Visit; - void Visit(const PrimFunc& prim_func, AccessPath path) override { + void Visit(const PrimFunc& prim_func, ffi::reflection::AccessPath path) override { Verifier::Visit(prim_func, path); redefine_allowed_within_function_.clear(); } - void EnterDef(const IterVar& iter_var, AccessPath path) override { + void EnterDef(const IterVar& iter_var, ffi::reflection::AccessPath path) override { Verifier::EnterDef(iter_var, path); if (iter_var->iter_type == IterVarType::kThreadIndex) { redefine_allowed_within_function_.insert(iter_var->var); } } - void EnterDef(const Var& var, AccessPath path) override { + void EnterDef(const Var& var, ffi::reflection::AccessPath path) override { bool redefine_is_allowed = redefine_allowed_within_function_.count(var); { auto it = currently_defined_.find(var); @@ -272,14 +271,14 @@ class UndefinedVarVerifier : public Verifier { currently_defined_.insert({var, path}); } - void ExitDef(const Var& var, AccessPath path) override { + void ExitDef(const Var& var, ffi::reflection::AccessPath path) override { auto active_def = currently_defined_.find(var); currently_defined_.erase(active_def); previously_defined_.insert({var, path}); } - void VisitExpr_(const VarNode* op, AccessPath path) override { + void VisitExpr_(const VarNode* op, ffi::reflection::AccessPath path) override { auto var = ffi::GetRef(op); auto active_def = currently_defined_.find(var); @@ -298,10 +297,10 @@ class UndefinedVarVerifier : public Verifier { } // Variables that are defined in the currently-visited scope. - std::unordered_map currently_defined_; + std::unordered_map currently_defined_; // Variables that were previously defined, and are now out of scope. - std::unordered_map previously_defined_; + std::unordered_map previously_defined_; // Special variables that are allowed to be re-defined, so long as // that re-definition occurs within the same PrimFunc. For example @@ -328,20 +327,20 @@ class UndefinedBufferVerifier : public Verifier { private: using Verifier::Visit; - void Visit(const PrimFunc& prim_func, AccessPath path) override { + void Visit(const PrimFunc& prim_func, ffi::reflection::AccessPath path) override { Verifier::Visit(prim_func, path); // Clear per-function state (buffers should not cross function boundaries). currently_defined_.clear(); previously_defined_.clear(); } - void EnterDef(const Buffer& buffer, AccessPath path) override { + void EnterDef(const Buffer& buffer, ffi::reflection::AccessPath path) override { // Call the base class to visit buffer's internal vars (shape, strides, etc.) Verifier::EnterDef(buffer, path); currently_defined_.insert({buffer, path}); } - void ExitDef(const Buffer& buffer, AccessPath path) override { + void ExitDef(const Buffer& buffer, ffi::reflection::AccessPath path) override { auto active_def = currently_defined_.find(buffer); if (active_def != currently_defined_.end()) { currently_defined_.erase(active_def); @@ -349,7 +348,7 @@ class UndefinedBufferVerifier : public Verifier { previously_defined_.insert({buffer, path}); } - void VisitBufferUse(const Buffer& buffer, AccessPath path) override { + void VisitBufferUse(const Buffer& buffer, ffi::reflection::AccessPath path) override { bool is_declared = currently_defined_.count(buffer); bool was_declared = previously_defined_.count(buffer); @@ -369,10 +368,10 @@ class UndefinedBufferVerifier : public Verifier { } // Buffers defined in the currently-visited scope. - std::unordered_map + std::unordered_map currently_defined_; // Buffers that were previously defined and are now out of scope. - std::unordered_map + std::unordered_map previously_defined_; }; @@ -389,12 +388,12 @@ class SingleEnvThreadVerifier : public Verifier { using Verifier::Verifier; private: - void Visit(const PrimFunc& prim_func, AccessPath path) override { + void Visit(const PrimFunc& prim_func, ffi::reflection::AccessPath path) override { Verifier::Visit(prim_func, path); env_thread_vars_.clear(); } - void EnterDef(const IterVar& iter_var, AccessPath path) override { + void EnterDef(const IterVar& iter_var, ffi::reflection::AccessPath path) override { if (iter_var->iter_type == IterVarType::kThreadIndex) { if (auto it = env_thread_vars_.find(iter_var->thread_tag); it != env_thread_vars_.end()) { const auto& [prev_var, prev_path] = it->second; @@ -413,7 +412,7 @@ class SingleEnvThreadVerifier : public Verifier { } } - std::unordered_map> env_thread_vars_; + std::unordered_map> env_thread_vars_; }; bool VerifyWellFormed(const PrimFunc& func, bool assert_mode) { diff --git a/src/tirx/ir/stmt.cc b/src/tirx/ir/stmt.cc index 983cdaa9c602..4e2cfd3f1474 100644 --- a/src/tirx/ir/stmt.cc +++ b/src/tirx/ir/stmt.cc @@ -510,7 +510,7 @@ MatchBufferRegion::MatchBufferRegion(Buffer buffer, BufferRegion source) { // Validate shape TVM_FFI_ICHECK(source->region.size() >= buffer->shape.size()) - << "Dimension of source Region expected to be larger or equal than target buffer shape, but " + << "Dimension of source ffi::Array expected to be larger or equal than target buffer shape, but " "got " << source->region.size() << " vs. " << buffer->shape.size(); size_t offset = source->region.size() - buffer->shape.size(); diff --git a/src/tirx/ir/tir_visitor_with_path.cc b/src/tirx/ir/tir_visitor_with_path.cc index 857ccca08eff..512225344959 100644 --- a/src/tirx/ir/tir_visitor_with_path.cc +++ b/src/tirx/ir/tir_visitor_with_path.cc @@ -35,9 +35,8 @@ namespace tvm { namespace tirx { -using AccessPath = ffi::reflection::AccessPath; -void TIRVisitorWithPath::Visit(const IRModule& mod, AccessPath path) { +void TIRVisitorWithPath::Visit(const IRModule& mod, ffi::reflection::AccessPath path) { // To ensure deterministic order of visits, sort the GlobalVar first // by visibility (public then private), then alphabetically by name. std::vector gvars; @@ -76,7 +75,7 @@ void TIRVisitorWithPath::Visit(const IRModule& mod, AccessPath path) { while (context.size()) context.pop_back(); } -void TIRVisitorWithPath::Visit(const PrimFunc& func, AccessPath path) { +void TIRVisitorWithPath::Visit(const PrimFunc& func, ffi::reflection::AccessPath path) { // The implicit definitions from a PrimFunc::buffer_map are pretty // weird. They only apply if no previous definition of that // variable has occurred. Therefore, to ensure that we only avoid @@ -115,25 +114,25 @@ void TIRVisitorWithPath::Visit(const PrimFunc& func, AccessPath path) { while (context.size()) context.pop_back(); } -void TIRVisitorWithPath::EnterDef(const IterVar& iter_var, AccessPath path) { +void TIRVisitorWithPath::EnterDef(const IterVar& iter_var, ffi::reflection::AccessPath path) { if (iter_var->dom.defined()) { Visit(iter_var->dom, path->Attr("dom")); } EnterDef(iter_var->var, path->Attr("var")); } -void TIRVisitorWithPath::ExitDef(const IterVar& iter_var, AccessPath path) { +void TIRVisitorWithPath::ExitDef(const IterVar& iter_var, ffi::reflection::AccessPath path) { ExitDef(iter_var->var, path->Attr("var")); } -void TIRVisitorWithPath::EnterDef(const Buffer& buffer, AccessPath path) { +void TIRVisitorWithPath::EnterDef(const Buffer& buffer, ffi::reflection::AccessPath path) { // Defining a buffer counts as using all parameters in the buffer // (e.g. shape/strides). VisitBufferDef(buffer, path); } -void TIRVisitorWithPath::ExitDef(const Buffer& buffer, AccessPath path) {} +void TIRVisitorWithPath::ExitDef(const Buffer& buffer, ffi::reflection::AccessPath path) {} -void TIRVisitorWithPath::VisitBufferDef(const Buffer& buffer, AccessPath path) { +void TIRVisitorWithPath::VisitBufferDef(const Buffer& buffer, ffi::reflection::AccessPath path) { Visit(buffer->data, path->Attr("data")); Visit(buffer->shape, path->Attr("shape")); Visit(buffer->strides, path->Attr("strides")); @@ -145,14 +144,14 @@ void TIRVisitorWithPath::VisitBufferDef(const Buffer& buffer, AccessPath path) { // VisitBufferDef/EnterDef. Re-visiting at use sites would require those // variables to be in scope at every use, which may not hold when buffers // are allocated in a different scope than where they are used. -void TIRVisitorWithPath::VisitBufferUse(const Buffer& buffer, AccessPath path) {} +void TIRVisitorWithPath::VisitBufferUse(const Buffer& buffer, ffi::reflection::AccessPath path) {} -void TIRVisitorWithPath::Visit(const BufferRegion& region, AccessPath path) { +void TIRVisitorWithPath::Visit(const BufferRegion& region, ffi::reflection::AccessPath path) { VisitBufferUse(region->buffer, path->Attr("buffer")); Visit(region->region, path->Attr("region")); } -void TIRVisitorWithPath::Visit(const MatchBufferRegion& match, AccessPath path) { +void TIRVisitorWithPath::Visit(const MatchBufferRegion& match, ffi::reflection::AccessPath path) { Visit(match->source, path->Attr("source")); // MatchBufferRegion define the match->buffer, but do not own the @@ -160,26 +159,26 @@ void TIRVisitorWithPath::Visit(const MatchBufferRegion& match, AccessPath path) // definitions are handled in the BlockNode visitor. } -void TIRVisitorWithPath::Visit(const IterVar& iter_var, AccessPath path) { +void TIRVisitorWithPath::Visit(const IterVar& iter_var, ffi::reflection::AccessPath path) { if (iter_var->dom.defined()) { Visit(iter_var->dom, path->Attr("dom")); } Visit(iter_var->var, path->Attr("var")); } -void TIRVisitorWithPath::Visit(const Range& range, AccessPath path) { +void TIRVisitorWithPath::Visit(const Range& range, ffi::reflection::AccessPath path) { Visit(range->min, path->Attr("min")); Visit(range->extent, path->Attr("extent")); } -void TIRVisitorWithPath::VisitStmt_(const BindNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const BindNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); // Push the Bind's var definition into the current scope. // The def lives until the enclosing scope (body-carrying stmt) exits. bind_scope_.Current().push_back(WithDef(op->var, path->Attr("var"))); } -void TIRVisitorWithPath::VisitStmt_(const AttrStmtNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const AttrStmtNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); std::vector, DefContext, DefContext>> context; @@ -200,19 +199,19 @@ void TIRVisitorWithPath::VisitStmt_(const AttrStmtNode* op, AccessPath path) { } } -void TIRVisitorWithPath::VisitStmt_(const ForNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const ForNode* op, ffi::reflection::AccessPath path) { Visit(op->min, path->Attr("min")); Visit(op->extent, path->Attr("extent")); auto context = WithDef(op->loop_var, path->Attr("loop_var")); bind_scope_.WithNewScope([&]() { Visit(op->body, path->Attr("body")); }); } -void TIRVisitorWithPath::VisitStmt_(const WhileNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const WhileNode* op, ffi::reflection::AccessPath path) { Visit(op->condition, path->Attr("condition")); bind_scope_.WithNewScope([&]() { Visit(op->body, path->Attr("body")); }); } -void TIRVisitorWithPath::VisitStmt_(const AllocBufferNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const AllocBufferNode* op, ffi::reflection::AccessPath path) { // AllocBuffer both allocates the data variable and declares the buffer. // Push definitions into the current scope so they are visible to subsequent siblings. auto buf_path = path->Attr("buffer"); @@ -220,41 +219,41 @@ void TIRVisitorWithPath::VisitStmt_(const AllocBufferNode* op, AccessPath path) bind_scope_.Current().push_back(WithDef(op->buffer, buf_path)); } -void TIRVisitorWithPath::VisitStmt_(const DeclBufferNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const DeclBufferNode* op, ffi::reflection::AccessPath path) { // Push buffer definition into the current scope so it is visible to subsequent siblings. bind_scope_.Current().push_back(WithDef(op->buffer, path->Attr("buffer"))); } -void TIRVisitorWithPath::VisitStmt_(const BufferStoreNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const BufferStoreNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); VisitBufferUse(op->buffer, path->Attr("buffer")); Visit(op->indices, path->Attr("indices")); } -void TIRVisitorWithPath::VisitStmt_(const IfThenElseNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const IfThenElseNode* op, ffi::reflection::AccessPath path) { Visit(op->condition, path->Attr("condition")); bind_scope_.WithNewScope([&]() { Visit(op->then_case, path->Attr("then_case")); }); bind_scope_.WithNewScope([&]() { Visit(op->else_case, path->Attr("else_case")); }); } -void TIRVisitorWithPath::VisitStmt_(const AssertStmtNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const AssertStmtNode* op, ffi::reflection::AccessPath path) { Visit(op->condition, path->Attr("condition")); Visit(op->error_kind, path->Attr("error_kind")); Visit(op->message_parts, path->Attr("message_parts")); } -void TIRVisitorWithPath::VisitStmt_(const SeqStmtNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const SeqStmtNode* op, ffi::reflection::AccessPath path) { auto seq_path = path->Attr("seq"); for (size_t i = 0; i < op->seq.size(); i++) { Visit(op->seq[i], seq_path->ArrayItem(i)); } } -void TIRVisitorWithPath::VisitStmt_(const EvaluateNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const EvaluateNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); } -void TIRVisitorWithPath::VisitStmt_(const SBlockNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const SBlockNode* op, ffi::reflection::AccessPath path) { std::vector, DefContext, DefContext>> context; { @@ -300,34 +299,34 @@ void TIRVisitorWithPath::VisitStmt_(const SBlockNode* op, AccessPath path) { while (context.size()) context.pop_back(); } -void TIRVisitorWithPath::VisitStmt_(const SBlockRealizeNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitStmt_(const SBlockRealizeNode* op, ffi::reflection::AccessPath path) { Visit(op->iter_values, path->Attr("iter_values")); Visit(op->predicate, path->Attr("predicate")); Visit(op->block, path->Attr("block")); } -void TIRVisitorWithPath::VisitExpr_(const VarNode* op, AccessPath path) {} +void TIRVisitorWithPath::VisitExpr_(const VarNode* op, ffi::reflection::AccessPath path) {} -void TIRVisitorWithPath::VisitExpr_(const SizeVarNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const SizeVarNode* op, ffi::reflection::AccessPath path) { VisitExpr_(static_cast(op), path); } -void TIRVisitorWithPath::VisitExpr_(const BufferLoadNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const BufferLoadNode* op, ffi::reflection::AccessPath path) { VisitBufferUse(op->buffer, path->Attr("buffer")); Visit(op->indices, path->Attr("indices")); } -void TIRVisitorWithPath::VisitExpr_(const ProducerLoadNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const ProducerLoadNode* op, ffi::reflection::AccessPath path) { Visit(op->indices, path->Attr("indices")); } -void TIRVisitorWithPath::VisitExpr_(const LetNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const LetNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); auto context = WithDef(op->var, path->Attr("var")); Visit(op->body, path->Attr("body")); } -void TIRVisitorWithPath::VisitExpr_(const CallNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const CallNode* op, ffi::reflection::AccessPath path) { if (auto gvar = op->op.as()) { Visit(gvar.value(), path->Attr("op")); } @@ -335,7 +334,7 @@ void TIRVisitorWithPath::VisitExpr_(const CallNode* op, AccessPath path) { } #define DEFINE_BINOP_VISIT_(OP) \ - void TIRVisitorWithPath::VisitExpr_(const OP* op, AccessPath path) { \ + void TIRVisitorWithPath::VisitExpr_(const OP* op, ffi::reflection::AccessPath path) { \ Visit(op->a, path->Attr("a")); \ Visit(op->b, path->Attr("b")); \ } @@ -360,43 +359,43 @@ DEFINE_BINOP_VISIT_(OrNode); #undef DEFINE_BINOP_VISIT_ -void TIRVisitorWithPath::VisitExpr_(const IntImmNode* op, AccessPath path) {} -void TIRVisitorWithPath::VisitExpr_(const FloatImmNode* op, AccessPath path) {} -void TIRVisitorWithPath::VisitExpr_(const StringImmNode* op, AccessPath path) {} +void TIRVisitorWithPath::VisitExpr_(const IntImmNode* op, ffi::reflection::AccessPath path) {} +void TIRVisitorWithPath::VisitExpr_(const FloatImmNode* op, ffi::reflection::AccessPath path) {} +void TIRVisitorWithPath::VisitExpr_(const StringImmNode* op, ffi::reflection::AccessPath path) {} -void TIRVisitorWithPath::VisitExpr_(const ReduceNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const ReduceNode* op, ffi::reflection::AccessPath path) { Visit(op->axis, path->Attr("axis")); Visit(op->source, path->Attr("source")); Visit(op->init, path->Attr("init")); Visit(op->condition, path->Attr("condition")); } -void TIRVisitorWithPath::VisitExpr_(const CastNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const CastNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); } -void TIRVisitorWithPath::VisitExpr_(const NotNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const NotNode* op, ffi::reflection::AccessPath path) { Visit(op->a, path->Attr("a")); } -void TIRVisitorWithPath::VisitExpr_(const SelectNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const SelectNode* op, ffi::reflection::AccessPath path) { Visit(op->condition, path->Attr("condition")); Visit(op->true_value, path->Attr("true_value")); Visit(op->false_value, path->Attr("false_value")); } -void TIRVisitorWithPath::VisitExpr_(const RampNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const RampNode* op, ffi::reflection::AccessPath path) { Visit(op->base, path->Attr("base")); Visit(op->stride, path->Attr("stride")); Visit(op->lanes, path->Attr("lanes")); } -void TIRVisitorWithPath::VisitExpr_(const ShuffleNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const ShuffleNode* op, ffi::reflection::AccessPath path) { Visit(op->indices, path->Attr("indices")); Visit(op->vectors, path->Attr("vectors")); } -void TIRVisitorWithPath::VisitExpr_(const BroadcastNode* op, AccessPath path) { +void TIRVisitorWithPath::VisitExpr_(const BroadcastNode* op, ffi::reflection::AccessPath path) { Visit(op->value, path->Attr("value")); Visit(op->lanes, path->Attr("lanes")); } diff --git a/src/tirx/op/op.cc b/src/tirx/op/op.cc index f1c9c8a9b507..91539c9e7c28 100644 --- a/src/tirx/op/op.cc +++ b/src/tirx/op/op.cc @@ -25,6 +25,7 @@ #include #include +#include #include #include #include diff --git a/src/tirx/script/builder/ir.cc b/src/tirx/script/builder/ir.cc index 7044cfe7e390..6dc316590c01 100644 --- a/src/tirx/script/builder/ir.cc +++ b/src/tirx/script/builder/ir.cc @@ -23,6 +23,7 @@ #include #include #include +#include #include "./utils.h" diff --git a/src/tirx/script/printer/expr.cc b/src/tirx/script/printer/expr.cc index 6d2e13cbd4b7..d9902eb3aab0 100644 --- a/src/tirx/script/printer/expr.cc +++ b/src/tirx/script/printer/expr.cc @@ -16,6 +16,7 @@ * specific language governing permissions and limitations * under the License. */ +#include #include #include "./utils.h" diff --git a/src/tirx/transform/lower_intrin.cc b/src/tirx/transform/lower_intrin.cc index 772bdd4d5b2a..a8c60d33b9d2 100644 --- a/src/tirx/transform/lower_intrin.cc +++ b/src/tirx/transform/lower_intrin.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include #include diff --git a/src/tirx/transform/lower_tvm_builtin.cc b/src/tirx/transform/lower_tvm_builtin.cc index 8e71d2f26f40..3ba72294bb2c 100644 --- a/src/tirx/transform/lower_tvm_builtin.cc +++ b/src/tirx/transform/lower_tvm_builtin.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include diff --git a/src/tirx/transform/storage_rewrite.cc b/src/tirx/transform/storage_rewrite.cc index 99d62ab02b20..f64d262f97c8 100644 --- a/src/tirx/transform/storage_rewrite.cc +++ b/src/tirx/transform/storage_rewrite.cc @@ -33,6 +33,7 @@ #include #include #include +#include #include #include diff --git a/src/tirx/transform/tvm_ffi_binder.cc b/src/tirx/transform/tvm_ffi_binder.cc index 4a0c2d3e124c..881d94ba61ab 100644 --- a/src/tirx/transform/tvm_ffi_binder.cc +++ b/src/tirx/transform/tvm_ffi_binder.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include "ir_utils.h" @@ -107,7 +108,7 @@ void TVMFFIABIBuilder::EmitTypeIndexCheck(int param_index, const PrimExpr& cond, // RenderAccessPath // ============================================================ -ffi::String TVMFFIABIBuilder::RenderAccessPath(const AccessPath& path) const { +ffi::String TVMFFIABIBuilder::RenderAccessPath(const ffi::reflection::AccessPath& path) const { ffi::Array steps = path->ToSteps(); std::ostringstream os; bool first_printed = false; @@ -148,7 +149,7 @@ ffi::String TVMFFIABIBuilder::RenderAccessPath(const AccessPath& path) const { // GetParamIndex // ============================================================ -int TVMFFIABIBuilder::GetParamIndex(const AccessPath& path) const { +int TVMFFIABIBuilder::GetParamIndex(const ffi::reflection::AccessPath& path) const { ffi::Array steps = path->ToSteps(); if (steps.size() >= 1 && steps[0]->kind == ffi::reflection::AccessKind::kArrayItem) { return static_cast(steps[0]->key.cast()); @@ -157,11 +158,11 @@ int TVMFFIABIBuilder::GetParamIndex(const AccessPath& path) const { } // ============================================================ -// BindScalar (scalar bind with AccessPath) +// BindScalar (scalar bind with ffi::reflection::AccessPath) // ============================================================ bool TVMFFIABIBuilder::BindScalar(const PrimExpr& arg, const PrimExpr& value, - const AccessPath& path, bool with_lets) { + const ffi::reflection::AccessPath& path, bool with_lets) { TVM_FFI_ICHECK_EQ(arg.dtype(), value.dtype()); if (arg.as()) { Var v_arg = Downcast(arg); @@ -225,7 +226,7 @@ bool TVMFFIABIBuilder::BindScalar(const PrimExpr& arg, const PrimExpr& value, // ============================================================ /*! - * \brief Render PrimExpr to string with variable names replaced by AccessPath names. + * \brief Render PrimExpr to string with variable names replaced by ffi::reflection::AccessPath names. * * Uses ExprFunctor for generic dispatch over all expression types. * The default TIR printer sanitizes Var name_hints (e.g. "B.shape[0]" -> "B_shape_0_") @@ -324,24 +325,24 @@ TVMFFIABIBuilder::Result TVMFFIABIBuilder::Finalize() { } // ============================================================ -// BindArray (array bind with AccessPath) +// BindArray (array bind with ffi::reflection::AccessPath) // ============================================================ void TVMFFIABIBuilder::BindArray(const ffi::Array& arg, const ffi::Array& value, - const AccessPath& base_path) { + const ffi::reflection::AccessPath& base_path) { TVM_FFI_ICHECK_EQ(arg.size(), value.size()) << "Array size mismatch at " << RenderAccessPath(base_path); for (size_t i = 0; i < arg.size(); ++i) { - AccessPath elem_path = base_path->ArrayItem(i); + ffi::reflection::AccessPath elem_path = base_path->ArrayItem(i); BindScalar(arg[i], value[i], elem_path, false); } } // ============================================================ -// BindBuffer (buffer-to-buffer bind with AccessPath) +// BindBuffer (buffer-to-buffer bind with ffi::reflection::AccessPath) // ============================================================ -void TVMFFIABIBuilder::BindBuffer(const Buffer& arg, const Buffer& value, AccessPath base_path, +void TVMFFIABIBuilder::BindBuffer(const Buffer& arg, const Buffer& value, ffi::reflection::AccessPath base_path, bool fuzzy_match) { TVM_FFI_ICHECK_EQ(arg.scope(), value.scope()) << "Argument " << arg->name << " Buffer bind scope mismatch"; @@ -360,9 +361,9 @@ void TVMFFIABIBuilder::BindBuffer(const Buffer& arg, const Buffer& value, Access << " required elem_offset=" << arg->elem_offset << ", provided elem_offset=" << value->elem_offset; } - AccessPath data_path = base_path->Attr(ffi::String("data")); + ffi::reflection::AccessPath data_path = base_path->Attr(ffi::String("data")); BindScalar(arg->data, value->data, data_path, false); - AccessPath offset_path = base_path->Attr(ffi::String("elem_offset")); + ffi::reflection::AccessPath offset_path = base_path->Attr(ffi::String("elem_offset")); if (BindScalar(arg->elem_offset, value->elem_offset, offset_path, false)) { if (arg->offset_factor > 1) { PrimExpr offset = value->elem_offset; @@ -385,8 +386,8 @@ void TVMFFIABIBuilder::BindBuffer(const Buffer& arg, const Buffer& value, Access } } - AccessPath shape_path = base_path->Attr(ffi::String("shape")); - AccessPath strides_path = base_path->Attr(ffi::String("strides")); + ffi::reflection::AccessPath shape_path = base_path->Attr(ffi::String("shape")); + ffi::reflection::AccessPath strides_path = base_path->Attr(ffi::String("strides")); if (arg->shape.size() < value->shape.size()) { TVM_FFI_ICHECK(fuzzy_match) << "Buffer size mismatch at " << RenderAccessPath(base_path); @@ -397,14 +398,14 @@ void TVMFFIABIBuilder::BindBuffer(const Buffer& arg, const Buffer& value, Access << " vs " << value->shape; } for (size_t i = 0; i < arg->shape.size(); ++i) { - AccessPath shape_k_path = shape_path->ArrayItem(i); + ffi::reflection::AccessPath shape_k_path = shape_path->ArrayItem(i); BindScalar(arg->shape[i], value->shape[i + diff], shape_k_path, false); } if (value->strides.size() != 0) { TVM_FFI_ICHECK_EQ(arg->strides.size(), arg->shape.size()); TVM_FFI_ICHECK_EQ(value->strides.size(), value->shape.size()); for (size_t i = 0; i < arg->strides.size(); ++i) { - AccessPath strides_k_path = strides_path->ArrayItem(i); + ffi::reflection::AccessPath strides_k_path = strides_path->ArrayItem(i); BindScalar(arg->strides[i], value->strides[i + diff], strides_k_path, false); } } @@ -513,7 +514,7 @@ void TVMFFIABIBuilder::DecodeParam(int param_index) { } // Bind scalar param to loaded value (defines vars before buffer binds reference them) - AccessPath param_path = AccessPath::Root()->Extend(AccessStep::ArrayItem(param_index)); + ffi::reflection::AccessPath param_path = ffi::reflection::AccessPath::Root()->Extend(AccessStep::ArrayItem(param_index)); BindScalar(param, arg_value, param_path, true); } @@ -535,8 +536,8 @@ void TVMFFIABIBuilder::DecodeAllParams() { Var param = params_[i]; if (buffer_map_.count(param)) { Buffer buffer = buffer_map_[param]; - AccessPath param_path = - AccessPath::Root()->Extend(AccessStep::ArrayItem(i))->Attr(ffi::String(buffer->name)); + ffi::reflection::AccessPath param_path = + ffi::reflection::AccessPath::Root()->Extend(AccessStep::ArrayItem(i))->Attr(ffi::String(buffer->name)); DecodeParamDLTensor(buffer, device_type_, device_id_, param, func_name_ + "." + param->name_hint, param_path); decl_buffers_.push_back(DeclBuffer(buffer)); @@ -571,7 +572,7 @@ PrimExpr TVMFFIABIBuilder::LoadInt64ArrayElem(const Var& ptr, int index) { void TVMFFIABIBuilder::BindCompactStrides(const Buffer& buffer, const Var& strides_ptr, const PrimExpr& v_strides_is_null, - const AccessPath& param_path) { + const ffi::reflection::AccessPath& param_path) { DataType stype = buffer->DefaultIndexType(); PrimExpr expect_stride = make_const(stype, 1); ffi::Array conds; @@ -598,7 +599,7 @@ void TVMFFIABIBuilder::BindCompactStrides(const Buffer& buffer, const Var& strid void TVMFFIABIBuilder::BindAutoBroadcastStrides(const Buffer& buffer, const Var& strides_ptr, const PrimExpr& v_strides_is_null, - const AccessPath& param_path) { + const ffi::reflection::AccessPath& param_path) { DataType stype = buffer->DefaultIndexType(); PrimExpr stride = make_const(stype, 1); for (size_t i = buffer->shape.size(); i != 0; --i) { @@ -606,7 +607,7 @@ void TVMFFIABIBuilder::BindAutoBroadcastStrides(const Buffer& buffer, const Var& PrimExpr value = cast(buffer->shape[k].dtype(), LoadInt64ArrayElem(strides_ptr, k)); value = tvm::if_then_else(v_strides_is_null, stride, value); value = tvm::if_then_else(buffer->shape[k] == 1, 0, value); - AccessPath strides_k_path = param_path->Attr(ffi::String("strides"))->ArrayItem(k); + ffi::reflection::AccessPath strides_k_path = param_path->Attr(ffi::String("strides"))->ArrayItem(k); BindScalar(buffer->strides[k], value, strides_k_path, true); stride = analyzer_.Simplify(stride * buffer->shape[k]); } @@ -614,11 +615,11 @@ void TVMFFIABIBuilder::BindAutoBroadcastStrides(const Buffer& buffer, const Var& void TVMFFIABIBuilder::BindRegularStrides(const Buffer& buffer, const Var& strides_ptr, const Var& shape_ptr, const PrimExpr& v_strides_is_null, - const AccessPath& param_path) { + const ffi::reflection::AccessPath& param_path) { PrimExpr stride_from_shape = 1; for (int k = buffer->strides.size() - 1; k >= 0; k--) { PrimExpr explicit_stride = cast(buffer->shape[k].dtype(), LoadInt64ArrayElem(strides_ptr, k)); - AccessPath strides_k_path = param_path->Attr(ffi::String("strides"))->ArrayItem(k); + ffi::reflection::AccessPath strides_k_path = param_path->Attr(ffi::String("strides"))->ArrayItem(k); BindScalar(buffer->strides[k], tvm::if_then_else(v_strides_is_null, stride_from_shape, explicit_stride), strides_k_path, true); @@ -632,11 +633,11 @@ void TVMFFIABIBuilder::BindRegularStrides(const Buffer& buffer, const Var& strid void TVMFFIABIBuilder::DecodeParamDLTensor(const Buffer& buffer, const PrimExpr& device_type, const PrimExpr& device_id, const Var& handle, - const std::string& arg_name, AccessPath base_path) { + const std::string& arg_name, ffi::reflection::AccessPath base_path) { const DataType tvm_ndim_type = DataType::Int(32); std::string buf_name = buffer->name; - AccessPath param_path = base_path; + ffi::reflection::AccessPath param_path = base_path; int param_index = GetParamIndex(base_path); // ── Section: Null pointer check ────────────────────────────── @@ -675,7 +676,7 @@ void TVMFFIABIBuilder::DecodeParamDLTensor(const Buffer& buffer, const PrimExpr& buffer->dtype == DataType::Int(1)) { break; } - AccessPath shape_k_path = param_path->Attr(ffi::String("shape"))->ArrayItem(k); + ffi::reflection::AccessPath shape_k_path = param_path->Attr(ffi::String("shape"))->ArrayItem(k); BindScalar(buffer->shape[k], cast(buffer->shape[k].dtype(), LoadInt64ArrayElem(shape_ptr, k)), shape_k_path, true); } @@ -693,7 +694,7 @@ void TVMFFIABIBuilder::DecodeParamDLTensor(const Buffer& buffer, const PrimExpr& // ── Section: byte_offset ───────────────────────────────────── int data_bytes = GetVectorBytes(buffer->dtype); - AccessPath byte_offset_path = param_path->Attr(ffi::String("byte_offset")); + ffi::reflection::AccessPath byte_offset_path = param_path->Attr(ffi::String("byte_offset")); if (const auto* const_offset = buffer->elem_offset.as()) { BindScalar(make_const(DataType::UInt(64), const_offset->value * data_bytes), TVMStructGet(DataType::UInt(64), handle, 0, builtin::kDLTensorByteOffset), @@ -739,17 +740,17 @@ void TVMFFIABIBuilder::DecodeParamDLTensor(const Buffer& buffer, const PrimExpr& device_name); } } else { - AccessPath device_type_path = param_path->Attr(ffi::String("device_type")); + ffi::reflection::AccessPath device_type_path = param_path->Attr(ffi::String("device_type")); BindScalar(device_type_, actual_device_type, device_type_path, true); } - AccessPath device_id_path = param_path->Attr(ffi::String("device_id")); + ffi::reflection::AccessPath device_id_path = param_path->Attr(ffi::String("device_id")); BindScalar(device_id_, TVMStructGet(DataType::Int(32), handle, 0, builtin::kDLTensorDeviceId), device_id_path, true); } // ── Section: data pointer ──────────────────────────────────── { - AccessPath data_path = param_path->Attr(ffi::String("data")); + ffi::reflection::AccessPath data_path = param_path->Attr(ffi::String("data")); if (BindScalar(buffer->data, TVMStructGet(DataType::Handle(), handle, 0, builtin::kDLTensorData), data_path, true)) { diff --git a/src/tirx/transform/tvm_ffi_binder.h b/src/tirx/transform/tvm_ffi_binder.h index 5f17d970dd8f..03ed0b77fede 100644 --- a/src/tirx/transform/tvm_ffi_binder.h +++ b/src/tirx/transform/tvm_ffi_binder.h @@ -49,7 +49,7 @@ namespace tirx { * generation for packed function parameters. The primary public method is * DecodeAllParams(), which handles everything: type index extraction, * type checking (TypeError), value loading, scalar binding, buffer - * binding, and rich error message generation with AccessPath. + * binding, and rich error message generation with ffi::reflection::AccessPath. * * ## Generated statement ordering * @@ -85,7 +85,7 @@ namespace tirx { */ class TVMFFIABIBuilder { public: - /*! \brief Variable definition info: bound value and the AccessPath where first defined. */ + /*! \brief Variable definition info: bound value and the ffi::reflection::AccessPath where first defined. */ struct VarDefInfo { PrimExpr value; ffi::reflection::AccessPath first_def_path; @@ -221,10 +221,10 @@ class TVMFFIABIBuilder { */ PrimExpr DecodeParamFloat(int param_index, const Var& type_index, DataType dtype); - // ── Private binding submethods (all take AccessPath) ─────────── + // ── Private binding submethods (all take ffi::reflection::AccessPath) ─────────── /*! - * \brief Internal scalar bind with AccessPath tracking and rich error messages. + * \brief Internal scalar bind with ffi::reflection::AccessPath tracking and rich error messages. * * Binds \p arg to \p value. If arg is a Var not yet in var_defs_, creates a * new definition (Bind to init_nest_); otherwise emits a rich assertion @@ -232,36 +232,36 @@ class TVMFFIABIBuilder { * * When arg is a non-Var expression (e.g. batch_size + 1), the assertion is * deferred to Finalize() so display-var substitution can render the expression - * using AccessPath names (e.g. "k.shape[0] + 1" instead of "batch_size + 1"). + * using ffi::reflection::AccessPath names (e.g. "k.shape[0] + 1" instead of "batch_size + 1"). * * \param arg The argument expression to bind (typically a Var or constant). * \param value The value expression to bind to the argument. * \param with_lets If true, emit Bind bindings into init_nest_. - * \param path AccessPath for rich error message rendering. + * \param path ffi::reflection::AccessPath for rich error message rendering. * \return True if this was the first bind (definition created), false otherwise. */ bool BindScalar(const PrimExpr& arg, const PrimExpr& value, const ffi::reflection::AccessPath& path, bool with_lets); /*! - * \brief Array bind: binds element-wise with AccessPath[k] for each element. + * \brief Array bind: binds element-wise with ffi::reflection::AccessPath[k] for each element. * * \param arg The expected array of expressions. * \param value The actual array of expressions to bind against. - * \param base_path Base AccessPath; each element appends ArrayItem(k). + * \param base_path Base ffi::reflection::AccessPath; each element appends ArrayItem(k). */ void BindArray(const ffi::Array& arg, const ffi::Array& value, const ffi::reflection::AccessPath& base_path); /*! - * \brief Buffer-to-buffer bind with AccessPath. + * \brief Buffer-to-buffer bind with ffi::reflection::AccessPath. * * Binds data, elem_offset, shape, and strides of \p arg against \p value, * emitting assertions for any mismatches. * * \param arg The expected buffer definition. * \param value The actual buffer to bind against. - * \param base_path Base AccessPath for the buffer parameter. + * \param base_path Base ffi::reflection::AccessPath for the buffer parameter. * \param fuzzy_match If true, allow value to have more dimensions than arg. */ void BindBuffer(const Buffer& arg, const Buffer& value, ffi::reflection::AccessPath base_path, @@ -275,7 +275,7 @@ class TVMFFIABIBuilder { * \param device_id The expected device id expression. * \param handle The variable holding the DLTensor handle. * \param arg_name Human-readable name for error messages. - * \param base_path Base AccessPath for the tensor parameter. + * \param base_path Base ffi::reflection::AccessPath for the tensor parameter. */ void DecodeParamDLTensor(const Buffer& buffer, const PrimExpr& device_type, const PrimExpr& device_id, const Var& handle, @@ -310,7 +310,7 @@ class TVMFFIABIBuilder { * \param buffer The expected buffer definition. * \param strides_ptr The strides pointer variable. * \param v_strides_is_null Expression checking if strides pointer is NULL. - * \param param_path AccessPath for the tensor parameter. + * \param param_path ffi::reflection::AccessPath for the tensor parameter. */ void BindCompactStrides(const Buffer& buffer, const Var& strides_ptr, const PrimExpr& v_strides_is_null, @@ -322,7 +322,7 @@ class TVMFFIABIBuilder { * \param buffer The expected buffer definition. * \param strides_ptr The strides pointer variable. * \param v_strides_is_null Expression checking if strides pointer is NULL. - * \param param_path AccessPath for the tensor parameter. + * \param param_path ffi::reflection::AccessPath for the tensor parameter. */ void BindAutoBroadcastStrides(const Buffer& buffer, const Var& strides_ptr, const PrimExpr& v_strides_is_null, @@ -335,7 +335,7 @@ class TVMFFIABIBuilder { * \param strides_ptr The strides pointer variable. * \param shape_ptr The shape pointer variable (for computing C-contiguous strides). * \param v_strides_is_null Expression checking if strides pointer is NULL. - * \param param_path AccessPath for the tensor parameter. + * \param param_path ffi::reflection::AccessPath for the tensor parameter. */ void BindRegularStrides(const Buffer& buffer, const Var& strides_ptr, const Var& shape_ptr, const PrimExpr& v_strides_is_null, @@ -356,15 +356,15 @@ class TVMFFIABIBuilder { void EmitTypeIndexCheck(int param_index, const PrimExpr& cond, const std::string& expected_type); /*! - * \brief Render an AccessPath as a human-readable string (e.g. "a.shape[0]"). - * \param path The AccessPath to render. + * \brief Render an ffi::reflection::AccessPath as a human-readable string (e.g. "a.shape[0]"). + * \param path The ffi::reflection::AccessPath to render. * \return A human-readable string representation of the path. */ ffi::String RenderAccessPath(const ffi::reflection::AccessPath& path) const; /*! * \brief Extract param_index from the root ArrayItem step of a path. - * \param path The AccessPath to extract the index from. + * \param path The ffi::reflection::AccessPath to extract the index from. * \return The param index, or -1 if not found. */ int GetParamIndex(const ffi::reflection::AccessPath& path) const; @@ -373,7 +373,7 @@ class TVMFFIABIBuilder { * \brief Render pending constant-expression assertions with display-var substitution. * * For each pending assertion, substitutes known variable names with their - * AccessPath-rendered names (e.g. batch_size → "k.shape[0]") so error messages + * ffi::reflection::AccessPath-rendered names (e.g. batch_size → "k.shape[0]") so error messages * show human-readable expressions like "k.shape[0] + 1" instead of "batch_size + 1". */ void RenderPendingAsserts(); @@ -416,7 +416,7 @@ class TVMFFIABIBuilder { PrimExpr device_type_; /*! \brief The device id variable. */ PrimExpr device_id_; - /*! \brief Map from param_index to param_name for AccessPath rendering. */ + /*! \brief Map from param_index to param_name for ffi::reflection::AccessPath rendering. */ std::unordered_map param_names_; // Pre-cached common message fragments for string sharing across assertions diff --git a/src/tirx/transform/vectorize_loop.cc b/src/tirx/transform/vectorize_loop.cc index 104c111c722b..45d6a5e118be 100644 --- a/src/tirx/transform/vectorize_loop.cc +++ b/src/tirx/transform/vectorize_loop.cc @@ -33,6 +33,7 @@ #include #include #include +#include #include #include diff --git a/tests/cpp/ndarray_test.cc b/tests/cpp/ndarray_test.cc index fdb064b4a46b..c02efecc5148 100644 --- a/tests/cpp/ndarray_test.cc +++ b/tests/cpp/ndarray_test.cc @@ -30,7 +30,7 @@ TEST(TensorTest, IsContiguous_ContiguousStride) { int64_t strides[] = {10, 1}; managed_tensor->dl_tensor.strides = strides; - TVM_FFI_ICHECK(runtime::IsContiguous(managed_tensor->dl_tensor)); + TVM_FFI_ICHECK(ffi::IsContiguous(managed_tensor->dl_tensor)); managed_tensor->deleter(managed_tensor); } @@ -41,7 +41,7 @@ TEST(TensorTest, IsContiguous_NullStride) { managed_tensor->dl_tensor.strides = nullptr; - TVM_FFI_ICHECK(runtime::IsContiguous(managed_tensor->dl_tensor)); + TVM_FFI_ICHECK(ffi::IsContiguous(managed_tensor->dl_tensor)); managed_tensor->deleter(managed_tensor); } @@ -53,7 +53,7 @@ TEST(TensorTest, IsContiguous_AnyStrideForSingular) { int64_t strides[] = {10, 1, 1}; // strides[1] is normalized to 1 because shape[1] == 1. managed_tensor->dl_tensor.strides = strides; - TVM_FFI_ICHECK(runtime::IsContiguous(managed_tensor->dl_tensor)); + TVM_FFI_ICHECK(ffi::IsContiguous(managed_tensor->dl_tensor)); managed_tensor->dl_tensor.strides = nullptr; managed_tensor->deleter(managed_tensor); @@ -66,7 +66,7 @@ TEST(TensorTest, IsContiguous_UncontiguousStride) { int64_t strides[] = {1, 1, 1}; managed_tensor->dl_tensor.strides = strides; - TVM_FFI_ICHECK(!runtime::IsContiguous(managed_tensor->dl_tensor)); + TVM_FFI_ICHECK(!ffi::IsContiguous(managed_tensor->dl_tensor)); managed_tensor->dl_tensor.strides = nullptr; managed_tensor->deleter(managed_tensor); diff --git a/tests/cpp/tir_scalable_datatype.cc b/tests/cpp/tir_scalable_datatype.cc index 9be9e8552e83..fd9f76eee366 100644 --- a/tests/cpp/tir_scalable_datatype.cc +++ b/tests/cpp/tir_scalable_datatype.cc @@ -83,7 +83,7 @@ TEST(ScalableDataType, TestIsScalar) { TEST(ScalableDataType, TestScalableDataTypeToString) { tvm::DataType scalable_type = tvm::DataType(kDLInt, 32, 4, true); - EXPECT_EQ(tvm::runtime::DLDataTypeToString(scalable_type), "int32xvscalex4"); + EXPECT_EQ(tvm::ffi::DLDataTypeToString(scalable_type), "int32xvscalex4"); } TEST(ScalableDataType, TestStringToScalableDataType) {