This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new cc9bd00d97 [FFI] Bump tvm-ffi dependency (#18370)
cc9bd00d97 is described below
commit cc9bd00d97e750615e0faaa34f49e89bb782b50d
Author: Tianqi Chen <[email protected]>
AuthorDate: Thu Oct 16 09:59:07 2025 -0400
[FFI] Bump tvm-ffi dependency (#18370)
* [FFI] Bump tvm-ffi dependency
This PR bumps tvm-ffi dependency to latest
Co-authored-by: Junru Shao <[email protected]>
Co-authored-by: Junru Shao <[email protected]>
---
3rdparty/tvm-ffi | 2 +-
include/tvm/ir/name_supply.h | 6 ++++++
include/tvm/meta_schedule/database.h | 2 ++
include/tvm/meta_schedule/feature_extractor.h | 5 ++++-
include/tvm/meta_schedule/measure_callback.h | 5 ++++-
include/tvm/meta_schedule/mutator.h | 3 ++-
include/tvm/meta_schedule/postproc.h | 5 ++++-
include/tvm/meta_schedule/profiler.h | 4 ++--
include/tvm/meta_schedule/runner.h | 9 +++++++++
include/tvm/meta_schedule/schedule_rule.h | 5 ++++-
include/tvm/meta_schedule/search_strategy.h | 2 ++
include/tvm/meta_schedule/space_generator.h | 2 ++
include/tvm/relax/nested_msg.h | 10 +++++++++-
include/tvm/runtime/data_type.h | 4 ++++
include/tvm/runtime/tensor.h | 1 +
include/tvm/tir/block_scope.h | 3 ++-
include/tvm/tir/schedule/schedule.h | 6 ++++--
python/tvm/testing/_ffi_api.py | 3 +++
src/arith/presburger_set.cc | 1 +
src/ir/name_supply.cc | 1 +
src/meta_schedule/measure_callback/add_to_database.cc | 7 +++++++
.../measure_callback/remove_build_artifact.cc | 7 +++++++
src/meta_schedule/measure_callback/update_cost_model.cc | 7 +++++++
.../postproc/disallow_async_strided_mem_copy.cc | 7 +++++++
src/meta_schedule/postproc/disallow_dynamic_loop.cc | 7 +++++++
src/meta_schedule/postproc/rewrite_cooperative_fetch.cc | 1 +
src/meta_schedule/postproc/rewrite_layout.cc | 7 +++++++
.../postproc/rewrite_parallel_vectorize_unroll.cc | 7 +++++++
src/meta_schedule/postproc/rewrite_reduction_block.cc | 4 ++--
src/meta_schedule/postproc/rewrite_tensorize.cc | 4 ++--
src/meta_schedule/postproc/verify_gpu_code.cc | 7 +++++++
src/meta_schedule/postproc/verify_vtcm_limit.cc | 7 +++++++
src/meta_schedule/runner/runner.cc | 1 +
.../schedule_rule/multi_level_tiling_tensor_core.cc | 14 ++++++++++++++
.../schedule_rule/multi_level_tiling_wide_vector.cc | 7 +++++++
.../schedule_rule/multi_level_tiling_with_intrin.cc | 6 ++++++
src/meta_schedule/search_strategy/replay_func.cc | 3 ++-
src/meta_schedule/space_generator/post_order_apply.cc | 3 ++-
src/meta_schedule/space_generator/schedule_fn.cc | 6 ++++--
src/relax/ir/block_builder.cc | 1 +
src/relax/ir/py_expr_functor.cc | 1 +
src/relax/transform/fold_constant.cc | 2 +-
src/runtime/contrib/random/mt_random_engine.cc | 9 +++++----
src/runtime/disco/distributed/socket_session.cc | 1 +
src/runtime/disco/process_session.cc | 1 +
src/runtime/disco/session.cc | 2 ++
src/runtime/disco/threaded_session.cc | 6 ++++++
src/runtime/profiling.cc | 3 +++
src/runtime/rpc/rpc_session.cc | 3 +++
src/runtime/vm/builtin.cc | 2 +-
src/tir/ir/py_functor.cc | 6 ++++--
src/tir/schedule/concrete_schedule.h | 3 ++-
src/tir/schedule/schedule.cc | 2 ++
src/tir/schedule/traced_schedule.h | 3 ++-
tests/python/tir-base/test_tir_ptx_cp_async.py | 3 ++-
.../test_tir_transform_inject_ptx_async_copy.py | 1 +
web/emcc/wasm_runtime.cc | 2 +-
57 files changed, 210 insertions(+), 32 deletions(-)
diff --git a/3rdparty/tvm-ffi b/3rdparty/tvm-ffi
index 4fefeb0f59..59c91c17eb 160000
--- a/3rdparty/tvm-ffi
+++ b/3rdparty/tvm-ffi
@@ -1 +1 @@
-Subproject commit 4fefeb0f5913fc41cf860f517b9320f1bf1d0e98
+Subproject commit 59c91c17eb7ef4f24cf00faedc82f1a8e0fc53a3
diff --git a/include/tvm/ir/name_supply.h b/include/tvm/ir/name_supply.h
index 2de0164eb2..d3139ea2c8 100644
--- a/include/tvm/ir/name_supply.h
+++ b/include/tvm/ir/name_supply.h
@@ -86,6 +86,12 @@ class NameSupplyNode : public Object {
std::string prefix_;
static constexpr const bool _type_mutable = true;
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<NameSupplyNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("ir.NameSupply", NameSupplyNode, Object);
private:
diff --git a/include/tvm/meta_schedule/database.h
b/include/tvm/meta_schedule/database.h
index 6ffd188319..ebd945482f 100644
--- a/include/tvm/meta_schedule/database.h
+++ b/include/tvm/meta_schedule/database.h
@@ -391,6 +391,8 @@ class PyDatabaseNode : public DatabaseNode {
// `f_query_schedule` is not registered
// `f_query_ir_module` is not registered
// `f_size` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyDatabaseNode>();
}
bool HasWorkload(const IRModule& mod) final {
diff --git a/include/tvm/meta_schedule/feature_extractor.h
b/include/tvm/meta_schedule/feature_extractor.h
index a2f7b90196..9a339d39e7 100644
--- a/include/tvm/meta_schedule/feature_extractor.h
+++ b/include/tvm/meta_schedule/feature_extractor.h
@@ -40,7 +40,8 @@ class FeatureExtractorNode : public runtime::Object {
virtual ~FeatureExtractorNode() = default;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<FeatureExtractorNode>();
}
/*!
@@ -79,6 +80,8 @@ class PyFeatureExtractorNode : public FeatureExtractorNode {
static void RegisterReflection() {
// `f_extract_from` is not registered
// `f_as_string` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyFeatureExtractorNode>();
}
ffi::Array<tvm::runtime::Tensor> ExtractFrom(
diff --git a/include/tvm/meta_schedule/measure_callback.h
b/include/tvm/meta_schedule/measure_callback.h
index 04c855e705..9e7d49a0c9 100644
--- a/include/tvm/meta_schedule/measure_callback.h
+++ b/include/tvm/meta_schedule/measure_callback.h
@@ -43,7 +43,8 @@ class MeasureCallbackNode : public runtime::Object {
virtual ~MeasureCallbackNode() = default;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<MeasureCallbackNode>();
}
/*!
@@ -95,6 +96,8 @@ class PyMeasureCallbackNode : public MeasureCallbackNode {
static void RegisterReflection() {
// `f_apply` is not registered
// `f_as_string` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyMeasureCallbackNode>();
}
void Apply(const TaskScheduler& task_scheduler, //
diff --git a/include/tvm/meta_schedule/mutator.h
b/include/tvm/meta_schedule/mutator.h
index a6522c23f3..05489c7552 100644
--- a/include/tvm/meta_schedule/mutator.h
+++ b/include/tvm/meta_schedule/mutator.h
@@ -41,7 +41,8 @@ class MutatorNode : public runtime::Object {
virtual ~MutatorNode() = default;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<MutatorNode>();
}
/*!
diff --git a/include/tvm/meta_schedule/postproc.h
b/include/tvm/meta_schedule/postproc.h
index fbf96fe990..948f752107 100644
--- a/include/tvm/meta_schedule/postproc.h
+++ b/include/tvm/meta_schedule/postproc.h
@@ -40,7 +40,8 @@ class PostprocNode : public runtime::Object {
virtual ~PostprocNode() = default;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PostprocNode>();
}
/*!
@@ -199,6 +200,8 @@ class PyPostprocNode : public PostprocNode {
// `f_apply` is not registered
// `f_clone` is not registered
// `f_as_string` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyPostprocNode>();
}
void InitializeWithTuneContext(const TuneContext& context) final;
diff --git a/include/tvm/meta_schedule/profiler.h
b/include/tvm/meta_schedule/profiler.h
index abad1ae54f..5b82e6606b 100644
--- a/include/tvm/meta_schedule/profiler.h
+++ b/include/tvm/meta_schedule/profiler.h
@@ -60,8 +60,8 @@ class ProfilerNode : public runtime::Object {
ffi::Function total_timer;
static void RegisterReflection() {
- // `stats_sec` is not registered
- // `total_timer` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ProfilerNode>();
}
static constexpr const bool _type_mutable = true;
diff --git a/include/tvm/meta_schedule/runner.h
b/include/tvm/meta_schedule/runner.h
index 9457167b30..a88ae5feac 100644
--- a/include/tvm/meta_schedule/runner.h
+++ b/include/tvm/meta_schedule/runner.h
@@ -128,6 +128,8 @@ class RunnerFutureNode : public runtime::Object {
static void RegisterReflection() {
// `f_done` is not registered
// `f_result` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<RunnerFutureNode>();
}
/*!
@@ -189,6 +191,11 @@ class RunnerNode : public runtime::Object {
*/
virtual ffi::Array<RunnerFuture> Run(ffi::Array<RunnerInput> runner_inputs)
= 0;
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<RunnerNode>();
+ }
+
static constexpr const bool _type_mutable = true;
TVM_FFI_DECLARE_OBJECT_INFO("meta_schedule.Runner", RunnerNode,
runtime::Object);
};
@@ -222,6 +229,8 @@ class PyRunnerNode : public RunnerNode {
static void RegisterReflection() {
// `f_run` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyRunnerNode>();
}
ffi::Array<RunnerFuture> Run(ffi::Array<RunnerInput> runner_inputs) final {
diff --git a/include/tvm/meta_schedule/schedule_rule.h
b/include/tvm/meta_schedule/schedule_rule.h
index d55d47373c..be9074acbd 100644
--- a/include/tvm/meta_schedule/schedule_rule.h
+++ b/include/tvm/meta_schedule/schedule_rule.h
@@ -43,7 +43,8 @@ class ScheduleRuleNode : public runtime::Object {
virtual ~ScheduleRuleNode() = default;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ScheduleRuleNode>();
}
/*!
@@ -337,6 +338,8 @@ class PyScheduleRuleNode : public ScheduleRuleNode {
// `f_apply` is not registered
// `f_as_string` is not registered
// `f_clone` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyScheduleRuleNode>();
}
void InitializeWithTuneContext(const TuneContext& context) final;
diff --git a/include/tvm/meta_schedule/search_strategy.h
b/include/tvm/meta_schedule/search_strategy.h
index aeb2a4da35..714c43470f 100644
--- a/include/tvm/meta_schedule/search_strategy.h
+++ b/include/tvm/meta_schedule/search_strategy.h
@@ -249,6 +249,8 @@ class PySearchStrategyNode : public SearchStrategyNode {
// `f_generate_measure_candidates` is not registered
// `f_notify_runner_results` is not registered
// `f_clone` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PySearchStrategyNode>();
}
void InitializeWithTuneContext(const TuneContext& context) final;
diff --git a/include/tvm/meta_schedule/space_generator.h
b/include/tvm/meta_schedule/space_generator.h
index 67d15ebe96..460a41e44a 100644
--- a/include/tvm/meta_schedule/space_generator.h
+++ b/include/tvm/meta_schedule/space_generator.h
@@ -227,6 +227,8 @@ class PySpaceGeneratorNode : public SpaceGeneratorNode {
// `f_initialize_with_tune_context` is not registered
// `f_generate_design_space` is not registered
// `f_clone` is not registered
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PySpaceGeneratorNode>();
}
void InitializeWithTuneContext(const TuneContext& context) final;
diff --git a/include/tvm/relax/nested_msg.h b/include/tvm/relax/nested_msg.h
index aac3175d72..77f001630f 100644
--- a/include/tvm/relax/nested_msg.h
+++ b/include/tvm/relax/nested_msg.h
@@ -639,7 +639,7 @@ struct TypeTraits<relax::NestedMsg<T>> : public
TypeTraitsBase {
}
TVM_FFI_INLINE static relax::NestedMsg<T> MoveFromAnyAfterCheck(TVMFFIAny*
src) {
- return
relax::NestedMsg<T>(details::AnyUnsafe::MoveTVMFFIAnyToAny(std::move(*src)));
+ return relax::NestedMsg<T>(details::AnyUnsafe::MoveTVMFFIAnyToAny(src));
}
static std::optional<relax::NestedMsg<T>> TryCastFromAnyView(const
TVMFFIAny* src) {
@@ -673,6 +673,14 @@ struct TypeTraits<relax::NestedMsg<T>> : public
TypeTraitsBase {
TVM_FFI_INLINE static std::string TypeStr() {
return "NestedMsg<" + details::Type2Str<T>::v() + ">";
}
+
+ TVM_FFI_INLINE static std::string TypeSchema() {
+ std::ostringstream oss;
+ oss << R"({"type":"NestedMsg","args":[)";
+ oss << details::TypeSchema<T>::v();
+ oss << "]}";
+ return oss.str();
+ }
};
} // namespace ffi
} // namespace tvm
diff --git a/include/tvm/runtime/data_type.h b/include/tvm/runtime/data_type.h
index 230f73747f..0af3022bbd 100644
--- a/include/tvm/runtime/data_type.h
+++ b/include/tvm/runtime/data_type.h
@@ -497,6 +497,10 @@ struct TypeTraits<runtime::DataType> : public
TypeTraitsBase {
}
TVM_FFI_INLINE static std::string TypeStr() { return
ffi::StaticTypeKey::kTVMFFIDataType; }
+
+ TVM_FFI_INLINE static std::string TypeSchema() {
+ return R"({"type":")" + std::string(ffi::StaticTypeKey::kTVMFFIDataType) +
R"("})";
+ }
};
} // namespace ffi
diff --git a/include/tvm/runtime/tensor.h b/include/tvm/runtime/tensor.h
index 3028723957..e32101aac2 100644
--- a/include/tvm/runtime/tensor.h
+++ b/include/tvm/runtime/tensor.h
@@ -74,6 +74,7 @@ class Tensor : public tvm::ffi::Tensor {
static Tensor FromDLPackVersioned(DLManagedTensorVersioned* tensor) {
return tvm::ffi::Tensor::FromDLPackVersioned(tensor, kAllocAlignment,
true);
}
+ inline const DLTensor* operator->() const { return this->get(); }
/*!
* \brief Copy data content from another array.
* \param other The source array to be copied from.
diff --git a/include/tvm/tir/block_scope.h b/include/tvm/tir/block_scope.h
index ae30613eb2..f1120c7837 100644
--- a/include/tvm/tir/block_scope.h
+++ b/include/tvm/tir/block_scope.h
@@ -264,7 +264,8 @@ class BlockScopeNode : public Object {
std::unordered_map<Buffer, ffi::Array<StmtSRef>, ObjectPtrHash,
ObjectPtrEqual> buffer_writers;
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<BlockScopeNode>();
}
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tir.BlockScope", BlockScopeNode, Object);
diff --git a/include/tvm/tir/schedule/schedule.h
b/include/tvm/tir/schedule/schedule.h
index c5695f62d9..60deae801f 100644
--- a/include/tvm/tir/schedule/schedule.h
+++ b/include/tvm/tir/schedule/schedule.h
@@ -51,7 +51,8 @@ enum class BufferIndexType : int32_t {
class BlockRVNode : public runtime::Object {
public:
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<BlockRVNode>();
}
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tir.BlockRV", BlockRVNode,
runtime::Object);
};
@@ -73,7 +74,8 @@ class BlockRV : public runtime::ObjectRef {
class LoopRVNode : public runtime::Object {
public:
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<LoopRVNode>();
}
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tir.LoopRV", LoopRVNode, runtime::Object);
};
diff --git a/python/tvm/testing/_ffi_api.py b/python/tvm/testing/_ffi_api.py
index 6cb0b9bac4..b7a0b59fd0 100644
--- a/python/tvm/testing/_ffi_api.py
+++ b/python/tvm/testing/_ffi_api.py
@@ -17,5 +17,8 @@
"""FFI APIs for tvm.testing"""
import tvm_ffi
+# must import testing before init_ffi_api
+import tvm_ffi.testing
+
tvm_ffi.init_ffi_api("testing", __name__)
diff --git a/src/arith/presburger_set.cc b/src/arith/presburger_set.cc
index 3722837830..a91b4e2117 100644
--- a/src/arith/presburger_set.cc
+++ b/src/arith/presburger_set.cc
@@ -277,6 +277,7 @@ PresburgerSet MakePresburgerSet(const PrimExpr& constraint)
{ return PresburgerS
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ PresburgerSetNode::RegisterReflection();
refl::GlobalDef().def("arith.PresburgerSet", MakePresburgerSet);
}
diff --git a/src/ir/name_supply.cc b/src/ir/name_supply.cc
index cc6db0c21f..e5b94dff5a 100644
--- a/src/ir/name_supply.cc
+++ b/src/ir/name_supply.cc
@@ -93,6 +93,7 @@ std::string NameSupplyNode::GetUniqueName(std::string name,
bool add_underscore)
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ NameSupplyNode::RegisterReflection();
refl::GlobalDef()
.def("ir.NameSupply", [](ffi::String prefix) { return
NameSupply(prefix); })
.def_method("ir.NameSupply_FreshName", &NameSupplyNode::FreshName)
diff --git a/src/meta_schedule/measure_callback/add_to_database.cc
b/src/meta_schedule/measure_callback/add_to_database.cc
index 76d5b1c7ce..a7b455eec7 100644
--- a/src/meta_schedule/measure_callback/add_to_database.cc
+++ b/src/meta_schedule/measure_callback/add_to_database.cc
@@ -56,6 +56,12 @@ class AddToDatabaseNode : public MeasureCallbackNode {
/*args_info=*/candidate->args_info));
}
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<AddToDatabaseNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.AddToDatabase",
AddToDatabaseNode,
MeasureCallbackNode);
};
@@ -67,6 +73,7 @@ MeasureCallback MeasureCallback::AddToDatabase() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<AddToDatabaseNode>();
refl::GlobalDef().def("meta_schedule.MeasureCallbackAddToDatabase",
MeasureCallback::AddToDatabase);
}
diff --git a/src/meta_schedule/measure_callback/remove_build_artifact.cc
b/src/meta_schedule/measure_callback/remove_build_artifact.cc
index bee5b0b03e..18f00efab5 100644
--- a/src/meta_schedule/measure_callback/remove_build_artifact.cc
+++ b/src/meta_schedule/measure_callback/remove_build_artifact.cc
@@ -37,6 +37,12 @@ class RemoveBuildArtifactNode : public MeasureCallbackNode {
}
}
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<RemoveBuildArtifactNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.RemoveBuildArtifact",
RemoveBuildArtifactNode,
MeasureCallbackNode);
};
@@ -48,6 +54,7 @@ MeasureCallback MeasureCallback::RemoveBuildArtifact() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ RemoveBuildArtifactNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.MeasureCallbackRemoveBuildArtifact",
MeasureCallback::RemoveBuildArtifact);
}
diff --git a/src/meta_schedule/measure_callback/update_cost_model.cc
b/src/meta_schedule/measure_callback/update_cost_model.cc
index 38f714b03a..845e14e1e7 100644
--- a/src/meta_schedule/measure_callback/update_cost_model.cc
+++ b/src/meta_schedule/measure_callback/update_cost_model.cc
@@ -54,6 +54,12 @@ class UpdateCostModelNode : public MeasureCallbackNode {
}
cost_model->Update(task->ctx, pruned_candidate, pruned_runner_result);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<UpdateCostModelNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.UpdateCostModel",
UpdateCostModelNode,
MeasureCallbackNode);
};
@@ -65,6 +71,7 @@ MeasureCallback MeasureCallback::UpdateCostModel() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ UpdateCostModelNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.MeasureCallbackUpdateCostModel",
MeasureCallback::UpdateCostModel);
}
diff --git a/src/meta_schedule/postproc/disallow_async_strided_mem_copy.cc
b/src/meta_schedule/postproc/disallow_async_strided_mem_copy.cc
index 94789ee402..9f59404de5 100644
--- a/src/meta_schedule/postproc/disallow_async_strided_mem_copy.cc
+++ b/src/meta_schedule/postproc/disallow_async_strided_mem_copy.cc
@@ -173,6 +173,12 @@ class DisallowAsyncStridedMemCopyNode : public
PostprocNode {
ffi::make_object<DisallowAsyncStridedMemCopyNode>(*this);
return Postproc(n);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<DisallowAsyncStridedMemCopyNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.DisallowAsyncStridedMemCopy",
DisallowAsyncStridedMemCopyNode,
PostprocNode);
@@ -188,6 +194,7 @@ Postproc Postproc::DisallowAsyncStridedMemCopy() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ DisallowAsyncStridedMemCopyNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocDisallowAsyncStridedMemCopy",
Postproc::DisallowAsyncStridedMemCopy);
}
diff --git a/src/meta_schedule/postproc/disallow_dynamic_loop.cc
b/src/meta_schedule/postproc/disallow_dynamic_loop.cc
index bd69b3a21a..df7344455e 100644
--- a/src/meta_schedule/postproc/disallow_dynamic_loop.cc
+++ b/src/meta_schedule/postproc/disallow_dynamic_loop.cc
@@ -74,6 +74,12 @@ class DisallowDynamicLoopNode : public PostprocNode {
ObjectPtr<DisallowDynamicLoopNode> n =
ffi::make_object<DisallowDynamicLoopNode>(*this);
return Postproc(n);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<DisallowDynamicLoopNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.DisallowDynamicLoop",
DisallowDynamicLoopNode,
PostprocNode);
};
@@ -85,6 +91,7 @@ Postproc Postproc::DisallowDynamicLoop() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ DisallowDynamicLoopNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocDisallowDynamicLoop",
Postproc::DisallowDynamicLoop);
}
diff --git a/src/meta_schedule/postproc/rewrite_cooperative_fetch.cc
b/src/meta_schedule/postproc/rewrite_cooperative_fetch.cc
index e0c4b5c8f1..ae7b693efd 100644
--- a/src/meta_schedule/postproc/rewrite_cooperative_fetch.cc
+++ b/src/meta_schedule/postproc/rewrite_cooperative_fetch.cc
@@ -139,6 +139,7 @@ class RewriteCooperativeFetchNode : public PostprocNode {
ObjectPtr<RewriteCooperativeFetchNode> n =
ffi::make_object<RewriteCooperativeFetchNode>(*this);
return Postproc(n);
}
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.RewriteCooperativeFetch",
RewriteCooperativeFetchNode, PostprocNode);
diff --git a/src/meta_schedule/postproc/rewrite_layout.cc
b/src/meta_schedule/postproc/rewrite_layout.cc
index 3712c77791..17acdcc9bf 100644
--- a/src/meta_schedule/postproc/rewrite_layout.cc
+++ b/src/meta_schedule/postproc/rewrite_layout.cc
@@ -264,6 +264,12 @@ class RewriteLayoutNode : public PostprocNode {
ObjectPtr<RewriteLayoutNode> n =
ffi::make_object<RewriteLayoutNode>(*this);
return Postproc(n);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<RewriteLayoutNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.RewriteLayout",
RewriteLayoutNode, PostprocNode);
};
@@ -274,6 +280,7 @@ Postproc Postproc::RewriteLayout() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ RewriteLayoutNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocRewriteLayout",
Postproc::RewriteLayout);
}
diff --git a/src/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc
b/src/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc
index 340211663b..d833af6142 100644
--- a/src/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc
+++ b/src/meta_schedule/postproc/rewrite_parallel_vectorize_unroll.cc
@@ -455,6 +455,12 @@ class RewriteParallelVectorizeUnrollNode : public
PostprocNode {
ffi::make_object<RewriteParallelVectorizeUnrollNode>(*this);
return Postproc(n);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<RewriteParallelVectorizeUnrollNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.RewriteParallelVectorizeUnroll",
RewriteParallelVectorizeUnrollNode,
PostprocNode);
};
@@ -467,6 +473,7 @@ Postproc Postproc::RewriteParallelVectorizeUnroll() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ RewriteParallelVectorizeUnrollNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocRewriteParallelVectorizeUnroll",
Postproc::RewriteParallelVectorizeUnroll);
}
diff --git a/src/meta_schedule/postproc/rewrite_reduction_block.cc
b/src/meta_schedule/postproc/rewrite_reduction_block.cc
index 74a80cf80b..fffef8ba68 100644
--- a/src/meta_schedule/postproc/rewrite_reduction_block.cc
+++ b/src/meta_schedule/postproc/rewrite_reduction_block.cc
@@ -125,6 +125,7 @@ class RewriteReductionBlockNode : public PostprocNode {
ObjectPtr<RewriteReductionBlockNode> n =
ffi::make_object<RewriteReductionBlockNode>(*this);
return Postproc(n);
}
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.RewriteReductionBlock",
RewriteReductionBlockNode, PostprocNode);
};
@@ -178,11 +179,10 @@ Postproc Postproc::RewriteReductionBlock() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ RewriteReductionBlockNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocRewriteReductionBlock",
Postproc::RewriteReductionBlock);
}
-TVM_FFI_STATIC_INIT_BLOCK() { RewriteReductionBlockNode::RegisterReflection();
}
-
} // namespace meta_schedule
} // namespace tvm
diff --git a/src/meta_schedule/postproc/rewrite_tensorize.cc
b/src/meta_schedule/postproc/rewrite_tensorize.cc
index 3a1024e410..473731b5a7 100644
--- a/src/meta_schedule/postproc/rewrite_tensorize.cc
+++ b/src/meta_schedule/postproc/rewrite_tensorize.cc
@@ -78,6 +78,7 @@ class RewriteTensorizeNode : public PostprocNode {
}
bool vectorize_init_loop = false;
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.RewriteTensorize",
RewriteTensorizeNode,
PostprocNode);
};
@@ -109,10 +110,9 @@ Postproc Postproc::RewriteTensorize(bool
vectorize_init_loop) {
return Postproc(n);
}
-TVM_FFI_STATIC_INIT_BLOCK() { RewriteTensorizeNode::RegisterReflection(); }
-
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ RewriteTensorizeNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocRewriteTensorize",
Postproc::RewriteTensorize);
}
diff --git a/src/meta_schedule/postproc/verify_gpu_code.cc
b/src/meta_schedule/postproc/verify_gpu_code.cc
index f02790cb49..04a9cf2ea7 100644
--- a/src/meta_schedule/postproc/verify_gpu_code.cc
+++ b/src/meta_schedule/postproc/verify_gpu_code.cc
@@ -206,6 +206,12 @@ class VerifyGPUCodeNode : public PostprocNode {
n->target_constraints_ = this->target_constraints_;
return Postproc(n);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<VerifyGPUCodeNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.VerifyGPUCode",
VerifyGPUCodeNode, PostprocNode);
};
@@ -216,6 +222,7 @@ Postproc Postproc::VerifyGPUCode() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ VerifyGPUCodeNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocVerifyGPUCode",
Postproc::VerifyGPUCode);
}
diff --git a/src/meta_schedule/postproc/verify_vtcm_limit.cc
b/src/meta_schedule/postproc/verify_vtcm_limit.cc
index 38234ef011..f0fe8be1c1 100644
--- a/src/meta_schedule/postproc/verify_vtcm_limit.cc
+++ b/src/meta_schedule/postproc/verify_vtcm_limit.cc
@@ -59,6 +59,12 @@ class VerifyVTCMLimitNode : public PostprocNode {
ObjectPtr<VerifyVTCMLimitNode> n =
ffi::make_object<VerifyVTCMLimitNode>(*this);
return Postproc(n);
}
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<VerifyVTCMLimitNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.VerifyVTCMLimit",
VerifyVTCMLimitNode,
PostprocNode);
};
@@ -70,6 +76,7 @@ Postproc Postproc::VerifyVTCMLimit() {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ VerifyVTCMLimitNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.PostprocVerifyVTCMLimit",
Postproc::VerifyVTCMLimit);
}
diff --git a/src/meta_schedule/runner/runner.cc
b/src/meta_schedule/runner/runner.cc
index 0d620fb3b3..1b9a3ea9a9 100644
--- a/src/meta_schedule/runner/runner.cc
+++ b/src/meta_schedule/runner/runner.cc
@@ -56,6 +56,7 @@ Runner Runner::PyRunner(Runner::FRun f_run) {
/******** FFI ********/
TVM_FFI_STATIC_INIT_BLOCK() {
+ RunnerNode::RegisterReflection();
RunnerInputNode::RegisterReflection();
RunnerResultNode::RegisterReflection();
RunnerFutureNode::RegisterReflection();
diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc
b/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc
index cdf69d8f11..c58e81dc33 100644
--- a/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc
+++ b/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc
@@ -90,6 +90,12 @@ class TensorCoreStateNode : public StateNode {
bool use_async;
State Copy() const final;
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<TensorCoreStateNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.TensorCoreState",
TensorCoreStateNode,
StateNode);
};
@@ -191,6 +197,12 @@ class MultiLevelTilingTensorCoreNode : public
MultiLevelTilingNode {
std::vector<TensorCoreIntrinGroup> intrin_groups;
/*! \brief Whether to use software pipeline */
bool use_software_pipeline = false;
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<MultiLevelTilingTensorCoreNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.MultiLevelTilingTensorCore",
MultiLevelTilingTensorCoreNode,
MultiLevelTilingNode);
@@ -927,6 +939,8 @@ ScheduleRule ScheduleRule::MultiLevelTilingTensorCore(
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ MultiLevelTilingTensorCoreNode::RegisterReflection();
+ TensorCoreStateNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.ScheduleRuleMultiLevelTilingTensorCore",
ScheduleRule::MultiLevelTilingTensorCore);
}
diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc
b/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc
index a09a38230d..080e1c9c0f 100644
--- a/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc
+++ b/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc
@@ -39,6 +39,12 @@ using tir::Schedule;
class MultiLevelTilingWideVectorNode : public MultiLevelTilingNode {
public:
size_t vector_length_in_bits;
+
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<MultiLevelTilingWideVectorNode>();
+ }
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.MultiLevelTilingWideVector",
MultiLevelTilingWideVectorNode,
MultiLevelTilingNode);
@@ -129,6 +135,7 @@ ScheduleRule ScheduleRule::MultiLevelTilingWideVector(
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ MultiLevelTilingWideVectorNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.ScheduleRuleMultiLevelTilingWideVector",
ScheduleRule::MultiLevelTilingWideVector);
}
diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc
b/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc
index 7b67823ad7..4a375689e4 100644
--- a/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc
+++ b/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc
@@ -86,6 +86,11 @@ class MultiLevelTilingWithIntrinNode : public
MultiLevelTilingNode {
}
public:
+ static void RegisterReflection() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<MultiLevelTilingWithIntrinNode>();
+ }
+
/*! \brief The name of a tensor intrinsic. */
ffi::String intrin_name;
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.MultiLevelTilingWithIntrin",
@@ -108,6 +113,7 @@ ScheduleRule ScheduleRule::MultiLevelTilingWithIntrin(
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ MultiLevelTilingWithIntrinNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.ScheduleRuleMultiLevelTilingWithIntrin",
ScheduleRule::MultiLevelTilingWithIntrin);
}
diff --git a/src/meta_schedule/search_strategy/replay_func.cc
b/src/meta_schedule/search_strategy/replay_func.cc
index 498857ad96..9082c6c3a9 100644
--- a/src/meta_schedule/search_strategy/replay_func.cc
+++ b/src/meta_schedule/search_strategy/replay_func.cc
@@ -63,7 +63,8 @@ class ReplayFuncNode : public SearchStrategyNode {
std::unique_ptr<State> state_ = nullptr;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ReplayFuncNode>();
}
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.ReplayFunc",
ReplayFuncNode, SearchStrategyNode);
diff --git a/src/meta_schedule/space_generator/post_order_apply.cc
b/src/meta_schedule/space_generator/post_order_apply.cc
index 26829356e5..e3786a4d61 100644
--- a/src/meta_schedule/space_generator/post_order_apply.cc
+++ b/src/meta_schedule/space_generator/post_order_apply.cc
@@ -37,7 +37,8 @@ class PostOrderApplyNode : public SpaceGeneratorNode {
TRandState rand_state_ = -1;
static void RegisterReflection() {
- // No fields to register
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PostOrderApplyNode>();
}
void InitializeWithTuneContext(const TuneContext& context) final {
diff --git a/src/meta_schedule/space_generator/schedule_fn.cc
b/src/meta_schedule/space_generator/schedule_fn.cc
index 687abef75f..7d22635b76 100644
--- a/src/meta_schedule/space_generator/schedule_fn.cc
+++ b/src/meta_schedule/space_generator/schedule_fn.cc
@@ -33,6 +33,8 @@ class ScheduleFnNode : public SpaceGeneratorNode {
static void RegisterReflection() {
// `schedule_fn_` is not registered.
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ScheduleFnNode>();
}
void InitializeWithTuneContext(const TuneContext& context) final {
@@ -80,6 +82,7 @@ class ScheduleFnNode : public SpaceGeneratorNode {
CloneRules(this, n.get());
return SpaceGenerator(n);
}
+
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("meta_schedule.ScheduleFn",
ScheduleFnNode, SpaceGeneratorNode);
};
@@ -95,10 +98,9 @@ SpaceGenerator SpaceGenerator::ScheduleFn(
return SpaceGenerator(n);
}
-TVM_FFI_STATIC_INIT_BLOCK() { ScheduleFnNode::RegisterReflection(); }
-
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ ScheduleFnNode::RegisterReflection();
refl::GlobalDef().def("meta_schedule.SpaceGeneratorScheduleFn",
SpaceGenerator::ScheduleFn);
}
diff --git a/src/relax/ir/block_builder.cc b/src/relax/ir/block_builder.cc
index 00be02270b..09f404d29c 100644
--- a/src/relax/ir/block_builder.cc
+++ b/src/relax/ir/block_builder.cc
@@ -1055,6 +1055,7 @@ BlockBuilder BlockBuilder::Create(ffi::Optional<IRModule>
mod,
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<BlockBuilderNode>();
refl::GlobalDef()
.def("relax.BlockBuilderCreate",
[](ffi::Optional<IRModule> mod) { return BlockBuilder::Create(mod);
})
diff --git a/src/relax/ir/py_expr_functor.cc b/src/relax/ir/py_expr_functor.cc
index 73f41f185d..b7d61bfda8 100644
--- a/src/relax/ir/py_expr_functor.cc
+++ b/src/relax/ir/py_expr_functor.cc
@@ -554,6 +554,7 @@ class PyExprMutator : public ObjectRef {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyExprVisitorNode>();
refl::GlobalDef()
.def("relax.MakePyExprVisitor", PyExprVisitor::MakePyExprVisitor)
.def("relax.PyExprVisitorVisitExpr",
diff --git a/src/relax/transform/fold_constant.cc
b/src/relax/transform/fold_constant.cc
index 0892adcc1a..b714d49243 100644
--- a/src/relax/transform/fold_constant.cc
+++ b/src/relax/transform/fold_constant.cc
@@ -272,7 +272,7 @@ class ConstantFolder : public ExprMutator {
Constant constant = Downcast<Constant>(arg);
runtime::Tensor ndarray = constant->data;
ICHECK_EQ(ndarray->device.device_type, kDLCPU);
- ICHECK(ffi::IsContiguous(*ndarray.get()));
+ ICHECK(ndarray.IsContiguous());
ICHECK_EQ(ndarray->byte_offset, 0);
ICHECK_EQ(ndarray->ndim, 1);
const int64_t* data = static_cast<const int64_t*>(ndarray->data);
diff --git a/src/runtime/contrib/random/mt_random_engine.cc
b/src/runtime/contrib/random/mt_random_engine.cc
index ce9b959a53..0158a66be5 100644
--- a/src/runtime/contrib/random/mt_random_engine.cc
+++ b/src/runtime/contrib/random/mt_random_engine.cc
@@ -124,8 +124,9 @@ class RandomEngine {
} else {
runtime::Tensor local = runtime::Tensor::Empty(
std::vector<int64_t>{data->shape, data->shape + data->ndim},
data->dtype, {kDLCPU, 0});
- DLTensor* tensor = const_cast<ffi::TensorObj*>(local.operator->());
- FillData(tensor);
+
+ const DLTensor* tensor = local.GetDLTensorPtr();
+ FillData(const_cast<DLTensor*>(tensor));
runtime::Tensor::CopyFromTo(tensor, data);
}
}
@@ -136,8 +137,8 @@ class RandomEngine {
} else {
runtime::Tensor local = runtime::Tensor::Empty(
std::vector<int64_t>{data->shape, data->shape + data->ndim},
data->dtype, {kDLCPU, 0});
- DLTensor* tensor = const_cast<ffi::TensorObj*>(local.operator->());
- FillDataForMeasure(tensor);
+ const DLTensor* tensor = local.GetDLTensorPtr();
+ FillDataForMeasure(const_cast<DLTensor*>(tensor));
runtime::Tensor::CopyFromTo(tensor, data);
}
}
diff --git a/src/runtime/disco/distributed/socket_session.cc
b/src/runtime/disco/distributed/socket_session.cc
index b1845bdcfe..99c54933bf 100644
--- a/src/runtime/disco/distributed/socket_session.cc
+++ b/src/runtime/disco/distributed/socket_session.cc
@@ -307,6 +307,7 @@ Session SocketSession(int num_nodes, int
num_workers_per_node, int num_groups,
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<SocketSessionObj>();
refl::GlobalDef()
.def("runtime.disco.SocketSession", SocketSession)
.def("runtime.disco.socket_session_init_workers",
diff --git a/src/runtime/disco/process_session.cc
b/src/runtime/disco/process_session.cc
index aca1fef90c..c13cd9e60e 100644
--- a/src/runtime/disco/process_session.cc
+++ b/src/runtime/disco/process_session.cc
@@ -194,6 +194,7 @@ void WorkerProcess(int worker_id, int num_workers, int
num_group, int64_t read_f
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ProcessSessionObj>();
refl::GlobalDef()
.def("runtime.disco.SessionProcess", Session::ProcessSession)
.def("runtime.disco.WorkerProcess", WorkerProcess);
diff --git a/src/runtime/disco/session.cc b/src/runtime/disco/session.cc
index ab8505d169..2bf132f362 100644
--- a/src/runtime/disco/session.cc
+++ b/src/runtime/disco/session.cc
@@ -32,6 +32,8 @@ struct SessionObj::FFI {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<DRefObj>();
+ refl::ObjectDef<SessionObj>();
refl::GlobalDef()
.def("runtime.disco.SessionThreaded", Session::ThreadedSession)
.def_method("runtime.disco.DRefDebugGetFromRemote",
&DRefObj::DebugGetFromRemote)
diff --git a/src/runtime/disco/threaded_session.cc
b/src/runtime/disco/threaded_session.cc
index 89245000a5..029038625f 100644
--- a/src/runtime/disco/threaded_session.cc
+++ b/src/runtime/disco/threaded_session.cc
@@ -17,6 +17,7 @@
* under the License.
*/
#include <dmlc/io.h>
+#include <tvm/ffi/reflection/registry.h>
#include <tvm/runtime/base.h>
#include <tvm/runtime/disco/disco_worker.h>
#include <tvm/runtime/object.h>
@@ -193,5 +194,10 @@ Session Session::ThreadedSession(int num_workers, int
num_group) {
return Session(std::move(n));
}
+TVM_FFI_STATIC_INIT_BLOCK() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ThreadedSessionObj>();
+}
+
} // namespace runtime
} // namespace tvm
diff --git a/src/runtime/profiling.cc b/src/runtime/profiling.cc
index 6e25fb5d34..21c5681176 100644
--- a/src/runtime/profiling.cc
+++ b/src/runtime/profiling.cc
@@ -784,6 +784,9 @@ Report Report::FromJSON(ffi::String json) {
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<MetricCollectorNode>();
+ refl::ObjectDef<DeviceWrapperNode>();
+
refl::GlobalDef()
.def_method("runtime.profiling.AsTable", &ReportNode::AsTable)
.def("runtime.profiling.AsCSV", [](Report n) { return n->AsCSV(); })
diff --git a/src/runtime/rpc/rpc_session.cc b/src/runtime/rpc/rpc_session.cc
index ace9cf9b94..1fee1424ea 100644
--- a/src/runtime/rpc/rpc_session.cc
+++ b/src/runtime/rpc/rpc_session.cc
@@ -24,6 +24,7 @@
#include "rpc_session.h"
#include <tvm/ffi/function.h>
+#include <tvm/ffi/reflection/registry.h>
#include <tvm/runtime/device_api.h>
#include <array>
@@ -127,5 +128,7 @@ void
RPCSession::InsertToSessionTable(std::shared_ptr<RPCSession> sess) {
sess->table_index_ = RPCSessTable::Global()->Insert(sess);
}
+TVM_FFI_STATIC_INIT_BLOCK() {
tvm::ffi::reflection::ObjectDef<RPCObjectRefObj>(); }
+
} // namespace runtime
} // namespace tvm
diff --git a/src/runtime/vm/builtin.cc b/src/runtime/vm/builtin.cc
index d94d5676bb..13446a158f 100644
--- a/src/runtime/vm/builtin.cc
+++ b/src/runtime/vm/builtin.cc
@@ -735,7 +735,7 @@ int TVMBackendAnyListMoveFromPackedReturn(void* anylist,
int index, TVMFFIAny* a
using namespace tvm::runtime;
TVM_FFI_SAFE_CALL_BEGIN();
auto* list = static_cast<tvm::ffi::Any*>(anylist);
- list[index] =
tvm::ffi::details::AnyUnsafe::MoveTVMFFIAnyToAny(std::move(args[ret_offset]));
+ list[index] =
tvm::ffi::details::AnyUnsafe::MoveTVMFFIAnyToAny(&args[ret_offset]);
TVM_FFI_SAFE_CALL_END();
}
} // extern "C"
diff --git a/src/tir/ir/py_functor.cc b/src/tir/ir/py_functor.cc
index d2cf81eae7..61bdfb15e7 100644
--- a/src/tir/ir/py_functor.cc
+++ b/src/tir/ir/py_functor.cc
@@ -215,7 +215,8 @@ class PyStmtExprVisitorNode : public Object, public
StmtExprVisitor {
}
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyStmtExprVisitorNode>();
}
static constexpr const bool _type_mutable = true;
@@ -581,7 +582,8 @@ class PyStmtExprMutatorNode : public Object, public
StmtExprMutator {
}
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<PyStmtExprMutatorNode>();
}
static constexpr const bool _type_mutable = true;
diff --git a/src/tir/schedule/concrete_schedule.h
b/src/tir/schedule/concrete_schedule.h
index f19fb3143e..b6f87a3aae 100644
--- a/src/tir/schedule/concrete_schedule.h
+++ b/src/tir/schedule/concrete_schedule.h
@@ -51,7 +51,8 @@ class ConcreteScheduleNode : public ScheduleNode {
public:
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ConcreteScheduleNode>();
}
virtual ~ConcreteScheduleNode() = default;
diff --git a/src/tir/schedule/schedule.cc b/src/tir/schedule/schedule.cc
index 9648154289..845bbb5cc2 100644
--- a/src/tir/schedule/schedule.cc
+++ b/src/tir/schedule/schedule.cc
@@ -23,6 +23,8 @@ namespace tvm {
namespace tir {
TVM_FFI_STATIC_INIT_BLOCK() {
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<ScheduleNode>();
BlockRVNode::RegisterReflection();
LoopRVNode::RegisterReflection();
}
diff --git a/src/tir/schedule/traced_schedule.h
b/src/tir/schedule/traced_schedule.h
index cf9e53a3a7..0b91dc2833 100644
--- a/src/tir/schedule/traced_schedule.h
+++ b/src/tir/schedule/traced_schedule.h
@@ -32,7 +32,8 @@ class TracedScheduleNode : public ConcreteScheduleNode {
public:
static void RegisterReflection() {
- // No fields to register as they are not visited
+ namespace refl = tvm::ffi::reflection;
+ refl::ObjectDef<TracedScheduleNode>();
}
~TracedScheduleNode() = default;
diff --git a/tests/python/tir-base/test_tir_ptx_cp_async.py
b/tests/python/tir-base/test_tir_ptx_cp_async.py
index f3255bd257..8b15e385d2 100644
--- a/tests/python/tir-base/test_tir_ptx_cp_async.py
+++ b/tests/python/tir-base/test_tir_ptx_cp_async.py
@@ -18,6 +18,7 @@ import tvm
from tvm.script import tir as T
import numpy as np
import tvm.testing
+import pytest
@T.prim_func
@@ -94,7 +95,7 @@ def ptx_cp_async_barrier(
B[tx, i] = A_shared[tx, i]
[email protected]_cuda_compute_version(8)
[email protected](reason="temp skip test due to cuda env update")
def test_ptx_cp_async_barrier():
f = ptx_cp_async_barrier
diff --git
a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
index aa4f5138a1..bcec1d4843 100644
--- a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
+++ b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py
@@ -214,6 +214,7 @@ def ptx_global_to_shared_copy_fp32x1_barrier(
B[tx, i] = A_shared[tx, i]
[email protected](reason="temp skip test due to cuda env update")
@tvm.testing.requires_cuda
def test_inject_async_copy_barrier():
dtype = "float32"
diff --git a/web/emcc/wasm_runtime.cc b/web/emcc/wasm_runtime.cc
index 31547269e1..d787c295e0 100644
--- a/web/emcc/wasm_runtime.cc
+++ b/web/emcc/wasm_runtime.cc
@@ -55,10 +55,10 @@
#include "3rdparty/tvm-ffi/src/ffi/extra/library_module.cc"
#include "3rdparty/tvm-ffi/src/ffi/extra/library_module_system_lib.cc"
#include "3rdparty/tvm-ffi/src/ffi/extra/module.cc"
-#include "3rdparty/tvm-ffi/src/ffi/extra/testing.cc"
#include "3rdparty/tvm-ffi/src/ffi/function.cc"
#include "3rdparty/tvm-ffi/src/ffi/object.cc"
#include "3rdparty/tvm-ffi/src/ffi/tensor.cc"
+#include "3rdparty/tvm-ffi/src/ffi/testing/testing.cc"
#include "src/runtime/memory/memory_manager.cc"
#include "src/runtime/nvtx.cc"
#include "src/runtime/vm/attn_backend.cc"
