You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by gi...@apache.org on 2026/09/15 15:15:07 UTC
(tvm-site) branch asf-site updated: deploying ffi docs
This is an automated email from the ASF dual-hosted git repository.
github-actions[bot] pushed a commit to branch asf-site
in repository https://gitbox.apache.org/repos/asf/tvm-site.git
The following commit(s) were added to refs/heads/asf-site by this push:
new e350a141a23 deploying ffi docs
e350a141a23 is described below
commit e350a141a23cdcb3f4aa0f3a4becb758827648f3
Author: tvm-bot <95...@users.noreply.github.com>
AuthorDate: Tue Sep 15 15:14:58 2026 +0000
deploying ffi docs
---
ffi/_modules/index.html | 329 +--
ffi/_modules/tvm_ffi/_convert.html | 369 +--
ffi/_modules/tvm_ffi/_dtype.html | 503 ++--
ffi/_modules/tvm_ffi/_tensor.html | 357 +--
ffi/_modules/tvm_ffi/access_path.html | 395 +--
ffi/_modules/tvm_ffi/container.html | 529 ++--
ffi/_modules/tvm_ffi/cpp/extension.html | 1017 ++++---
ffi/_modules/tvm_ffi/cpp/nvrtc.html | 387 +--
ffi/_modules/tvm_ffi/error.html | 397 ++-
ffi/_modules/tvm_ffi/libinfo.html | 505 ++--
ffi/_modules/tvm_ffi/module.html | 463 ++--
ffi/_modules/tvm_ffi/registry.html | 499 ++--
ffi/_modules/tvm_ffi/serialization.html | 347 +--
ffi/_modules/tvm_ffi/stream.html | 377 +--
ffi/_modules/tvm_ffi/structural.html | 952 ++++---
ffi/_sources/concepts/structural_eq_hash.rst | 482 ++--
ffi/_sources/guides/python_lang_guide.md | 21 +
ffi/_sources/guides/rust_lang_guide.md | 438 +++
ffi/_sources/packaging/stubgen.rst | 16 +
.../cpp/generated/classtvm_1_1ffi_1_1BigInt.rst | 17 +
.../cpp/generated/classtvm_1_1ffi_1_1Object.rst | 2 +
.../cpp/generated/classtvm_1_1ffi_1_1ObjectRef.rst | 3 +
.../classtvm_1_1ffi_1_1StructuralMapDynEngine.rst | 26 +
.../classtvm_1_1ffi_1_1StructuralMapEngine.rst | 26 +
.../classtvm_1_1ffi_1_1StructuralMapEngineBase.rst | 26 +
.../classtvm_1_1ffi_1_1StructuralMutateEngine.rst | 26 +
.../classtvm_1_1ffi_1_1StructuralMutator.rst | 26 +
.../classtvm_1_1ffi_1_1StructuralMutatorObj.rst | 32 +
.../classtvm_1_1ffi_1_1StructuralVisitEngine.rst | 26 +
.../classtvm_1_1ffi_1_1StructuralVisitorObj.rst | 6 -
.../classtvm_1_1ffi_1_1StructuralWalkEngine.rst | 26 +
.../generated/classtvm_1_1ffi_1_1UnchangedOr.rst | 27 +
...__api_8h_1ae5d3f69eb7caba8077fbc18dbeab494b.rst | 14 +
...ected_8h_1a12814e63007046dc2eeeea807301b18e.rst | 14 +
...ected_8h_1a1bdcc4b0013750cda4330687cef726ba.rst | 14 +
...ected_8h_1a222bb766cd3ed50251e570fba8eb6707.rst | 14 +
...ected_8h_1a47048735aaf27fd8eb9063006b8d3ba1.rst | 14 +
...ected_8h_1a4fb262f9e930a0be433b549a70f6f2fa.rst | 14 +
...ected_8h_1a5cba5c185a18ce67f3f97474d6b1a033.rst | 14 +
...ected_8h_1a7ee294a645a49ccf4c8247dc4027d9f6.rst | 14 +
...ected_8h_1a8415b9392b980d719d9bff1a796cc2bd.rst | 14 +
...ected_8h_1a8b6abe1ba48a00e2a3a80ccb914e8489.rst | 14 +
...ected_8h_1a90c343b32e301472b7e576ba2fb17bd5.rst | 14 +
...ected_8h_1a97ad19e683ea80f3c9f5ce36629ff9a2.rst | 14 +
...ected_8h_1aa3db947e07dd51d2e5eeef8cc89e6dc7.rst | 14 +
...ected_8h_1aac4c9af29b71ebdf3cb0918255cf08e1.rst | 14 +
...ected_8h_1af4272b23a03c44b439162406f9162784.rst | 14 +
...ected_8h_1af7ef9130393bd0034392c72491f31fc1.rst | 14 +
...utate_8h_1a5338fd615444f86fceae424d26a5baa7.rst | 14 +
...utate_8h_1a58cd32be70a3aa6902f09556fec4cf24.rst | 14 +
...visit_8h_1ac74b688a86e26ee7dc5518e277d58393.rst | 14 +
.../reference/cpp/generated/dir_tvm_ffi.rst | 1 +
.../reference/cpp/generated/dir_tvm_ffi_extra.rst | 1 +
...utate_8h_1a45478eb8e9652bf4c18ec5d9bfc195a2.rst | 14 +
.../reference/cpp/generated/file_tvm_ffi_any.h.rst | 2 +
.../cpp/generated/file_tvm_ffi_big_int.h.rst | 87 +
.../cpp/generated/file_tvm_ffi_c_api.h.rst | 8 +
.../cpp/generated/file_tvm_ffi_cast.h.rst | 2 +
.../generated/file_tvm_ffi_container_array.h.rst | 2 +
.../generated/file_tvm_ffi_container_tensor.h.rst | 8 +-
.../generated/file_tvm_ffi_container_tuple.h.rst | 2 +
.../generated/file_tvm_ffi_container_variant.h.rst | 2 +
.../cpp/generated/file_tvm_ffi_error.h.rst | 2 +
.../cpp/generated/file_tvm_ffi_expected.h.rst | 41 +
.../cpp/generated/file_tvm_ffi_extra_base64.h.rst | 4 +-
.../file_tvm_ffi_extra_structural_mutate.h.rst | 157 ++
.../file_tvm_ffi_extra_structural_visit.h.rst | 21 +
.../file_tvm_ffi_extra_visit_error_context.h.rst | 2 +
.../cpp/generated/file_tvm_ffi_function.h.rst | 2 +
.../cpp/generated/file_tvm_ffi_memory.h.rst | 2 +
.../cpp/generated/file_tvm_ffi_optional.h.rst | 4 +-
.../file_tvm_ffi_reflection_accessor.h.rst | 8 +
.../cpp/generated/file_tvm_ffi_type_traits.h.rst | 2 +
...__api_8h_1a7306ae6d3ae66f958ce6a777a8b515fe.rst | 14 +
...__api_8h_1a7d1e6bea055fa1e4e9422342f8f096ac.rst | 14 +
...utate_8h_1a1f946d2f6dd55c80fc81d35e6fd841a2.rst | 14 +
...utate_8h_1a781fd0c5b30ec9a988812b47580a30cf.rst | 14 +
...utate_8h_1a84c3f7a1a45ea13f2ba233c4ad6cbf31.rst | 14 +
...utate_8h_1aefadde05e52e44ac89aea2a1db8ec75c.rst | 14 +
...visit_8h_1ad632463d20e2bee04ca4fc38ec95df23.rst | 14 +
...visit_8h_1aece5d69ac3a9b3ee2c0c2bc2210c57b3.rst | 14 +
.../reference/cpp/generated/namespace_tvm__ffi.rst | 50 +-
.../namespace_tvm__ffi__reflection__type_attr.rst | 6 +
.../program_listing_file_tvm_ffi_any.h.rst | 61 +-
.../program_listing_file_tvm_ffi_big_int.h.rst | 1295 +++++++++
.../program_listing_file_tvm_ffi_c_api.h.rst | 25 +-
...gram_listing_file_tvm_ffi_container_array.h.rst | 10 -
...ogram_listing_file_tvm_ffi_container_dict.h.rst | 4 +
...rogram_listing_file_tvm_ffi_container_map.h.rst | 4 +
...m_listing_file_tvm_ffi_container_seq_base.h.rst | 16 +
...am_listing_file_tvm_ffi_container_variant.h.rst | 13 +-
.../program_listing_file_tvm_ffi_expected.h.rst | 248 +-
...ting_file_tvm_ffi_extra_structural_mutate.h.rst | 1318 ++++++++++
...sting_file_tvm_ffi_extra_structural_visit.h.rst | 453 ++--
...ng_file_tvm_ffi_extra_visit_error_context.h.rst | 15 +
.../program_listing_file_tvm_ffi_memory.h.rst | 21 +-
.../program_listing_file_tvm_ffi_object.h.rst | 80 +-
.../program_listing_file_tvm_ffi_optional.h.rst | 353 ++-
..._listing_file_tvm_ffi_reflection_accessor.h.rst | 5 +
..._listing_file_tvm_ffi_reflection_registry.h.rst | 10 +-
.../program_listing_file_tvm_ffi_type_traits.h.rst | 2 +
...structtvm_1_1ffi_1_1StructuralMutatorVTable.rst | 17 +
.../generated/structtvm_1_1ffi_1_1Unchanged.rst | 17 +
...ructtvm_1_1ffi_1_1UnchangedOr_1_1UnsafeInit.rst | 23 +
...utate_8h_1a060cef1ff53c97235144336dac227d74.rst | 14 +
...utate_8h_1a074bf7a4d1b8b09e4ff3f3a189e77c2b.rst | 14 +
...utate_8h_1ae69b95653133acacec6825c62bc41191.rst | 14 +
.../reference/cpp/generated/unabridged_orphan.rst | 10 +
...pe__attr_1a09a99f5f4a3b5ec6730fafa33309461f.rst | 14 +
...pe__attr_1a0e928fa6d490779c06ed997dc8af2f35.rst | 14 +
...pe__attr_1aa78a5fd8e7af94d0540d67c278128960.rst | 14 +
.../python/generated/tvm_ffi.DefRegionKind.rst | 30 +
.../python/generated/tvm_ffi.StructuralMutator.rst | 39 +
.../python/generated/tvm_ffi.StructuralVisitor.rst | 1 +
.../python/generated/tvm_ffi.structural_map.rst | 6 +
.../python/generated/tvm_ffi.structural_visit.rst | 6 +
ffi/_sources/reference/python/index.rst | 4 +
ffi/_static/documentation_options.js | 2 +-
ffi/_static/downloads/d3.min.js | 2 -
ffi/_static/locales/hu/LC_MESSAGES/booktheme.mo | Bin 0 -> 1422 bytes
ffi/_static/locales/hu/LC_MESSAGES/booktheme.po | 75 +
ffi/_static/scripts/bootstrap.js | 2 +-
ffi/_static/scripts/bootstrap.js.map | 2 +-
ffi/_static/scripts/fontawesome.js | 2 +-
ffi/_static/scripts/fontawesome.js.LICENSE.txt | 4 +-
ffi/_static/scripts/fontawesome.js.map | 2 +-
ffi/_static/scripts/pydata-sphinx-theme.js | 2 +-
ffi/_static/scripts/pydata-sphinx-theme.js.map | 2 +-
ffi/_static/styles/pydata-sphinx-theme.css | 27 +-
ffi/_static/styles/pydata-sphinx-theme.css.map | 2 +-
ffi/_static/styles/sphinx-book-theme.css | 7 +-
ffi/_static/styles/sphinx-book-theme.css.map | 2 +-
.../vendor/fontawesome/webfonts/fa-brands-400.ttf | Bin 209128 -> 0 bytes
.../fontawesome/webfonts/fa-brands-400.woff2 | Bin 117852 -> 110088 bytes
.../vendor/fontawesome/webfonts/fa-regular-400.ttf | Bin 67860 -> 0 bytes
.../fontawesome/webfonts/fa-regular-400.woff2 | Bin 25392 -> 18924 bytes
.../vendor/fontawesome/webfonts/fa-solid-900.ttf | Bin 420332 -> 0 bytes
.../vendor/fontawesome/webfonts/fa-solid-900.woff2 | Bin 156400 -> 114740 bytes
ffi/_static/webpack-macros.html | 14 +-
ffi/concepts/abi_overview.html | 144 +-
ffi/concepts/any.html | 146 +-
ffi/concepts/containers.html | 132 +-
ffi/concepts/exception_handling.html | 152 +-
ffi/concepts/func_module.html | 152 +-
ffi/concepts/object_and_class.html | 166 +-
ffi/concepts/structural_eq_hash.html | 955 +++----
ffi/concepts/tensor.html | 130 +-
ffi/dev/ci_cd.html | 124 +-
ffi/dev/doc_build.html | 114 +-
ffi/dev/release_process.html | 248 +-
ffi/dev/source_build.html | 114 +-
ffi/genindex.html | 629 +++--
ffi/get_started/quickstart.html | 168 +-
ffi/get_started/stable_c_abi.html | 152 +-
ffi/guides/compiler_integration.html | 122 +-
ffi/guides/cpp_lang_guide.html | 188 +-
ffi/guides/cubin_launcher.html | 232 +-
ffi/guides/dataclass_reflection.html | 204 +-
ffi/guides/export_func_cls.html | 148 +-
ffi/guides/kernel_library_guide.html | 154 +-
ffi/guides/python_lang_guide.html | 202 +-
ffi/guides/rust_lang_guide.html | 559 +++-
ffi/index.html | 118 +-
ffi/objects.inv | Bin 195497 -> 214303 bytes
ffi/packaging/cpp_tooling.html | 130 +-
ffi/packaging/python_packaging.html | 148 +-
ffi/packaging/stubgen.html | 161 +-
ffi/py-modindex.html | 331 +--
.../cpp/generated/classtvm_1_1ffi_1_1Any.html | 129 +-
.../cpp/generated/classtvm_1_1ffi_1_1AnyView.html | 132 +-
.../cpp/generated/classtvm_1_1ffi_1_1Arc.html | 129 +-
.../cpp/generated/classtvm_1_1ffi_1_1Array.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1ArrayObj.html | 120 +-
...EnumObj.html => classtvm_1_1ffi_1_1BigInt.html} | 322 ++-
.../cpp/generated/classtvm_1_1ffi_1_1Bytes.html | 116 +-
.../generated/classtvm_1_1ffi_1_1CubinKernel.html | 118 +-
.../generated/classtvm_1_1ffi_1_1CubinModule.html | 120 +-
.../classtvm_1_1ffi_1_1DenseMapBaseObj.html | 124 +-
.../cpp/generated/classtvm_1_1ffi_1_1Dict.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1DictObj.html | 147 +-
.../cpp/generated/classtvm_1_1ffi_1_1Enum.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1EnumObj.html | 120 +-
.../generated/classtvm_1_1ffi_1_1EnumState.html | 120 +-
.../generated/classtvm_1_1ffi_1_1EnumStateObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1Error.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1ErrorObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1Expected.html | 197 +-
.../classtvm_1_1ffi_1_1Expected_3_01void_01_4.html | 123 +-
.../cpp/generated/classtvm_1_1ffi_1_1Function.html | 126 +-
.../generated/classtvm_1_1ffi_1_1FunctionObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1IntEnum.html | 120 +-
.../generated/classtvm_1_1ffi_1_1IntEnumObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1List.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1ListObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1Map.html | 120 +-
.../generated/classtvm_1_1ffi_1_1MapBaseObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1MapObj.html | 147 +-
.../cpp/generated/classtvm_1_1ffi_1_1Module.html | 122 +-
.../generated/classtvm_1_1ffi_1_1ModuleObj.html | 126 +-
.../cpp/generated/classtvm_1_1ffi_1_1Object.html | 141 +-
.../generated/classtvm_1_1ffi_1_1ObjectPtr.html | 124 +-
.../generated/classtvm_1_1ffi_1_1ObjectRef.html | 132 +-
.../cpp/generated/classtvm_1_1ffi_1_1Optional.html | 118 +-
.../generated/classtvm_1_1ffi_1_1PackedArgs.html | 116 +-
.../generated/classtvm_1_1ffi_1_1RValueRef.html | 116 +-
.../generated/classtvm_1_1ffi_1_1SeqBaseObj.html | 159 +-
.../cpp/generated/classtvm_1_1ffi_1_1Shape.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1ShapeObj.html | 120 +-
.../generated/classtvm_1_1ffi_1_1ShapeView.html | 116 +-
.../classtvm_1_1ffi_1_1SmallMapBaseObj.html | 120 +-
.../cpp/generated/classtvm_1_1ffi_1_1StrEnum.html | 120 +-
.../generated/classtvm_1_1ffi_1_1StrEnumObj.html | 120 +-
.../generated/classtvm_1_1ffi_1_1StrictBool.html | 116 +-
.../cpp/generated/classtvm_1_1ffi_1_1String.html | 116 +-
.../classtvm_1_1ffi_1_1StructuralEqual.html | 116 +-
.../classtvm_1_1ffi_1_1StructuralHash.html | 116 +-
.../classtvm_1_1ffi_1_1StructuralKey.html | 120 +-
.../classtvm_1_1ffi_1_1StructuralKeyObj.html | 120 +-
...classtvm_1_1ffi_1_1StructuralMapDynEngine.html} | 203 +-
...=> classtvm_1_1ffi_1_1StructuralMapEngine.html} | 211 +-
...lasstvm_1_1ffi_1_1StructuralMapEngineBase.html} | 321 ++-
...classtvm_1_1ffi_1_1StructuralMutateEngine.html} | 182 +-
...l => classtvm_1_1ffi_1_1StructuralMutator.html} | 156 +-
...> classtvm_1_1ffi_1_1StructuralMutatorObj.html} | 345 ++-
... classtvm_1_1ffi_1_1StructuralVisitEngine.html} | 179 +-
.../classtvm_1_1ffi_1_1StructuralVisitor.html | 120 +-
.../classtvm_1_1ffi_1_1StructuralVisitorObj.html | 168 +-
...> classtvm_1_1ffi_1_1StructuralWalkEngine.html} | 211 +-
.../cpp/generated/classtvm_1_1ffi_1_1Tensor.html | 126 +-
.../generated/classtvm_1_1ffi_1_1TensorObj.html | 120 +-
.../generated/classtvm_1_1ffi_1_1TensorView.html | 116 +-
.../cpp/generated/classtvm_1_1ffi_1_1Tuple.html | 120 +-
.../classtvm_1_1ffi_1_1TypedFunction.html | 116 +-
...i_1_1TypedFunction_3_01R_07Args_8_8_8_08_4.html | 118 +-
.../generated/classtvm_1_1ffi_1_1UnchangedOr.html | 863 ++++++
.../generated/classtvm_1_1ffi_1_1Unexpected.html | 123 +-
.../cpp/generated/classtvm_1_1ffi_1_1Variant.html | 123 +-
.../classtvm_1_1ffi_1_1VisitErrorContext.html | 120 +-
.../classtvm_1_1ffi_1_1VisitErrorContextObj.html | 120 +-
.../classtvm_1_1ffi_1_1VisitInterrupt.html | 120 +-
.../classtvm_1_1ffi_1_1VisitInterruptObj.html | 120 +-
.../generated/classtvm_1_1ffi_1_1WalkResult.html | 120 +-
.../classtvm_1_1ffi_1_1WeakObjectPtr.html | 116 +-
...lasstvm_1_1ffi_1_1reflection_1_1AccessPath.html | 120 +-
...stvm_1_1ffi_1_1reflection_1_1AccessPathObj.html | 120 +-
...lasstvm_1_1ffi_1_1reflection_1_1AccessStep.html | 120 +-
...stvm_1_1ffi_1_1reflection_1_1AccessStepObj.html | 120 +-
...vm_1_1ffi_1_1reflection_1_1AttachFieldFlag.html | 144 +-
...tvm_1_1ffi_1_1reflection_1_1DefaultFactory.html | 120 +-
...sstvm_1_1ffi_1_1reflection_1_1DefaultValue.html | 120 +-
.../classtvm_1_1ffi_1_1reflection_1_1EnumDef.html | 120 +-
...asstvm_1_1ffi_1_1reflection_1_1FieldGetter.html | 116 +-
...asstvm_1_1ffi_1_1reflection_1_1FieldSetter.html | 116 +-
...classtvm_1_1ffi_1_1reflection_1_1GlobalDef.html | 122 +-
.../classtvm_1_1ffi_1_1reflection_1_1Metadata.html | 120 +-
...stvm_1_1ffi_1_1reflection_1_1ObjectCreator.html | 116 +-
...classtvm_1_1ffi_1_1reflection_1_1ObjectDef.html | 124 +-
..._1_1ffi_1_1reflection_1_1OverloadObjectDef.html | 128 +-
...tvm_1_1ffi_1_1reflection_1_1TypeAttrColumn.html | 116 +-
...asstvm_1_1ffi_1_1reflection_1_1TypeAttrDef.html | 122 +-
.../classtvm_1_1ffi_1_1reflection_1_1compare.html | 120 +-
...tvm_1_1ffi_1_1reflection_1_1default__value.html | 120 +-
.../classtvm_1_1ffi_1_1reflection_1_1hash.html | 120 +-
.../classtvm_1_1ffi_1_1reflection_1_1kw__only.html | 120 +-
.../classtvm_1_1ffi_1_1reflection_1_1repr.html | 120 +-
...ails_8h_1ad96d0c78609aaf2660689e5408ddc40b.html | 116 +-
..._api_8h_1a01871644eed04162005df7faa5442f92.html | 116 +-
..._api_8h_1a4241f838270253b885e1de02534bc273.html | 116 +-
..._api_8h_1a42634cb2690dce4d9af6c2a402fc903c.html | 116 +-
..._api_8h_1a717daddad98d127813ba8ff5a715e929.html | 116 +-
...api_8h_1ae5d3f69eb7caba8077fbc18dbeab494b.html} | 136 +-
...cher_8h_1a55832b50e83cf39108f1e306f031433d.html | 118 +-
...cher_8h_1a99736a44462543179cb434ebd4512ade.html | 124 +-
...cher_8h_1ae8d64fa1cc7db9d38632e32054df72fc.html | 120 +-
...base_8h_1a8b7e871d0d7206eacc60705f11827ee1.html | 116 +-
...dian_8h_1a2ebd9542b313b3b45b2892f2b3d4f26d.html | 116 +-
...dian_8h_1a912dbc7a68913612c723f877acaa8105.html | 116 +-
...rror_8h_1a07c8a4b3cbde3471d290fdbc73e14b0c.html | 116 +-
...rror_8h_1a0b440ab7013c78c61ab5dd5d8f27bd91.html | 116 +-
...rror_8h_1a16f0dbfb27520f4d044d82e32fd5bbec.html | 116 +-
...rror_8h_1a1a94c8b78e74adec05051abb9055e7b8.html | 116 +-
...rror_8h_1a26827bb8bc9e39255060f6be1d2a37fb.html | 116 +-
...rror_8h_1a27374a2cd94c6211ce2b754bd87e3fa3.html | 116 +-
...rror_8h_1a288653475dd42db1c715a11ab2972578.html | 116 +-
...rror_8h_1a2bce4014203aa455160771f1e8e698e5.html | 116 +-
...rror_8h_1a34350f477b8479a26ec5a2683100bc04.html | 116 +-
...rror_8h_1a386641a33bb14802830471a2442ddd7d.html | 116 +-
...rror_8h_1a394ec1781a34c49e29d47db5c3dfad9c.html | 116 +-
...rror_8h_1a4e04117882ed9f1469ec36f7979ace27.html | 116 +-
...rror_8h_1a64c939463f90816ffc2b70bc89ea3e43.html | 116 +-
...rror_8h_1a6c01470f76c6437edf93ac668ea6242a.html | 116 +-
...rror_8h_1a7205bbaef1c45e57f939af6deefb6449.html | 116 +-
...rror_8h_1a7d7a510c235963a0f1c7b7ab575c15fb.html | 116 +-
...rror_8h_1a8df4e0ff8a6195274c6f7c031a4f0db4.html | 118 +-
...rror_8h_1a927d4ae2bbf7a08723643993e780f183.html | 116 +-
...rror_8h_1a990e2467aea416d9b8b3a338df07cb32.html | 116 +-
...rror_8h_1aa7a5d2005aa58b2c7d9fa80c64f26e81.html | 116 +-
...rror_8h_1aaab66d9de933f4134006b0ad3d29e74a.html | 116 +-
...rror_8h_1aade00b2e0a9789e3e8bf33441d651455.html | 116 +-
...rror_8h_1aba8087204dace0473072f763f47ca943.html | 116 +-
...rror_8h_1aca415693e572f4602ca4e3a961fcb1f1.html | 116 +-
...rror_8h_1ad4674e1cb9c051f47127e5331a4bcadc.html | 116 +-
...rror_8h_1ae5be195e80d32413cbe5e2ad11e146f7.html | 116 +-
...rror_8h_1aeec11597fea71f1db8072324fe6d8b92.html | 116 +-
...rror_8h_1af1f84686ce5635b3a378567368e89e4d.html | 116 +-
...rror_8h_1af23d9224afd96bc1db250ab98d79ab45.html | 116 +-
...rror_8h_1af88efba89f134754f31e52c7024acda4.html | 116 +-
...rror_8h_1afcd5778e656fa2c20aa0b81fbcebc497.html | 116 +-
...ted_8h_1a12814e63007046dc2eeeea807301b18e.html} | 140 +-
...ted_8h_1a1bdcc4b0013750cda4330687cef726ba.html} | 155 +-
...ted_8h_1a222bb766cd3ed50251e570fba8eb6707.html} | 153 +-
...ted_8h_1a47048735aaf27fd8eb9063006b8d3ba1.html} | 172 +-
...ted_8h_1a4fb262f9e930a0be433b549a70f6f2fa.html} | 180 +-
...ted_8h_1a5cba5c185a18ce67f3f97474d6b1a033.html} | 153 +-
...ted_8h_1a7ee294a645a49ccf4c8247dc4027d9f6.html} | 140 +-
...ted_8h_1a8415b9392b980d719d9bff1a796cc2bd.html} | 140 +-
...ted_8h_1a8b6abe1ba48a00e2a3a80ccb914e8489.html} | 140 +-
...ted_8h_1a90c343b32e301472b7e576ba2fb17bd5.html} | 155 +-
...ted_8h_1a97ad19e683ea80f3c9f5ce36629ff9a2.html} | 140 +-
...ted_8h_1aa3db947e07dd51d2e5eeef8cc89e6dc7.html} | 153 +-
...ted_8h_1aac4c9af29b71ebdf3cb0918255cf08e1.html} | 145 +-
...ted_8h_1af4272b23a03c44b439162406f9162784.html} | 155 +-
...ted_8h_1af7ef9130393bd0034392c72491f31fc1.html} | 140 +-
...tion_8h_1a01777f12f8f99771331d679a0c1d2ed4.html | 120 +-
...tion_8h_1a10bd932564c8e71a3978a061fbb82360.html | 116 +-
...tion_8h_1a14328d95ccda4b1691b9874645aef4bd.html | 116 +-
...tion_8h_1a78c6d17c063fba9ca07175f150838046.html | 116 +-
...tion_8h_1ae3c016c4bf383aa205b4485edd7ca17c.html | 116 +-
...tion_8h_1af84cdf62b6c193275db7e63e3bffc8bd.html | 120 +-
...ject_8h_1a0fc0cc2d8f4318bc41a896c117e8a15c.html | 116 +-
...ject_8h_1a438ff7ee1fc503216d35e156081280c0.html | 116 +-
...ject_8h_1a4ac33631a9b51877808ec7edf010f0d2.html | 116 +-
...ject_8h_1a6ad2a76d53c8c695743f7a275c8c0523.html | 116 +-
...ject_8h_1a9d1c42071de0440f70c7bb5e0bd0f325.html | 116 +-
...ject_8h_1ac593f4065568c8f61cf38d00d585befd.html | 116 +-
...ate_8h_1a5338fd615444f86fceae424d26a5baa7.html} | 150 +-
...ate_8h_1a58cd32be70a3aa6902f09556fec4cf24.html} | 209 +-
...sit_8h_1ac74b688a86e26ee7dc5518e277d58393.html} | 180 +-
...aits_8h_1a63cd913de4d021d2e201531d5c675a88.html | 116 +-
..._api_8h_1a160bb722922e17fa17c938fcf8561732.html | 116 +-
..._api_8h_1a6dcf7c0dec497922c66282f1162ee081.html | 116 +-
..._api_8h_1a9e03c2471638d2243c3c3b79d800efe1.html | 116 +-
...text_8h_1a2c7a0445bca80cbe5572c1210ece5d9a.html | 116 +-
...text_8h_1a83081a2550b85038ac6195ee80dbc3d5.html | 118 +-
...text_8h_1a869b3127e19530e45aa7f2052dd290f7.html | 116 +-
...text_8h_1ad7a104799c7bc592fb7f7bb33b15549c.html | 116 +-
ffi/reference/cpp/generated/dir_tvm.html | 112 +-
ffi/reference/cpp/generated/dir_tvm_ffi.html | 113 +-
.../cpp/generated/dir_tvm_ffi_container.html | 112 +-
ffi/reference/cpp/generated/dir_tvm_ffi_extra.html | 113 +-
.../cpp/generated/dir_tvm_ffi_extra_cuda.html | 112 +-
.../generated/dir_tvm_ffi_extra_cuda_internal.html | 112 +-
.../cpp/generated/dir_tvm_ffi_reflection.html | 112 +-
..._api_8h_1a1925bb5d568a3f5c92a6c28934c9bcc2.html | 136 +-
..._api_8h_1a8eb3fb55f6ec3c46b6e57740af4ca58a.html | 116 +-
..._api_8h_1ad5a608f648c4f86cafdc8dc09f870bc5.html | 136 +-
..._api_8h_1ad6a441d28d4a6406418489609e975a19.html | 116 +-
..._api_8h_1aea0f30bf60e26c1322e525e246ef5be9.html | 116 +-
..._api_8h_1aff6130e5bc350339d7e71153e396cf4a.html | 142 +-
...type_8h_1a16f66b97fa44d7bd6114324c8b8e1bed.html | 116 +-
...lection_1a9c2cd4d683e7c7065942e36e628a2517.html | 116 +-
...ate_8h_1a45478eb8e9652bf4c18ec5d9bfc195a2.html} | 178 +-
...isit_8h_1ab164fdc59c01d20b5fa07b88ef5ac09a.html | 116 +-
.../cpp/generated/file_tvm_ffi_any.h.html | 113 +-
.../cpp/generated/file_tvm_ffi_base_details.h.html | 112 +-
...overload.h.html => file_tvm_ffi_big_int.h.html} | 177 +-
.../cpp/generated/file_tvm_ffi_c_api.h.html | 116 +-
.../cpp/generated/file_tvm_ffi_cast.h.html | 113 +-
.../generated/file_tvm_ffi_container_array.h.html | 113 +-
.../generated/file_tvm_ffi_container_dict.h.html | 112 +-
.../generated/file_tvm_ffi_container_list.h.html | 112 +-
.../generated/file_tvm_ffi_container_map.h.html | 112 +-
.../file_tvm_ffi_container_map_base.h.html | 112 +-
.../file_tvm_ffi_container_seq_base.h.html | 112 +-
.../generated/file_tvm_ffi_container_shape.h.html | 112 +-
.../generated/file_tvm_ffi_container_tensor.h.html | 116 +-
.../generated/file_tvm_ffi_container_tuple.h.html | 113 +-
.../file_tvm_ffi_container_variant.h.html | 113 +-
.../cpp/generated/file_tvm_ffi_device.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_dtype.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_endian.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_enum.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_error.h.html | 113 +-
.../cpp/generated/file_tvm_ffi_expected.h.html | 138 +-
.../cpp/generated/file_tvm_ffi_extra_base.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_extra_base64.h.html | 114 +-
.../generated/file_tvm_ffi_extra_c_env_api.h.html | 112 +-
.../generated/file_tvm_ffi_extra_cuda_base.h.html | 112 +-
.../file_tvm_ffi_extra_cuda_cubin_launcher.h.html | 112 +-
.../file_tvm_ffi_extra_cuda_device_guard.h.html | 112 +-
..._tvm_ffi_extra_cuda_internal_unified_api.h.html | 112 +-
.../generated/file_tvm_ffi_extra_dataclass.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_extra_dtype.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_extra_json.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_extra_module.h.html | 112 +-
.../file_tvm_ffi_extra_serialization.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_extra_stl.h.html | 112 +-
.../file_tvm_ffi_extra_structural_equal.h.html | 112 +-
.../file_tvm_ffi_extra_structural_hash.h.html | 112 +-
.../file_tvm_ffi_extra_structural_key.h.html | 112 +-
...=> file_tvm_ffi_extra_structural_mutate.h.html} | 180 +-
.../file_tvm_ffi_extra_structural_visit.h.html | 144 +-
.../file_tvm_ffi_extra_visit_error_context.h.html | 113 +-
.../cpp/generated/file_tvm_ffi_function.h.html | 113 +-
.../cpp/generated/file_tvm_ffi_memory.h.html | 113 +-
.../cpp/generated/file_tvm_ffi_object.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_optional.h.html | 115 +-
.../file_tvm_ffi_reflection_access_path.h.html | 112 +-
.../file_tvm_ffi_reflection_accessor.h.html | 116 +-
.../file_tvm_ffi_reflection_creator.h.html | 112 +-
.../file_tvm_ffi_reflection_enum_def.h.html | 112 +-
.../file_tvm_ffi_reflection_overload.h.html | 112 +-
.../file_tvm_ffi_reflection_registry.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_rvalue_ref.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_string.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_tvm_ffi.h.html | 112 +-
.../cpp/generated/file_tvm_ffi_type_traits.h.html | 113 +-
...tion_8h_1a111e2ed7444eb220e19c7624d0932d89.html | 116 +-
...rray_8h_1a3b56317c7cfb2760cf6c09c303840967.html | 116 +-
...rray_8h_1af5983fd2a31e79635307875b3a216472.html | 116 +-
...se64_8h_1a279b4ba2f4c7ea16291ef24b73e43963.html | 116 +-
...se64_8h_1aa5b2a4531449585b9f079cb170a79f0b.html | 116 +-
...se64_8h_1acb849cdb1d9b9620be10bbcda0257afa.html | 116 +-
...se64_8h_1ae1d4736f7dc56b224acd303ba7bf7a45.html | 116 +-
..._api_8h_1a0d3dcd6ceff356614a6924201c02cf97.html | 116 +-
..._api_8h_1a0ee6d2da9a72749561838bf503f62f93.html | 116 +-
..._api_8h_1a1162f2cbaf1a7628a03ed9488cff858b.html | 116 +-
..._api_8h_1a145f2694254410fd4814541e2cde99c1.html | 116 +-
..._api_8h_1a147c7c5616be4acf5a71623dbb7dd680.html | 116 +-
..._api_8h_1a177d080127ef8f3e338beb7dba4c04fd.html | 116 +-
..._api_8h_1a1a617dd18aa7dd47aa51a984611b67a1.html | 116 +-
..._api_8h_1a1a99289a314c6c3de865721bd3efb525.html | 116 +-
..._api_8h_1a1ac76d88e0d79843403ce2e0239d7fe2.html | 116 +-
..._api_8h_1a1dd5584d94f6e00a18f4cfa003dfb93a.html | 116 +-
..._api_8h_1a268d0b5dc89146020c09f77fa0eb8f67.html | 116 +-
..._api_8h_1a31a317a32bfa641180eed1bfb3d5cf01.html | 116 +-
..._api_8h_1a3253bbf35e2ad1361b2ce1de13d2c9b2.html | 116 +-
..._api_8h_1a3479c144d165e717d1464b7b4df84897.html | 116 +-
..._api_8h_1a3bf47f6e1229627902c80217d674ed29.html | 116 +-
..._api_8h_1a3f859b60d99c85f43fce8226c69f3c5e.html | 116 +-
..._api_8h_1a402f64dec7a665f281e55cf946a7481b.html | 116 +-
..._api_8h_1a4152d18590b7967eca4a528cbe33a84e.html | 116 +-
..._api_8h_1a44604b1d0db208ed82cbc7cf0fd0f47c.html | 116 +-
..._api_8h_1a4ce84fc667040bba425664c78b5f2c69.html | 116 +-
..._api_8h_1a4e8b3caff4da72205ec1ac9a811748e5.html | 116 +-
..._api_8h_1a53c5ec2971806966ed96a6add32b5238.html | 116 +-
..._api_8h_1a56ce10d1e6fc2d34a1a496bdace4af33.html | 116 +-
..._api_8h_1a5a7e3fe8a94e066bc1bf897fbf736651.html | 116 +-
..._api_8h_1a63149ed43a6ecffa2146f92914c0bc7b.html | 116 +-
..._api_8h_1a6415c088928d62af8749f3dc57dcf23b.html | 116 +-
..._api_8h_1a724b0d386cf3d50c7d141ca37be19092.html | 116 +-
...api_8h_1a7306ae6d3ae66f958ce6a777a8b515fe.html} | 153 +-
..._api_8h_1a74c26efdd74840abc3e1581c73a38bfa.html | 116 +-
..._api_8h_1a762fc5c881dc5f66066eb45b65a5781f.html | 116 +-
...api_8h_1a7d1e6bea055fa1e4e9422342f8f096ac.html} | 142 +-
..._api_8h_1aa26fa16c024c6f2fa15e46aa1333fdb9.html | 116 +-
..._api_8h_1aa2a81275b69ba0f61ce556df85517ef9.html | 116 +-
..._api_8h_1aa7d102c19d2ac12f730935e6fb210bf1.html | 116 +-
..._api_8h_1ab15b1e535d645969fde869a5e515a885.html | 116 +-
..._api_8h_1ab31a093c3633c793dc2b4ffed0825c31.html | 116 +-
..._api_8h_1ab691fd5bc4442c631c1f70985a962f50.html | 116 +-
..._api_8h_1aba50ba2cee84b1ef7a718211ab349031.html | 116 +-
..._api_8h_1ac0e2d59d3a8ba7f2b32eac47b1c6fc5d.html | 116 +-
..._api_8h_1ac5a4f48eacde3c69ad3154b6aa39fa3c.html | 116 +-
..._api_8h_1ac64df2793c84fc8ba5d24968d60bcada.html | 116 +-
..._api_8h_1ac66c1d142240be0cdf65c8ba38770eb6.html | 116 +-
..._api_8h_1adaf9a5ab3d30bb3eca58b430f2592fc0.html | 116 +-
..._api_8h_1adec9f8a8f0683caeb9dd194d9a3c62d5.html | 116 +-
..._api_8h_1aedd89f80f73d54dd5904d943dd0690c6.html | 116 +-
..._api_8h_1aeee9dab8873b5b2f327f4dfe50f802ab.html | 116 +-
..._api_8h_1af108ca4da596c634b4cb9e82c8bec1ce.html | 116 +-
..._api_8h_1af1b8a29eb3f9405982c344f5e572dd65.html | 116 +-
..._api_8h_1afee5a803e22f9baf5ee2ff1f0ace6429.html | 123 +-
..._api_8h_1a0e01ad3ed44c0240f1b76dbed7014234.html | 116 +-
..._api_8h_1a36c15024744584acceaf0233947f9505.html | 116 +-
..._api_8h_1a683191d48d3fdebc98b68f3b717baf6a.html | 116 +-
..._api_8h_1a7451ab18593d70d30a6214d622e200d2.html | 116 +-
..._api_8h_1a8ce2ba89ca1d010a084b4c36ed6b2965.html | 116 +-
..._api_8h_1a9b142b59ab84bc06e4a0f66a3510e8f7.html | 116 +-
..._api_8h_1aaf8960c33546788950b8e35ac18efe39.html | 116 +-
..._api_8h_1abcbbccfc02c57418e757033afce42ece.html | 116 +-
..._api_8h_1ac8f7894c4f205756b7aee01108ab8e84.html | 116 +-
..._api_8h_1ae24b47945440d1c0aca13af13bdaf916.html | 116 +-
...cast_8h_1a3c5fed9bb23b4804b61037c8ac4874f1.html | 116 +-
...cast_8h_1a4dba749dd934d8a7891ad126a8fdcc1f.html | 116 +-
...ator_8h_1a9f39a70dd0c30ef28f11b745a6753222.html | 116 +-
...ator_8h_1ad292023f85eb05de65c40b53a8126b0c.html | 116 +-
...lass_8h_1a0f53d4334768d5c1c11ace3013a40c48.html | 116 +-
...lass_8h_1a2ca288baccb3a4e2305a4f0e20e43b54.html | 116 +-
...lass_8h_1a35400e2a96bcc305d505add70359b77f.html | 116 +-
...lass_8h_1a5a1645610efa2f94feb1b0653839eb7d.html | 116 +-
...lass_8h_1aa2c22b4f25336e2bc4326d08c6654080.html | 116 +-
...lass_8h_1ac0c56b2059600296bc93a03d340d3637.html | 116 +-
...lass_8h_1ae661fadd13fbc7aef6e15cddba7607e8.html | 116 +-
...lass_8h_1af7642a00508a88b93e7103a109529a79.html | 116 +-
...type_8h_1a16f5ec0ba28cdd73f61247f4a662115a.html | 116 +-
...type_8h_1a74a70941a707682598886f8475f86848.html | 116 +-
...type_8h_1a89696eef7917af8ff34285cb6fae4c63.html | 116 +-
...type_8h_1a8a71c64237efaa13a8d40dfacd6c3165.html | 116 +-
...type_8h_1a94787950883e5a7e97b764ddd342afc3.html | 116 +-
...type_8h_1a95190c80d6dfa260c156b00ce5336c83.html | 116 +-
...dian_8h_1ab4e0ed8000c632cd2c7cd089c0a0d1af.html | 116 +-
...rror_8h_1af4cf9e94df0f183e9c5c0a3ccb05fda5.html | 116 +-
...json_8h_1a0a0dc3202335cbaa265618821764fa2d.html | 116 +-
...json_8h_1a86217bda8bfebd85f8265878bc8a59b8.html | 116 +-
..._map_8h_1a611844db41026b7a95846dd316b435af.html | 116 +-
...mory_8h_1a6001462e18d71614c0f11a4c2e064f7e.html | 116 +-
...mory_8h_1a801e8ecc2e3adbc1dac52c714e7cef11.html | 116 +-
...mory_8h_1a8892f2a3cad8dbc51930c464f2ab863b.html | 116 +-
...lection_1a32d91001c82e10566c69e237ab0de50a.html | 116 +-
...lection_1a441e2af5c0ecf1cfa2860cc2fb672dc6.html | 116 +-
...lection_1a4d5446d7f6ef60852e177084e3d57743.html | 116 +-
...lection_1a6b1f837cf822ded72a31f343ce987f4c.html | 116 +-
...lection_1a735285a70c7c4308816450a5b27f03f2.html | 116 +-
...lection_1a80c4984a43d14113b20a67c966833709.html | 116 +-
...lection_1a9321975ab375d35fcf6beb55f1d7d52a.html | 116 +-
...lection_1aba4d7fd771bea88a32803dda191b7044.html | 116 +-
...lection_1ac902a725008e348254ee7a2d07c8a95a.html | 116 +-
...lection_1acc0ded3ead1cc6ffc836005d3bb2cc9c.html | 116 +-
...tion_8h_1a24d1febeb5c1026e7fd66dee9dd9915f.html | 122 +-
...tion_8h_1a56baa5d1ca415dbdebf19ac1f54bd5b7.html | 116 +-
...ring_8h_1a462b67a1ca03cb444893ec950bf5e916.html | 116 +-
...ring_8h_1a88e8350835739df34d48ab4e9d776535.html | 116 +-
...ring_8h_1aa190eed2516ff9a633009cc62f660502.html | 116 +-
...ring_8h_1ad045700080c053c50d68cde298b2b4b2.html | 116 +-
...ate_8h_1a1f946d2f6dd55c80fc81d35e6fd841a2.html} | 151 +-
...ate_8h_1a781fd0c5b30ec9a988812b47580a30cf.html} | 160 +-
...ate_8h_1a84c3f7a1a45ea13f2ba233c4ad6cbf31.html} | 161 +-
...ate_8h_1aefadde05e52e44ac89aea2a1db8ec75c.html} | 167 +-
...isit_8h_1a22c788e5e785d50c51b588526601f9cb.html | 118 +-
...isit_8h_1ab94d5197b2b27e397f0bcdef5a63048b.html | 155 +-
...sit_8h_1ad632463d20e2bee04ca4fc38ec95df23.html} | 155 +-
...sit_8h_1aece5d69ac3a9b3ee2c0c2bc2210c57b3.html} | 162 +-
...nsor_8h_1a011d0a2dcbe5702156124d7a7223b3dd.html | 116 +-
...nsor_8h_1a226723176642cde5996498e7da6a6368.html | 116 +-
...nsor_8h_1a65b3a61c540c72d9410150a3a7a76dde.html | 116 +-
...nsor_8h_1a71e4a492b5a03a19a95087771a7e9ed1.html | 116 +-
...nsor_8h_1a93094b6b3f20c2cd5aa1c1c2522bc3e4.html | 116 +-
...nsor_8h_1aa79cf4e1e3873d8947050b3cee41153a.html | 116 +-
...nsor_8h_1aab70a15ca1dde710bc3052c5ce1f6fbd.html | 116 +-
...aits_8h_1a0d6e075f219808c4ff7adebd99210a7b.html | 116 +-
..._api_8h_1a08e79db6318f80253199c3eb6c3a61d9.html | 116 +-
..._api_8h_1a1d6ae21c397f551f1f2f12d5d7cdb88a.html | 116 +-
..._api_8h_1a3016d84c6ad28cc7213440c0e2513737.html | 116 +-
..._api_8h_1a309be0e2b6f41e8f602d16931cc4ca38.html | 116 +-
..._api_8h_1a4d4f73e984897c9a4da2d0dccc11e1a1.html | 116 +-
..._api_8h_1a8d6ac296b5142517449c57c387ec8c21.html | 116 +-
..._api_8h_1a8dfc6904c4424309f4a490daaec7ed3e.html | 116 +-
..._api_8h_1a9d158620f8ae2da73b9d5d53eb6151de.html | 116 +-
..._api_8h_1a9f15e6a91d49c1a9d373b9f031edfe69.html | 116 +-
..._api_8h_1aaf09051ae19acf959b1521f9495f1beb.html | 116 +-
..._api_8h_1aaf633b14f4acc686a2f0dc23f6b3367e.html | 116 +-
..._api_8h_1afe299c48605bbb76772042acae7dc6de.html | 116 +-
ffi/reference/cpp/generated/index.html | 803 +++++-
ffi/reference/cpp/generated/namespace_tvm.html | 112 +-
.../cpp/generated/namespace_tvm__ffi.html | 137 +-
.../generated/namespace_tvm__ffi__cuda_api.html | 112 +-
.../cpp/generated/namespace_tvm__ffi__json.html | 112 +-
.../generated/namespace_tvm__ffi__reflection.html | 112 +-
.../namespace_tvm__ffi__reflection__type_attr.html | 115 +-
.../cpp/generated/namespace_tvm__ffi__symbol.html | 112 +-
ffi/reference/cpp/generated/namespace_tvm_ffi.html | 104 +-
ffi/reference/cpp/generated/page_deprecated.html | 104 +-
.../program_listing_file_tvm_ffi_any.h.html | 243 +-
...rogram_listing_file_tvm_ffi_base_details.h.html | 112 +-
.../program_listing_file_tvm_ffi_big_int.h.html | 1848 +++++++++++++
.../program_listing_file_tvm_ffi_c_api.h.html | 139 +-
.../program_listing_file_tvm_ffi_cast.h.html | 116 +-
...ram_listing_file_tvm_ffi_container_array.h.html | 160 +-
...gram_listing_file_tvm_ffi_container_dict.h.html | 122 +-
...gram_listing_file_tvm_ffi_container_list.h.html | 142 +-
...ogram_listing_file_tvm_ffi_container_map.h.html | 120 +-
..._listing_file_tvm_ffi_container_map_base.h.html | 156 +-
..._listing_file_tvm_ffi_container_seq_base.h.html | 152 +-
...ram_listing_file_tvm_ffi_container_shape.h.html | 118 +-
...am_listing_file_tvm_ffi_container_tensor.h.html | 148 +-
...ram_listing_file_tvm_ffi_container_tuple.h.html | 134 +-
...m_listing_file_tvm_ffi_container_variant.h.html | 137 +-
.../program_listing_file_tvm_ffi_device.h.html | 144 +-
.../program_listing_file_tvm_ffi_dtype.h.html | 200 +-
.../program_listing_file_tvm_ffi_endian.h.html | 112 +-
.../program_listing_file_tvm_ffi_enum.h.html | 120 +-
.../program_listing_file_tvm_ffi_error.h.html | 134 +-
.../program_listing_file_tvm_ffi_expected.h.html | 396 ++-
.../program_listing_file_tvm_ffi_extra_base.h.html | 108 +-
...rogram_listing_file_tvm_ffi_extra_base64.h.html | 134 +-
...ram_listing_file_tvm_ffi_extra_c_env_api.h.html | 112 +-
...ram_listing_file_tvm_ffi_extra_cuda_base.h.html | 112 +-
...g_file_tvm_ffi_extra_cuda_cubin_launcher.h.html | 124 +-
...ing_file_tvm_ffi_extra_cuda_device_guard.h.html | 108 +-
..._tvm_ffi_extra_cuda_internal_unified_api.h.html | 116 +-
...ram_listing_file_tvm_ffi_extra_dataclass.h.html | 108 +-
...program_listing_file_tvm_ffi_extra_dtype.h.html | 110 +-
.../program_listing_file_tvm_ffi_extra_json.h.html | 108 +-
...rogram_listing_file_tvm_ffi_extra_module.h.html | 126 +-
...listing_file_tvm_ffi_extra_serialization.h.html | 108 +-
.../program_listing_file_tvm_ffi_extra_stl.h.html | 196 +-
...ting_file_tvm_ffi_extra_structural_equal.h.html | 108 +-
...sting_file_tvm_ffi_extra_structural_hash.h.html | 108 +-
...isting_file_tvm_ffi_extra_structural_key.h.html | 110 +-
...ing_file_tvm_ffi_extra_structural_mutate.h.html | 1871 +++++++++++++
...ting_file_tvm_ffi_extra_structural_visit.h.html | 567 ++--
...g_file_tvm_ffi_extra_visit_error_context.h.html | 125 +-
.../program_listing_file_tvm_ffi_function.h.html | 146 +-
.../program_listing_file_tvm_ffi_memory.h.html | 139 +-
.../program_listing_file_tvm_ffi_object.h.html | 232 +-
.../program_listing_file_tvm_ffi_optional.h.html | 473 +++-
...ting_file_tvm_ffi_reflection_access_path.h.html | 126 +-
...listing_file_tvm_ffi_reflection_accessor.h.html | 165 +-
..._listing_file_tvm_ffi_reflection_creator.h.html | 134 +-
...listing_file_tvm_ffi_reflection_enum_def.h.html | 122 +-
...listing_file_tvm_ffi_reflection_overload.h.html | 148 +-
...listing_file_tvm_ffi_reflection_registry.h.html | 166 +-
.../program_listing_file_tvm_ffi_rvalue_ref.h.html | 116 +-
.../program_listing_file_tvm_ffi_string.h.html | 218 +-
.../program_listing_file_tvm_ffi_tvm_ffi.h.html | 108 +-
...program_listing_file_tvm_ffi_type_traits.h.html | 182 +-
ffi/reference/cpp/generated/structTVMFFIAny.html | 116 +-
.../cpp/generated/structTVMFFIByteArray.html | 120 +-
.../cpp/generated/structTVMFFICustomAllocator.html | 116 +-
.../cpp/generated/structTVMFFIErrorCell.html | 120 +-
.../cpp/generated/structTVMFFIFieldInfo.html | 120 +-
.../cpp/generated/structTVMFFIFunctionCell.html | 120 +-
.../cpp/generated/structTVMFFIMethodInfo.html | 120 +-
.../cpp/generated/structTVMFFIObject.html | 116 +-
.../generated/structTVMFFIObjectAllocHeader.html | 116 +-
.../generated/structTVMFFIOpaqueObjectCell.html | 116 +-
.../cpp/generated/structTVMFFISeqCell.html | 120 +-
.../cpp/generated/structTVMFFIShapeCell.html | 120 +-
.../cpp/generated/structTVMFFITypeAttrColumn.html | 116 +-
.../cpp/generated/structTVMFFITypeInfo.html | 116 +-
.../cpp/generated/structTVMFFITypeMetadata.html | 116 +-
.../cpp/generated/structTVMFFIVersion.html | 116 +-
.../generated/structtvm_1_1ffi_1_1AnyEqual.html | 116 +-
.../cpp/generated/structtvm_1_1ffi_1_1AnyHash.html | 116 +-
.../structtvm_1_1ffi_1_1CUDADeviceGuard.html | 116 +-
...ructtvm_1_1ffi_1_1DenseMapBaseObj_1_1Block.html | 116 +-
...ttvm_1_1ffi_1_1DenseMapBaseObj_1_1ItemType.html | 116 +-
...ttvm_1_1ffi_1_1DenseMapBaseObj_1_1ListNode.html | 116 +-
...structtvm_1_1ffi_1_1FallbackOnlyTraitsBase.html | 120 +-
.../structtvm_1_1ffi_1_1ObjectPtrEqual.html | 116 +-
.../structtvm_1_1ffi_1_1ObjectPtrHash.html | 116 +-
..._1_1ffi_1_1ObjectRefWithFallbackTraitsBase.html | 120 +-
.../structtvm_1_1ffi_1_1StaticTypeKey.html | 132 +-
...ructtvm_1_1ffi_1_1StructuralMutatorVTable.html} | 224 +-
...tructtvm_1_1ffi_1_1StructuralVisitorVTable.html | 116 +-
...cttvm_1_1ffi_1_1TypeToFieldStaticTypeIndex.html | 116 +-
...structtvm_1_1ffi_1_1TypeToRuntimeTypeIndex.html | 116 +-
...der.html => structtvm_1_1ffi_1_1Unchanged.html} | 170 +-
...cttvm_1_1ffi_1_1UnchangedOr_1_1UnsafeInit.html} | 155 +-
.../generated/structtvm_1_1ffi_1_1UnsafeInit.html | 116 +-
.../cpp/generated/structtvm_1_1ffi_1_1dim3.html | 116 +-
.../structtvm_1_1ffi_1_1is__valid__iterator.html | 120 +-
...valid__iterator_3_01Any_00_01IterType_01_4.html | 120 +-
...3_01Optional_3_01T_01_4_00_01IterType_01_4.html | 120 +-
...m_1_1ffi_1_1reflection_1_1FieldInfoBuilder.html | 120 +-
...tructtvm_1_1ffi_1_1reflection_1_1InfoTrait.html | 120 +-
..._1_1ffi_1_1reflection_1_1MethodInfoBuilder.html | 120 +-
...n_1_1OverloadObjectDef_1_1RegisteredMethod.html | 116 +-
.../structtvm_1_1ffi_1_1reflection_1_1init.html | 120 +-
...structtvm_1_1ffi_1_1reflection_1_1init_3_4.html | 120 +-
..._api_8h_1a19fdd3c500a70abb3def842c73694548.html | 116 +-
..._api_8h_1a9eee5afddb1b39723847d67d109c03fe.html | 116 +-
..._api_8h_1aa0abef2444fde82d76337218e5c13b3f.html | 116 +-
..._api_8h_1aba6bf585b03799369fb535db4a5d2e79.html | 116 +-
..._api_8h_1ae22f1a87cbaf7f744d9201948c05d2a5.html | 116 +-
..._api_8h_1aa7abccefa41ac9d3040df689c0c72fa4.html | 116 +-
...json_8h_1a74ca72d0b686c114af683d63e96f2998.html | 116 +-
...json_8h_1a8e808548bb6a99d0a6093ab3e2229ad3.html | 116 +-
...json_8h_1aeb9cdccfbc28034e94e46f9f6bd8d92a.html | 116 +-
...mory_8h_1a65bf625279df21dec5c8097264415113.html | 116 +-
...lection_1a32b044a71a094de492867144ca299c15.html | 116 +-
...lection_1a96513314d45b43eada132d9b96653d9a.html | 116 +-
...ate_8h_1a060cef1ff53c97235144336dac227d74.html} | 184 +-
...ate_8h_1a074bf7a4d1b8b09e4ff3f3a189e77c2b.html} | 155 +-
...ate_8h_1ae69b95653133acacec6825c62bc41191.html} | 150 +-
...isit_8h_1a745d053678181bd38e0fee17b772656b.html | 116 +-
...aits_8h_1a7ba296d84188e5970ba854332a1f04a6.html | 116 +-
...aits_8h_1ae125f5f5a8a94c7340577c1959805295.html | 116 +-
..._api_8h_1a17eed3020dde14e09827a83036143d07.html | 116 +-
..._api_8h_1a3a921bfefc753492061768fced304aa7.html | 116 +-
..._api_8h_1a630f9293b2557ea7cf02fc5dc8975608.html | 116 +-
..._api_8h_1a6378bff570cca29703bf0a888e62c2e1.html | 116 +-
..._api_8h_1a6c2e323543c694ac2fc8e68c40a5fa3d.html | 116 +-
..._api_8h_1a92973dc82316f80c6ad935aafba39418.html | 116 +-
..._api_8h_1ab6a530cbc289dbb52b705f00c8095fc7.html | 116 +-
..._api_8h_1ae60e2c3b6b2591a181c9b0ff8905c2fb.html | 116 +-
ffi/reference/cpp/generated/unabridged_orphan.html | 155 +-
...rray_8h_1a3bbf4c2bf19b72b026ad9e19a23143d5.html | 116 +-
...dule_8h_1a07576e172e0788da4982148edcf53b05.html | 116 +-
...dule_8h_1a5a3a8ad10620c05ded5876f8e607e0eb.html | 116 +-
...dule_8h_1a76f8a8fd1e38fd98cc501623d959ee58.html | 116 +-
...dule_8h_1aae0405c7615139afd4f68b72a862cffe.html | 116 +-
...dule_8h_1ac261d9ce6ffb49dd9878610ce217075e.html | 116 +-
...dule_8h_1ad9c186f1359987cef9278993037b3f66.html | 116 +-
...e__attr_1a096d266d5a2b4523f98f044de310489d.html | 116 +-
...__attr_1a09a99f5f4a3b5ec6730fafa33309461f.html} | 139 +-
...__attr_1a0e928fa6d490779c06ed997dc8af2f35.html} | 150 +-
...e__attr_1a121b81bcda7ce56913cd5cee586ab85f.html | 116 +-
...e__attr_1a18dd88667c1998578f7fdbbf25c9a33d.html | 116 +-
...e__attr_1a214da76aef89b240bf763409cc7ce8f4.html | 116 +-
...e__attr_1a23d22935800d0c4ccefe7a3a0b9a44c0.html | 116 +-
...e__attr_1a3a4fa9be3e1bd6c285c731f0b805fdd2.html | 116 +-
...e__attr_1a4e67838e846b587f2f82127c53b6ef86.html | 116 +-
...e__attr_1a8473fde721d924e9ffb48fc818f6b609.html | 116 +-
...e__attr_1a877b4867a278823be18e6c7f7bff378e.html | 116 +-
...e__attr_1a8ad904131880910d82adfd8e695fe259.html | 116 +-
...e__attr_1a9a2230acd6ebfab403460ad574a1c53b.html | 116 +-
...__attr_1aa78a5fd8e7af94d0540d67c278128960.html} | 150 +-
...e__attr_1ab8b34ca1b301024a7f48d226906bf5c1.html | 116 +-
...e__attr_1abfb8210e9d25dea8ac19ae403e4d09f7.html | 116 +-
...e__attr_1ac49c0e8607e2a8fe80b5a02250330955.html | 116 +-
...e__attr_1ac92fdfe731cfda2d51324cea574c7e08.html | 116 +-
...e__attr_1ade6e6dd8a1abc985d20d7e855a30b7ea.html | 116 +-
...e__attr_1ae970768fd777e6c1dd312de67bfba39c.html | 116 +-
...lection_1aab5d2ed0a499b8237fedeead671a3fa0.html | 118 +-
...ject_8h_1a31dc4b1b71941b3e4217a0e389f807a4.html | 116 +-
...ject_8h_1a5f52cb03b4b23d5d7738ff0386335845.html | 116 +-
...aits_8h_1a62ef22ceb232bddd3e9540895af38cc1.html | 116 +-
..._api_8h_1acd118a980ee3e7e77dd76858652255a3.html | 116 +-
ffi/reference/cpp/index.html | 112 +-
.../python/cpp/generated/tvm_ffi.cpp.build.html | 128 +-
.../cpp/generated/tvm_ffi.cpp.build_inline.html | 130 +-
.../python/cpp/generated/tvm_ffi.cpp.load.html | 128 +-
.../cpp/generated/tvm_ffi.cpp.load_inline.html | 130 +-
.../generated/tvm_ffi.cpp.nvrtc.nvrtc_compile.html | 118 +-
.../generated/tvm_ffi.libinfo.load_lib_module.html | 112 +-
ffi/reference/python/generated/tvm_ffi.Array.html | 116 +-
.../python/generated/tvm_ffi.DLDeviceType.html | 118 +-
....WalkResult.html => tvm_ffi.DefRegionKind.html} | 196 +-
.../python/generated/tvm_ffi.Device_class.html | 122 +-
.../python/generated/tvm_ffi.Function.html | 120 +-
ffi/reference/python/generated/tvm_ffi.Map.html | 122 +-
ffi/reference/python/generated/tvm_ffi.Module.html | 136 +-
ffi/reference/python/generated/tvm_ffi.Object.html | 128 +-
.../generated/tvm_ffi.ObjectConvertible.html | 116 +-
ffi/reference/python/generated/tvm_ffi.Shape.html | 118 +-
.../python/generated/tvm_ffi.StreamContext.html | 116 +-
.../python/generated/tvm_ffi.StructuralKey.html | 122 +-
...Visitor.html => tvm_ffi.StructuralMutator.html} | 287 +-
.../generated/tvm_ffi.StructuralVisitor.html | 167 +-
ffi/reference/python/generated/tvm_ffi.Tensor.html | 122 +-
.../python/generated/tvm_ffi.VisitInterrupt.html | 130 +-
.../python/generated/tvm_ffi.WalkOrder.html | 126 +-
.../python/generated/tvm_ffi.WalkResult.html | 116 +-
.../generated/tvm_ffi.access_path.AccessKind.html | 116 +-
.../generated/tvm_ffi.access_path.AccessPath.html | 118 +-
.../generated/tvm_ffi.access_path.AccessStep.html | 116 +-
.../python/generated/tvm_ffi.convert.html | 120 +-
.../python/generated/tvm_ffi.convert_func.html | 112 +-
.../python/generated/tvm_ffi.device_function.html | 116 +-
ffi/reference/python/generated/tvm_ffi.dtype.html | 148 +-
.../python/generated/tvm_ffi.from_dlpack.html | 114 +-
.../tvm_ffi.get_first_structural_mismatch.html | 112 +-
.../python/generated/tvm_ffi.get_global_func.html | 116 +-
.../tvm_ffi.get_global_func_metadata.html | 114 +-
.../python/generated/tvm_ffi.get_raw_stream.html | 114 +-
.../python/generated/tvm_ffi.init_ffi_api.html | 114 +-
.../python/generated/tvm_ffi.load_module.html | 116 +-
.../python/generated/tvm_ffi.register_error.html | 120 +-
.../generated/tvm_ffi.register_global_func.html | 120 +-
.../python/generated/tvm_ffi.register_object.html | 118 +-
.../generated/tvm_ffi.remove_global_func.html | 120 +-
.../tvm_ffi.serialization.from_json_graph_str.html | 112 +-
.../tvm_ffi.serialization.to_json_graph_str.html | 112 +-
.../python/generated/tvm_ffi.structural_equal.html | 112 +-
.../python/generated/tvm_ffi.structural_hash.html | 112 +-
...tural_walk.html => tvm_ffi.structural_map.html} | 191 +-
...ral_walk.html => tvm_ffi.structural_visit.html} | 190 +-
.../python/generated/tvm_ffi.structural_walk.html | 130 +-
.../python/generated/tvm_ffi.system_lib.html | 120 +-
.../python/generated/tvm_ffi.use_raw_stream.html | 114 +-
.../python/generated/tvm_ffi.use_torch_stream.html | 112 +-
ffi/reference/python/index.html | 132 +-
ffi/reference/rust/generated/help.html | 2 +-
.../rust/generated/search.index/03772cb4fb.js | 1 -
.../rust/generated/search.index/0387b32beb.js | 1 +
.../rust/generated/search.index/03bc5952ae.js | 1 -
.../rust/generated/search.index/04dabe67a7.js | 1 +
.../rust/generated/search.index/0551af7a87.js | 1 +
.../rust/generated/search.index/05ec30a22d.js | 1 +
.../rust/generated/search.index/0613f7a71f.js | 1 +
.../rust/generated/search.index/08543f9576.js | 1 -
.../rust/generated/search.index/08a44370c6.js | 1 +
.../rust/generated/search.index/091bd5731c.js | 1 +
.../rust/generated/search.index/092bd42091.js | 1 -
.../rust/generated/search.index/0a705ec80b.js | 1 +
.../rust/generated/search.index/0aee35cae5.js | 1 +
.../rust/generated/search.index/0c9486ff92.js | 1 +
.../rust/generated/search.index/0e031b0f52.js | 1 -
.../rust/generated/search.index/0e2b0abab6.js | 1 +
.../rust/generated/search.index/0e77f34418.js | 1 -
.../rust/generated/search.index/0e8ab3f5ae.js | 1 +
.../rust/generated/search.index/0ed80fe8fd.js | 1 +
.../rust/generated/search.index/0fff3b064e.js | 1 +
.../rust/generated/search.index/1036ca6021.js | 1 -
.../rust/generated/search.index/11d45689a1.js | 1 -
.../rust/generated/search.index/13232f30cb.js | 1 -
.../rust/generated/search.index/13527bb9b7.js | 1 +
.../rust/generated/search.index/14299f2acb.js | 1 -
.../rust/generated/search.index/1533176707.js | 1 -
.../rust/generated/search.index/16cb798e3d.js | 1 +
.../rust/generated/search.index/17e2c4c655.js | 1 -
.../rust/generated/search.index/18bf5f8bb7.js | 1 -
.../rust/generated/search.index/1a0313691b.js | 1 -
.../rust/generated/search.index/1a687f1f76.js | 1 +
.../rust/generated/search.index/1aaf784bda.js | 1 -
.../rust/generated/search.index/1b7dc9fe56.js | 1 -
.../rust/generated/search.index/1cb38efd5e.js | 1 -
.../rust/generated/search.index/1db7d586fd.js | 1 +
.../rust/generated/search.index/1f15f942e1.js | 1 +
.../rust/generated/search.index/1fc4df4ac8.js | 1 -
.../rust/generated/search.index/20c4d978f8.js | 1 -
.../rust/generated/search.index/20f35cdde0.js | 1 +
.../rust/generated/search.index/2258d9da13.js | 1 -
.../rust/generated/search.index/22bfd35153.js | 1 -
.../rust/generated/search.index/23b4ba2ede.js | 1 -
.../rust/generated/search.index/23b5a66ad7.js | 1 +
.../rust/generated/search.index/23f38ead30.js | 1 +
.../rust/generated/search.index/2468612cf9.js | 1 -
.../rust/generated/search.index/2668a9a546.js | 1 +
.../rust/generated/search.index/26c114fd90.js | 1 -
.../rust/generated/search.index/26c98bef1f.js | 1 +
.../rust/generated/search.index/272b381929.js | 1 +
.../rust/generated/search.index/273a9fc45e.js | 1 -
.../rust/generated/search.index/283af4234a.js | 1 +
.../rust/generated/search.index/28a920b910.js | 1 +
.../rust/generated/search.index/28b73a9a37.js | 1 +
.../rust/generated/search.index/28d446a0c5.js | 1 +
.../rust/generated/search.index/296f03e1a6.js | 1 -
.../rust/generated/search.index/2ae332f173.js | 1 +
.../rust/generated/search.index/2afcf3e2c0.js | 1 +
.../rust/generated/search.index/2b3ddc5f6c.js | 1 -
.../rust/generated/search.index/2e55d77e4e.js | 1 +
.../rust/generated/search.index/2e7f1438bd.js | 1 +
.../rust/generated/search.index/2f8621037e.js | 1 -
.../rust/generated/search.index/2ffd7d5f9d.js | 1 -
.../rust/generated/search.index/3019f0b0b0.js | 1 +
.../rust/generated/search.index/32111108bd.js | 1 +
.../rust/generated/search.index/3244f78789.js | 1 +
.../rust/generated/search.index/325d654d6e.js | 1 +
.../rust/generated/search.index/3382ae3a12.js | 1 +
.../rust/generated/search.index/3476c2fba8.js | 1 +
.../rust/generated/search.index/35dc255970.js | 1 +
.../rust/generated/search.index/37c14bdabf.js | 1 +
.../rust/generated/search.index/390d853c11.js | 1 +
.../rust/generated/search.index/39bd945585.js | 1 +
.../rust/generated/search.index/39c1d329d8.js | 1 +
.../rust/generated/search.index/39e26d185e.js | 1 -
.../rust/generated/search.index/3bd83b5127.js | 1 +
.../rust/generated/search.index/3cdf437706.js | 1 +
.../rust/generated/search.index/3cf1a281b8.js | 1 +
.../rust/generated/search.index/3d2682bf6c.js | 1 +
.../rust/generated/search.index/3d57734fbc.js | 1 +
.../rust/generated/search.index/3d617a3065.js | 1 +
.../rust/generated/search.index/3db5bc2670.js | 1 +
.../rust/generated/search.index/3ecaf0f832.js | 1 +
.../rust/generated/search.index/3f608780fc.js | 1 +
.../rust/generated/search.index/3f7a9f6a12.js | 1 +
.../rust/generated/search.index/3f9d9bd309.js | 1 -
.../rust/generated/search.index/3fc2e587eb.js | 1 +
.../rust/generated/search.index/4119d20f31.js | 1 +
.../rust/generated/search.index/44b1366631.js | 1 -
.../rust/generated/search.index/4571cc8cb1.js | 1 +
.../rust/generated/search.index/46ae1a4f7b.js | 1 -
.../rust/generated/search.index/4744fff1aa.js | 1 -
.../rust/generated/search.index/4790434571.js | 1 +
.../rust/generated/search.index/47b0ffce47.js | 1 -
.../rust/generated/search.index/47e321df3f.js | 1 -
.../rust/generated/search.index/48d41c4333.js | 1 -
.../rust/generated/search.index/4934900d2d.js | 1 +
.../rust/generated/search.index/4945196acd.js | 1 -
.../rust/generated/search.index/4b00601928.js | 1 +
.../rust/generated/search.index/4b06847d19.js | 1 -
.../rust/generated/search.index/4bb02f4dde.js | 1 +
.../rust/generated/search.index/4d9472c4bf.js | 1 -
.../rust/generated/search.index/4dba48dc5a.js | 1 -
.../rust/generated/search.index/4dda778ba5.js | 1 -
.../rust/generated/search.index/4df0218d79.js | 1 -
.../rust/generated/search.index/4e21669ead.js | 1 -
.../rust/generated/search.index/4e7efbd415.js | 1 +
.../rust/generated/search.index/4fb20af1f1.js | 1 +
.../rust/generated/search.index/50bd30183a.js | 1 +
.../rust/generated/search.index/50f0d76875.js | 1 +
.../rust/generated/search.index/513b5fda61.js | 1 -
.../rust/generated/search.index/52606d314f.js | 1 +
.../rust/generated/search.index/52da59aadc.js | 1 -
.../rust/generated/search.index/534072635a.js | 1 +
.../rust/generated/search.index/53ee7a4b56.js | 1 -
.../rust/generated/search.index/5508c37d6d.js | 1 -
.../rust/generated/search.index/55770e70dd.js | 1 +
.../rust/generated/search.index/55ca2081c2.js | 1 +
.../rust/generated/search.index/56d5b3fa2f.js | 1 -
.../rust/generated/search.index/574f82d70b.js | 1 -
.../rust/generated/search.index/58a41244be.js | 1 +
.../rust/generated/search.index/58b127a285.js | 1 -
.../rust/generated/search.index/58f399de3b.js | 1 +
.../rust/generated/search.index/5986a37e9c.js | 1 -
.../rust/generated/search.index/599d3ac125.js | 1 -
.../rust/generated/search.index/59f9ea7b52.js | 1 +
.../rust/generated/search.index/5a4ca175d1.js | 1 +
.../rust/generated/search.index/5a849e915c.js | 1 +
.../rust/generated/search.index/5a9ea03967.js | 1 -
.../rust/generated/search.index/5b190e5149.js | 1 +
.../rust/generated/search.index/5bb2c54844.js | 1 +
.../rust/generated/search.index/5c1d6b1af5.js | 1 +
.../rust/generated/search.index/5c4182f53e.js | 1 +
.../rust/generated/search.index/5c723a7fda.js | 1 +
.../rust/generated/search.index/5ce8e9adf4.js | 1 +
.../rust/generated/search.index/5ee5730130.js | 1 -
.../rust/generated/search.index/5f292ecfa5.js | 1 +
.../rust/generated/search.index/5f314f495e.js | 1 +
.../rust/generated/search.index/5fc75d2cfd.js | 1 -
.../rust/generated/search.index/60adf77e8e.js | 1 +
.../rust/generated/search.index/62f6a3c72f.js | 1 -
.../rust/generated/search.index/6332c66ead.js | 1 +
.../rust/generated/search.index/636cd9caca.js | 1 -
.../rust/generated/search.index/64ed5a664e.js | 1 -
.../rust/generated/search.index/65292c3d6f.js | 1 +
.../rust/generated/search.index/6600abf96e.js | 1 -
.../rust/generated/search.index/661cb6f2ab.js | 1 -
.../rust/generated/search.index/66b75c3c37.js | 1 -
.../rust/generated/search.index/66b9ef2e90.js | 1 +
.../rust/generated/search.index/676405c4b6.js | 1 +
.../rust/generated/search.index/686b27ffce.js | 1 -
.../rust/generated/search.index/68f453e8cb.js | 1 +
.../rust/generated/search.index/68fce82d5f.js | 1 +
.../rust/generated/search.index/6941ddede2.js | 1 +
.../rust/generated/search.index/6aadacf08a.js | 1 +
.../rust/generated/search.index/6c7df64ad9.js | 1 +
.../rust/generated/search.index/6c9f09339a.js | 1 -
.../rust/generated/search.index/6e17953dd3.js | 1 -
.../rust/generated/search.index/6e3327e61e.js | 1 +
.../rust/generated/search.index/6e99ac1c32.js | 1 -
.../rust/generated/search.index/6f4a766c5e.js | 1 -
.../rust/generated/search.index/6f54bbb0a3.js | 1 +
.../rust/generated/search.index/701529078d.js | 1 +
.../rust/generated/search.index/701c0c9623.js | 1 +
.../rust/generated/search.index/70701df4d4.js | 1 -
.../rust/generated/search.index/7072584f1a.js | 1 +
.../rust/generated/search.index/71150d4e92.js | 1 +
.../rust/generated/search.index/71328c525c.js | 1 -
.../rust/generated/search.index/71402f23b9.js | 1 +
.../rust/generated/search.index/71435cc324.js | 1 -
.../rust/generated/search.index/71e919f6e1.js | 1 +
.../rust/generated/search.index/727d128249.js | 1 +
.../rust/generated/search.index/72ad31a590.js | 1 -
.../rust/generated/search.index/736a0526de.js | 1 +
.../rust/generated/search.index/73b8599a39.js | 1 -
.../rust/generated/search.index/74bc080ebb.js | 1 +
.../rust/generated/search.index/770349ccf7.js | 1 +
.../rust/generated/search.index/770912ee5b.js | 1 +
.../rust/generated/search.index/779c46eea8.js | 1 -
.../rust/generated/search.index/7883dc797f.js | 1 +
.../rust/generated/search.index/78a48e7a9f.js | 1 +
.../rust/generated/search.index/793b905fed.js | 1 +
.../rust/generated/search.index/7a57311828.js | 1 -
.../rust/generated/search.index/7b30d60253.js | 1 -
.../rust/generated/search.index/7c06187aee.js | 1 +
.../rust/generated/search.index/7c0ba185a6.js | 1 +
.../rust/generated/search.index/7cdff861b3.js | 1 +
.../rust/generated/search.index/7d4326a6ae.js | 1 -
.../rust/generated/search.index/7d619004d1.js | 1 +
.../rust/generated/search.index/7d967506f2.js | 1 -
.../rust/generated/search.index/desc/ae7fd24bb7.js | 1 -
.../rust/generated/search.index/desc/b06ba070a6.js | 1 +
.../generated/search.index/entry/610372671e.js | 1 +
.../generated/search.index/entry/ff540c65e6.js | 1 -
.../generated/search.index/function/0429def3d0.js | 1 -
.../generated/search.index/function/3dd890be48.js | 1 +
.../generic_inverted_index/bb24b1b0f5.js | 1 -
.../generic_inverted_index/efc4a0d194.js | 1 +
.../rust/generated/search.index/name/2d61dfb249.js | 1 +
.../rust/generated/search.index/name/8e6e820f4c.js | 1 -
.../search.index/normalizedName/7460098516.js | 1 -
.../search.index/normalizedName/bad9900b24.js | 1 +
.../rust/generated/search.index/path/37434c85b9.js | 1 -
.../rust/generated/search.index/path/4afdbdf55f.js | 1 +
ffi/reference/rust/generated/search.index/root.js | 2 +-
.../rust/generated/search.index/type/274bec8097.js | 1 -
.../rust/generated/search.index/type/5ffa5e98b4.js | 1 +
ffi/reference/rust/generated/settings.html | 2 +-
ffi/reference/rust/generated/src-files.js | 4 +-
.../rust/generated/src/tvm_ffi/any.rs.html | 698 ++---
.../src/tvm_ffi/collections/array.rs.html | 548 ++--
.../generated/src/tvm_ffi/collections/map.rs.html | 993 +++----
.../generated/src/tvm_ffi/collections/mod.rs.html | 2 +-
.../src/tvm_ffi/collections/shape.rs.html | 2 +-
.../src/tvm_ffi/collections/tensor.rs.html | 2 +-
.../rust/generated/src/tvm_ffi/derive.rs.html | 2 +-
.../rust/generated/src/tvm_ffi/device.rs.html | 2 +-
.../rust/generated/src/tvm_ffi/dtype.rs.html | 2 +-
.../rust/generated/src/tvm_ffi/error.rs.html | 2 +-
.../generated/src/tvm_ffi/extra/dispatch.rs.html | 74 +
.../rust/generated/src/tvm_ffi/extra/mod.rs.html | 9 +-
.../generated/src/tvm_ffi/extra/module.rs.html | 2 +-
.../src/tvm_ffi/extra/structural_common.rs.html | 193 ++
.../src/tvm_ffi/extra/structural_mutate.rs.html | 2779 ++++++++++++++++++++
.../src/tvm_ffi/extra/structural_visit.rs.html | 2013 ++++++++++++++
.../generated/src/tvm_ffi/extra/unchanged.rs.html | 295 +++
.../rust/generated/src/tvm_ffi/function.rs.html | 608 +++--
.../src/tvm_ffi/function_internal.rs.html | 538 ++--
.../rust/generated/src/tvm_ffi/lib.rs.html | 73 +-
.../rust/generated/src/tvm_ffi/macros.rs.html | 777 +++---
.../src/tvm_ffi/match_any_internal.rs.html | 311 ++-
.../rust/generated/src/tvm_ffi/object.rs.html | 1158 ++++----
.../rust/generated/src/tvm_ffi/optional.rs.html | 396 +--
.../rust/generated/src/tvm_ffi/reflection.rs.html | 235 ++
.../rust/generated/src/tvm_ffi/rvalue_ref.rs.html | 184 ++
.../rust/generated/src/tvm_ffi/string.rs.html | 2 +-
.../rust/generated/src/tvm_ffi/type_traits.rs.html | 767 +++---
.../generated/src/tvm_ffi_macros/dispatch.rs.html | 508 ++++
.../rust/generated/src/tvm_ffi_macros/lib.rs.html | 64 +-
.../generated/src/tvm_ffi_macros/match_any.rs.html | 892 ++++---
.../src/tvm_ffi_macros/object_macros.rs.html | 306 +--
.../generated/src/tvm_ffi_macros/utils.rs.html | 2 +-
.../rust/generated/src/tvm_ffi_sys/c_api.rs.html | 878 ++++---
.../generated/src/tvm_ffi_sys/c_env_api.rs.html | 2 +-
.../rust/generated/src/tvm_ffi_sys/dlpack.rs.html | 2 +-
.../rust/generated/src/tvm_ffi_sys/lib.rs.html | 2 +-
.../{main-fcd733ba.js => main-417ebf0a.js} | 6 +-
.../{rustdoc-17e0aaed.css => rustdoc-1ed95124.css} | 2 +-
.../{search-fd9372ac.js => search-9944d736.js} | 2 +-
.../generated/trait.impl/core/clone/trait.Clone.js | 4 +-
.../rust/generated/trait.impl/core/cmp/trait.Eq.js | 4 +-
.../generated/trait.impl/core/cmp/trait.Ord.js | 2 +-
.../trait.impl/core/cmp/trait.PartialEq.js | 4 +-
.../trait.impl/core/cmp/trait.PartialOrd.js | 2 +-
.../trait.impl/core/convert/trait.From.js | 4 +-
.../trait.impl/core/convert/trait.TryFrom.js | 4 +-
.../trait.impl/core/default/trait.Default.js | 4 +-
.../generated/trait.impl/core/error/trait.Error.js | 2 +-
.../generated/trait.impl/core/fmt/trait.Debug.js | 4 +-
.../generated/trait.impl/core/fmt/trait.Display.js | 2 +-
.../generated/trait.impl/core/hash/trait.Hash.js | 4 +-
.../core/iter/traits/collect/trait.FromIterator.js | 4 +-
.../core/iter/traits/collect/trait.IntoIterator.js | 4 +-
.../traits/exact_size/trait.ExactSizeIterator.js | 2 +-
.../core/iter/traits/iterator/trait.Iterator.js | 4 +-
.../generated/trait.impl/core/marker/trait.Copy.js | 4 +-
.../trait.impl/core/marker/trait.Freeze.js | 4 +-
.../generated/trait.impl/core/marker/trait.Send.js | 4 +-
.../core/marker/trait.StructuralPartialEq.js | 4 +-
.../generated/trait.impl/core/marker/trait.Sync.js | 4 +-
.../trait.impl/core/marker/trait.Unpin.js | 4 +-
.../trait.impl/core/marker/trait.UnsafeUnpin.js | 4 +-
.../trait.impl/core/ops/deref/trait.Deref.js | 4 +-
.../trait.impl/core/ops/deref/trait.DerefMut.js | 2 +-
.../trait.impl/core/ops/drop/trait.Drop.js | 4 +-
.../trait.impl/core/ops/index/trait.Index.js | 4 +-
.../core/panic/unwind_safe/trait.RefUnwindSafe.js | 4 +-
.../core/panic/unwind_safe/trait.UnwindSafe.js | 4 +-
.../tvm_ffi/extra/dispatch/trait.WalkDispatch.js | 9 +
.../extra/structural_mutate/trait.IntoMapper.js | 9 +
.../extra/structural_mutate/trait.IntoMutator.js | 9 +
.../extra/structural_mutate/trait.MapDispatch.js | 9 +
.../structural_mutate/trait.StructuralMutator.js | 9 +
.../extra/structural_visit/trait.IntoVisitor.js | 9 +
.../extra/structural_visit/trait.IntoWalkResult.js | 9 +
.../extra/structural_visit/trait.IntoWalker.js | 9 +
.../structural_visit/trait.StructuralVisitor.js | 9 +
.../tvm_ffi/optional/trait.OptionalCompatible.js | 9 +
ffi/reference/rust/generated/tvm_ffi/all.html | 2 +-
.../rust/generated/tvm_ffi/any/index.html | 2 +-
.../rust/generated/tvm_ffi/any/struct.Any.html | 38 +-
.../rust/generated/tvm_ffi/any/struct.AnyView.html | 38 +-
.../generated/tvm_ffi/any/struct.TryFromTemp.html | 40 +-
.../rust/generated/tvm_ffi/attr.dispatch.html | 2 +
.../generated/tvm_ffi/collections/array/index.html | 2 +-
.../tvm_ffi/collections/array/struct.Array.html | 48 +-
.../collections/array/struct.ArrayIterator.html | 408 +--
.../tvm_ffi/collections/array/struct.ArrayObj.html | 36 +-
.../rust/generated/tvm_ffi/collections/index.html | 2 +-
.../generated/tvm_ffi/collections/map/index.html | 2 +-
.../tvm_ffi/collections/map/struct.Map.html | 100 +-
.../tvm_ffi/collections/map/struct.MapIter.html | 406 +--
.../tvm_ffi/collections/map/struct.MapObj.html | 36 +-
.../tvm_ffi/collections/map/type.MapItems.html | 2 +-
.../tvm_ffi/collections/map/type.MapKeys.html | 2 +-
.../tvm_ffi/collections/map/type.MapValues.html | 2 +-
.../generated/tvm_ffi/collections/shape/index.html | 2 +-
.../tvm_ffi/collections/shape/struct.Shape.html | 388 +--
.../tvm_ffi/collections/shape/struct.ShapeObj.html | 22 +-
.../tvm_ffi/collections/tensor/index.html | 2 +-
.../collections/tensor/struct.CPUNDAlloc.html | 22 +-
.../tvm_ffi/collections/tensor/struct.Tensor.html | 46 +-
.../collections/tensor/struct.TensorObj.html | 22 +-
.../collections/tensor/trait.DLTensorExt.html | 8 +-
.../collections/tensor/trait.NDAllocator.html | 16 +-
.../generated/tvm_ffi/derive/derive.Object.html | 2 +-
.../generated/tvm_ffi/derive/derive.ObjectRef.html | 2 +-
.../rust/generated/tvm_ffi/derive/index.html | 2 +-
.../tvm_ffi/device/fn.current_stream.html | 2 +-
.../generated/tvm_ffi/device/fn.with_stream.html | 4 +-
.../rust/generated/tvm_ffi/device/index.html | 2 +-
.../rust/generated/tvm_ffi/dtype/index.html | 2 +-
.../tvm_ffi/dtype/trait.AsDLDataType.html | 4 +-
.../tvm_ffi/dtype/trait.DLDataTypeExt.html | 8 +-
.../generated/tvm_ffi/enum.DLDataTypeCode.html | 28 +-
.../rust/generated/tvm_ffi/enum.DLDeviceType.html | 28 +-
.../rust/generated/tvm_ffi/enum.TypeIndex.html | 44 +-
.../tvm_ffi/error/constant.ATTRIBUTE_ERROR.html | 2 +-
.../tvm_ffi/error/constant.INDEX_ERROR.html | 2 +-
.../tvm_ffi/error/constant.KEY_ERROR.html | 2 +-
.../tvm_ffi/error/constant.RUNTIME_ERROR.html | 2 +-
.../tvm_ffi/error/constant.TYPE_ERROR.html | 2 +-
.../tvm_ffi/error/constant.VALUE_ERROR.html | 2 +-
.../rust/generated/tvm_ffi/error/index.html | 2 +-
.../rust/generated/tvm_ffi/error/struct.Error.html | 48 +-
.../generated/tvm_ffi/error/struct.ErrorKind.html | 30 +-
.../generated/tvm_ffi/error/struct.ErrorObj.html | 24 +-
.../rust/generated/tvm_ffi/error/type.Result.html | 4 +-
.../generated/tvm_ffi/extra/dispatch/index.html | 2 +
.../tvm_ffi/extra/dispatch/sidebar-items.js | 1 +
.../tvm_ffi/extra/dispatch/trait.WalkDispatch.html | 19 +
.../rust/generated/tvm_ffi/extra/index.html | 2 +-
.../rust/generated/tvm_ffi/extra/module/index.html | 2 +-
.../tvm_ffi/extra/module/struct.Module.html | 36 +-
.../tvm_ffi/extra/module/struct.ModuleObj.html | 24 +-
.../rust/generated/tvm_ffi/extra/sidebar-items.js | 2 +-
.../structural_common/struct.StructuralValue.html | 11 +
.../extra/structural_mutate/fn.structural_map.html | 18 +
.../structural_mutate/fn.structural_mutate.html | 9 +
.../tvm_ffi/extra/structural_mutate/index.html | 5 +
.../extra/structural_mutate/sidebar-items.js | 1 +
.../structural_mutate/struct.InplaceValue.html | 23 +
.../extra/structural_mutate/struct.MapValue.html | 20 +
.../structural_mutate/struct.MutateCallbacks.html | 77 +
.../structural_mutate/struct.MutateContext.html | 58 +
.../extra/structural_mutate/struct.Mutator.html | 78 +
.../struct.StructuralVarRemap.html | 16 +
.../extra/structural_mutate/trait.IntoMapper.html | 85 +
.../extra/structural_mutate/trait.IntoMutator.html | 6 +
.../structural_mutate/trait.MapChainLink.html | 4 +
.../extra/structural_mutate/trait.MapDispatch.html | 20 +
.../structural_mutate/trait.MutateChainLink.html | 2 +
.../structural_mutate/trait.MutateDispatch.html | 17 +
.../structural_mutate/trait.StructuralMutator.html | 151 ++
.../structural_mutate/type.CallbackMutator.html | 4 +
.../extra/structural_visit/enum.DefRegionKind.html | 25 +
.../extra/structural_visit/enum.WalkOrder.html | 20 +
.../extra/structural_visit/enum.WalkResult.html | 22 +
.../structural_visit/fn.structural_visit.html | 8 +
.../extra/structural_visit/fn.structural_walk.html | 17 +
.../tvm_ffi/extra/structural_visit/index.html | 23 +
.../extra/structural_visit/sidebar-items.js | 1 +
.../structural_visit/struct.VisitCallbacks.html | 38 +
.../structural_visit/struct.VisitContext.html | 28 +
.../structural_visit/struct.VisitInterrupt.html | 20 +
.../extra/structural_visit/struct.VisitValue.html | 23 +
.../extra/structural_visit/trait.IntoVisitor.html | 4 +
.../structural_visit/trait.IntoWalkResult.html | 7 +
.../extra/structural_visit/trait.IntoWalker.html | 82 +
.../structural_visit/trait.StructuralVisitor.html | 42 +
.../structural_visit/trait.VisitChainLink.html | 2 +
.../structural_visit/trait.WalkChainLink.html | 5 +
.../generated/tvm_ffi/extra/unchanged/index.html | 2 +
.../tvm_ffi/extra/unchanged/sidebar-items.js | 1 +
.../tvm_ffi/extra/unchanged/struct.Unchanged.html | 19 +
.../extra/unchanged/struct.UnchangedOr.html | 43 +
.../rust/generated/tvm_ffi/function/index.html | 2 +-
.../tvm_ffi/function/struct.Function.html | 103 +-
.../tvm_ffi/function/struct.FunctionObj.html | 24 +-
.../function_internal/fn.call_packed_callable.html | 2 +-
.../generated/tvm_ffi/function_internal/index.html | 2 +-
.../function_internal/trait.ArgIntoRef.html | 5 +-
.../function_internal/trait.AsPackedCallable.html | 54 +-
.../function_internal/trait.IntoArgHolder.html | 5 +-
.../trait.IntoArgHolderTuple.html | 38 +-
.../function_internal/trait.TupleAsPackedArgs.html | 102 +-
ffi/reference/rust/generated/tvm_ffi/index.html | 8 +-
.../generated/tvm_ffi/macro.attach_context.html | 2 +-
.../rust/generated/tvm_ffi/macro.bail.html | 2 +-
.../tvm_ffi/macro.cached_global_func.html | 2 +-
.../generated/tvm_ffi/macro.check_safe_call.html | 2 +-
.../rust/generated/tvm_ffi/macro.ensure.html | 2 +-
.../generated/tvm_ffi/macro.function_name.html | 2 +-
.../generated/tvm_ffi/macro.impl_arg_into_ref.html | 2 +-
.../macro.impl_into_arg_holder_default.html | 2 +-
.../tvm_ffi/macro.impl_object_upcast!.html | 11 +
.../tvm_ffi/macro.impl_object_upcast.html | 9 +
.../generated/tvm_ffi/macro.impl_try_from_any.html | 2 +-
.../macro.impl_try_from_any_for_parametric.html | 2 +-
.../generated/tvm_ffi/macro.into_typed_fn.html | 2 +-
.../rust/generated/tvm_ffi/macro.match_any.html | 2 +-
.../macro.tvm_ffi_dll_export_typed_func.html | 2 +-
.../rust/generated/tvm_ffi/macros/index.html | 2 +-
.../generated/tvm_ffi/object/enum.TypeIndex.html | 44 +-
.../rust/generated/tvm_ffi/object/index.html | 2 +-
.../rust/generated/tvm_ffi/object/sidebar-items.js | 2 +-
.../generated/tvm_ffi/object/struct.Object.html | 24 +-
.../generated/tvm_ffi/object/struct.ObjectArc.html | 52 +-
.../tvm_ffi/object/struct.ObjectIdentity.html | 21 +
.../generated/tvm_ffi/object/struct.ObjectRef.html | 34 +-
.../generated/tvm_ffi/object/trait.ObjectCore.html | 16 +-
.../object/trait.ObjectCoreWithExtraItems.html | 14 +-
.../tvm_ffi/object/trait.ObjectRefCast.html | 6 +-
.../tvm_ffi/object/trait.ObjectRefCore.html | 21 +-
.../rust/generated/tvm_ffi/optional/index.html | 30 +-
.../generated/tvm_ffi/optional/sidebar-items.js | 2 +-
.../tvm_ffi/optional/struct.Optional.html | 52 +-
.../tvm_ffi/optional/trait.OptionalCompatible.html | 8 +
.../tvm_ffi/reflection/fn.get_type_attr.html | 2 +
.../rust/generated/tvm_ffi/reflection/index.html | 2 +
.../generated/tvm_ffi/reflection/sidebar-items.js | 1 +
.../tvm_ffi/reflection/struct.FieldGetter.html | 20 +
.../tvm_ffi/reflection/struct.TypeAttrColumn.html | 18 +
.../rust/generated/tvm_ffi/rvalue_ref/index.html | 2 +
.../generated/tvm_ffi/rvalue_ref/sidebar-items.js | 1 +
.../tvm_ffi/rvalue_ref/struct.RValueRef.html | 27 +
.../rust/generated/tvm_ffi/sidebar-items.js | 2 +-
.../rust/generated/tvm_ffi/string/index.html | 2 +-
.../generated/tvm_ffi/string/struct.Bytes.html | 438 +--
.../generated/tvm_ffi/string/struct.String.html | 657 +++--
.../rust/generated/tvm_ffi/struct.DLDataType.html | 42 +-
.../rust/generated/tvm_ffi/struct.DLDevice.html | 34 +-
.../rust/generated/tvm_ffi/struct.TVMFFIAny.html | 40 +-
.../generated/tvm_ffi/struct.TVMFFIObject.html | 30 +-
.../generated/tvm_ffi/type.TVMFFIStreamHandle.html | 2 +-
.../rust/generated/tvm_ffi/type_traits/index.html | 2 +-
.../tvm_ffi/type_traits/trait.AnyCompatible.html | 47 +-
.../rust/generated/tvm_ffi_macros/all.html | 2 +-
.../generated/tvm_ffi_macros/attr.dispatch.html | 2 +
.../generated/tvm_ffi_macros/derive.Object.html | 2 +-
.../generated/tvm_ffi_macros/derive.ObjectRef.html | 2 +-
.../rust/generated/tvm_ffi_macros/index.html | 2 +-
.../generated/tvm_ffi_macros/macro.match_any.html | 2 +-
.../rust/generated/tvm_ffi_macros/sidebar-items.js | 2 +-
ffi/reference/rust/generated/tvm_ffi_sys/all.html | 2 +-
.../constant.COMBINED_REF_COUNT_BOTH_ONE.html | 2 +-
.../constant.COMBINED_REF_COUNT_MASK_U32.html | 2 +-
.../constant.COMBINED_REF_COUNT_STRONG_ONE.html | 2 +-
.../constant.COMBINED_REF_COUNT_WEAK_ONE.html | 2 +-
.../c_api/enum.TVMFFIBacktraceUpdateMode.html | 28 +-
.../c_api/enum.TVMFFIDefRegionKind.html | 19 +
.../c_api/enum.TVMFFIFieldFlagBitMask.html | 29 +
.../c_api/enum.TVMFFIObjectDeleterFlagBitMask.html | 28 +-
.../tvm_ffi_sys/c_api/enum.TVMFFISEqHashKind.html | 22 +
.../tvm_ffi_sys/c_api/enum.TVMFFITypeIndex.html | 44 +-
.../c_api/fn.TVMFFIAnyViewToOwnedAny.html | 8 +-
.../c_api/fn.TVMFFIBytesFromByteArray.html | 8 +-
.../c_api/fn.TVMFFIDataTypeFromString.html | 8 +-
.../c_api/fn.TVMFFIDataTypeToString.html | 8 +-
.../tvm_ffi_sys/c_api/fn.TVMFFIErrorCreate.html | 12 +-
.../c_api/fn.TVMFFIErrorMoveFromRaised.html | 4 +-
.../tvm_ffi_sys/c_api/fn.TVMFFIErrorSetRaised.html | 2 +-
.../c_api/fn.TVMFFIErrorSetRaisedFromCStr.html | 6 +-
.../tvm_ffi_sys/c_api/fn.TVMFFIFunctionCall.html | 10 +-
.../tvm_ffi_sys/c_api/fn.TVMFFIFunctionCreate.html | 10 +-
.../c_api/fn.TVMFFIFunctionGetGlobal.html | 8 +-
.../c_api/fn.TVMFFIFunctionSetGlobal.html | 8 +-
.../c_api/fn.TVMFFIGetCustomAllocator.html | 2 +-
...nedAny.html => fn.TVMFFIGetTypeAttrColumn.html} | 7 +-
.../tvm_ffi_sys/c_api/fn.TVMFFIGetTypeInfo.html | 6 +-
.../c_api/fn.TVMFFISetCustomAllocator.html | 6 +-
.../c_api/fn.TVMFFIStringFromByteArray.html | 8 +-
.../c_api/fn.TVMFFITensorFromDLPack.html | 12 +-
.../c_api/fn.TVMFFITensorFromDLPackVersioned.html | 12 +-
.../tvm_ffi_sys/c_api/fn.TVMFFITensorToDLPack.html | 6 +-
.../c_api/fn.TVMFFITensorToDLPackVersioned.html | 6 +-
.../c_api/fn.TVMFFITestingDummyTarget.html | 2 +-
.../tvm_ffi_sys/c_api/fn.TVMFFITraceback.html | 12 +-
.../tvm_ffi_sys/c_api/fn.TVMFFITypeKeyToIndex.html | 8 +-
...oString.html => fn.TVMFFITypeRegisterAttr.html} | 9 +-
.../rust/generated/tvm_ffi_sys/c_api/index.html | 4 +-
.../generated/tvm_ffi_sys/c_api/sidebar-items.js | 2 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIAny.html | 36 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIByteArray.html | 30 +-
.../c_api/struct.TVMFFICustomAllocator.html | 26 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIErrorCell.html | 24 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIFieldInfo.html | 50 +-
.../c_api/struct.TVMFFIFunctionCell.html | 24 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIMethodInfo.html | 26 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIObject.html | 30 +-
.../c_api/struct.TVMFFIObjectAllocHeader.html | 24 +-
.../tvm_ffi_sys/c_api/struct.TVMFFIShapeCell.html | 26 +-
.../c_api/struct.TVMFFITypeAttrColumn.html | 34 +-
.../tvm_ffi_sys/c_api/struct.TVMFFITypeInfo.html | 58 +-
.../c_api/struct.TVMFFITypeMetadata.html | 34 +-
.../tvm_ffi_sys/c_api/type.TVMFFIFieldGetter.html | 2 +-
.../tvm_ffi_sys/c_api/type.TVMFFIFieldSetter.html | 2 +-
.../c_api/type.TVMFFIObjectCreator.html | 2 +-
.../c_api/type.TVMFFIObjectDeleter.html | 2 +-
.../tvm_ffi_sys/c_api/type.TVMFFIObjectHandle.html | 2 +-
.../tvm_ffi_sys/c_api/type.TVMFFISafeCallType.html | 2 +-
.../c_api/union.TVMFFIAnyDataUnion.html | 54 +-
.../c_env_api/fn.TVMFFIEnvCheckSignals.html | 2 +-
...n.TVMFFIEnvGetDLPackManagedTensorAllocator.html | 2 +-
.../c_env_api/fn.TVMFFIEnvGetStream.html | 6 +-
.../fn.TVMFFIEnvModLookupFromImports.html | 8 +-
.../fn.TVMFFIEnvModRegisterContextSymbol.html | 8 +-
.../fn.TVMFFIEnvModRegisterSystemLibSymbol.html | 8 +-
.../c_env_api/fn.TVMFFIEnvRegisterCAPI.html | 8 +-
...n.TVMFFIEnvSetDLPackManagedTensorAllocator.html | 8 +-
.../c_env_api/fn.TVMFFIEnvSetStream.html | 10 +-
.../generated/tvm_ffi_sys/c_env_api/index.html | 2 +-
.../type.DLPackManagedTensorAllocator.html | 2 +-
.../c_env_api/type.TVMFFIStreamHandle.html | 2 +-
.../tvm_ffi_sys/dlpack/enum.DLDataTypeCode.html | 28 +-
.../tvm_ffi_sys/dlpack/enum.DLDeviceType.html | 28 +-
.../rust/generated/tvm_ffi_sys/dlpack/index.html | 2 +-
.../tvm_ffi_sys/dlpack/struct.DLDataType.html | 34 +-
.../tvm_ffi_sys/dlpack/struct.DLDevice.html | 30 +-
.../tvm_ffi_sys/dlpack/struct.DLTensor.html | 48 +-
.../rust/generated/tvm_ffi_sys/index.html | 2 +-
.../generated/type.impl/core/result/enum.Result.js | 4 +-
.../generated/type.impl/std/primitive.pointer.js | 4 +-
.../tvm_ffi/collections/map/struct.MapIter.js | 2 +-
.../structural_mutate/struct.MutateContext.js | 9 +
ffi/reference/rust/index.html | 112 +-
ffi/search.html | 326 +--
ffi/searchindex.js | 2 +-
1312 files changed, 91272 insertions(+), 32687 deletions(-)
diff --git a/ffi/_modules/index.html b/ffi/_modules/index.html
index 6aa81aea97e..88302b2205a 100644
--- a/ffi/_modules/index.html
+++ b/ffi/_modules/index.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../_static/documentation_options.js?v=bd121c97"></script>
<script src="../_static/doctools.js?v=9bcbadda"></script>
<script src="../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../_static/downloads/require.min.js"></script>
- <script src="../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/index';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../genindex.html" />
<link rel="search" title="Search" href="../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../_static/searchtools.js"></script>
+ <script src="../_static/language_data.js"></script>
+ <script src="../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -930,7 +827,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -978,8 +875,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/_convert.html b/ffi/_modules/tvm_ffi/_convert.html
index 580c9afc23b..3d749580a6b 100644
--- a/ffi/_modules/tvm_ffi/_convert.html
+++ b/ffi/_modules/tvm_ffi/_convert.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/_convert';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,18 +777,18 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""Conversion utilities to convert Python objects into TVM FFI values."""</span>
+<span class="sd">"""Conversion utilities to convert Python objects into TVM FFI values."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -917,7 +814,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="convert">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.convert.html#tvm_ffi.convert">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">convert</span><span class="p">(</span><span class="n">value</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="n">Any</span><span class="p">:</span> <span class="c1"># noqa: PLR0911,PLR0912</span>
-<span class="w"> </span><span class="sd">"""Convert a Python object into TVM FFI values.</span>
+<span class="w"> </span><span class="sd">"""Convert a Python object into TVM FFI values.</span>
<span class="sd"> This helper mirrors the automatic argument conversion that happens when</span>
<span class="sd"> calling FFI functions. It is primarily useful in tests or places where</span>
@@ -944,12 +841,12 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> assert isinstance(a, tvm_ffi.Array)</span>
<span class="sd"> # Dicts become tvm_ffi.Map</span>
-<span class="sd"> m = tvm_ffi.convert({"a": 1, "b": 2})</span>
+<span class="sd"> m = tvm_ffi.convert({"a": 1, "b": 2})</span>
<span class="sd"> assert isinstance(m, tvm_ffi.Map)</span>
<span class="sd"> # Strings and bytes become zero-copy FFI-aware types</span>
-<span class="sd"> s = tvm_ffi.convert("hello")</span>
-<span class="sd"> b = tvm_ffi.convert(b"bytes")</span>
+<span class="sd"> s = tvm_ffi.convert("hello")</span>
+<span class="sd"> b = tvm_ffi.convert(b"bytes")</span>
<span class="sd"> assert isinstance(s, tvm_ffi.core.String)</span>
<span class="sd"> assert isinstance(b, tvm_ffi.core.Bytes)</span>
@@ -960,7 +857,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # Array libraries that support DLPack export can be converted to Tensor</span>
<span class="sd"> import numpy as np</span>
-<span class="sd"> x = tvm_ffi.convert(np.arange(4, dtype="int32"))</span>
+<span class="sd"> x = tvm_ffi.convert(np.arange(4, dtype="int32"))</span>
<span class="sd"> assert isinstance(x, tvm_ffi.Tensor)</span>
<span class="sd"> Note</span>
@@ -969,7 +866,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> automatically converted. So this function is mainly</span>
<span class="sd"> only used in internal or testing scenarios.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span>
<span class="n">value</span><span class="p">,</span> <span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="n">Object</span><span class="p">,</span> <span class="n">core</span><span class="o">.</span><span class="n">PyNativeObject</span><span class="p">,</span> <span class="nb">bool</span><span class="p">,</span> <span class="n">Number</span><span class="p">,</span> <span class="n">ctypes</span><span class="o">.</span><span class="n">c_void_p</spa [...]
<span class="p">):</span>
@@ -988,31 +885,31 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">_convert_to_ffi_func</span><span class="p">(</span><span class="n">value</span><span class="p">)</span>
<span class="k">elif</span> <span class="n">value</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
<span class="k">return</span> <span class="kc">None</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__dlpack__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__dlpack__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">from_dlpack</span><span class="p">(</span><span class="n">value</span><span class="p">)</span>
<span class="k">elif</span> <span class="n">torch</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span> <span class="ow">and</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="n">torch</span><span class="o">.</span><span class="n">dtype</span><span class="p">):</span>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">_convert_torch_dtype_to_ffi_dtype</span><span class="p">(</span><span class="n">value</span><span class="p">)</span>
<span class="k">elif</span> <span class="n">numpy</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span> <span class="ow">and</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="n">numpy</span><span class="o">.</span><span class="n">dtype</span><span class="p">):</span>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">_convert_numpy_dtype_to_ffi_dtype</span><span class="p">(</span><span class="n">value</span><span class="p">)</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__dlpack_data_type__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__dlpack_data_type__"</span><span class="p">):</span>
<span class="n">cdtype</span> <span class="o">=</span> <span class="n">core</span><span class="o">.</span><span class="n">_create_cdtype_from_tuple</span><span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="n">DataType</span><span class="p">,</span> <span class="o">*</span><span class="n">value</span><span class="o">.</span><span class="n">__dlpack_data_type__</span><span class="p">())</span>
<span class="n">dtype</span> <span class="o">=</span> <span class="nb">str</span><span class="o">.</span><span class="fm">__new__</span><span class="p">(</span><span class="n">_dtype</span><span class="o">.</span><span class="n">dtype</span><span class="p">,</span> <span class="nb">str</span><span class="p">(</span><span class="n">cdtype</span><span class="p">))</span>
<span class="n">dtype</span><span class="o">.</span><span class="n">_tvm_ffi_dtype</span> <span class="o">=</span> <span class="n">cdtype</span>
<span class="k">return</span> <span class="n">dtype</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="ne">Exception</span><span class="p">):</span>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">_convert_to_ffi_error</span><span class="p">(</span><span class="n">value</span><span class="p">)</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_object__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_object__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">value</span><span class="o">.</span><span class="n">__tvm_ffi_object__</span><span class="p">()</span>
<span class="c1"># keep rest protocol values as it is as they can be handled by ffi function</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__cuda_stream__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__cuda_stream__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">value</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_opaque_ptr__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_opaque_ptr__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">value</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__dlpack_device__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__dlpack_device__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">value</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_int__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_int__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">value</span>
- <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_float__"</span><span class="p">):</span>
+ <span class="k">elif</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="s2">"__tvm_ffi_float__"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">value</span>
<span class="k">else</span><span class="p">:</span>
<span class="c1"># in this case, it is an opaque python object</span>
@@ -1026,7 +923,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">pyfunc</span><span class="p">:</span> <span class="n">Callable</span><span class="p">[</span><span class="o">...</span><span class="p">,</span> <span class="n">Any</span><span class="p">],</span>
<span class="n">tensor_cls</span><span class="p">:</span> <span class="nb">type</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">Any</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Convert a Python callable to an FFI :py:class:`~tvm_ffi.Function`.</span>
+<span class="w"> </span><span class="sd">"""Convert a Python callable to an FFI :py:class:`~tvm_ffi.Function`.</span>
<span class="sd"> This is the callable-specific sibling of :py:func:`tvm_ffi.convert`.</span>
<span class="sd"> It accepts one extra argument, ``tensor_cls``, that lets the caller</span>
@@ -1082,10 +979,10 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> See Also</span>
<span class="sd"> --------</span>
<span class="sd"> :py:func:`tvm_ffi.convert` :</span>
-<span class="sd"> Generic value-to-FFI conversion. Use this when you don't need to</span>
+<span class="sd"> Generic value-to-FFI conversion. Use this when you don't need to</span>
<span class="sd"> specify ``tensor_cls``.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">_convert_to_ffi_func</span><span class="p">(</span><span class="n">pyfunc</span><span class="p">,</span> <span class="n">tensor_cls</span><span class="o">=</span><span class="n">tensor_cls</span><span class="p">)</span></div>
</pre></div>
@@ -1128,7 +1025,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1176,8 +1073,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/_dtype.html b/ffi/_modules/tvm_ffi/_dtype.html
index 23495dc0bcf..6531316316e 100644
--- a/ffi/_modules/tvm_ffi/_dtype.html
+++ b/ffi/_modules/tvm_ffi/_dtype.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/_dtype';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,18 +777,18 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""Lightweight dtype wrapper for TVM FFI."""</span>
+<span class="sd">"""Lightweight dtype wrapper for TVM FFI."""</span>
<span class="c1"># pylint: disable=invalid-name</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -904,7 +801,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">class</span><span class="w"> </span><span class="nc">DataTypeCode</span><span class="p">(</span><span class="n">IntEnum</span><span class="p">):</span>
-<span class="w"> </span><span class="sd">"""DLDataTypeCode code in DLTensor."""</span>
+<span class="w"> </span><span class="sd">"""DLDataTypeCode code in DLTensor."""</span>
<span class="n">INT</span> <span class="o">=</span> <span class="mi">0</span>
<span class="n">UINT</span> <span class="o">=</span> <span class="mi">1</span>
@@ -928,7 +825,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="dtype">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.dtype.html#tvm_ffi.dtype">[docs]</a>
<span class="k">class</span><span class="w"> </span><span class="nc">dtype</span><span class="p">(</span><span class="nb">str</span><span class="p">):</span>
-<span class="w"> </span><span class="sd">"""Lightweight data type in TVM FFI.</span>
+<span class="w"> </span><span class="sd">"""Lightweight data type in TVM FFI.</span>
<span class="sd"> It behaves like a Python :class:`str` but also carries an internal FFI</span>
<span class="sd"> representation. You can construct it from strings, NumPy/ML dtypes, or</span>
@@ -946,26 +843,26 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
<span class="sd"> # Create from string</span>
-<span class="sd"> f32 = tvm_ffi.dtype("float32")</span>
+<span class="sd"> f32 = tvm_ffi.dtype("float32")</span>
<span class="sd"> assert f32.bits == 32</span>
<span class="sd"> assert f32.itemsize == 4</span>
<span class="sd"> # Adjust lanes to create vector types</span>
<span class="sd"> v4f32 = f32.with_lanes(4)</span>
-<span class="sd"> assert v4f32 == "float32x4"</span>
+<span class="sd"> assert v4f32 == "float32x4"</span>
<span class="sd"> # Round-trip from a DLPack (code, bits, lanes) triple</span>
<span class="sd"> f16 = tvm_ffi.dtype.from_dlpack_data_type((2, 16, 1))</span>
-<span class="sd"> assert f16 == "float16"</span>
+<span class="sd"> assert f16 == "float16"</span>
<span class="sd"> Note</span>
<span class="sd"> ----</span>
<span class="sd"> This class subclasses str so it can be directly passed</span>
-<span class="sd"> into other array api's dtype arguments.</span>
+<span class="sd"> into other array api's dtype arguments.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
- <span class="vm">__slots__</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"_tvm_ffi_dtype"</span><span class="p">]</span>
+ <span class="vm">__slots__</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"_tvm_ffi_dtype"</span><span class="p">]</span>
<span class="n">_tvm_ffi_dtype</span><span class="p">:</span> <span class="n">core</span><span class="o">.</span><span class="n">DataType</span>
<span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">:</span> <span class="n">ClassVar</span><span class="p">[</span><span class="nb">dict</span><span class="p">[</span><span class="n">Any</span><span class="p">,</span> <span class="nb">str</span><span class="p">]]</span> <span class="o">=</span> <span class="p">{}</span>
@@ -996,7 +893,7 @@ window.runMermaid = runMermaid;</script>
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.dtype.html#tvm_ffi.dtype.from_dlpack_data_type">[docs]</a>
<span class="nd">@staticmethod</span>
<span class="k">def</span><span class="w"> </span><span class="nf">from_dlpack_data_type</span><span class="p">(</span><span class="n">dltype_data_type</span><span class="p">:</span> <span class="nb">tuple</span><span class="p">[</span><span class="nb">int</span><span class="p">,</span> <span class="nb">int</span><span class="p">,</span> <span class="nb">int</span><span class="p">])</span> <span class="o">-></span> <span class="n">dtype</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create a dtype from a DLPack data type tuple.</span>
+<span class="w"> </span><span class="sd">"""Create a dtype from a DLPack data type tuple.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1017,8 +914,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # Create float16 and int8 directly from DLPack triples</span>
<span class="sd"> f16 = tvm_ffi.dtype.from_dlpack_data_type((2, 16, 1))</span>
<span class="sd"> i8 = tvm_ffi.dtype.from_dlpack_data_type((0, 8, 1))</span>
-<span class="sd"> assert f16 == "float16"</span>
-<span class="sd"> assert i8 == "int8"</span>
+<span class="sd"> assert f16 == "float16"</span>
+<span class="sd"> assert i8 == "int8"</span>
<span class="sd"> See Also</span>
<span class="sd"> --------</span>
@@ -1027,7 +924,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> :py:meth:`tvm_ffi.dtype.with_lanes`</span>
<span class="sd"> Create vector dtypes from a scalar base.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="n">cdtype</span> <span class="o">=</span> <span class="n">core</span><span class="o">.</span><span class="n">_create_cdtype_from_tuple</span><span class="p">(</span>
<span class="n">core</span><span class="o">.</span><span class="n">DataType</span><span class="p">,</span>
<span class="n">dltype_data_type</span><span class="p">[</span><span class="mi">0</span><span class="p">],</span>
@@ -1040,12 +937,12 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__repr__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
- <span class="k">return</span> <span class="sa">f</span><span class="s2">"dtype('</span><span class="si">{</span><span class="bp">self</span><span class="si">}</span><span class="s2">')"</span>
+ <span class="k">return</span> <span class="sa">f</span><span class="s2">"dtype('</span><span class="si">{</span><span class="bp">self</span><span class="si">}</span><span class="s2">')"</span>
<div class="viewcode-block" id="dtype.with_lanes">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.dtype.html#tvm_ffi.dtype.with_lanes">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">with_lanes</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">lanes</span><span class="p">:</span> <span class="nb">int</span><span class="p">)</span> <span class="o">-></span> <span class="n">dtype</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create a new dtype with the given number of lanes.</span>
+<span class="w"> </span><span class="sd">"""Create a new dtype with the given number of lanes.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1063,9 +960,9 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> f32 = tvm_ffi.dtype("float32")</span>
+<span class="sd"> f32 = tvm_ffi.dtype("float32")</span>
<span class="sd"> v4f32 = f32.with_lanes(4)</span>
-<span class="sd"> assert v4f32 == "float32x4"</span>
+<span class="sd"> assert v4f32 == "float32x4"</span>
<span class="sd"> assert v4f32.bits == f32.bits and v4f32.lanes == 4</span>
<span class="sd"> See Also</span>
@@ -1073,7 +970,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> :py:meth:`tvm_ffi.dtype.from_dlpack_data_type`</span>
<span class="sd"> Construct from a DLPack ``(code, bits, lanes)`` triple.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="n">cdtype</span> <span class="o">=</span> <span class="n">core</span><span class="o">.</span><span class="n">_create_cdtype_from_tuple</span><span class="p">(</span>
<span class="n">core</span><span class="o">.</span><span class="n">DataType</span><span class="p">,</span>
<span class="bp">self</span><span class="o">.</span><span class="n">_tvm_ffi_dtype</span><span class="o">.</span><span class="n">type_code</span><span class="p">,</span>
@@ -1087,7 +984,7 @@ window.runMermaid = runMermaid;</script>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">itemsize</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Size of one element in bytes.</span>
+<span class="w"> </span><span class="sd">"""Size of one element in bytes.</span>
<span class="sd"> The size is computed as ``bits * lanes // 8``. When the number of</span>
<span class="sd"> lanes is greater than 1, the ``itemsize`` represents the byte size</span>
@@ -1099,8 +996,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> assert tvm_ffi.dtype("float32").itemsize == 4</span>
-<span class="sd"> assert tvm_ffi.dtype("float32").with_lanes(4).itemsize == 16</span>
+<span class="sd"> assert tvm_ffi.dtype("float32").itemsize == 4</span>
+<span class="sd"> assert tvm_ffi.dtype("float32").with_lanes(4).itemsize == 16</span>
<span class="sd"> See Also</span>
<span class="sd"> --------</span>
@@ -1111,12 +1008,12 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> :py:meth:`tvm_ffi.dtype.with_lanes`</span>
<span class="sd"> Create a vector dtype from a scalar base.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_tvm_ffi_dtype</span><span class="o">.</span><span class="n">itemsize</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">type_code</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Integer DLDataTypeCode of the scalar base type.</span>
+<span class="w"> </span><span class="sd">"""Integer DLDataTypeCode of the scalar base type.</span>
<span class="sd"> Examples</span>
<span class="sd"> --------</span>
@@ -1124,7 +1021,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> f32 = tvm_ffi.dtype("float32")</span>
+<span class="sd"> f32 = tvm_ffi.dtype("float32")</span>
<span class="sd"> # The type code is an integer following DLPack conventions</span>
<span class="sd"> assert isinstance(f32.type_code, int)</span>
<span class="sd"> # Consistent with constructing from an explicit (code, bits, lanes)</span>
@@ -1137,12 +1034,12 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> :py:meth:`tvm_ffi.dtype.from_dlpack_data_type`</span>
<span class="sd"> Construct a dtype from a DLPack ``(code, bits, lanes)`` triple.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_tvm_ffi_dtype</span><span class="o">.</span><span class="n">type_code</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">bits</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Number of bits of the scalar base type.</span>
+<span class="w"> </span><span class="sd">"""Number of bits of the scalar base type.</span>
<span class="sd"> Examples</span>
<span class="sd"> --------</span>
@@ -1150,8 +1047,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> assert tvm_ffi.dtype("int8").bits == 8</span>
-<span class="sd"> v4f32 = tvm_ffi.dtype("float32").with_lanes(4)</span>
+<span class="sd"> assert tvm_ffi.dtype("int8").bits == 8</span>
+<span class="sd"> v4f32 = tvm_ffi.dtype("float32").with_lanes(4)</span>
<span class="sd"> assert v4f32.bits == 32 # per-lane bit width</span>
<span class="sd"> See Also</span>
@@ -1161,12 +1058,12 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> :py:attr:`tvm_ffi.dtype.lanes`</span>
<span class="sd"> Number of lanes for vector types.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_tvm_ffi_dtype</span><span class="o">.</span><span class="n">bits</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">lanes</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Number of lanes (for vector types).</span>
+<span class="w"> </span><span class="sd">"""Number of lanes (for vector types).</span>
<span class="sd"> Returns ``1`` for scalar dtypes and the lane count for vector dtypes</span>
<span class="sd"> created via :py:meth:`tvm_ffi.dtype.with_lanes`.</span>
@@ -1177,8 +1074,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> assert tvm_ffi.dtype("float32").lanes == 1</span>
-<span class="sd"> assert tvm_ffi.dtype("float32").with_lanes(4).lanes == 4</span>
+<span class="sd"> assert tvm_ffi.dtype("float32").lanes == 1</span>
+<span class="sd"> assert tvm_ffi.dtype("float32").with_lanes(4).lanes == 4</span>
<span class="sd"> See Also</span>
<span class="sd"> --------</span>
@@ -1187,34 +1084,34 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> :py:attr:`tvm_ffi.dtype.itemsize`</span>
<span class="sd"> Byte size accounting for lanes.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_tvm_ffi_dtype</span><span class="o">.</span><span class="n">lanes</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">is_bool</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">builtins</span><span class="o">.</span><span class="n">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Whether this dtype stores boolean values."""</span>
+<span class="w"> </span><span class="sd">"""Whether this dtype stores boolean values."""</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">type_code</span> <span class="o">==</span> <span class="n">DataTypeCode</span><span class="o">.</span><span class="n">BOOL</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">is_integer</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">builtins</span><span class="o">.</span><span class="n">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Whether this dtype stores signed or unsigned integer values."""</span>
+<span class="w"> </span><span class="sd">"""Whether this dtype stores signed or unsigned integer values."""</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">type_code</span> <span class="ow">in</span> <span class="bp">self</span><span class="o">.</span><span class="n">_INTEGER_CODES</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">is_float</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">builtins</span><span class="o">.</span><span class="n">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Whether this dtype stores floating-point values."""</span>
+<span class="w"> </span><span class="sd">"""Whether this dtype stores floating-point values."""</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">type_code</span> <span class="ow">in</span> <span class="bp">self</span><span class="o">.</span><span class="n">_FLOAT_CODES</span>
<span class="nd">@property</span>
<span class="k">def</span><span class="w"> </span><span class="nf">is_handle</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">builtins</span><span class="o">.</span><span class="n">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Whether this dtype stores opaque handle values."""</span>
+<span class="w"> </span><span class="sd">"""Whether this dtype stores opaque handle values."""</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">type_code</span> <span class="o">==</span> <span class="n">DataTypeCode</span><span class="o">.</span><span class="n">HANDLE</span>
<div class="viewcode-block" id="dtype.is_dtype">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.dtype.html#tvm_ffi.dtype.is_dtype">[docs]</a>
<span class="nd">@staticmethod</span>
<span class="k">def</span><span class="w"> </span><span class="nf">is_dtype</span><span class="p">(</span><span class="n">arg</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="n">dtype</span><span class="p">)</span> <span class="o">-></span> <span class="n">builtins</span><span class="o">.</span><span class="n">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Check if the given string is a valid dtype.</span>
+<span class="w"> </span><span class="sd">"""Check if the given string is a valid dtype.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1232,10 +1129,10 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> assert tvm_ffi.dtype.is_dtype("float32")</span>
-<span class="sd"> assert not tvm_ffi.dtype.is_dtype("not_a_dtype")</span>
+<span class="sd"> assert tvm_ffi.dtype.is_dtype("float32")</span>
+<span class="sd"> assert not tvm_ffi.dtype.is_dtype("not_a_dtype")</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">arg</span><span class="p">,</span> <span class="n">dtype</span><span class="p">):</span>
<span class="k">return</span> <span class="kc">True</span>
<span class="k">try</span><span class="p">:</span>
@@ -1252,31 +1149,31 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># although almost in all cases we want numpy</span>
<span class="kn">import</span><span class="w"> </span><span class="nn">numpy</span><span class="w"> </span><span class="k">as</span><span class="w"> </span><span class="nn">np</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">bool_</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"bool"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int8</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int8"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int16"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int32</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int32"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int64"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint8</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint8"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint16"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint32</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint32"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint64"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float16"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float32"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float64"</span>
- <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">np</span><span class="p">,</span> <span class="s2">"float_"</span><span class="p">):</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float64"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">bool_</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"bool"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int8</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int8"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int16"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int32</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int32"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">int64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"int64"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint8</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint8"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint16"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint32</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint32"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">uint64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"uint64"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float16"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float32"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float64"</span>
+ <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">np</span><span class="p">,</span> <span class="s2">"float_"</span><span class="p">):</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float64</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float64"</span>
<span class="k">except</span> <span class="ne">ImportError</span><span class="p">:</span>
<span class="k">pass</span>
<span class="k">try</span><span class="p">:</span>
<span class="kn">import</span><span class="w"> </span><span class="nn">ml_dtypes</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">bfloat16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"bfloat16"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">float8_e4m3fn</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float8_e4m3fn"</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">float8_e5m2</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float8_e5m2"</span>
- <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="p">,</span> <span class="s2">"float4_e2m1fn"</span><span class="p">):</span> <span class="c1"># ml_dtypes >= 0.5.0</span>
- <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">float4_e2m1fn</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float4_e2m1fn"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">bfloat16</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"bfloat16"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">float8_e4m3fn</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float8_e4m3fn"</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">float8_e5m2</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float8_e5m2"</span>
+ <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="p">,</span> <span class="s2">"float4_e2m1fn"</span><span class="p">):</span> <span class="c1"># ml_dtypes >= 0.5.0</span>
+ <span class="n">dtype</span><span class="o">.</span><span class="n">_NUMPY_DTYPE_TO_STR</span><span class="p">[</span><span class="n">np</span><span class="o">.</span><span class="n">dtype</span><span class="p">(</span><span class="n">ml_dtypes</span><span class="o">.</span><span class="n">float4_e2m1fn</span><span class="p">)]</span> <span class="o">=</span> <span class="s2">"float4_e2m1fn"</span>
<span class="k">except</span> <span class="ne">ImportError</span><span class="p">:</span>
<span class="k">pass</span>
@@ -1286,35 +1183,35 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># note that we can always cover more dtypes via explicit construction</span>
<span class="c1"># from dlpack data type tuple</span>
<span class="c1"># align with choice of numpy 2.0, which moved away from bool_ to bool</span>
-<span class="nb">bool</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"bool"</span><span class="p">)</span>
-<span class="n">int8</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int8"</span><span class="p">)</span>
-<span class="n">int16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int16"</span><span class="p">)</span>
-<span class="n">int32</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int32"</span><span class="p">)</span>
-<span class="n">int64</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int64"</span><span class="p">)</span>
-<span class="n">uint8</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint8"</span><span class="p">)</span>
-<span class="n">uint16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint16"</span><span class="p">)</span>
-<span class="n">uint32</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint32"</span><span class="p">)</span>
-<span class="n">uint64</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint64"</span><span class="p">)</span>
-<span class="n">float64</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float64"</span><span class="p">)</span>
-<span class="n">float32</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float32"</span><span class="p">)</span>
-<span class="n">float16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float16"</span><span class="p">)</span>
-<span class="n">bfloat16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"bfloat16"</span><span class="p">)</span>
+<span class="nb">bool</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"bool"</span><span class="p">)</span>
+<span class="n">int8</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int8"</span><span class="p">)</span>
+<span class="n">int16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int16"</span><span class="p">)</span>
+<span class="n">int32</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int32"</span><span class="p">)</span>
+<span class="n">int64</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"int64"</span><span class="p">)</span>
+<span class="n">uint8</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint8"</span><span class="p">)</span>
+<span class="n">uint16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint16"</span><span class="p">)</span>
+<span class="n">uint32</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint32"</span><span class="p">)</span>
+<span class="n">uint64</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"uint64"</span><span class="p">)</span>
+<span class="n">float64</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float64"</span><span class="p">)</span>
+<span class="n">float32</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float32"</span><span class="p">)</span>
+<span class="n">float16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float16"</span><span class="p">)</span>
+<span class="n">bfloat16</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"bfloat16"</span><span class="p">)</span>
<span class="c1"># float8 dtypes</span>
-<span class="n">float8_e3m4</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e3m4"</span><span class="p">)</span>
-<span class="n">float8_e4m3</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3"</span><span class="p">)</span>
-<span class="n">float8_e4m3b11fnuz</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3b11fnuz"</span><span class="p">)</span>
-<span class="n">float8_e4m3fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3fn"</span><span class="p">)</span>
-<span class="n">float8_e4m3fnuz</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3fnuz"</span><span class="p">)</span>
-<span class="n">float8_e5m2</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e5m2"</span><span class="p">)</span>
-<span class="n">float8_e5m2fnuz</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e5m2fnuz"</span><span class="p">)</span>
-<span class="n">float8_e8m0fnu</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e8m0fnu"</span><span class="p">)</span>
+<span class="n">float8_e3m4</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e3m4"</span><span class="p">)</span>
+<span class="n">float8_e4m3</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3"</span><span class="p">)</span>
+<span class="n">float8_e4m3b11fnuz</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3b11fnuz"</span><span class="p">)</span>
+<span class="n">float8_e4m3fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3fn"</span><span class="p">)</span>
+<span class="n">float8_e4m3fnuz</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e4m3fnuz"</span><span class="p">)</span>
+<span class="n">float8_e5m2</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e5m2"</span><span class="p">)</span>
+<span class="n">float8_e5m2fnuz</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e5m2fnuz"</span><span class="p">)</span>
+<span class="n">float8_e8m0fnu</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float8_e8m0fnu"</span><span class="p">)</span>
<span class="c1"># float6 dtypes</span>
-<span class="n">float6_e2m3fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float6_e2m3fn"</span><span class="p">)</span>
-<span class="n">float6_e3m2fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float6_e3m2fn"</span><span class="p">)</span>
+<span class="n">float6_e2m3fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float6_e2m3fn"</span><span class="p">)</span>
+<span class="n">float6_e3m2fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float6_e3m2fn"</span><span class="p">)</span>
<span class="c1"># float4 dtypes</span>
-<span class="n">float4_e2m1fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float4_e2m1fn"</span><span class="p">)</span>
+<span class="n">float4_e2m1fn</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float4_e2m1fn"</span><span class="p">)</span>
<span class="c1"># float4x2 dtypes</span>
-<span class="n">float4_e2m1fnx2</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float4_e2m1fnx2"</span><span class="p">)</span>
+<span class="n">float4_e2m1fnx2</span> <span class="o">=</span> <span class="n">dtype</span><span class="p">(</span><span class="s2">"float4_e2m1fnx2"</span><span class="p">)</span>
<span class="c1"># alias for torch naming pattern for f4x2</span>
<span class="n">float4_e2m1fn_x2</span> <span class="o">=</span> <span class="n">float4_e2m1fnx2</span>
</pre></div>
@@ -1357,7 +1254,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1405,8 +1302,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/_tensor.html b/ffi/_modules/tvm_ffi/_tensor.html
index 481529f8143..86319b8afb4 100644
--- a/ffi/_modules/tvm_ffi/_tensor.html
+++ b/ffi/_modules/tvm_ffi/_tensor.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/_tensor';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,18 +777,18 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""Tensor related objects and functions."""</span>
+<span class="sd">"""Tensor related objects and functions."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -913,9 +810,9 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="Shape">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Shape.html#tvm_ffi.Shape">[docs]</a>
-<span class="nd">@registry</span><span class="o">.</span><span class="n">register_object</span><span class="p">(</span><span class="s2">"ffi.Shape"</span><span class="p">)</span>
+<span class="nd">@registry</span><span class="o">.</span><span class="n">register_object</span><span class="p">(</span><span class="s2">"ffi.Shape"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">Shape</span><span class="p">(</span><span class="nb">tuple</span><span class="p">,</span> <span class="n">PyNativeObject</span><span class="p">):</span>
-<span class="w"> </span><span class="sd">"""Shape tuple that represents :cpp:class:`tvm::ffi::Shape` returned by an FFI call.</span>
+<span class="w"> </span><span class="sd">"""Shape tuple that represents :cpp:class:`tvm::ffi::Shape` returned by an FFI call.</span>
<span class="sd"> Notes</span>
<span class="sd"> -----</span>
@@ -929,10 +826,10 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import numpy as np</span>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> x = tvm_ffi.from_dlpack(np.arange(6, dtype="int32").reshape(2, 3))</span>
+<span class="sd"> x = tvm_ffi.from_dlpack(np.arange(6, dtype="int32").reshape(2, 3))</span>
<span class="sd"> assert x.shape == (2, 3)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="n">_tvm_ffi_cached_object</span><span class="p">:</span> <span class="n">Any</span>
@@ -943,14 +840,14 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__new__</span><span class="p">(</span><span class="bp">cls</span><span class="p">,</span> <span class="n">content</span><span class="p">:</span> <span class="nb">tuple</span><span class="p">[</span><span class="nb">int</span><span class="p">,</span> <span class="o">...</span><span class="p">])</span> <span class="o">-></span> <span class="n">Shape</span><span class="p">:</span>
<span class="k">if</span> <span class="nb">any</span><span class="p">(</span><span class="ow">not</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">x</span><span class="p">,</span> <span class="n">Integral</span><span class="p">)</span> <span class="k">for</span> <span class="n">x</span> <span class="ow">in</span> <span class="n">content</span><span class="p">):</span>
- <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="s2">"Shape must be a tuple of integers"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="s2">"Shape must be a tuple of integers"</span><span class="p">)</span>
<span class="n">val</span><span class="p">:</span> <span class="n">Shape</span> <span class="o">=</span> <span class="nb">tuple</span><span class="o">.</span><span class="fm">__new__</span><span class="p">(</span><span class="bp">cls</span><span class="p">,</span> <span class="n">content</span><span class="p">)</span>
<span class="n">val</span><span class="o">.</span><span class="n">__init_cached_object_by_constructor__</span><span class="p">(</span><span class="n">_ffi_api</span><span class="o">.</span><span class="n">Shape</span><span class="p">,</span> <span class="o">*</span><span class="n">content</span><span class="p">)</span>
<span class="k">return</span> <span class="n">val</span>
<span class="c1"># pylint: disable=no-self-argument</span>
<span class="k">def</span><span class="w"> </span><span class="nf">__from_tvm_ffi_object__</span><span class="p">(</span><span class="bp">cls</span><span class="p">,</span> <span class="n">obj</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="n">Shape</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Construct from a given tvm object."""</span>
+<span class="w"> </span><span class="sd">"""Construct from a given tvm object."""</span>
<span class="n">content</span> <span class="o">=</span> <span class="n">_shape_obj_get_py_tuple</span><span class="p">(</span><span class="n">obj</span><span class="p">)</span>
<span class="n">val</span><span class="p">:</span> <span class="n">Shape</span> <span class="o">=</span> <span class="nb">tuple</span><span class="o">.</span><span class="fm">__new__</span><span class="p">(</span><span class="bp">cls</span><span class="p">,</span> <span class="n">content</span><span class="p">)</span> <span class="c1"># ty: ignore[invalid-argument-type]</span>
<span class="n">val</span><span class="o">.</span><span class="n">_tvm_ffi_cached_object</span> <span class="o">=</span> <span class="n">obj</span>
@@ -961,7 +858,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="device">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.device_function.html#tvm_ffi.device">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">device</span><span class="p">(</span><span class="n">device_type</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="nb">int</span> <span class="o">|</span> <span class="n">DLDeviceType</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">int</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <sp [...]
-<span class="w"> </span><span class="sd">"""Construct a TVM FFI device with given device type and index.</span>
+<span class="w"> </span><span class="sd">"""Construct a TVM FFI device with given device type and index.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -984,16 +881,16 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> assert tvm_ffi.device("cuda:0") == tvm_ffi.device("cuda", 0)</span>
-<span class="sd"> assert tvm_ffi.device("cpu:0") == tvm_ffi.device("cpu", 0)</span>
+<span class="sd"> assert tvm_ffi.device("cuda:0") == tvm_ffi.device("cuda", 0)</span>
+<span class="sd"> assert tvm_ffi.device("cpu:0") == tvm_ffi.device("cpu", 0)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="c1"># must refer to core._CLASS_DEVICE so we pick up override here</span>
<span class="k">return</span> <span class="n">core</span><span class="o">.</span><span class="n">_CLASS_DEVICE</span><span class="p">(</span><span class="n">device_type</span><span class="p">,</span> <span class="n">index</span><span class="p">)</span></div>
-<span class="n">__all__</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"DLDeviceType"</span><span class="p">,</span> <span class="s2">"Device"</span><span class="p">,</span> <span class="s2">"Tensor"</span><span class="p">,</span> <span class="s2">"device"</span><span class="p">,</span> <span class="s2">"from_dlpack"</span><span class="p">]</span>
+<span class="n">__all__</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"DLDeviceType"</span><span class="p">,</span> <span class="s2">"Device"</span><span class="p">,</span> <span class="s2">"Tensor"</span><span class="p">,</span> <span class="s2">"device"</span><span class="p">,</span> <span class="s2">"from_dlpack"</span><span class="p">]</span>
</pre></div>
</article>
@@ -1034,7 +931,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1082,8 +979,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/access_path.html b/ffi/_modules/tvm_ffi/access_path.html
index 433f6b07ec1..9d15b3d5232 100644
--- a/ffi/_modules/tvm_ffi/access_path.html
+++ b/ffi/_modules/tvm_ffi/access_path.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/access_path';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,19 +777,19 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
<span class="c1"># pylint: disable=invalid-name</span>
-<span class="sd">"""Access path classes."""</span>
+<span class="sd">"""Access path classes."""</span>
<span class="c1"># tvm-ffi-stubgen(begin): import-section</span>
<span class="c1"># fmt: off</span>
@@ -915,7 +812,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="AccessKind">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessKind.html#tvm_ffi.access_path.AccessKind">[docs]</a>
<span class="k">class</span><span class="w"> </span><span class="nc">AccessKind</span><span class="p">(</span><span class="n">IntEnum</span><span class="p">):</span>
-<span class="w"> </span><span class="sd">"""Kinds of access steps in an access path."""</span>
+<span class="w"> </span><span class="sd">"""Kinds of access steps in an access path."""</span>
<span class="n">ATTR</span> <span class="o">=</span> <span class="mi">0</span>
<span class="n">ARRAY_ITEM</span> <span class="o">=</span> <span class="mi">1</span>
@@ -928,9 +825,9 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="AccessStep">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessStep.html#tvm_ffi.access_path.AccessStep">[docs]</a>
-<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.reflection.AccessStep"</span><span class="p">)</span>
+<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.reflection.AccessStep"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">AccessStep</span><span class="p">(</span><span class="n">Object</span><span class="p">):</span>
-<span class="w"> </span><span class="sd">"""Access step container."""</span>
+<span class="w"> </span><span class="sd">"""Access step container."""</span>
<span class="c1"># tvm-ffi-stubgen(ty-map): ffi.reflection.AccessStep -> ffi.access_path.AccessStep</span>
<span class="c1"># tvm-ffi-stubgen(begin): object/ffi.reflection.AccessStep</span>
@@ -944,9 +841,9 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="AccessPath">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath">[docs]</a>
-<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.reflection.AccessPath"</span><span class="p">)</span>
+<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.reflection.AccessPath"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">AccessPath</span><span class="p">(</span><span class="n">Object</span><span class="p">):</span>
-<span class="w"> </span><span class="sd">"""Access path container.</span>
+<span class="w"> </span><span class="sd">"""Access path container.</span>
<span class="sd"> It describes how to reach a nested attribute or item</span>
<span class="sd"> inside a complex FFI object by recording a sequence of steps</span>
@@ -961,10 +858,10 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> root = AccessPath.root()</span>
<span class="sd"> # Build a path equivalent to obj.layer.weight[2]</span>
-<span class="sd"> p = root.attr("layer").attr("weight").array_item(2)</span>
+<span class="sd"> p = root.attr("layer").attr("weight").array_item(2)</span>
<span class="sd"> assert isinstance(p, AccessPath)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="c1"># tvm-ffi-stubgen(ty-map): ffi.reflection.AccessPath -> ffi.access_path.AccessPath</span>
<span class="c1"># tvm-ffi-stubgen(begin): object/ffi.reflection.AccessPath</span>
@@ -991,10 +888,10 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="AccessPath.__init__">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.__init__">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Disallow direct construction; use `AccessPath.root()` instead."""</span>
+<span class="w"> </span><span class="sd">"""Disallow direct construction; use `AccessPath.root()` instead."""</span>
<span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span>
- <span class="s2">"AccessPath can't be initialized directly. "</span>
- <span class="s2">"Use AccessPath.root() to create a path to the root object"</span>
+ <span class="s2">"AccessPath can't be initialized directly. "</span>
+ <span class="s2">"Use AccessPath.root() to create a path to the root object"</span>
<span class="p">)</span></div>
@@ -1002,25 +899,25 @@ window.runMermaid = runMermaid;</script>
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.root">[docs]</a>
<span class="nd">@staticmethod</span>
<span class="k">def</span><span class="w"> </span><span class="nf">root</span><span class="p">()</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create a root access path.</span>
+<span class="w"> </span><span class="sd">"""Create a root access path.</span>
<span class="sd"> Returns</span>
<span class="sd"> -------</span>
<span class="sd"> AccessPath</span>
<span class="sd"> A path representing the root of an object graph.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="n">AccessPath</span><span class="o">.</span><span class="n">_root</span><span class="p">()</span></div>
<span class="k">def</span><span class="w"> </span><span class="fm">__eq__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return whether two access paths are equal."""</span>
+<span class="w"> </span><span class="sd">"""Return whether two access paths are equal."""</span>
<span class="k">if</span> <span class="ow">not</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="n">AccessPath</span><span class="p">):</span>
<span class="k">return</span> <span class="kc">False</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_path_equal</span><span class="p">(</span><span class="n">other</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__ne__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return whether two access paths are not equal."""</span>
+<span class="w"> </span><span class="sd">"""Return whether two access paths are not equal."""</span>
<span class="k">if</span> <span class="ow">not</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="n">AccessPath</span><span class="p">):</span>
<span class="k">return</span> <span class="kc">True</span>
<span class="k">return</span> <span class="ow">not</span> <span class="bp">self</span><span class="o">.</span><span class="n">_path_equal</span><span class="p">(</span><span class="n">other</span><span class="p">)</span>
@@ -1028,7 +925,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="AccessPath.is_prefix_of">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.is_prefix_of">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">is_prefix_of</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">AccessPath</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Check if this access path is a prefix of another access path.</span>
+<span class="w"> </span><span class="sd">"""Check if this access path is a prefix of another access path.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1040,14 +937,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> bool</span>
<span class="sd"> True if this access path is a prefix of the other access path, False otherwise</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_is_prefix_of</span><span class="p">(</span><span class="n">other</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.attr">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.attr">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">attr</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">attr_key</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create an access path to the attribute of the current object.</span>
+<span class="w"> </span><span class="sd">"""Create an access path to the attribute of the current object.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1059,14 +956,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> AccessPath</span>
<span class="sd"> The extended access path</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_attr</span><span class="p">(</span><span class="n">attr_key</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.attr_missing">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.attr_missing">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">attr_missing</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">attr_key</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create an access path that indicate an attribute is missing.</span>
+<span class="w"> </span><span class="sd">"""Create an access path that indicate an attribute is missing.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1078,14 +975,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> AccessPath</span>
<span class="sd"> The extended access path</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_attr_missing</span><span class="p">(</span><span class="n">attr_key</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.array_item">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.array_item">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">array_item</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">int</span><span class="p">)</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create an access path to the item of the current array.</span>
+<span class="w"> </span><span class="sd">"""Create an access path to the item of the current array.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1097,14 +994,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> AccessPath</span>
<span class="sd"> The extended access path</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_array_item</span><span class="p">(</span><span class="n">index</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.array_item_missing">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.array_item_missing">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">array_item_missing</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">int</span><span class="p">)</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create an access path that indicate an array item is missing.</span>
+<span class="w"> </span><span class="sd">"""Create an access path that indicate an array item is missing.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1116,14 +1013,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> AccessPath</span>
<span class="sd"> The extended access path</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_array_item_missing</span><span class="p">(</span><span class="n">index</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.map_item">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.map_item">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">map_item</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create an access path to the item of the current map.</span>
+<span class="w"> </span><span class="sd">"""Create an access path to the item of the current map.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1135,14 +1032,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> AccessPath</span>
<span class="sd"> The extended access path</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_map_item</span><span class="p">(</span><span class="n">key</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.map_item_missing">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.map_item_missing">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">map_item_missing</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="n">AccessPath</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Create an access path that indicate a map item is missing.</span>
+<span class="w"> </span><span class="sd">"""Create an access path that indicate a map item is missing.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1154,21 +1051,21 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> AccessPath</span>
<span class="sd"> The extended access path</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_map_item_missing</span><span class="p">(</span><span class="n">key</span><span class="p">)</span></div>
<div class="viewcode-block" id="AccessPath.to_steps">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.access_path.AccessPath.html#tvm_ffi.access_path.AccessPath.to_steps">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">to_steps</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">Sequence</span><span class="p">[</span><span class="n">AccessStep</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Convert the access path to a list of access steps.</span>
+<span class="w"> </span><span class="sd">"""Convert the access path to a list of access steps.</span>
<span class="sd"> Returns</span>
<span class="sd"> -------</span>
<span class="sd"> access_steps</span>
<span class="sd"> The list of access steps</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_to_steps</span><span class="p">()</span></div>
@@ -1214,7 +1111,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1262,8 +1159,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/container.html b/ffi/_modules/tvm_ffi/container.html
index 4f3b81bf0bb..3cd75548056 100644
--- a/ffi/_modules/tvm_ffi/container.html
+++ b/ffi/_modules/tvm_ffi/container.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/container';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,18 +777,18 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""Container classes."""</span>
+<span class="sd">"""Container classes."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -920,16 +817,16 @@ window.runMermaid = runMermaid;</script>
<span class="kn">from</span><span class="w"> </span><span class="nn">.</span><span class="w"> </span><span class="kn">import</span> <span class="n">_ffi_api</span><span class="p">,</span> <span class="n">core</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">.registry</span><span class="w"> </span><span class="kn">import</span> <span class="n">register_object</span>
-<span class="n">__all__</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"Array"</span><span class="p">,</span> <span class="s2">"Dict"</span><span class="p">,</span> <span class="s2">"List"</span><span class="p">,</span> <span class="s2">"Map"</span><span class="p">]</span>
+<span class="n">__all__</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"Array"</span><span class="p">,</span> <span class="s2">"Dict"</span><span class="p">,</span> <span class="s2">"List"</span><span class="p">,</span> <span class="s2">"Map"</span><span class="p">]</span>
-<span class="n">T</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"T"</span><span class="p">)</span>
-<span class="n">K</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"K"</span><span class="p">)</span>
-<span class="n">V</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"V"</span><span class="p">)</span>
+<span class="n">T</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"T"</span><span class="p">)</span>
+<span class="n">K</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"K"</span><span class="p">)</span>
+<span class="n">V</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"V"</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_sequence_compare_other</span><span class="p">(</span><span class="n">this</span><span class="p">:</span> <span class="nb">object</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">object</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Normalize plain Python sequences for structural container equality."""</span>
+<span class="w"> </span><span class="sd">"""Normalize plain Python sequences for structural container equality."""</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="p">(</span><span class="nb">str</span><span class="p">,</span> <span class="nb">bytes</span><span class="p">,</span> <span class="n">Mapping</span><span class="p">)):</span>
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="n">Sequence</span><span class="p">):</span>
@@ -945,7 +842,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">value</span><span class="p">:</span> <span class="nb">object</span><span class="p">,</span>
<span class="n">ffi_contains</span><span class="p">:</span> <span class="n">Callable</span><span class="p">[[</span><span class="n">Any</span><span class="p">,</span> <span class="nb">object</span><span class="p">],</span> <span class="nb">bool</span><span class="p">],</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Containment with a Python-level structural fallback for nested sequences."""</span>
+<span class="w"> </span><span class="sd">"""Containment with a Python-level structural fallback for nested sequences."""</span>
<span class="k">if</span> <span class="n">ffi_contains</span><span class="p">(</span><span class="n">this</span><span class="p">,</span> <span class="n">value</span><span class="p">):</span>
<span class="k">return</span> <span class="kc">True</span>
<span class="k">if</span> <span class="ow">not</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="n">Sequence</span><span class="p">)</span> <span class="ow">or</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">value</span><span class="p">,</span> <span class="p">(</span><span class="nb">str</span><span class="p">,</span> <span class="nb">bytes</span><span class="p">)):</span>
@@ -960,7 +857,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="kc">False</span>
-<span class="n">_DefaultT</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"_DefaultT"</span><span class="p">)</span>
+<span class="n">_DefaultT</span> <span class="o">=</span> <span class="n">TypeVar</span><span class="p">(</span><span class="s2">"_DefaultT"</span><span class="p">)</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">.core</span><span class="w"> </span><span class="kn">import</span> <span class="n">MISSING</span>
@@ -971,7 +868,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">length</span><span class="p">:</span> <span class="nb">int</span><span class="p">,</span>
<span class="n">idx</span><span class="p">:</span> <span class="n">SupportsIndex</span> <span class="o">|</span> <span class="nb">slice</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">T</span> <span class="o">|</span> <span class="nb">list</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Implement a pythonic __getitem__ helper.</span>
+<span class="w"> </span><span class="sd">"""Implement a pythonic __getitem__ helper.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -992,7 +889,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> result</span>
<span class="sd"> The element for integer indices or a :class:`list` for slices.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">idx</span><span class="p">,</span> <span class="nb">slice</span><span class="p">):</span>
<span class="n">start</span><span class="p">,</span> <span class="n">stop</span><span class="p">,</span> <span class="n">step</span> <span class="o">=</span> <span class="n">idx</span><span class="o">.</span><span class="n">indices</span><span class="p">(</span><span class="n">length</span><span class="p">)</span>
<span class="k">return</span> <span class="p">[</span><span class="n">elem_getter</span><span class="p">(</span><span class="n">obj</span><span class="p">,</span> <span class="n">i</span><span class="p">)</span> <span class="k">for</span> <span class="n">i</span> <span class="ow">in</span> <span class="nb">range</span><span class="p">(</span><span class="n">start</span><span class="p">,</span> <span class="n">stop</span><span class="p">,</span> <span class="n">step</span><span cl [...]
@@ -1002,13 +899,13 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">normalize_index</span><span class="p">(</span><span class="n">length</span><span class="p">:</span> <span class="nb">int</span><span class="p">,</span> <span class="n">idx</span><span class="p">:</span> <span class="n">SupportsIndex</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Normalize and bounds-check a Python index."""</span>
+<span class="w"> </span><span class="sd">"""Normalize and bounds-check a Python index."""</span>
<span class="k">try</span><span class="p">:</span>
<span class="n">index</span> <span class="o">=</span> <span class="n">operator</span><span class="o">.</span><span class="n">index</span><span class="p">(</span><span class="n">idx</span><span class="p">)</span>
<span class="k">except</span> <span class="ne">TypeError</span> <span class="k">as</span> <span class="n">exc</span><span class="p">:</span> <span class="c1"># pragma: no cover - defensive, matches list behaviour</span>
- <span class="k">raise</span> <span class="ne">TypeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"indices must be integers or slices, not </span><span class="si">{</span><span class="nb">type</span><span class="p">(</span><span class="n">idx</span><span class="p">)</span><span class="o">.</span><span class="vm">__name__</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span> <span class="kn">from</span><span class= [...]
+ <span class="k">raise</span> <span class="ne">TypeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"indices must be integers or slices, not </span><span class="si">{</span><span class="nb">type</span><span class="p">(</span><span class="n">idx</span><span class="p">)</span><span class="o">.</span><span class="vm">__name__</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span> <span class="kn">from</span><span class="w"> </spa [...]
<span class="k">if</span> <span class="n">index</span> <span class="o"><</span> <span class="o">-</span><span class="n">length</span> <span class="ow">or</span> <span class="n">index</span> <span class="o">>=</span> <span class="n">length</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">IndexError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Index out of range. size: </span><span class="si">{</span><span class="n">length</span><span class="si">}</span><span class="s2">, got index </span><span class="si">{</span><span class="n">index</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">IndexError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Index out of range. size: </span><span class="si">{</span><span class="n">length</span><span class="si">}</span><span class="s2">, got index </span><span class="si">{</span><span class="n">index</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">index</span> <span class="o"><</span> <span class="mi">0</span><span class="p">:</span>
<span class="n">index</span> <span class="o">+=</span> <span class="n">length</span>
<span class="k">return</span> <span class="n">index</span>
@@ -1016,9 +913,9 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="Array">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Array.html#tvm_ffi.Array">[docs]</a>
-<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.Array"</span><span class="p">)</span>
+<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.Array"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">Array</span><span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="n">CContainerBase</span><span class="p">,</span> <span class="n">core</span><span class="o">.</span><span class="n">Object</span><span class="p">,</span> <span class="n">Sequence</span><span class="p">[</span><span class="n">T</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Array container that represents a sequence of values in the FFI.</span>
+<span class="w"> </span><span class="sd">"""Array container that represents a sequence of values in the FFI.</span>
<span class="sd"> :py:func:`tvm_ffi.convert` will map python list/tuple to this class.</span>
@@ -1044,7 +941,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> --------</span>
<span class="sd"> :py:func:`tvm_ffi.convert`</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="c1"># tvm-ffi-stubgen(begin): object/ffi.Array</span>
<span class="c1"># fmt: off</span>
@@ -1057,7 +954,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="Array.__init__">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Array.html#tvm_ffi.Array.__init__">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">input_list</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">])</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Construct an Array from a Python sequence."""</span>
+<span class="w"> </span><span class="sd">"""Construct an Array from a Python sequence."""</span>
<span class="bp">self</span><span class="o">.</span><span class="n">__init_handle_by_constructor__</span><span class="p">(</span><span class="n">_ffi_api</span><span class="o">.</span><span class="n">Array</span><span class="p">,</span> <span class="o">*</span><span class="n">input_list</span><span class="p">)</span></div>
@@ -1068,27 +965,27 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__getitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">idx</span><span class="p">:</span> <span class="nb">slice</span><span class="p">,</span> <span class="o">/</span><span class="p">)</span> <span class="o">-></span> <span class="nb">list</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span> <span class="o">...</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__getitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">idx</span><span class="p">:</span> <span class="n">SupportsIndex</span> <span class="o">|</span> <span class="nb">slice</span><span class="p">,</span> <span class="o">/</span><span class="p">)</span> <span class="o">-></span> <span class="n">T</span> <span class="o">|</span> <span class="nb">list</span><sp [...]
-<span class="w"> </span><span class="sd">"""Return one element or a list for a slice."""</span>
+<span class="w"> </span><span class="sd">"""Return one element or a list for a slice."""</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="n">result</span> <span class="o">=</span> <span class="n">getitem_helper</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ArrayGetItem</span><span class="p">,</span> <span class="n">length</span><span class="p">,</span> <span class="n">idx</span><span class="p">)</span>
<span class="k">return</span> <span class="n">result</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__len__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return the number of elements in the array."""</span>
+<span class="w"> </span><span class="sd">"""Return the number of elements in the array."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ArraySize</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__iter__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">Iterator</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Iterate over the elements in the array."""</span>
+<span class="w"> </span><span class="sd">"""Iterate over the elements in the array."""</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">for</span> <span class="n">i</span> <span class="ow">in</span> <span class="nb">range</span><span class="p">(</span><span class="n">length</span><span class="p">):</span>
<span class="k">yield</span> <span class="bp">self</span><span class="p">[</span><span class="n">i</span><span class="p">]</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__contains__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">value</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Check if the array contains a value."""</span>
+<span class="w"> </span><span class="sd">"""Check if the array contains a value."""</span>
<span class="k">return</span> <span class="n">_sequence_contains</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">value</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ArrayContains</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__eq__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural equality."""</span>
+<span class="w"> </span><span class="sd">"""Structural equality."""</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">))</span> <span class="ow">or</span> <span class="nb">isinstance</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="n">other</span><span class="p">)):</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveEq</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
<span class="n">other</span> <span class="o">=</span> <span class="n">_sequence_compare_other</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
@@ -1097,33 +994,33 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveEq</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__ne__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
+<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
<span class="n">result</span> <span class="o">=</span> <span class="bp">self</span><span class="o">.</span><span class="fm">__eq__</span><span class="p">(</span><span class="n">other</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="ow">is</span> <span class="bp">NotImplemented</span><span class="p">:</span>
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">return</span> <span class="ow">not</span> <span class="n">result</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__hash__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural hash."""</span>
+<span class="w"> </span><span class="sd">"""Structural hash."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveHash</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__bool__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return True if the array is non-empty."""</span>
+<span class="w"> </span><span class="sd">"""Return True if the array is non-empty."""</span>
<span class="k">return</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">></span> <span class="mi">0</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__add__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">])</span> <span class="o">-></span> <span class="n">Array</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Concatenate two arrays."""</span>
+<span class="w"> </span><span class="sd">"""Concatenate two arrays."""</span>
<span class="k">return</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">)(</span><span class="n">itertools</span><span class="o">.</span><span class="n">chain</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">))</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__radd__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">])</span> <span class="o">-></span> <span class="n">Array</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Concatenate two arrays."""</span>
+<span class="w"> </span><span class="sd">"""Concatenate two arrays."""</span>
<span class="k">return</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">)(</span><span class="n">itertools</span><span class="o">.</span><span class="n">chain</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="bp">self</span><span class="p">))</span></div>
-<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.List"</span><span class="p">)</span>
+<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.List"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">List</span><span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="n">CContainerBase</span><span class="p">,</span> <span class="n">core</span><span class="o">.</span><span class="n">Object</span><span class="p">,</span> <span class="n">MutableSequence</span><span class="p">[</span><span class="n">T</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Mutable list container that represents a mutable sequence in the FFI."""</span>
+<span class="w"> </span><span class="sd">"""Mutable list container that represents a mutable sequence in the FFI."""</span>
<span class="c1"># tvm-ffi-stubgen(begin): object/ffi.List</span>
<span class="c1"># fmt: off</span>
@@ -1134,7 +1031,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DeepCopy</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">input_list</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">]</span> <span class="o">=</span> <span class="p">())</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Construct a List from a Python sequence."""</span>
+<span class="w"> </span><span class="sd">"""Construct a List from a Python sequence."""</span>
<span class="bp">self</span><span class="o">.</span><span class="n">__init_handle_by_constructor__</span><span class="p">(</span><span class="n">_ffi_api</span><span class="o">.</span><span class="n">List</span><span class="p">,</span> <span class="o">*</span><span class="n">input_list</span><span class="p">)</span>
<span class="nd">@overload</span>
@@ -1144,7 +1041,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__getitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">idx</span><span class="p">:</span> <span class="nb">slice</span><span class="p">,</span> <span class="o">/</span><span class="p">)</span> <span class="o">-></span> <span class="nb">list</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span> <span class="o">...</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__getitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">idx</span><span class="p">:</span> <span class="n">SupportsIndex</span> <span class="o">|</span> <span class="nb">slice</span><span class="p">,</span> <span class="o">/</span><span class="p">)</span> <span class="o">-></span> <span class="n">T</span> <span class="o">|</span> <span class="nb">list</span><sp [...]
-<span class="w"> </span><span class="sd">"""Return one element or a list for a slice."""</span>
+<span class="w"> </span><span class="sd">"""Return one element or a list for a slice."""</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">return</span> <span class="n">getitem_helper</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListGetItem</span><span class="p">,</span> <span class="n">length</span><span class="p">,</span> <span class="n">idx</span><span class="p">)</span>
@@ -1155,7 +1052,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__setitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">slice</span><span class="p">[</span><span class="nb">int</span> <span class="o">|</span> <span class="kc">None</span><span class="p">],</span> <span class="n">value</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span [...]
<span class="k">def</span><span class="w"> </span><span class="fm">__setitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="n">SupportsIndex</span> <span class="o">|</span> <span class="nb">slice</span><span class="p">[</span><span class="nb">int</span> <span class="o">|</span> <span class="kc">None</span><span class="p">],</span> <span class="n">value</span><span class="p">:</span [...]
-<span class="w"> </span><span class="sd">"""Set one element or assign a slice."""</span>
+<span class="w"> </span><span class="sd">"""Set one element or assign a slice."""</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">index</span><span class="p">,</span> <span class="nb">slice</span><span class="p">):</span>
<span class="n">replacement</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span><span class="n">cast</span><span class="p">(</span><span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">],</span> <span class="n">value</span><span class="p">))</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
@@ -1164,8 +1061,8 @@ window.runMermaid = runMermaid;</script>
<span class="n">target_indices</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span><span class="nb">range</span><span class="p">(</span><span class="n">start</span><span class="p">,</span> <span class="n">stop</span><span class="p">,</span> <span class="n">step</span><span class="p">))</span>
<span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">replacement</span><span class="p">)</span> <span class="o">!=</span> <span class="nb">len</span><span class="p">(</span><span class="n">target_indices</span><span class="p">):</span>
<span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span>
- <span class="s2">"attempt to assign sequence of size "</span>
- <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="nb">len</span><span class="p">(</span><span class="n">replacement</span><span class="p">)</span><span class="si">}</span><span class="s2"> to extended slice of size </span><span class="si">{</span><span class="nb">len</span><span class="p">(</span><span class="n">target_indices</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span>
+ <span class="s2">"attempt to assign sequence of size "</span>
+ <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="nb">len</span><span class="p">(</span><span class="n">replacement</span><span class="p">)</span><span class="si">}</span><span class="s2"> to extended slice of size </span><span class="si">{</span><span class="nb">len</span><span class="p">(</span><span class="n">target_indices</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span>
<span class="p">)</span>
<span class="k">for</span> <span class="n">i</span><span class="p">,</span> <span class="n">item</span> <span class="ow">in</span> <span class="nb">zip</span><span class="p">(</span><span class="n">target_indices</span><span class="p">,</span> <span class="n">replacement</span><span class="p">):</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListSetItem</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">i</span><span class="p">,</span> <span class="n">item</span><span class="p">)</span>
@@ -1184,7 +1081,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__delitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">slice</span><span class="p">[</span><span class="nb">int</span> <span class="o">|</span> <span class="kc">None</span><span class="p">])</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span> <span class="o">...</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__delitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="n">SupportsIndex</span> <span class="o">|</span> <span class="nb">slice</span><span class="p">[</span><span class="nb">int</span> <span class="o">|</span> <span class="kc">None</span><span class="p">])</span> <span class="o">-></span> <span class="kc">None< [...]
-<span class="w"> </span><span class="sd">"""Delete one element or a slice."""</span>
+<span class="w"> </span><span class="sd">"""Delete one element or a slice."""</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">index</span><span class="p">,</span> <span class="nb">slice</span><span class="p">):</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="n">start</span><span class="p">,</span> <span class="n">stop</span><span class="p">,</span> <span class="n">step</span> <span class="o">=</span> <span class="n">index</span><span class="o">.</span><span class="n">indices</span><span class="p">(</span><span class="n">length</span><span class="p">)</span>
@@ -1204,7 +1101,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListErase</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">normalized_index</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">insert</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">int</span><span class="p">,</span> <span class="n">value</span><span class="p">:</span> <span class="n">T</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Insert value before index."""</span>
+<span class="w"> </span><span class="sd">"""Insert value before index."""</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">if</span> <span class="n">index</span> <span class="o"><</span> <span class="mi">0</span><span class="p">:</span>
<span class="n">index</span> <span class="o">=</span> <span class="nb">max</span><span class="p">(</span><span class="mi">0</span><span class="p">,</span> <span class="n">index</span> <span class="o">+</span> <span class="n">length</span><span class="p">)</span>
@@ -1213,46 +1110,46 @@ window.runMermaid = runMermaid;</script>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListInsert</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">,</span> <span class="n">value</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">append</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">value</span><span class="p">:</span> <span class="n">T</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Append one value to the tail."""</span>
+<span class="w"> </span><span class="sd">"""Append one value to the tail."""</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListAppend</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">value</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">clear</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Remove all elements from the list."""</span>
+<span class="w"> </span><span class="sd">"""Remove all elements from the list."""</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListClear</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">reverse</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Reverse the list in-place."""</span>
+<span class="w"> </span><span class="sd">"""Reverse the list in-place."""</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListReverse</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">pop</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">index</span><span class="p">:</span> <span class="nb">int</span> <span class="o">=</span> <span class="o">-</span><span class="mi">1</span><span class="p">)</span> <span class="o">-></span> <span class="n">T</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Remove and return item at index (default last)."""</span>
+<span class="w"> </span><span class="sd">"""Remove and return item at index (default last)."""</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">if</span> <span class="n">length</span> <span class="o">==</span> <span class="mi">0</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">IndexError</span><span class="p">(</span><span class="s2">"pop from empty list"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">IndexError</span><span class="p">(</span><span class="s2">"pop from empty list"</span><span class="p">)</span>
<span class="n">normalized_index</span> <span class="o">=</span> <span class="n">normalize_index</span><span class="p">(</span><span class="n">length</span><span class="p">,</span> <span class="n">index</span><span class="p">)</span>
<span class="k">return</span> <span class="n">cast</span><span class="p">(</span><span class="n">T</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListPop</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">normalized_index</span><span class="p">))</span>
<span class="k">def</span><span class="w"> </span><span class="nf">extend</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">values</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">])</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Append elements from an iterable."""</span>
+<span class="w"> </span><span class="sd">"""Append elements from an iterable."""</span>
<span class="n">end</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="bp">self</span><span class="p">[</span><span class="n">end</span><span class="p">:</span><span class="n">end</span><span class="p">]</span> <span class="o">=</span> <span class="n">values</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__len__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return the number of elements in the list."""</span>
+<span class="w"> </span><span class="sd">"""Return the number of elements in the list."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListSize</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__iter__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">Iterator</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Iterate over the elements in the list."""</span>
+<span class="w"> </span><span class="sd">"""Iterate over the elements in the list."""</span>
<span class="n">length</span> <span class="o">=</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">for</span> <span class="n">i</span> <span class="ow">in</span> <span class="nb">range</span><span class="p">(</span><span class="n">length</span><span class="p">):</span>
<span class="k">yield</span> <span class="n">cast</span><span class="p">(</span><span class="n">T</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListGetItem</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">i</span><span class="p">))</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__contains__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">value</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Check if the list contains a value."""</span>
+<span class="w"> </span><span class="sd">"""Check if the list contains a value."""</span>
<span class="k">return</span> <span class="n">_sequence_contains</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">value</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">ListContains</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__eq__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural equality."""</span>
+<span class="w"> </span><span class="sd">"""Structural equality."""</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">))</span> <span class="ow">or</span> <span class="nb">isinstance</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="n">other</span><span class="p">)):</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveEq</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
<span class="n">other</span> <span class="o">=</span> <span class="n">_sequence_compare_other</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
@@ -1261,31 +1158,31 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveEq</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__ne__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
+<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
<span class="n">result</span> <span class="o">=</span> <span class="bp">self</span><span class="o">.</span><span class="fm">__eq__</span><span class="p">(</span><span class="n">other</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="ow">is</span> <span class="bp">NotImplemented</span><span class="p">:</span>
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">return</span> <span class="ow">not</span> <span class="n">result</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__hash__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural hash."""</span>
+<span class="w"> </span><span class="sd">"""Structural hash."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveHash</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__bool__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return True if the list is non-empty."""</span>
+<span class="w"> </span><span class="sd">"""Return True if the list is non-empty."""</span>
<span class="k">return</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">></span> <span class="mi">0</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__add__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">])</span> <span class="o">-></span> <span class="n">List</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Concatenate two lists."""</span>
+<span class="w"> </span><span class="sd">"""Concatenate two lists."""</span>
<span class="k">return</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">)(</span><span class="n">itertools</span><span class="o">.</span><span class="n">chain</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">))</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__radd__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Iterable</span><span class="p">[</span><span class="n">T</span><span class="p">])</span> <span class="o">-></span> <span class="n">List</span><span class="p">[</span><span class="n">T</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Concatenate two lists."""</span>
+<span class="w"> </span><span class="sd">"""Concatenate two lists."""</span>
<span class="k">return</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">)(</span><span class="n">itertools</span><span class="o">.</span><span class="n">chain</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="bp">self</span><span class="p">))</span>
<span class="k">class</span><span class="w"> </span><span class="nc">KeysView</span><span class="p">(</span><span class="n">KeysViewBase</span><span class="p">[</span><span class="n">K</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Helper class to return keys view."""</span>
+<span class="w"> </span><span class="sd">"""Helper class to return keys view."""</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span>
<span class="bp">self</span><span class="p">,</span>
@@ -1312,7 +1209,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">class</span><span class="w"> </span><span class="nc">ValuesView</span><span class="p">(</span><span class="n">ValuesViewBase</span><span class="p">[</span><span class="n">V</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Helper class to return values view."""</span>
+<span class="w"> </span><span class="sd">"""Helper class to return values view."""</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span>
<span class="bp">self</span><span class="p">,</span>
@@ -1336,7 +1233,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">class</span><span class="w"> </span><span class="nc">ItemsView</span><span class="p">(</span><span class="n">ItemsViewBase</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Helper class to return items view."""</span>
+<span class="w"> </span><span class="sd">"""Helper class to return items view."""</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span>
<span class="bp">self</span><span class="p">,</span>
@@ -1372,9 +1269,9 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="Map">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Map.html#tvm_ffi.Map">[docs]</a>
-<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.Map"</span><span class="p">)</span>
+<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.Map"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">Map</span><span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="n">CContainerBase</span><span class="p">,</span> <span class="n">core</span><span class="o">.</span><span class="n">Object</span><span class="p">,</span> <span class="n">Mapping</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Map container.</span>
+<span class="w"> </span><span class="sd">"""Map container.</span>
<span class="sd"> :py:func:`tvm_ffi.convert` will map python dict to this class.</span>
@@ -1389,10 +1286,10 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> amap = tvm_ffi.Map({"a": 1, "b": 2})</span>
+<span class="sd"> amap = tvm_ffi.Map({"a": 1, "b": 2})</span>
<span class="sd"> assert len(amap) == 2</span>
-<span class="sd"> assert amap["a"] == 1</span>
-<span class="sd"> assert amap["b"] == 2</span>
+<span class="sd"> assert amap["a"] == 1</span>
+<span class="sd"> assert amap["b"] == 2</span>
<span class="sd"> Notes</span>
<span class="sd"> -----</span>
@@ -1402,7 +1299,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> --------</span>
<span class="sd"> :py:func:`tvm_ffi.convert`</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="c1"># tvm-ffi-stubgen(begin): object/ffi.Map</span>
<span class="c1"># fmt: off</span>
@@ -1415,7 +1312,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="Map.__init__">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Map.html#tvm_ffi.Map.__init__">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">input_dict</span><span class="p">:</span> <span class="n">Mapping</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">])</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Construct a Map from a Python mapping."""</span>
+<span class="w"> </span><span class="sd">"""Construct a Map from a Python mapping."""</span>
<span class="n">list_kvs</span><span class="p">:</span> <span class="nb">list</span><span class="p">[</span><span class="n">Any</span><span class="p">]</span> <span class="o">=</span> <span class="p">[]</span>
<span class="k">for</span> <span class="n">k</span><span class="p">,</span> <span class="n">v</span> <span class="ow">in</span> <span class="n">input_dict</span><span class="o">.</span><span class="n">items</span><span class="p">():</span>
<span class="n">list_kvs</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">k</span><span class="p">)</span>
@@ -1424,61 +1321,61 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="fm">__getitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">:</span> <span class="n">K</span><span class="p">)</span> <span class="o">-></span> <span class="n">V</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return the value for key `k` or raise KeyError."""</span>
+<span class="w"> </span><span class="sd">"""Return the value for key `k` or raise KeyError."""</span>
<span class="k">return</span> <span class="n">cast</span><span class="p">(</span><span class="n">V</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">MapGetItem</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">))</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__contains__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return True if the map contains key `k`."""</span>
+<span class="w"> </span><span class="sd">"""Return True if the map contains key `k`."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">MapCount</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">)</span> <span class="o">!=</span> <span class="mi">0</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__eq__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural equality."""</span>
+<span class="w"> </span><span class="sd">"""Structural equality."""</span>
<span class="k">if</span> <span class="ow">not</span> <span class="p">(</span><span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">))</span> <span class="ow">or</span> <span class="nb">isinstance</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span [...]
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveEq</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__ne__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
+<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
<span class="n">result</span> <span class="o">=</span> <span class="bp">self</span><span class="o">.</span><span class="fm">__eq__</span><span class="p">(</span><span class="n">other</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="ow">is</span> <span class="bp">NotImplemented</span><span class="p">:</span>
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">return</span> <span class="ow">not</span> <span class="n">result</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__hash__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural hash."""</span>
+<span class="w"> </span><span class="sd">"""Structural hash."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveHash</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<div class="viewcode-block" id="Map.keys">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Map.html#tvm_ffi.Map.keys">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">keys</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">KeysView</span><span class="p">[</span><span class="n">K</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Return a dynamic view of the map's keys."""</span>
+<span class="w"> </span><span class="sd">"""Return a dynamic view of the map's keys."""</span>
<span class="k">return</span> <span class="n">KeysView</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span></div>
<div class="viewcode-block" id="Map.values">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Map.html#tvm_ffi.Map.values">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">values</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">ValuesView</span><span class="p">[</span><span class="n">V</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Return a dynamic view of the map's values."""</span>
+<span class="w"> </span><span class="sd">"""Return a dynamic view of the map's values."""</span>
<span class="k">return</span> <span class="n">ValuesView</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span></div>
<div class="viewcode-block" id="Map.items">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Map.html#tvm_ffi.Map.items">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">items</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">ItemsView</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Get the items from the map."""</span>
+<span class="w"> </span><span class="sd">"""Get the items from the map."""</span>
<span class="k">return</span> <span class="n">ItemsView</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span></div>
<span class="k">def</span><span class="w"> </span><span class="fm">__len__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return the number of items in the map."""</span>
+<span class="w"> </span><span class="sd">"""Return the number of items in the map."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">MapSize</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__bool__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return True if the map is non-empty."""</span>
+<span class="w"> </span><span class="sd">"""Return True if the map is non-empty."""</span>
<span class="k">return</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">></span> <span class="mi">0</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__iter__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">Iterator</span><span class="p">[</span><span class="n">K</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Iterate over the map's keys."""</span>
+<span class="w"> </span><span class="sd">"""Iterate over the map's keys."""</span>
<span class="k">return</span> <span class="nb">iter</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">keys</span><span class="p">())</span>
<span class="nd">@overload</span>
@@ -1490,7 +1387,7 @@ window.runMermaid = runMermaid;</script>
<div class="viewcode-block" id="Map.get">
<a class="viewcode-back" href="../../reference/python/generated/tvm_ffi.Map.html#tvm_ffi.Map.get">[docs]</a>
<span class="k">def</span><span class="w"> </span><span class="nf">get</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">:</span> <span class="n">K</span><span class="p">,</span> <span class="n">default</span><span class="p">:</span> <span class="n">V</span> <span class="o">|</span> <span class="n">_DefaultT</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc"> [...]
-<span class="w"> </span><span class="sd">"""Get an element with a default value.</span>
+<span class="w"> </span><span class="sd">"""Get an element with a default value.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1505,7 +1402,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> value</span>
<span class="sd"> The result value.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="n">ret</span> <span class="o">=</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">MapGetItemOrMissing</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">)</span>
<span class="k">if</span> <span class="n">MISSING</span><span class="o">.</span><span class="n">same_as</span><span class="p">(</span><span class="n">ret</span><span class="p">):</span>
<span class="k">return</span> <span class="n">default</span>
@@ -1514,9 +1411,9 @@ window.runMermaid = runMermaid;</script>
-<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.Dict"</span><span class="p">)</span>
+<span class="nd">@register_object</span><span class="p">(</span><span class="s2">"ffi.Dict"</span><span class="p">)</span>
<span class="k">class</span><span class="w"> </span><span class="nc">Dict</span><span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="n">CContainerBase</span><span class="p">,</span> <span class="n">core</span><span class="o">.</span><span class="n">Object</span><span class="p">,</span> <span class="n">MutableMapping</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">]):</span>
-<span class="w"> </span><span class="sd">"""Mutable dictionary container with shared reference semantics.</span>
+<span class="w"> </span><span class="sd">"""Mutable dictionary container with shared reference semantics.</span>
<span class="sd"> Unlike :class:`Map`, ``Dict`` does NOT implement copy-on-write.</span>
<span class="sd"> Mutations happen directly on the underlying shared object.</span>
@@ -1533,17 +1430,17 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi</span>
-<span class="sd"> d = tvm_ffi.Dict({"a": 1, "b": 2})</span>
-<span class="sd"> d["c"] = 3</span>
+<span class="sd"> d = tvm_ffi.Dict({"a": 1, "b": 2})</span>
+<span class="sd"> d["c"] = 3</span>
<span class="sd"> assert len(d) == 3</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">def</span><span class="w"> </span><span class="nf">__deepcopy__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">memo</span><span class="p">:</span> <span class="n">Any</span> <span class="o">=</span> <span class="kc">None</span><span class="p">)</span> <span class="o">-></span> <span class="n">Any</span><span class="p">:</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DeepCopy</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">input_dict</span><span class="p">:</span> <span class="n">Mapping</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">]</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class= [...]
-<span class="w"> </span><span class="sd">"""Construct a Dict from a Python mapping."""</span>
+<span class="w"> </span><span class="sd">"""Construct a Dict from a Python mapping."""</span>
<span class="n">list_kvs</span><span class="p">:</span> <span class="nb">list</span><span class="p">[</span><span class="n">Any</span><span class="p">]</span> <span class="o">=</span> <span class="p">[]</span>
<span class="k">if</span> <span class="n">input_dict</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="k">for</span> <span class="n">k</span><span class="p">,</span> <span class="n">v</span> <span class="ow">in</span> <span class="n">input_dict</span><span class="o">.</span><span class="n">items</span><span class="p">():</span>
@@ -1552,62 +1449,62 @@ window.runMermaid = runMermaid;</script>
<span class="bp">self</span><span class="o">.</span><span class="n">__init_handle_by_constructor__</span><span class="p">(</span><span class="n">_ffi_api</span><span class="o">.</span><span class="n">Dict</span><span class="p">,</span> <span class="o">*</span><span class="n">list_kvs</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__getitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">:</span> <span class="n">K</span><span class="p">)</span> <span class="o">-></span> <span class="n">V</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return the value for key `k` or raise KeyError."""</span>
+<span class="w"> </span><span class="sd">"""Return the value for key `k` or raise KeyError."""</span>
<span class="k">return</span> <span class="n">cast</span><span class="p">(</span><span class="n">V</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictGetItem</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">))</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__setitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">:</span> <span class="n">K</span><span class="p">,</span> <span class="n">v</span><span class="p">:</span> <span class="n">V</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Set the value for key `k`."""</span>
+<span class="w"> </span><span class="sd">"""Set the value for key `k`."""</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictSetItem</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">,</span> <span class="n">v</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__delitem__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">:</span> <span class="n">K</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Delete the entry for key `k`."""</span>
+<span class="w"> </span><span class="sd">"""Delete the entry for key `k`."""</span>
<span class="k">if</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictCount</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">)</span> <span class="o">==</span> <span class="mi">0</span><span class="p">:</span>
<span class="k">raise</span> <span class="ne">KeyError</span><span class="p">(</span><span class="n">k</span><span class="p">)</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictErase</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__contains__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return True if the dict contains key `k`."""</span>
+<span class="w"> </span><span class="sd">"""Return True if the dict contains key `k`."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictCount</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">k</span><span class="p">)</span> <span class="o">!=</span> <span class="mi">0</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__eq__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural equality."""</span>
+<span class="w"> </span><span class="sd">"""Structural equality."""</span>
<span class="k">if</span> <span class="ow">not</span> <span class="p">(</span><span class="nb">isinstance</span><span class="p">(</span><span class="n">other</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span><span class="bp">self</span><span class="p">))</span> <span class="ow">or</span> <span class="nb">isinstance</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="nb">type</span><span class="p">(</span [...]
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveEq</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__ne__</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="nb">object</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
+<span class="w"> </span><span class="sd">"""Structural inequality."""</span>
<span class="n">result</span> <span class="o">=</span> <span class="bp">self</span><span class="o">.</span><span class="fm">__eq__</span><span class="p">(</span><span class="n">other</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="ow">is</span> <span class="bp">NotImplemented</span><span class="p">:</span>
<span class="k">return</span> <span class="bp">NotImplemented</span>
<span class="k">return</span> <span class="ow">not</span> <span class="n">result</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__hash__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Structural hash."""</span>
+<span class="w"> </span><span class="sd">"""Structural hash."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">RecursiveHash</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__len__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">int</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return the number of items in the dict."""</span>
+<span class="w"> </span><span class="sd">"""Return the number of items in the dict."""</span>
<span class="k">return</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictSize</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__bool__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="nb">bool</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Return True if the dict is non-empty."""</span>
+<span class="w"> </span><span class="sd">"""Return True if the dict is non-empty."""</span>
<span class="k">return</span> <span class="nb">len</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">></span> <span class="mi">0</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__iter__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">Iterator</span><span class="p">[</span><span class="n">K</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Iterate over the dict's keys."""</span>
+<span class="w"> </span><span class="sd">"""Iterate over the dict's keys."""</span>
<span class="k">return</span> <span class="nb">iter</span><span class="p">(</span><span class="bp">self</span><span class="o">.</span><span class="n">keys</span><span class="p">())</span>
<span class="k">def</span><span class="w"> </span><span class="nf">keys</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">KeysView</span><span class="p">[</span><span class="n">K</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Return a dynamic view of the dict's keys."""</span>
+<span class="w"> </span><span class="sd">"""Return a dynamic view of the dict's keys."""</span>
<span class="k">return</span> <span class="n">KeysView</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictForwardIterFunctor</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">values</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">ValuesView</span><span class="p">[</span><span class="n">V</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Return a dynamic view of the dict's values."""</span>
+<span class="w"> </span><span class="sd">"""Return a dynamic view of the dict's values."""</span>
<span class="k">return</span> <span class="n">ValuesView</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictForwardIterFunctor</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">items</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="n">ItemsView</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Get the items from the dict."""</span>
+<span class="w"> </span><span class="sd">"""Get the items from the dict."""</span>
<span class="k">return</span> <span class="n">ItemsView</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictForwardIterFunctor</span><span class="p">)</span>
<span class="nd">@overload</span>
@@ -1617,16 +1514,16 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">get</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">:</span> <span class="n">K</span><span class="p">,</span> <span class="n">default</span><span class="p">:</span> <span class="n">V</span> <span class="o">|</span> <span class="n">_DefaultT</span><span class="p">)</span> <span class="o">-></span> <span class="n">V</span> <span class="o">|< [...]
<span class="k">def</span><span class="w"> </span><span class="nf">get</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">:</span> <span class="n">K</span><span class="p">,</span> <span class="n">default</span><span class="p">:</span> <span class="n">V</span> <span class="o">|</span> <span class="n">_DefaultT</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc"> [...]
-<span class="w"> </span><span class="sd">"""Get an element with a default value."""</span>
+<span class="w"> </span><span class="sd">"""Get an element with a default value."""</span>
<span class="n">ret</span> <span class="o">=</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictGetItemOrMissing</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">)</span>
<span class="k">if</span> <span class="n">MISSING</span><span class="o">.</span><span class="n">same_as</span><span class="p">(</span><span class="n">ret</span><span class="p">):</span>
<span class="k">return</span> <span class="n">default</span>
<span class="k">return</span> <span class="n">ret</span>
<span class="k">def</span><span class="w"> </span><span class="nf">pop</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">:</span> <span class="n">K</span><span class="p">,</span> <span class="o">*</span><span class="n">args</span><span class="p">:</span> <span class="n">V</span> <span class="o">|</span> <span class="n">_DefaultT</span><span class="p">)</span> <span class="o">-></span> <span class="n">V</spa [...]
-<span class="w"> </span><span class="sd">"""Remove and return value for key, or default if not present."""</span>
+<span class="w"> </span><span class="sd">"""Remove and return value for key, or default if not present."""</span>
<span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">args</span><span class="p">)</span> <span class="o">></span> <span class="mi">1</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">TypeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"pop expected at most 2 arguments, got </span><span class="si">{</span><span class="mi">1</span><span class="w"> </span><span class="o">+</span><span class="w"> </span><span class="nb">len</span><span class="p">(</span><span class="n">args</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">TypeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"pop expected at most 2 arguments, got </span><span class="si">{</span><span class="mi">1</span><span class="w"> </span><span class="o">+</span><span class="w"> </span><span class="nb">len</span><span class="p">(</span><span class="n">args</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="n">ret</span> <span class="o">=</span> <span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictGetItemOrMissing</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">key</span><span class="p">)</span>
<span class="k">if</span> <span class="n">MISSING</span><span class="o">.</span><span class="n">same_as</span><span class="p">(</span><span class="n">ret</span><span class="p">):</span>
<span class="k">if</span> <span class="n">args</span><span class="p">:</span>
@@ -1636,11 +1533,11 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="n">cast</span><span class="p">(</span><span class="n">V</span><span class="p">,</span> <span class="n">ret</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">clear</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Remove all elements from the dict."""</span>
+<span class="w"> </span><span class="sd">"""Remove all elements from the dict."""</span>
<span class="n">_ffi_api</span><span class="o">.</span><span class="n">DictClear</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">update</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">other</span><span class="p">:</span> <span class="n">Mapping</span><span class="p">[</span><span class="n">K</span><span class="p">,</span> <span class="n">V</span><span class="p">])</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span> <span class="c1"># type: ignore[override]</span>
-<span class="w"> </span><span class="sd">"""Update the dict from a mapping."""</span>
+<span class="w"> </span><span class="sd">"""Update the dict from a mapping."""</span>
<span class="k">for</span> <span class="n">k</span><span class="p">,</span> <span class="n">v</span> <span class="ow">in</span> <span class="n">other</span><span class="o">.</span><span class="n">items</span><span class="p">():</span>
<span class="bp">self</span><span class="p">[</span><span class="n">k</span><span class="p">]</span> <span class="o">=</span> <span class="n">v</span>
</pre></div>
@@ -1683,7 +1580,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1731,8 +1628,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/cpp/extension.html b/ffi/_modules/tvm_ffi/cpp/extension.html
index 20362704b11..4e9f66025de 100644
--- a/ffi/_modules/tvm_ffi/cpp/extension.html
+++ b/ffi/_modules/tvm_ffi/cpp/extension.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../../_static/downloads/require.min.js"></script>
- <script src="../../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/cpp/extension';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../../genindex.html" />
<link rel="search" title="Search" href="../../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../../_static/searchtools.js"></script>
+ <script src="../../../_static/language_data.js"></script>
+ <script src="../../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,18 +777,18 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""Build and load C++/CUDA sources into a tvm_ffi Module using Ninja."""</span>
+<span class="sd">"""Build and load C++/CUDA sources into a tvm_ffi Module using Ninja."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -911,21 +808,21 @@ window.runMermaid = runMermaid;</script>
<span class="kn">from</span><span class="w"> </span><span class="nn">tvm_ffi.module</span><span class="w"> </span><span class="kn">import</span> <span class="n">Module</span><span class="p">,</span> <span class="n">load_module</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">tvm_ffi.utils</span><span class="w"> </span><span class="kn">import</span> <span class="n">FileLock</span>
-<span class="n">IS_WINDOWS</span> <span class="o">=</span> <span class="n">sys</span><span class="o">.</span><span class="n">platform</span> <span class="o">==</span> <span class="s2">"win32"</span>
-<span class="n">BACKEND_STR</span> <span class="o">=</span> <span class="n">Literal</span><span class="p">[</span><span class="s2">"cuda"</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">]</span>
+<span class="n">IS_WINDOWS</span> <span class="o">=</span> <span class="n">sys</span><span class="o">.</span><span class="n">platform</span> <span class="o">==</span> <span class="s2">"win32"</span>
+<span class="n">BACKEND_STR</span> <span class="o">=</span> <span class="n">Literal</span><span class="p">[</span><span class="s2">"cuda"</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">]</span>
<span class="n">logger</span> <span class="o">=</span> <span class="n">logging</span><span class="o">.</span><span class="n">getLogger</span><span class="p">(</span><span class="vm">__name__</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_find_compiler</span><span class="p">(</span><span class="n">name</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span> <span class="o">*</span><span class="n">home_vars</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Locate a compiler on ``PATH`` or under an env-specified toolkit home's ``bin``."""</span>
+<span class="w"> </span><span class="sd">"""Locate a compiler on ``PATH`` or under an env-specified toolkit home's ``bin``."""</span>
<span class="n">found</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="n">name</span><span class="p">)</span>
<span class="k">if</span> <span class="n">found</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="k">return</span> <span class="n">found</span>
<span class="k">for</span> <span class="n">var</span> <span class="ow">in</span> <span class="n">home_vars</span><span class="p">:</span>
<span class="n">home</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="n">var</span><span class="p">)</span>
<span class="k">if</span> <span class="n">home</span><span class="p">:</span>
- <span class="n">found</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="n">name</span><span class="p">,</span> <span class="n">path</span><span class="o">=</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">home</span><span class="p">)</span> <span class="o">/</span> <span class="s2">"bin"</span> [...]
+ <span class="n">found</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="n">name</span><span class="p">,</span> <span class="n">path</span><span class="o">=</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">home</span><span class="p">)</span> <span class="o">/</span> <span class="s2">"bin"</span><span clas [...]
<span class="k">if</span> <span class="n">found</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="k">return</span> <span class="n">found</span>
<span class="k">return</span> <span class="kc">None</span>
@@ -933,7 +830,7 @@ window.runMermaid = runMermaid;</script>
<span class="nd">@functools</span><span class="o">.</span><span class="n">lru_cache</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_detect_gpu_backend</span><span class="p">()</span> <span class="o">-></span> <span class="n">BACKEND_STR</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Auto-detect whether to use CUDA or HIP (ROCm), defaulting to CUDA.</span>
+<span class="w"> </span><span class="sd">"""Auto-detect whether to use CUDA or HIP (ROCm), defaulting to CUDA.</span>
<span class="sd"> Resolution order: the ``TVM_FFI_GPU_BACKEND`` override, then PyTorch build</span>
<span class="sd"> signals (``torch.version.hip`` / ``torch.version.cuda``), then a discoverable</span>
@@ -942,32 +839,32 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> toolkit homes, so an env-specified CUDA install is honored even when ``nvcc``</span>
<span class="sd"> is off ``PATH``. The mere existence of ``/opt/rocm`` is not treated as</span>
<span class="sd"> evidence of HIP.</span>
-<span class="sd"> """</span>
- <span class="n">backend</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_GPU_BACKEND"</span><span class="p">,</span> <span class="s2">""</span><span class="p">)</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
- <span class="k">if</span> <span class="n">backend</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"cuda"</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">):</span>
+<span class="sd"> """</span>
+ <span class="n">backend</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_GPU_BACKEND"</span><span class="p">,</span> <span class="s2">""</span><span class="p">)</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
+ <span class="k">if</span> <span class="n">backend</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"cuda"</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">backend</span> <span class="c1"># type: ignore[return-value]</span>
<span class="k">try</span><span class="p">:</span>
<span class="kn">import</span><span class="w"> </span><span class="nn">torch</span> <span class="c1"># noqa: PLC0415</span>
- <span class="n">version</span> <span class="o">=</span> <span class="nb">getattr</span><span class="p">(</span><span class="n">torch</span><span class="p">,</span> <span class="s2">"version"</span><span class="p">,</span> <span class="kc">None</span><span class="p">)</span>
- <span class="k">if</span> <span class="nb">getattr</span><span class="p">(</span><span class="n">version</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">,</span> <span class="kc">None</span><span class="p">):</span>
- <span class="k">return</span> <span class="s2">"hip"</span>
- <span class="k">if</span> <span class="nb">getattr</span><span class="p">(</span><span class="n">version</span><span class="p">,</span> <span class="s2">"cuda"</span><span class="p">,</span> <span class="kc">None</span><span class="p">):</span>
- <span class="k">return</span> <span class="s2">"cuda"</span>
+ <span class="n">version</span> <span class="o">=</span> <span class="nb">getattr</span><span class="p">(</span><span class="n">torch</span><span class="p">,</span> <span class="s2">"version"</span><span class="p">,</span> <span class="kc">None</span><span class="p">)</span>
+ <span class="k">if</span> <span class="nb">getattr</span><span class="p">(</span><span class="n">version</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">,</span> <span class="kc">None</span><span class="p">):</span>
+ <span class="k">return</span> <span class="s2">"hip"</span>
+ <span class="k">if</span> <span class="nb">getattr</span><span class="p">(</span><span class="n">version</span><span class="p">,</span> <span class="s2">"cuda"</span><span class="p">,</span> <span class="kc">None</span><span class="p">):</span>
+ <span class="k">return</span> <span class="s2">"cuda"</span>
<span class="k">except</span> <span class="ne">Exception</span><span class="p">:</span>
<span class="k">pass</span>
- <span class="k">if</span> <span class="n">_find_compiler</span><span class="p">(</span><span class="s2">"nvcc"</span><span class="p">,</span> <span class="s2">"CUDA_HOME"</span><span class="p">,</span> <span class="s2">"CUDA_PATH"</span><span class="p">)</span> <span class="ow">is</span> <span class="kc">None</span> <span class="ow">and</span> <span class="p">(</span>
- <span class="n">_find_compiler</span><span class="p">(</span><span class="s2">"hipcc"</span><span class="p">,</span> <span class="s2">"ROCM_HOME"</span><span class="p">,</span> <span class="s2">"ROCM_PATH"</span><span class="p">)</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span>
+ <span class="k">if</span> <span class="n">_find_compiler</span><span class="p">(</span><span class="s2">"nvcc"</span><span class="p">,</span> <span class="s2">"CUDA_HOME"</span><span class="p">,</span> <span class="s2">"CUDA_PATH"</span><span class="p">)</span> <span class="ow">is</span> <span class="kc">None</span> <span class="ow">and</span> <span class="p">(</span>
+ <span class="n">_find_compiler</span><span class="p">(</span><span class="s2">"hipcc"</span><span class="p">,</span> <span class="s2">"ROCM_HOME"</span><span class="p">,</span> <span class="s2">"ROCM_PATH"</span><span class="p">)</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span>
<span class="p">):</span>
- <span class="k">return</span> <span class="s2">"hip"</span>
- <span class="k">return</span> <span class="s2">"cuda"</span>
+ <span class="k">return</span> <span class="s2">"hip"</span>
+ <span class="k">return</span> <span class="s2">"cuda"</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_resolve_gpu_backend</span><span class="p">(</span><span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span><span class="p">)</span> <span class="o">-></span> <span class="n">BACKEND_STR</span><span class="p">:</span>
<span class="k">if</span> <span class="n">backend</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
- <span class="k">if</span> <span class="n">backend</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"cuda"</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">):</span>
+ <span class="k">if</span> <span class="n">backend</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"cuda"</span><span class="p">,</span> <span class="s2">"hip"</span><span class="p">):</span>
<span class="k">return</span> <span class="n">backend</span> <span class="c1"># type: ignore[return-value]</span>
- <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Invalid backend: </span><span class="si">{</span><span class="n">backend</span><span class="si">}</span><span class="s2">. Supported backends are 'cuda' and 'hip'."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Invalid backend: </span><span class="si">{</span><span class="n">backend</span><span class="si">}</span><span class="s2">. Supported backends are 'cuda' and 'hip'."</span><span class="p">)</span>
<span class="k">return</span> <span class="n">_detect_gpu_backend</span><span class="p">()</span>
@@ -983,30 +880,30 @@ window.runMermaid = runMermaid;</script>
<span class="n">extra_include_paths</span><span class="p">:</span> <span class="n">Sequence</span><span class="p">[</span><span class="nb">str</span><span class="p">],</span>
<span class="n">embed_cubin</span><span class="p">:</span> <span class="n">Mapping</span><span class="p">[</span><span class="nb">str</span><span class="p">,</span> <span class="nb">bytes</span><span class="p">]</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Generate a unique hash for the given sources and functions."""</span>
+<span class="w"> </span><span class="sd">"""Generate a unique hash for the given sources and functions."""</span>
<span class="n">m</span> <span class="o">=</span> <span class="n">hashlib</span><span class="o">.</span><span class="n">sha256</span><span class="p">()</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_hash</span><span class="p">(</span><span class="n">obj</span><span class="p">:</span> <span class="n">Any</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
<span class="k">if</span> <span class="n">obj</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
- <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"None"</span><span class="p">)</span>
+ <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"None"</span><span class="p">)</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">obj</span><span class="p">,</span> <span class="nb">str</span><span class="p">):</span>
- <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"str"</span><span class="p">)</span>
- <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="n">obj</span><span class="o">.</span><span class="n">encode</span><span class="p">(</span><span class="s2">"utf-8"</span><span class="p">))</span>
+ <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"str"</span><span class="p">)</span>
+ <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="n">obj</span><span class="o">.</span><span class="n">encode</span><span class="p">(</span><span class="s2">"utf-8"</span><span class="p">))</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">obj</span><span class="p">,</span> <span class="nb">bytes</span><span class="p">):</span>
- <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"bytes"</span><span class="p">)</span>
+ <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"bytes"</span><span class="p">)</span>
<span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="n">obj</span><span class="p">)</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">obj</span><span class="p">,</span> <span class="n">Mapping</span><span class="p">):</span>
- <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"Mapping"</span><span class="p">)</span>
+ <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"Mapping"</span><span class="p">)</span>
<span class="k">for</span> <span class="n">key</span> <span class="ow">in</span> <span class="nb">sorted</span><span class="p">(</span><span class="n">obj</span><span class="o">.</span><span class="n">keys</span><span class="p">()):</span>
<span class="n">item</span> <span class="o">=</span> <span class="n">obj</span><span class="p">[</span><span class="n">key</span><span class="p">]</span>
<span class="n">_hash</span><span class="p">(</span><span class="n">key</span><span class="p">)</span>
<span class="n">_hash</span><span class="p">(</span><span class="n">item</span><span class="p">)</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">obj</span><span class="p">,</span> <span class="n">Sequence</span><span class="p">):</span>
- <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"Sequence"</span><span class="p">)</span>
+ <span class="n">m</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="sa">b</span><span class="s2">"Sequence"</span><span class="p">)</span>
<span class="k">for</span> <span class="n">item</span> <span class="ow">in</span> <span class="n">obj</span><span class="p">:</span>
<span class="n">_hash</span><span class="p">(</span><span class="n">item</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Unsupported type: </span><span class="si">{</span><span class="nb">type</span><span class="p">(</span><span class="n">obj</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Unsupported type: </span><span class="si">{</span><span class="nb">type</span><span class="p">(</span><span class="n">obj</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="n">_hash</span><span class="p">(</span>
<span class="p">(</span>
@@ -1027,154 +924,154 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">_maybe_write</span><span class="p">(</span><span class="n">path</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span> <span class="n">content</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Write content to path if it does not already exist with the same content."""</span>
+<span class="w"> </span><span class="sd">"""Write content to path if it does not already exist with the same content."""</span>
<span class="n">p</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">path</span><span class="p">)</span>
<span class="k">if</span> <span class="n">p</span><span class="o">.</span><span class="n">exists</span><span class="p">():</span>
<span class="k">with</span> <span class="n">p</span><span class="o">.</span><span class="n">open</span><span class="p">()</span> <span class="k">as</span> <span class="n">f</span><span class="p">:</span>
<span class="n">existing_content</span> <span class="o">=</span> <span class="n">f</span><span class="o">.</span><span class="n">read</span><span class="p">()</span>
<span class="k">if</span> <span class="n">existing_content</span> <span class="o">==</span> <span class="n">content</span><span class="p">:</span>
<span class="k">return</span>
- <span class="k">with</span> <span class="n">p</span><span class="o">.</span><span class="n">open</span><span class="p">(</span><span class="s2">"w"</span><span class="p">)</span> <span class="k">as</span> <span class="n">f</span><span class="p">:</span>
+ <span class="k">with</span> <span class="n">p</span><span class="o">.</span><span class="n">open</span><span class="p">(</span><span class="s2">"w"</span><span class="p">)</span> <span class="k">as</span> <span class="n">f</span><span class="p">:</span>
<span class="n">f</span><span class="o">.</span><span class="n">write</span><span class="p">(</span><span class="n">content</span><span class="p">)</span>
<span class="nd">@functools</span><span class="o">.</span><span class="n">lru_cache</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_find_cuda_home</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find the CUDA install path."""</span>
+<span class="w"> </span><span class="sd">"""Find the CUDA install path."""</span>
<span class="c1"># Guess #1</span>
- <span class="n">cuda_home</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CUDA_HOME"</span><span class="p">)</span> <span class="ow">or</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CUDA [...]
+ <span class="n">cuda_home</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CUDA_HOME"</span><span class="p">)</span> <span class="ow">or</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CUDA_PATH"</span><s [...]
<span class="k">if</span> <span class="n">cuda_home</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
<span class="c1"># Guess #2</span>
- <span class="n">nvcc_path</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="s2">"nvcc"</span><span class="p">)</span>
+ <span class="n">nvcc_path</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="s2">"nvcc"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">nvcc_path</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="n">cuda_home</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">nvcc_path</span><span class="p">)</span><span class="o">.</span><span class="n">parent</span><span class="o">.</span><span class="n">parent</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="c1"># Guess #3</span>
<span class="k">if</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="n">cuda_root</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="s2">"C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA"</span><span class="p">)</span>
- <span class="n">cuda_homes</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span><span class="n">cuda_root</span><span class="o">.</span><span class="n">glob</span><span class="p">(</span><span class="s2">"v*.*"</span><span class="p">))</span>
+ <span class="n">cuda_root</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="s2">"C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA"</span><span class="p">)</span>
+ <span class="n">cuda_homes</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span><span class="n">cuda_root</span><span class="o">.</span><span class="n">glob</span><span class="p">(</span><span class="s2">"v*.*"</span><span class="p">))</span>
<span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">cuda_homes</span><span class="p">)</span> <span class="o">==</span> <span class="mi">0</span><span class="p">:</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"Could not find CUDA installation. Please set CUDA_HOME environment variable."</span>
+ <span class="s2">"Could not find CUDA installation. Please set CUDA_HOME environment variable."</span>
<span class="p">)</span>
<span class="n">cuda_home</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">cuda_homes</span><span class="p">[</span><span class="mi">0</span><span class="p">])</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">cuda_home</span> <span class="o">=</span> <span class="s2">"/usr/local/cuda"</span>
+ <span class="n">cuda_home</span> <span class="o">=</span> <span class="s2">"/usr/local/cuda"</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">Path</span><span class="p">(</span><span class="n">cuda_home</span><span class="p">)</span><span class="o">.</span><span class="n">exists</span><span class="p">():</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"Could not find CUDA installation. Please set CUDA_HOME environment variable."</span>
+ <span class="s2">"Could not find CUDA installation. Please set CUDA_HOME environment variable."</span>
<span class="p">)</span>
<span class="k">return</span> <span class="n">cuda_home</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_get_cuda_target</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Get the CUDA target architecture flag."""</span>
- <span class="k">if</span> <span class="s2">"TVM_FFI_CUDA_ARCH_LIST"</span> <span class="ow">in</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">:</span>
- <span class="n">arch_list</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">[</span><span class="s2">"TVM_FFI_CUDA_ARCH_LIST"</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">()</span> <span class="c1"># e.g., "8.9 9.0a"</span>
+<span class="w"> </span><span class="sd">"""Get the CUDA target architecture flag."""</span>
+ <span class="k">if</span> <span class="s2">"TVM_FFI_CUDA_ARCH_LIST"</span> <span class="ow">in</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">:</span>
+ <span class="n">arch_list</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">[</span><span class="s2">"TVM_FFI_CUDA_ARCH_LIST"</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">()</span> <span class="c1"># e.g., "8.9 9.0a"</span>
<span class="n">flags</span> <span class="o">=</span> <span class="p">[]</span>
<span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">arch_list</span><span class="p">:</span>
- <span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">arch</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"."</span><span class="p">))</span> <span class="o">!=</span> <span class="mi">2</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Invalid CUDA architecture: </span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">major</span><span class="p">,</span> <span class="n">minor</span> <span class="o">=</span> <span class="n">arch</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"."</span><span class="p">)</span>
- <span class="n">flags</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"-gencode=arch=compute_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">,code=sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span>< [...]
- <span class="k">return</span> <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">flags</span><span class="p">)</span>
+ <span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">arch</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"."</span><span class="p">))</span> <span class="o">!=</span> <span class="mi">2</span><span class="p">:</span>
+ <span class="k">raise</span> <span class="ne">ValueError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Invalid CUDA architecture: </span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">major</span><span class="p">,</span> <span class="n">minor</span> <span class="o">=</span> <span class="n">arch</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"."</span><span class="p">)</span>
+ <span class="n">flags</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"-gencode=arch=compute_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">,code=sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span [...]
+ <span class="k">return</span> <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">flags</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="k">try</span><span class="p">:</span>
<span class="n">status</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span>
- <span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="s2">"nvidia-smi"</span><span class="p">,</span> <span class="s2">"--query-gpu=compute_cap"</span><span class="p">,</span> <span class="s2">"--format=csv,noheader"</span><span class="p">],</span>
+ <span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="s2">"nvidia-smi"</span><span class="p">,</span> <span class="s2">"--query-gpu=compute_cap"</span><span class="p">,</span> <span class="s2">"--format=csv,noheader"</span><span class="p">],</span>
<span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
<span class="n">check</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
<span class="p">)</span>
- <span class="n">compute_cap</span> <span class="o">=</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="s2">"utf-8"</span><span class="p">)</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class=" [...]
- <span class="n">major</span><span class="p">,</span> <span class="n">minor</span> <span class="o">=</span> <span class="n">compute_cap</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"."</span><span class="p">)</span>
- <span class="k">return</span> <span class="sa">f</span><span class="s2">"-gencode=arch=compute_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">,code=sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">"</span>
+ <span class="n">compute_cap</span> <span class="o">=</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="s2">"utf-8"</span><span class="p">)</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><s [...]
+ <span class="n">major</span><span class="p">,</span> <span class="n">minor</span> <span class="o">=</span> <span class="n">compute_cap</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"."</span><span class="p">)</span>
+ <span class="k">return</span> <span class="sa">f</span><span class="s2">"-gencode=arch=compute_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">,code=sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">"</span>
<span class="k">except</span> <span class="ne">Exception</span><span class="p">:</span>
<span class="k">try</span><span class="p">:</span>
<span class="c1"># For old drivers, there is no compute_cap, but we can use the GPU name to determine the architecture.</span>
<span class="n">ampere_arch_map</span> <span class="o">=</span> <span class="p">{</span>
- <span class="s2">"A100"</span><span class="p">:</span> <span class="p">(</span><span class="s2">"8"</span><span class="p">,</span> <span class="s2">"0"</span><span class="p">),</span>
- <span class="s2">"A10"</span><span class="p">:</span> <span class="p">(</span><span class="s2">"8"</span><span class="p">,</span> <span class="s2">"6"</span><span class="p">),</span>
+ <span class="s2">"A100"</span><span class="p">:</span> <span class="p">(</span><span class="s2">"8"</span><span class="p">,</span> <span class="s2">"0"</span><span class="p">),</span>
+ <span class="s2">"A10"</span><span class="p">:</span> <span class="p">(</span><span class="s2">"8"</span><span class="p">,</span> <span class="s2">"6"</span><span class="p">),</span>
<span class="p">}</span>
<span class="n">status</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span>
- <span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="s2">"nvidia-smi"</span><span class="p">,</span> <span class="s2">"--query-gpu=name"</span><span class="p">,</span> <span class="s2">"--format=csv,noheader"</span><span class="p">],</span>
+ <span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="s2">"nvidia-smi"</span><span class="p">,</span> <span class="s2">"--query-gpu=name"</span><span class="p">,</span> <span class="s2">"--format=csv,noheader"</span><span class="p">],</span>
<span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
<span class="n">check</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
<span class="n">text</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
<span class="p">)</span>
- <span class="n">gpu_name</span> <span class="o">=</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">)[</span><span class="mi">0</span><span class="p">]</span>
+ <span class="n">gpu_name</span> <span class="o">=</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">)[</span><span class="mi">0</span><span class="p">]</span>
<span class="k">for</span> <span class="n">gpu_key</span><span class="p">,</span> <span class="p">(</span><span class="n">major</span><span class="p">,</span> <span class="n">minor</span><span class="p">)</span> <span class="ow">in</span> <span class="n">ampere_arch_map</span><span class="o">.</span><span class="n">items</span><span class="p">():</span>
<span class="k">if</span> <span class="n">gpu_key</span> <span class="ow">in</span> <span class="n">gpu_name</span><span class="p">:</span>
- <span class="k">return</span> <span class="sa">f</span><span class="s2">"-gencode=arch=compute_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">,code=sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">"</span>
+ <span class="k">return</span> <span class="sa">f</span><span class="s2">"-gencode=arch=compute_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">,code=sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">"</span>
<span class="k">except</span> <span class="p">(</span><span class="n">subprocess</span><span class="o">.</span><span class="n">CalledProcessError</span><span class="p">,</span> <span class="ne">FileNotFoundError</span><span class="p">):</span>
<span class="k">pass</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"Could not detect CUDA compute_cap automatically. Please set TVM_FFI_CUDA_ARCH_LIST environment variable."</span>
+ <span class="s2">"Could not detect CUDA compute_cap automatically. Please set TVM_FFI_CUDA_ARCH_LIST environment variable."</span>
<span class="p">)</span>
<span class="nd">@functools</span><span class="o">.</span><span class="n">lru_cache</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_find_rocm_home</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find the ROCm install path."""</span>
+<span class="w"> </span><span class="sd">"""Find the ROCm install path."""</span>
<span class="c1"># Guess #1: check environment variables</span>
- <span class="n">rocm_home</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"ROCM_HOME"</span><span class="p">)</span> <span class="ow">or</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"ROCM [...]
+ <span class="n">rocm_home</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"ROCM_HOME"</span><span class="p">)</span> <span class="ow">or</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"ROCM_PATH"</span><s [...]
<span class="k">if</span> <span class="n">rocm_home</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
- <span class="n">hipcc_path</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="s2">"hipcc"</span><span class="p">)</span>
+ <span class="n">hipcc_path</span> <span class="o">=</span> <span class="n">shutil</span><span class="o">.</span><span class="n">which</span><span class="p">(</span><span class="s2">"hipcc"</span><span class="p">)</span>
<span class="c1"># Guess #2: find hipcc in PATH and resolve ROCm home from it</span>
<span class="k">if</span> <span class="n">hipcc_path</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="n">rocm_home</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">hipcc_path</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">()</span><span class="o">.</span><span class="n">parent</span><span class="o">.</span><span class="n">parent</span><span class="p">)</span>
- <span class="k">if</span> <span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">"hip"</span><span class="p">:</span>
+ <span class="k">if</span> <span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span><span class="o">.</span><span class="n">name</span> <span class="o">==</span> <span class="s2">"hip"</span><span class="p">:</span>
<span class="n">rocm_home</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span><span class="o">.</span><span class="n">parent</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="c1"># Guess #3: use default installation path</span>
- <span class="n">rocm_home</span> <span class="o">=</span> <span class="s2">"/opt/rocm"</span>
+ <span class="n">rocm_home</span> <span class="o">=</span> <span class="s2">"/opt/rocm"</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span><span class="o">.</span><span class="n">exists</span><span class="p">():</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"Could not find ROCm installation. Please set ROCM_HOME environment variable."</span>
+ <span class="s2">"Could not find ROCm installation. Please set ROCM_HOME environment variable."</span>
<span class="p">)</span>
<span class="k">return</span> <span class="n">rocm_home</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_get_rocm_target</span><span class="p">()</span> <span class="o">-></span> <span class="nb">list</span><span class="p">[</span><span class="nb">str</span><span class="p">]:</span>
-<span class="w"> </span><span class="sd">"""Get the ROCm target architecture flags (--offload-arch=gfxXXXX)."""</span>
- <span class="k">if</span> <span class="s2">"TVM_FFI_ROCM_ARCH_LIST"</span> <span class="ow">in</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">:</span>
- <span class="n">arch_list</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">[</span><span class="s2">"TVM_FFI_ROCM_ARCH_LIST"</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">()</span> <span class="c1"># e.g., "gfx90a gfx942"</span>
- <span class="k">return</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"--offload-arch=</span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span> <span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">arch_list</span><span class="p">]</span>
+<span class="w"> </span><span class="sd">"""Get the ROCm target architecture flags (--offload-arch=gfxXXXX)."""</span>
+ <span class="k">if</span> <span class="s2">"TVM_FFI_ROCM_ARCH_LIST"</span> <span class="ow">in</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">:</span>
+ <span class="n">arch_list</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="p">[</span><span class="s2">"TVM_FFI_ROCM_ARCH_LIST"</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">()</span> <span class="c1"># e.g., "gfx90a gfx942"</span>
+ <span class="k">return</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"--offload-arch=</span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span> <span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">arch_list</span><span class="p">]</span>
<span class="c1"># Try rocm_agent_enumerator</span>
<span class="k">try</span><span class="p">:</span>
- <span class="n">agent_enum</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_rocm_home</span><span class="p">())</span> <span class="o">/</span> <span class="s2">"bin"</span> <span class="o">/</span> <span class="s2">"rocm_agent_enumerator"</span><span class="p">)</span>
+ <span class="n">agent_enum</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_rocm_home</span><span class="p">())</span> <span class="o">/</span> <span class="s2">"bin"</span> <span class="o">/</span> <span class="s2">"rocm_agent_enumerator"</span><span class="p">)</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">Path</span><span class="p">(</span><span class="n">agent_enum</span><span class="p">)</span><span class="o">.</span><span class="n">exists</span><span class="p">():</span>
- <span class="n">agent_enum</span> <span class="o">=</span> <span class="s2">"rocm_agent_enumerator"</span>
+ <span class="n">agent_enum</span> <span class="o">=</span> <span class="s2">"rocm_agent_enumerator"</span>
<span class="n">status</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="n">agent_enum</span><span class="p">],</span> <span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span> <span class="n">check</span><span class="o">=</span><span class="kc">True< [...]
<span class="n">archs</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span>
<span class="nb">dict</span><span class="o">.</span><span class="n">fromkeys</span><span class="p">(</span>
<span class="n">line</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span>
- <span class="k">for</span> <span class="n">line</span> <span class="ow">in</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">)</span>
- <span class="k">if</span> <span class="n">line</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="ow">and</span> <span class="n">line</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="o">!=</span> <span class="s2">"gfx000"</span>
+ <span class="k">for</span> <span class="n">line</span> <span class="ow">in</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">if</span> <span class="n">line</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="ow">and</span> <span class="n">line</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="o">!=</span> <span class="s2">"gfx000"</span>
<span class="p">)</span>
<span class="p">)</span>
<span class="k">if</span> <span class="n">archs</span><span class="p">:</span>
- <span class="k">return</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"--offload-arch=</span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span> <span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">archs</span><span class="p">]</span>
+ <span class="k">return</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"--offload-arch=</span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span> <span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">archs</span><span class="p">]</span>
<span class="k">except</span> <span class="p">(</span><span class="n">subprocess</span><span class="o">.</span><span class="n">CalledProcessError</span><span class="p">,</span> <span class="ne">FileNotFoundError</span><span class="p">):</span>
<span class="k">pass</span>
<span class="c1"># Try rocminfo</span>
<span class="k">try</span><span class="p">:</span>
- <span class="n">status</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="s2">"rocminfo"</span><span class="p">],</span> <span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span> <span class="n">check</span><span class="o">=</span><span class [...]
+ <span class="n">status</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">args</span><span class="o">=</span><span class="p">[</span><span class="s2">"rocminfo"</span><span class="p">],</span> <span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span> <span class="n">check</span><span class="o">=</span><span class="kc">True [...]
<span class="n">archs</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span>
<span class="nb">dict</span><span class="o">.</span><span class="n">fromkeys</span><span class="p">(</span>
- <span class="n">line</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">":"</span><span class="p">)[</span><span class="o">-</span><span class="mi">1</span><span class="p">]</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span>
- <span class="k">for</span> <span class="n">line</span> <span class="ow">in</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">)</span>
- <span class="k">if</span> <span class="s2">"Name:"</span> <span class="ow">in</span> <span class="n">line</span>
- <span class="ow">and</span> <span class="s2">"gfx"</span> <span class="ow">in</span> <span class="n">line</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
- <span class="ow">and</span> <span class="n">line</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">":"</span><span class="p">)[</span><span class="o">-</span><span class="mi">1</span><span class="p">]</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="o">!=</span> <span class="s2">"gfx000"</span>
+ <span class="n">line</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">":"</span><span class="p">)[</span><span class="o">-</span><span class="mi">1</span><span class="p">]</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span>
+ <span class="k">for</span> <span class="n">line</span> <span class="ow">in</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">if</span> <span class="s2">"Name:"</span> <span class="ow">in</span> <span class="n">line</span>
+ <span class="ow">and</span> <span class="s2">"gfx"</span> <span class="ow">in</span> <span class="n">line</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
+ <span class="ow">and</span> <span class="n">line</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">":"</span><span class="p">)[</span><span class="o">-</span><span class="mi">1</span><span class="p">]</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="o">!=</span> <span class="s2">"gfx000"</span>
<span class="p">)</span>
<span class="p">)</span>
<span class="k">if</span> <span class="n">archs</span><span class="p">:</span>
- <span class="k">return</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"--offload-arch=</span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span> <span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">archs</span><span class="p">]</span>
+ <span class="k">return</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"--offload-arch=</span><span class="si">{</span><span class="n">arch</span><span class="si">}</span><span class="s2">"</span> <span class="k">for</span> <span class="n">arch</span> <span class="ow">in</span> <span class="n">archs</span><span class="p">]</span>
<span class="k">except</span> <span class="p">(</span><span class="n">subprocess</span><span class="o">.</span><span class="n">CalledProcessError</span><span class="p">,</span> <span class="ne">FileNotFoundError</span><span class="p">):</span>
<span class="k">pass</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"Could not detect ROCm GPU architecture automatically. "</span>
- <span class="s2">"Please set TVM_FFI_ROCM_ARCH_LIST environment variable (e.g. 'gfx90a gfx942')."</span>
+ <span class="s2">"Could not detect ROCm GPU architecture automatically. "</span>
+ <span class="s2">"Please set TVM_FFI_ROCM_ARCH_LIST environment variable (e.g. 'gfx90a gfx942')."</span>
<span class="p">)</span>
@@ -1183,29 +1080,29 @@ window.runMermaid = runMermaid;</script>
<span class="n">cwd</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="n">os</span><span class="o">.</span><span class="n">PathLike</span><span class="p">[</span><span class="nb">str</span><span class="p">],</span>
<span class="n">capture_output</span><span class="p">:</span> <span class="nb">bool</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">subprocess</span><span class="o">.</span><span class="n">CompletedProcess</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Locates the Developer Command Prompt and runs a command within its environment."""</span>
+<span class="w"> </span><span class="sd">"""Locates the Developer Command Prompt and runs a command within its environment."""</span>
<span class="k">try</span><span class="p">:</span>
<span class="c1"># Path to vswhere.exe</span>
<span class="n">vswhere_path</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span>
- <span class="n">Path</span><span class="p">(</span><span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"ProgramFiles(x86)"</span><span class="p">,</span> <span class="s2">"C:</span><span class="se">\\</span><span class="s2">Program Files (x86)"</span><span class="p">))</span>
- <span class="o">/</span> <span class="s2">"Microsoft Visual Studio"</span>
- <span class="o">/</span> <span class="s2">"Installer"</span>
- <span class="o">/</span> <span class="s2">"vswhere.exe"</span>
+ <span class="n">Path</span><span class="p">(</span><span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"ProgramFiles(x86)"</span><span class="p">,</span> <span class="s2">"C:</span><span class="se">\\</span><span class="s2">Program Files (x86)"</span><span class="p">))</span>
+ <span class="o">/</span> <span class="s2">"Microsoft Visual Studio"</span>
+ <span class="o">/</span> <span class="s2">"Installer"</span>
+ <span class="o">/</span> <span class="s2">"vswhere.exe"</span>
<span class="p">)</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">Path</span><span class="p">(</span><span class="n">vswhere_path</span><span class="p">)</span><span class="o">.</span><span class="n">exists</span><span class="p">():</span>
- <span class="k">raise</span> <span class="ne">FileNotFoundError</span><span class="p">(</span><span class="s2">"vswhere.exe not found."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">FileNotFoundError</span><span class="p">(</span><span class="s2">"vswhere.exe not found."</span><span class="p">)</span>
<span class="c1"># Find the Visual Studio installation path</span>
<span class="n">vs_install_path</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span>
<span class="p">[</span>
<span class="n">vswhere_path</span><span class="p">,</span>
- <span class="s2">"-latest"</span><span class="p">,</span>
- <span class="s2">"-prerelease"</span><span class="p">,</span>
- <span class="s2">"-products"</span><span class="p">,</span>
- <span class="s2">"*"</span><span class="p">,</span>
- <span class="s2">"-property"</span><span class="p">,</span>
- <span class="s2">"installationPath"</span><span class="p">,</span>
+ <span class="s2">"-latest"</span><span class="p">,</span>
+ <span class="s2">"-prerelease"</span><span class="p">,</span>
+ <span class="s2">"-products"</span><span class="p">,</span>
+ <span class="s2">"*"</span><span class="p">,</span>
+ <span class="s2">"-property"</span><span class="p">,</span>
+ <span class="s2">"installationPath"</span><span class="p">,</span>
<span class="p">],</span>
<span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
<span class="n">text</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span>
@@ -1213,19 +1110,19 @@ window.runMermaid = runMermaid;</script>
<span class="p">)</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">vs_install_path</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">FileNotFoundError</span><span class="p">(</span><span class="s2">"No Visual Studio installation found."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">FileNotFoundError</span><span class="p">(</span><span class="s2">"No Visual Studio installation found."</span><span class="p">)</span>
<span class="c1"># Construct the path to the VsDevCmd.bat file</span>
- <span class="n">vsdevcmd_path</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">vs_install_path</span><span class="p">)</span> <span class="o">/</span> <span class="s2">"Common7"</span> <span class="o">/</span> <span class="s2">"Tools"</span> <span class="o">/</span> <span class="s2">"VsDevCmd.bat"</span><span class="p">)</span>
+ <span class="n">vsdevcmd_path</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">vs_install_path</span><span class="p">)</span> <span class="o">/</span> <span class="s2">"Common7"</span> <span class="o">/</span> <span class="s2">"Tools"</span> <span class="o">/</span> <span class="s2">"VsDevCmd.bat"</span><span class="p">)</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">Path</span><span class="p">(</span><span class="n">vsdevcmd_path</span><span class="p">)</span><span class="o">.</span><span class="n">exists</span><span class="p">():</span>
- <span class="k">raise</span> <span class="ne">FileNotFoundError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"VsDevCmd.bat not found at: </span><span class="si">{</span><span class="n">vsdevcmd_path</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">FileNotFoundError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"VsDevCmd.bat not found at: </span><span class="si">{</span><span class="n">vsdevcmd_path</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="c1"># Use cmd.exe to run the batch file and then your command.</span>
<span class="c1"># The /k flag keeps the command prompt open after the batch file runs.</span>
- <span class="c1"># The "&" symbol chains the commands.</span>
- <span class="n">cmd_command</span> <span class="o">=</span> <span class="s1">'"</span><span class="si">{vsdevcmd_path}</span><span class="s1">" -arch=x64 & </span><span class="si">{command}</span><span class="s1">'</span><span class="o">.</span><span class="n">format</span><span class="p">(</span>
- <span class="n">vsdevcmd_path</span><span class="o">=</span><span class="n">vsdevcmd_path</span><span class="p">,</span> <span class="n">command</span><span class="o">=</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">args</span><span class="p">)</span>
+ <span class="c1"># The "&" symbol chains the commands.</span>
+ <span class="n">cmd_command</span> <span class="o">=</span> <span class="s1">'"</span><span class="si">{vsdevcmd_path}</span><span class="s1">" -arch=x64 & </span><span class="si">{command}</span><span class="s1">'</span><span class="o">.</span><span class="n">format</span><span class="p">(</span>
+ <span class="n">vsdevcmd_path</span><span class="o">=</span><span class="n">vsdevcmd_path</span><span class="p">,</span> <span class="n">command</span><span class="o">=</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">args</span><span class="p">)</span>
<span class="p">)</span>
<span class="c1"># Execute the command in a new shell</span>
@@ -1235,8 +1132,8 @@ window.runMermaid = runMermaid;</script>
<span class="k">except</span> <span class="p">(</span><span class="ne">FileNotFoundError</span><span class="p">,</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">CalledProcessError</span><span class="p">)</span> <span class="k">as</span> <span class="n">e</span><span class="p">:</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"Failed to run the following command in MSVC developer environment: </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span>
- <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">args</span><span class="p">)</span>
+ <span class="s2">"Failed to run the following command in MSVC developer environment: </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span>
+ <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">args</span><span class="p">)</span>
<span class="p">)</span>
<span class="p">)</span> <span class="kn">from</span><span class="w"> </span><span class="nn">e</span>
@@ -1252,73 +1149,73 @@ window.runMermaid = runMermaid;</script>
<span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="n">output</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Generate the content of build.ninja for building the module."""</span>
+<span class="w"> </span><span class="sd">"""Generate the content of build.ninja for building the module."""</span>
<span class="c1"># Determine output format from extension</span>
<span class="k">if</span> <span class="n">output</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="n">out_ext</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">output</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
- <span class="n">object_mode</span> <span class="o">=</span> <span class="n">out_ext</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">".o"</span><span class="p">,</span> <span class="s2">".obj"</span><span class="p">)</span>
+ <span class="n">object_mode</span> <span class="o">=</span> <span class="n">out_ext</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">".o"</span><span class="p">,</span> <span class="s2">".obj"</span><span class="p">)</span>
<span class="n">output_name</span> <span class="o">=</span> <span class="n">output</span>
<span class="k">else</span><span class="p">:</span>
<span class="n">object_mode</span> <span class="o">=</span> <span class="kc">False</span>
- <span class="n">output_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}{</span><span class="s1">'.dll'</span><span class="w"> </span><span class="k">if</span><span class="w"> </span><span class="n">IS_WINDOWS</span><span class="w"> </span><span class="k">else</span><span class="w"> </span><span class="s1">'.so'</span><span class="si">}</span><span clas [...]
- <span class="n">has_cuda_sources</span> <span class="o">=</span> <span class="nb">any</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">s</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="o">==</span> <span class="s2">".cu"</span> <span class="k">for</span> <span class="n">s</span> <span class="ow">in</spa [...]
- <span class="n">with_hip</span> <span class="o">=</span> <span class="n">backend</span> <span class="o">==</span> <span class="s2">"hip"</span>
- <span class="n">with_cuda</span> <span class="o">=</span> <span class="n">backend</span> <span class="o">==</span> <span class="s2">"cuda"</span>
+ <span class="n">output_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}{</span><span class="s1">'.dll'</span><span class="w"> </span><span class="k">if</span><span class="w"> </span><span class="n">IS_WINDOWS</span><span class="w"> </span><span class="k">else</span><span class="w"> </span><span class="s1">'.so'</span><span class="si">}</span><span class="s2">"</span>
+ <span class="n">has_cuda_sources</span> <span class="o">=</span> <span class="nb">any</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">s</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="o">==</span> <span class="s2">".cu"</span> <span class="k">for</span> <span class="n">s</span> <span class="ow">in</span> <span c [...]
+ <span class="n">with_hip</span> <span class="o">=</span> <span class="n">backend</span> <span class="o">==</span> <span class="s2">"hip"</span>
+ <span class="n">with_cuda</span> <span class="o">=</span> <span class="n">backend</span> <span class="o">==</span> <span class="s2">"cuda"</span>
<span class="n">with_backend</span> <span class="o">=</span> <span class="n">with_hip</span> <span class="ow">or</span> <span class="n">with_cuda</span> <span class="ow">or</span> <span class="n">has_cuda_sources</span>
<span class="k">if</span> <span class="n">has_cuda_sources</span> <span class="ow">and</span> <span class="ow">not</span> <span class="p">(</span><span class="n">with_hip</span> <span class="ow">or</span> <span class="n">with_cuda</span><span class="p">):</span>
<span class="c1"># Auto-detect backend from available GPU</span>
<span class="n">detected</span> <span class="o">=</span> <span class="n">_resolve_gpu_backend</span><span class="p">(</span><span class="kc">None</span><span class="p">)</span>
- <span class="n">with_hip</span> <span class="o">=</span> <span class="n">detected</span> <span class="o">==</span> <span class="s2">"hip"</span>
- <span class="n">with_cuda</span> <span class="o">=</span> <span class="n">detected</span> <span class="o">==</span> <span class="s2">"cuda"</span>
+ <span class="n">with_hip</span> <span class="o">=</span> <span class="n">detected</span> <span class="o">==</span> <span class="s2">"hip"</span>
+ <span class="n">with_cuda</span> <span class="o">=</span> <span class="n">detected</span> <span class="o">==</span> <span class="s2">"cuda"</span>
<span class="n">default_include_paths</span> <span class="o">=</span> <span class="p">[</span><span class="n">find_include_path</span><span class="p">(),</span> <span class="n">find_dlpack_include_path</span><span class="p">()]</span>
<span class="n">tvm_ffi_lib</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">find_libtvm_ffi</span><span class="p">())</span>
<span class="n">tvm_ffi_lib_path</span> <span class="o">=</span> <span class="nb">str</span><span class="p">(</span><span class="n">tvm_ffi_lib</span><span class="o">.</span><span class="n">parent</span><span class="p">)</span>
<span class="n">tvm_ffi_lib_name</span> <span class="o">=</span> <span class="n">tvm_ffi_lib</span><span class="o">.</span><span class="n">stem</span>
<span class="k">if</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="n">default_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"/O2"</span><span class="p">,</span> <span class="s2">"/MD"</span><span class="p">]</span>
- <span class="n">default_cxxflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"/std:c++17"</span><span class="p">,</span> <span class="s2">"/MD"</span><span class="p">,</span> <span class="s2">"/EHsc"</span><span class="p">]</span>
+ <span class="n">default_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"/O2"</span><span class="p">,</span> <span class="s2">"/MD"</span><span class="p">]</span>
+ <span class="n">default_cxxflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"/std:c++17"</span><span class="p">,</span> <span class="s2">"/MD"</span><span class="p">,</span> <span class="s2">"/EHsc"</span><span class="p">]</span>
<span class="n">_win_warnings</span> <span class="o">=</span> <span class="p">[</span>
- <span class="s2">"/wd4819"</span><span class="p">,</span>
- <span class="s2">"/wd4251"</span><span class="p">,</span>
- <span class="s2">"/wd4244"</span><span class="p">,</span>
- <span class="s2">"/wd4267"</span><span class="p">,</span>
- <span class="s2">"/wd4275"</span><span class="p">,</span>
- <span class="s2">"/wd4018"</span><span class="p">,</span>
- <span class="s2">"/wd4190"</span><span class="p">,</span>
- <span class="s2">"/wd4624"</span><span class="p">,</span>
- <span class="s2">"/wd4067"</span><span class="p">,</span>
- <span class="s2">"/wd4068"</span><span class="p">,</span>
+ <span class="s2">"/wd4819"</span><span class="p">,</span>
+ <span class="s2">"/wd4251"</span><span class="p">,</span>
+ <span class="s2">"/wd4244"</span><span class="p">,</span>
+ <span class="s2">"/wd4267"</span><span class="p">,</span>
+ <span class="s2">"/wd4275"</span><span class="p">,</span>
+ <span class="s2">"/wd4018"</span><span class="p">,</span>
+ <span class="s2">"/wd4190"</span><span class="p">,</span>
+ <span class="s2">"/wd4624"</span><span class="p">,</span>
+ <span class="s2">"/wd4067"</span><span class="p">,</span>
+ <span class="s2">"/wd4068"</span><span class="p">,</span>
<span class="p">]</span>
<span class="n">default_cflags</span> <span class="o">+=</span> <span class="n">_win_warnings</span>
<span class="n">default_cxxflags</span> <span class="o">+=</span> <span class="n">_win_warnings</span>
- <span class="n">default_cuda_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-Xcompiler"</span><span class="p">,</span> <span class="s2">"/std:c++17"</span><span class="p">,</span> <span class="s2">"/O2"</span><span class="p">]</span>
+ <span class="n">default_cuda_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-Xcompiler"</span><span class="p">,</span> <span class="s2">"/std:c++17"</span><span class="p">,</span> <span class="s2">"/O2"</span><span class="p">]</span>
<span class="n">default_ldflags</span> <span class="o">=</span> <span class="p">[</span>
- <span class="s2">"/DLL"</span><span class="p">,</span>
- <span class="sa">f</span><span class="s2">"/LIBPATH:</span><span class="si">{</span><span class="n">tvm_ffi_lib_path</span><span class="si">}</span><span class="s2">"</span><span class="p">,</span>
- <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">tvm_ffi_lib_name</span><span class="si">}</span><span class="s2">.lib"</span><span class="p">,</span>
+ <span class="s2">"/DLL"</span><span class="p">,</span>
+ <span class="sa">f</span><span class="s2">"/LIBPATH:</span><span class="si">{</span><span class="n">tvm_ffi_lib_path</span><span class="si">}</span><span class="s2">"</span><span class="p">,</span>
+ <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">tvm_ffi_lib_name</span><span class="si">}</span><span class="s2">.lib"</span><span class="p">,</span>
<span class="p">]</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">default_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-fPIC"</span><span class="p">,</span> <span class="s2">"-O2"</span><span class="p">]</span>
- <span class="n">default_cxxflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-std=c++17"</span><span class="p">,</span> <span class="s2">"-fPIC"</span><span class="p">,</span> <span class="s2">"-O2"</span><span class="p">]</span>
- <span class="n">default_cuda_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-std=c++17"</span><span class="p">,</span> <span class="s2">"-O2"</span><span class="p">]</span>
- <span class="n">default_ldflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-shared"</span><span class="p">,</span> <span class="sa">f</span><span class="s2">"-L</span><span class="si">{</span><span class="n">tvm_ffi_lib_path</span><span class="si">}</span><span class="s2">"</span><span class="p">,</span> <span class="s2">"-ltvm_ffi"</span><span class="p">]</span>
+ <span class="n">default_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-fPIC"</span><span class="p">,</span> <span class="s2">"-O2"</span><span class="p">]</span>
+ <span class="n">default_cxxflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-std=c++17"</span><span class="p">,</span> <span class="s2">"-fPIC"</span><span class="p">,</span> <span class="s2">"-O2"</span><span class="p">]</span>
+ <span class="n">default_cuda_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-std=c++17"</span><span class="p">,</span> <span class="s2">"-O2"</span><span class="p">]</span>
+ <span class="n">default_ldflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-shared"</span><span class="p">,</span> <span class="sa">f</span><span class="s2">"-L</span><span class="si">{</span><span class="n">tvm_ffi_lib_path</span><span class="si">}</span><span class="s2">"</span><span class="p">,</span> <span class="s2">"-ltvm_ffi"</span><span class="p">]</span>
<span class="k">if</span> <span class="n">with_hip</span><span class="p">:</span>
<span class="n">rocm_home</span> <span class="o">=</span> <span class="n">_find_rocm_home</span><span class="p">()</span>
- <span class="n">default_cuda_cflags</span> <span class="o">+=</span> <span class="p">[</span><span class="s2">"-fPIC"</span><span class="p">,</span> <span class="s2">"-D__HIP_PLATFORM_AMD__=1"</span><span class="p">,</span> <span class="s2">"-fno-gpu-rdc"</span><span class="p">]</span>
+ <span class="n">default_cuda_cflags</span> <span class="o">+=</span> <span class="p">[</span><span class="s2">"-fPIC"</span><span class="p">,</span> <span class="s2">"-D__HIP_PLATFORM_AMD__=1"</span><span class="p">,</span> <span class="s2">"-fno-gpu-rdc"</span><span class="p">]</span>
<span class="n">default_cuda_cflags</span> <span class="o">+=</span> <span class="n">_get_rocm_target</span><span class="p">()</span>
- <span class="n">default_include_paths</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">))</span>
+ <span class="n">default_include_paths</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">))</span>
<span class="n">default_ldflags</span> <span class="o">+=</span> <span class="p">[</span>
- <span class="sa">f</span><span class="s2">"-L</span><span class="si">{</span><span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span><span class="w"> </span><span class="o">/</span><span class="w"> </span><span class="s1">'lib'</span><span class="si">!s}</span><span class="s2">"</span><span class="p">,</span>
- <span class="s2">"-lamdhip64"</span><span class="p">,</span>
+ <span class="sa">f</span><span class="s2">"-L</span><span class="si">{</span><span class="n">Path</span><span class="p">(</span><span class="n">rocm_home</span><span class="p">)</span><span class="w"> </span><span class="o">/</span><span class="w"> </span><span class="s1">'lib'</span><span class="si">!s}</span><span class="s2">"</span><span class="p">,</span>
+ <span class="s2">"-lamdhip64"</span><span class="p">,</span>
<span class="p">]</span>
<span class="k">if</span> <span class="n">with_cuda</span><span class="p">:</span>
- <span class="n">default_cuda_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-Xcompiler"</span><span class="p">,</span> <span class="s2">"-fPIC"</span><span class="p">,</span> <span class="o">*</span><span class="n">default_cuda_cflags</span><span class="p">]</span>
+ <span class="n">default_cuda_cflags</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"-Xcompiler"</span><span class="p">,</span> <span class="s2">"-fPIC"</span><span class="p">,</span> <span class="o">*</span><span class="n">default_cuda_cflags</span><span class="p">]</span>
<span class="n">default_cuda_cflags</span> <span class="o">+=</span> <span class="p">[</span><span class="n">_get_cuda_target</span><span class="p">()]</span>
<span class="n">default_ldflags</span> <span class="o">+=</span> <span class="p">[</span>
- <span class="s2">"-L</span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_cuda_home</span><span class="p">())</span> <span class="o">/</span> <span class="s2">"lib64"</span><span class="p">)),</span>
- <span class="s2">"-lcudart"</span><span class="p">,</span> <span class="c1"># cuda runtime library</span>
+ <span class="s2">"-L</span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_cuda_home</span><span class="p">())</span> <span class="o">/</span> <span class="s2">"lib64"</span><span class="p">)),</span>
+ <span class="s2">"-lcudart"</span><span class="p">,</span> <span class="c1"># cuda runtime library</span>
<span class="p">]</span>
<span class="n">extra_cflags_list</span> <span class="o">=</span> <span class="p">[</span><span class="n">flag</span><span class="o">.</span><span class="n">strip</span><span class="p">()</span> <span class="k">for</span> <span class="n">flag</span> <span class="ow">in</span> <span class="n">extra_cflags</span><span class="p">]</span>
@@ -1332,117 +1229,117 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># append include paths</span>
<span class="k">for</span> <span class="n">path</span> <span class="ow">in</span> <span class="n">include_paths</span><span class="p">:</span>
- <span class="n">inc</span> <span class="o">=</span> <span class="s2">"-I</span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="n">path</span><span class="o">.</span><span class="n">replace</span><span class="p">(</span><span class="s2">":"</span><span class="p">,</span> <span class="s2">"$:"</span><span class="p">))</span>
+ <span class="n">inc</span> <span class="o">=</span> <span class="s2">"-I</span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="n">path</span><span class="o">.</span><span class="n">replace</span><span class="p">(</span><span class="s2">":"</span><span class="p">,</span> <span class="s2">"$:"</span><span class="p">))</span>
<span class="n">cflags</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">inc</span><span class="p">)</span>
<span class="n">cxxflags</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">inc</span><span class="p">)</span>
<span class="n">cuda_cflags</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">inc</span><span class="p">)</span>
<span class="c1"># Classify sources by extension to determine which rules are needed</span>
- <span class="n">with_c</span> <span class="o">=</span> <span class="nb">any</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">s</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="o">==</span> <span class="s2">".c"</span> <span class="k">for</span> <span class="n">s</span> <span class="ow">in</span> <span cl [...]
+ <span class="n">with_c</span> <span class="o">=</span> <span class="nb">any</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">s</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="o">==</span> <span class="s2">".c"</span> <span class="k">for</span> <span class="n">s</span> <span class="ow">in</span> <span class="n">so [...]
<span class="n">ninja</span><span class="p">:</span> <span class="nb">list</span><span class="p">[</span><span class="nb">str</span><span class="p">]</span> <span class="o">=</span> <span class="p">[]</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"ninja_required_version = 1.3"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cxx = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CXX"</span><span class="p">,</spa [...]
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cxxflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cxxflags</span><span class="p">)))</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"ninja_required_version = 1.3"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cxx = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CXX"</span><span class="p">,</span> <span class="s2"> [...]
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cxxflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cxxflags</span><span class="p">)))</span>
<span class="k">if</span> <span class="n">with_c</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cc = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CC"</span><span class="p">,</s [...]
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cflags</span><span class="p">)))</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cc = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"CC"</span><span class="p">,</span> <span class="s2 [...]
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cflags</span><span class="p">)))</span>
<span class="k">if</span> <span class="n">with_backend</span><span class="p">:</span>
<span class="k">if</span> <span class="n">with_hip</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"nvcc = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_rocm_home</span><span class="p">())</span> <span class="o">/</span> <span class=" [...]
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"nvcc = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_rocm_home</span><span class="p">())</span> <span class="o">/</span> <span class="s2">"bin"< [...]
<span class="k">if</span> <span class="n">with_cuda</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"nvcc = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_cuda_home</span><span class="p">())</span> <span class="o">/</span> <span class=" [...]
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cuda_cflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cuda_cflags</span><span class="p">)))</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"ldflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">ldflags</span><span class="p">)))</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"nvcc = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">_find_cuda_home</span><span class="p">())</span> <span class="o">/</span> <span class="s2">"bin"< [...]
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"cuda_cflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cuda_cflags</span><span class="p">)))</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"ldflags = </span><span class="si">{}</span><span class="s2">"</span><span class="o">.</span><span class="n">format</span><span class="p">(</span><span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">ldflags</span><span class="p">)))</span>
<span class="c1"># rules</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule compile"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule compile"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx /showIncludes $cxxflags -c $in /Fo$out"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = msvc"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx /showIncludes $cxxflags -c $in /Fo$out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = msvc"</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" depfile = $out.d"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = gcc"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx -MMD -MF $out.d $cxxflags -c $in -o $out"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" depfile = $out.d"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = gcc"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx -MMD -MF $out.d $cxxflags -c $in -o $out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="k">if</span> <span class="n">with_c</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule c_compile"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule c_compile"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cc /showIncludes $cflags -c $in /Fo$out"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = msvc"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cc /showIncludes $cflags -c $in /Fo$out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = msvc"</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" depfile = $out.d"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = gcc"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cc -MMD -MF $out.d $cflags -c $in -o $out"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" depfile = $out.d"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = gcc"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cc -MMD -MF $out.d $cflags -c $in -o $out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="k">if</span> <span class="n">with_backend</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule compile_cuda"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" depfile = $out.d"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = gcc"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule compile_cuda"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" depfile = $out.d"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" deps = gcc"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">with_hip</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $nvcc $cuda_cflags -c $in -o $out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $nvcc $cuda_cflags -c $in -o $out"</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span>
- <span class="s2">" command = $nvcc --generate-dependencies-with-compile --dependency-output $out.d $cuda_cflags -c $in -o $out"</span>
+ <span class="s2">" command = $nvcc --generate-dependencies-with-compile --dependency-output $out.d $cuda_cflags -c $in -o $out"</span>
<span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="c1"># Add rules for object merging and cubin embedding (Unix only)</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule merge_objects"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = ld -r -o $out $in"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule merge_objects"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = ld -r -o $out $in"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="c1"># Copy a single object straight through, bypassing `ld -r`. A partial</span>
<span class="c1"># link of one input is a pointless round-trip, and some system linkers</span>
<span class="c1"># (e.g. GNU ld 2.45) assign a nonzero sh_addr to zero-size, non-SHF_ALLOC</span>
<span class="c1"># sections like `.note.GNU-stack` during it. JITLink then materializes a</span>
<span class="c1"># zero-length block at that address, which collides with a real block in</span>
- <span class="c1"># EHFrameEdgeFixer's block map and aborts C++ (eh_frame-bearing) loads.</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule copy_object"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = cp -f $in $out"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="c1"># EHFrameEdgeFixer's block map and aborts C++ (eh_frame-bearing) loads.</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule copy_object"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = cp -f $in $out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="k">if</span> <span class="n">embed_cubin</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule embed_cubin"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule embed_cubin"</span><span class="p">)</span>
<span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span>
- <span class="sa">f</span><span class="s2">" command = </span><span class="si">{</span><span class="n">sys</span><span class="o">.</span><span class="n">executable</span><span class="si">}</span><span class="s2"> -m tvm_ffi.utils.embed_cubin --output-obj $out --input-obj $in --cubin $cubin --name $name"</span>
+ <span class="sa">f</span><span class="s2">" command = </span><span class="si">{</span><span class="n">sys</span><span class="o">.</span><span class="n">executable</span><span class="si">}</span><span class="s2"> -m tvm_ffi.utils.embed_cubin --output-obj $out --input-obj $in --cubin $cubin --name $name"</span>
<span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">object_mode</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule link"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">"rule link"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx $in /link $ldflags /out:$out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx $in /link $ldflags /out:$out"</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx $in $ldflags -o $out"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">" command = $cxx $in $ldflags -o $out"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="c1"># build targets — dispatch by file extension</span>
<span class="n">obj_files</span><span class="p">:</span> <span class="nb">list</span><span class="p">[</span><span class="nb">str</span><span class="p">]</span> <span class="o">=</span> <span class="p">[]</span>
<span class="n">c_idx</span> <span class="o">=</span> <span class="n">cpp_idx</span> <span class="o">=</span> <span class="n">cuda_idx</span> <span class="o">=</span> <span class="mi">0</span>
<span class="k">for</span> <span class="n">src</span> <span class="ow">in</span> <span class="nb">sorted</span><span class="p">(</span><span class="n">sources</span><span class="p">):</span>
<span class="n">ext</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">src</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
- <span class="n">escaped</span> <span class="o">=</span> <span class="n">src</span><span class="o">.</span><span class="n">replace</span><span class="p">(</span><span class="s2">":"</span><span class="p">,</span> <span class="s2">"$:"</span><span class="p">)</span>
- <span class="k">if</span> <span class="n">ext</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">".o"</span><span class="p">,</span> <span class="s2">".obj"</span><span class="p">):</span>
+ <span class="n">escaped</span> <span class="o">=</span> <span class="n">src</span><span class="o">.</span><span class="n">replace</span><span class="p">(</span><span class="s2">":"</span><span class="p">,</span> <span class="s2">"$:"</span><span class="p">)</span>
+ <span class="k">if</span> <span class="n">ext</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">".o"</span><span class="p">,</span> <span class="s2">".obj"</span><span class="p">):</span>
<span class="c1"># Pre-compiled object file: pass directly to linker</span>
<span class="n">obj_files</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">escaped</span><span class="p">)</span>
- <span class="k">elif</span> <span class="n">ext</span> <span class="o">==</span> <span class="s2">".c"</span><span class="p">:</span>
- <span class="n">obj_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"c_</span><span class="si">{</span><span class="n">c_idx</span><span class="si">}</span><span class="s2">.o"</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">obj_name</span><span class="si">}</span><span class="s2">: c_compile </span><span class="si">{</span><span class="n">escaped</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">elif</span> <span class="n">ext</span> <span class="o">==</span> <span class="s2">".c"</span><span class="p">:</span>
+ <span class="n">obj_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"c_</span><span class="si">{</span><span class="n">c_idx</span><span class="si">}</span><span class="s2">.o"</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">obj_name</span><span class="si">}</span><span class="s2">: c_compile </span><span class="si">{</span><span class="n">escaped</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="n">obj_files</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">obj_name</span><span class="p">)</span>
<span class="n">c_idx</span> <span class="o">+=</span> <span class="mi">1</span>
- <span class="k">elif</span> <span class="n">ext</span> <span class="o">==</span> <span class="s2">".cu"</span><span class="p">:</span>
- <span class="n">obj_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"cuda_</span><span class="si">{</span><span class="n">cuda_idx</span><span class="si">}</span><span class="s2">.o"</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">obj_name</span><span class="si">}</span><span class="s2">: compile_cuda </span><span class="si">{</span><span class="n">escaped</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">elif</span> <span class="n">ext</span> <span class="o">==</span> <span class="s2">".cu"</span><span class="p">:</span>
+ <span class="n">obj_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"cuda_</span><span class="si">{</span><span class="n">cuda_idx</span><span class="si">}</span><span class="s2">.o"</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">obj_name</span><span class="si">}</span><span class="s2">: compile_cuda </span><span class="si">{</span><span class="n">escaped</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="n">obj_files</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">obj_name</span><span class="p">)</span>
<span class="n">cuda_idx</span> <span class="o">+=</span> <span class="mi">1</span>
<span class="k">else</span><span class="p">:</span>
<span class="c1"># .cc, .cpp, .cxx — default to C++ compilation</span>
- <span class="n">obj_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"cpp_</span><span class="si">{</span><span class="n">cpp_idx</span><span class="si">}</span><span class="s2">.o"</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">obj_name</span><span class="si">}</span><span class="s2">: compile </span><span class="si">{</span><span class="n">escaped</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">obj_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"cpp_</span><span class="si">{</span><span class="n">cpp_idx</span><span class="si">}</span><span class="s2">.o"</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">obj_name</span><span class="si">}</span><span class="s2">: compile </span><span class="si">{</span><span class="n">escaped</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="n">obj_files</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="n">obj_name</span><span class="p">)</span>
<span class="n">cpp_idx</span> <span class="o">+=</span> <span class="mi">1</span>
@@ -1451,15 +1348,15 @@ window.runMermaid = runMermaid;</script>
<span class="k">if</span> <span class="ow">not</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
<span class="c1"># A single object needs no partial link — copy it through to avoid</span>
<span class="c1"># `ld -r` rewriting section addresses (see the copy_object rule).</span>
- <span class="n">rule</span> <span class="o">=</span> <span class="s2">"copy_object"</span> <span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">obj_files</span><span class="p">)</span> <span class="o">==</span> <span class="mi">1</span> <span class="k">else</span> <span class="s2">"merge_objects"</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">: </span><span class="si">{</span><span class="n">rule</span><span class="si">}</span><span class="s2"> </span><span class="si">{</span><span class="s1">' '</span><span class="o">.</span><span class="n">jo [...]
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"default </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">rule</span> <span class="o">=</span> <span class="s2">"copy_object"</span> <span class="k">if</span> <span class="nb">len</span><span class="p">(</span><span class="n">obj_files</span><span class="p">)</span> <span class="o">==</span> <span class="mi">1</span> <span class="k">else</span> <span class="s2">"merge_objects"</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">: </span><span class="si">{</span><span class="n">rule</span><span class="si">}</span><span class="s2"> </span><span class="si">{</span><span class="s1">' '</span><span class="o">.</span><span class="n">join</span><spa [...]
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"default </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="c1"># Windows: no ld -r available; default to the first intermediate object</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"default </span><span class="si">{</span><span class="n">obj_files</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
- <span class="k">return</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">ninja</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"default </span><span class="si">{</span><span class="n">obj_files</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="k">return</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">ninja</span><span class="p">)</span>
<span class="c1"># For Unix systems with embed_cubin, use a 3-step process:</span>
<span class="c1"># 1. Merge all object files into a unified object file</span>
@@ -1467,108 +1364,108 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># 3. Link the final object file into a shared library</span>
<span class="k">if</span> <span class="ow">not</span> <span class="n">IS_WINDOWS</span> <span class="ow">and</span> <span class="n">embed_cubin</span><span class="p">:</span>
<span class="c1"># Step 1: Merge object files into unified.o</span>
- <span class="n">unified_obj</span> <span class="o">=</span> <span class="s2">"unified.o"</span>
- <span class="n">obj_files_str</span> <span class="o">=</span> <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">obj_files</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">unified_obj</span><span class="si">}</span><span class="s2">: merge_objects </span><span class="si">{</span><span class="n">obj_files_str</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">unified_obj</span> <span class="o">=</span> <span class="s2">"unified.o"</span>
+ <span class="n">obj_files_str</span> <span class="o">=</span> <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">obj_files</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">unified_obj</span><span class="si">}</span><span class="s2">: merge_objects </span><span class="si">{</span><span class="n">obj_files_str</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="c1"># Step 2: Chain embed_cubin operations for each cubin</span>
<span class="n">current_obj</span> <span class="o">=</span> <span class="n">unified_obj</span>
<span class="k">for</span> <span class="n">cubin_name</span> <span class="ow">in</span> <span class="nb">sorted</span><span class="p">(</span><span class="n">embed_cubin</span><span class="o">.</span><span class="n">keys</span><span class="p">()):</span>
<span class="c1"># Create next object file name</span>
- <span class="n">next_obj</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"unified_with_</span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">.o"</span>
- <span class="n">cubin_file</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">.cubin"</span>
+ <span class="n">next_obj</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"unified_with_</span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">.o"</span>
+ <span class="n">cubin_file</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">.cubin"</span>
<span class="c1"># Add ninja build rule</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">next_obj</span><span class="si">}</span><span class="s2">: embed_cubin </span><span class="si">{</span><span class="n">current_obj</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">" cubin = </span><span class="si">{</span><span class="n">cubin_file</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">" name = </span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">next_obj</span><span class="si">}</span><span class="s2">: embed_cubin </span><span class="si">{</span><span class="n">current_obj</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">" cubin = </span><span class="si">{</span><span class="n">cubin_file</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">" name = </span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="n">current_obj</span> <span class="o">=</span> <span class="n">next_obj</span>
<span class="c1"># Step 3: Link the final object file</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">: link </span><span class="si">{</span><span class="n">current_obj</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">: link </span><span class="si">{</span><span class="n">current_obj</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="c1"># Directly link object files (for Windows or no cubin embedding)</span>
- <span class="n">link_files_str</span> <span class="o">=</span> <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">obj_files</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">: link </span><span class="si">{</span><span class="n">link_files_str</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">link_files_str</span> <span class="o">=</span> <span class="s2">" "</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">obj_files</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"build </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">: link </span><span class="si">{</span><span class="n">link_files_str</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
<span class="c1"># default target</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"default </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
- <span class="k">return</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">ninja</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"default </span><span class="si">{</span><span class="n">output_name</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="n">ninja</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="k">return</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">ninja</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">build_ninja</span><span class="p">(</span><span class="n">build_dir</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Build the module in the given build directory using ninja."""</span>
- <span class="n">command</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"ninja"</span><span class="p">,</span> <span class="s2">"-v"</span><span class="p">]</span>
- <span class="n">num_workers</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"MAX_JOBS"</span><span class="p">,</span> <span class="kc">None</span><span class="p">)</span>
+<span class="w"> </span><span class="sd">"""Build the module in the given build directory using ninja."""</span>
+ <span class="n">command</span> <span class="o">=</span> <span class="p">[</span><span class="s2">"ninja"</span><span class="p">,</span> <span class="s2">"-v"</span><span class="p">]</span>
+ <span class="n">num_workers</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"MAX_JOBS"</span><span class="p">,</span> <span class="kc">None</span><span class="p">)</span>
<span class="k">if</span> <span class="n">num_workers</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
- <span class="n">command</span> <span class="o">+=</span> <span class="p">[</span><span class="s2">"-j"</span><span class="p">,</span> <span class="n">num_workers</span><span class="p">]</span>
+ <span class="n">command</span> <span class="o">+=</span> <span class="p">[</span><span class="s2">"-j"</span><span class="p">,</span> <span class="n">num_workers</span><span class="p">]</span>
<span class="k">if</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
<span class="n">status</span> <span class="o">=</span> <span class="n">_run_command_in_dev_prompt</span><span class="p">(</span><span class="n">args</span><span class="o">=</span><span class="n">command</span><span class="p">,</span> <span class="n">cwd</span><span class="o">=</span><span class="n">build_dir</span><span class="p">,</span> <span class="n">capture_output</span><span class="o">=</span><span class="kc">True</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
<span class="n">status</span> <span class="o">=</span> <span class="n">subprocess</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">check</span><span class="o">=</span><span class="kc">False</span><span class="p">,</span> <span class="n">args</span><span class="o">=</span><span class="n">command</span><span class="p">,</span> <span class="n">cwd</span><span class="o">=</span><span class="n">build_dir</span><span class="p">,</span> <sp [...]
- <span class="n">encoding</span> <span class="o">=</span> <span class="s2">"oem"</span> <span class="k">if</span> <span class="n">IS_WINDOWS</span> <span class="k">else</span> <span class="s2">"utf-8"</span>
+ <span class="n">encoding</span> <span class="o">=</span> <span class="s2">"oem"</span> <span class="k">if</span> <span class="n">IS_WINDOWS</span> <span class="k">else</span> <span class="s2">"utf-8"</span>
<span class="k">if</span> <span class="n">status</span><span class="o">.</span><span class="n">returncode</span> <span class="o">!=</span> <span class="mi">0</span><span class="p">:</span>
- <span class="n">msg</span> <span class="o">=</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"ninja exited with status </span><span class="si">{</span><span class="n">status</span><span class="o">.</span><span class="n">returncode</span><span class="si">}</span><span class="s2">"</span><span class="p">]</span>
+ <span class="n">msg</span> <span class="o">=</span> <span class="p">[</span><span class="sa">f</span><span class="s2">"ninja exited with status </span><span class="si">{</span><span class="n">status</span><span class="o">.</span><span class="n">returncode</span><span class="si">}</span><span class="s2">"</span><span class="p">]</span>
<span class="k">if</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="p">:</span>
- <span class="n">msg</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"stdout:</span><span class="se">\n</span><span class="si">{</span><span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">)</span><span class="si">}</span><span class="s2">"</ [...]
+ <span class="n">msg</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"stdout:</span><span class="se">\n</span><span class="si">{</span><span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span><span [...]
<span class="k">if</span> <span class="n">status</span><span class="o">.</span><span class="n">stderr</span><span class="p">:</span>
- <span class="n">msg</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"stderr:</span><span class="se">\n</span><span class="si">{</span><span class="n">status</span><span class="o">.</span><span class="n">stderr</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">)</span><span class="si">}</span><span class="s2">"</ [...]
+ <span class="n">msg</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"stderr:</span><span class="se">\n</span><span class="si">{</span><span class="n">status</span><span class="o">.</span><span class="n">stderr</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span><span [...]
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">msg</span><span class="p">))</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">msg</span><span class="p">))</span>
- <span class="n">LOG_BUILD</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_CPP_EXTENSION_LOG_BUILD"</span><span class="p">,</span> <span class="s2">"0"</span><span class="p">)</span>
- <span class="k">if</span> <span class="n">LOG_BUILD</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"1"</span><span class="p">,</span> <span class="s2">"stdout"</span><span class="p">):</span>
- <span class="n">logger</span><span class="o">.</span><span class="n">info</span><span class="p">(</span><span class="s2">"ninja build stdout:</span><span class="se">\n</span><span class="si">%s</span><span class="s2">"</span><span class="p">,</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">))</span>
- <span class="k">if</span> <span class="n">LOG_BUILD</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"1"</span><span class="p">,</span> <span class="s2">"stderr"</span><span class="p">):</span>
- <span class="n">logger</span><span class="o">.</span><span class="n">info</span><span class="p">(</span><span class="s2">"ninja build stderr:</span><span class="se">\n</span><span class="si">%s</span><span class="s2">"</span><span class="p">,</span> <span class="n">status</span><span class="o">.</span><span class="n">stderr</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">))</span>
+ <span class="n">LOG_BUILD</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_CPP_EXTENSION_LOG_BUILD"</span><span class="p">,</span> <span class="s2">"0"</span><span class="p">)</span>
+ <span class="k">if</span> <span class="n">LOG_BUILD</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"1"</span><span class="p">,</span> <span class="s2">"stdout"</span><span class="p">):</span>
+ <span class="n">logger</span><span class="o">.</span><span class="n">info</span><span class="p">(</span><span class="s2">"ninja build stdout:</span><span class="se">\n</span><span class="si">%s</span><span class="s2">"</span><span class="p">,</span> <span class="n">status</span><span class="o">.</span><span class="n">stdout</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">))</span>
+ <span class="k">if</span> <span class="n">LOG_BUILD</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">"1"</span><span class="p">,</span> <span class="s2">"stderr"</span><span class="p">):</span>
+ <span class="n">logger</span><span class="o">.</span><span class="n">info</span><span class="p">(</span><span class="s2">"ninja build stderr:</span><span class="se">\n</span><span class="si">%s</span><span class="s2">"</span><span class="p">,</span> <span class="n">status</span><span class="o">.</span><span class="n">stderr</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="n">encoding</span><span class="p">))</span>
<span class="c1"># Translation table for escaping C++ string literals</span>
<span class="n">_CPP_ESCAPE_TABLE</span> <span class="o">=</span> <span class="nb">str</span><span class="o">.</span><span class="n">maketrans</span><span class="p">(</span>
<span class="p">{</span>
- <span class="s2">"</span><span class="se">\\</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\\\</span><span class="s2">"</span><span class="p">,</span>
- <span class="s1">'"'</span><span class="p">:</span> <span class="s1">'</span><span class="se">\\</span><span class="s1">"'</span><span class="p">,</span>
- <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\</span><span class="s2">n"</span><span class="p">,</span>
- <span class="s2">"</span><span class="se">\r</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\</span><span class="s2">r"</span><span class="p">,</span>
- <span class="s2">"</span><span class="se">\t</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\</span><span class="s2">t"</span><span class="p">,</span>
+ <span class="s2">"</span><span class="se">\\</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\\\</span><span class="s2">"</span><span class="p">,</span>
+ <span class="s1">'"'</span><span class="p">:</span> <span class="s1">'</span><span class="se">\\</span><span class="s1">"'</span><span class="p">,</span>
+ <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\</span><span class="s2">n"</span><span class="p">,</span>
+ <span class="s2">"</span><span class="se">\r</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\</span><span class="s2">r"</span><span class="p">,</span>
+ <span class="s2">"</span><span class="se">\t</span><span class="s2">"</span><span class="p">:</span> <span class="s2">"</span><span class="se">\\</span><span class="s2">t"</span><span class="p">,</span>
<span class="p">}</span>
<span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_escape_cpp_string_literal</span><span class="p">(</span><span class="n">s</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Escape special characters for C++ string literals."""</span>
+<span class="w"> </span><span class="sd">"""Escape special characters for C++ string literals."""</span>
<span class="k">return</span> <span class="n">s</span><span class="o">.</span><span class="n">translate</span><span class="p">(</span><span class="n">_CPP_ESCAPE_TABLE</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_decorate_with_tvm_ffi</span><span class="p">(</span><span class="n">source</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span> <span class="n">functions</span><span class="p">:</span> <span class="n">Mapping</span><span class="p">[</span><span class="nb">str</span><span class="p">,</span> <span class="nb">str</span><span class="p">])</span> <span class="o">-></span> <span class="nb">str< [...]
-<span class="w"> </span><span class="sd">"""Decorate the given source code with TVM FFI export macros."""</span>
+<span class="w"> </span><span class="sd">"""Decorate the given source code with TVM FFI export macros."""</span>
<span class="n">sources</span> <span class="o">=</span> <span class="p">[</span>
- <span class="s2">"#include <tvm/ffi/container/tensor.h>"</span><span class="p">,</span>
- <span class="s2">"#include <tvm/ffi/dtype.h>"</span><span class="p">,</span>
- <span class="s2">"#include <tvm/ffi/error.h>"</span><span class="p">,</span>
- <span class="s2">"#include <tvm/ffi/extra/c_env_api.h>"</span><span class="p">,</span>
- <span class="s2">"#include <tvm/ffi/function.h>"</span><span class="p">,</span>
- <span class="s2">""</span><span class="p">,</span>
+ <span class="s2">"#include <tvm/ffi/container/tensor.h>"</span><span class="p">,</span>
+ <span class="s2">"#include <tvm/ffi/dtype.h>"</span><span class="p">,</span>
+ <span class="s2">"#include <tvm/ffi/error.h>"</span><span class="p">,</span>
+ <span class="s2">"#include <tvm/ffi/extra/c_env_api.h>"</span><span class="p">,</span>
+ <span class="s2">"#include <tvm/ffi/function.h>"</span><span class="p">,</span>
+ <span class="s2">""</span><span class="p">,</span>
<span class="n">source</span><span class="p">,</span>
<span class="p">]</span>
<span class="k">for</span> <span class="n">func_name</span><span class="p">,</span> <span class="n">func_doc</span> <span class="ow">in</span> <span class="n">functions</span><span class="o">.</span><span class="n">items</span><span class="p">():</span>
- <span class="n">sources</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"TVM_FFI_DLL_EXPORT_TYPED_FUNC(</span><span class="si">{</span><span class="n">func_name</span><span class="si">}</span><span class="s2">, </span><span class="si">{</span><span class="n">func_name</span><span class="si">}</span><span class="s2">);"</span><span class="p">)</span>
+ <span class="n">sources</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s2">"TVM_FFI_DLL_EXPORT_TYPED_FUNC(</span><span class="si">{</span><span class="n">func_name</span><span class="si">}</span><span class="s2">, </span><span class="si">{</span><span class="n">func_name</span><span class="si">}</span><span class="s2">);"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">func_doc</span><span class="p">:</span>
<span class="c1"># Escape the docstring for C++ string literal</span>
<span class="n">escaped_doc</span> <span class="o">=</span> <span class="n">_escape_cpp_string_literal</span><span class="p">(</span><span class="n">func_doc</span><span class="p">)</span>
- <span class="n">sources</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s1">'TVM_FFI_DLL_EXPORT_TYPED_FUNC_DOC(</span><span class="si">{</span><span class="n">func_name</span><span class="si">}</span><span class="s1">, "</span><span class="si">{</span><span class="n">escaped_doc</span><span class="si">}</span><span class="s1">");'</span><span class="p">)</span>
+ <span class="n">sources</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s1">'TVM_FFI_DLL_EXPORT_TYPED_FUNC_DOC(</span><span class="si">{</span><span class="n">func_name</span><span class="si">}</span><span class="s1">, "</span><span class="si">{</span><span class="n">escaped_doc</span><span class="si">}</span><span class="s1">");'</span><span class="p">)</span>
- <span class="n">sources</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
+ <span class="n">sources</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="s2">""</span><span class="p">)</span>
- <span class="k">return</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">sources</span><span class="p">)</span>
+ <span class="k">return</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">sources</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_str_seq2list</span><span class="p">(</span><span class="n">seq</span><span class="p">:</span> <span class="n">Sequence</span><span class="p">[</span><span class="nb">str</span><span class="p">]</span> <span class="o">|</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span><span class="p">)</span> <span class="o">-></span> <span class="nb">list</span><span class="p">[</span><span class [...]
@@ -1593,12 +1490,12 @@ window.runMermaid = runMermaid;</script>
<span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="n">output</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Real implementation of build function."""</span>
+<span class="w"> </span><span class="sd">"""Real implementation of build function."""</span>
<span class="c1"># need to resolve the path to make it unique</span>
<span class="n">source_path_list</span> <span class="o">=</span> <span class="p">[</span><span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">p</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span> <span class="k">for</span> <span class="n">p</span> <span class="ow">in</span> <span class="n">_str_seq2list</span><span class="p">(</span><span class="n">sources</span [...]
- <span class="k">assert</span> <span class="n">source_path_list</span><span class="p">,</span> <span class="s2">"sources must be provided."</span>
+ <span class="k">assert</span> <span class="n">source_path_list</span><span class="p">,</span> <span class="s2">"sources must be provided."</span>
- <span class="n">has_cuda</span> <span class="o">=</span> <span class="nb">any</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">p</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="o">==</span> <span class="s2">".cu"</span> <span class="k">for</span> <span class="n">p</span> <span class="ow">in</span> <span [...]
+ <span class="n">has_cuda</span> <span class="o">=</span> <span class="nb">any</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">p</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="o">==</span> <span class="s2">".cu"</span> <span class="k">for</span> <span class="n">p</span> <span class="ow">in</span> <span class="n" [...]
<span class="n">resolved_backend</span> <span class="o">=</span> <span class="n">_resolve_gpu_backend</span><span class="p">(</span><span class="n">backend</span><span class="p">)</span> <span class="k">if</span> <span class="n">has_cuda</span> <span class="k">else</span> <span class="kc">None</span>
<span class="n">extra_ldflags_list</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span><span class="n">extra_ldflags</span><span class="p">)</span> <span class="k">if</span> <span class="n">extra_ldflags</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span> <span class="k">else</span> <span class="p">[]</span>
<span class="n">extra_cflags_list</span> <span class="o">=</span> <span class="nb">list</span><span class="p">(</span><span class="n">extra_cflags</span><span class="p">)</span> <span class="k">if</span> <span class="n">extra_cflags</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span> <span class="k">else</span> <span class="p">[]</span>
@@ -1607,7 +1504,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">build_dir</span><span class="p">:</span> <span class="n">Path</span>
<span class="k">if</span> <span class="n">build_directory</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
- <span class="n">cache_dir</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_CACHE_DIR"</span><span class="p">,</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="s2">"~/.cache/tvm-ffi"</span><span class="p">)</span><span class="o">. [...]
+ <span class="n">cache_dir</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_CACHE_DIR"</span><span class="p">,</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="s2">"~/.cache/tvm-ffi"</span><span class="p">)</span><span class="o">.</span><span class=" [...]
<span class="n">source_hash</span><span class="p">:</span> <span class="nb">str</span> <span class="o">=</span> <span class="n">_hash_sources</span><span class="p">(</span>
<span class="kc">None</span><span class="p">,</span>
<span class="kc">None</span><span class="p">,</span>
@@ -1620,20 +1517,20 @@ window.runMermaid = runMermaid;</script>
<span class="n">extra_include_paths_list</span><span class="p">,</span>
<span class="n">embed_cubin</span><span class="p">,</span>
<span class="p">)</span>
- <span class="n">build_dir</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">cache_dir</span><span class="p">)</span><span class="o">.</span><span class="n">expanduser</span><span class="p">()</span> <span class="o">/</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}</span><span class="s2">_</span><span class="si">{</span><span class="n">source_hash</ [...]
+ <span class="n">build_dir</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">cache_dir</span><span class="p">)</span><span class="o">.</span><span class="n">expanduser</span><span class="p">()</span> <span class="o">/</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}</span><span class="s2">_</span><span class="si">{</span><span class="n">source_hash</span> [...]
<span class="k">else</span><span class="p">:</span>
<span class="n">build_dir</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">build_directory</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">()</span>
<span class="n">build_dir</span><span class="o">.</span><span class="n">mkdir</span><span class="p">(</span><span class="n">parents</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span> <span class="n">exist_ok</span><span class="o">=</span><span class="kc">True</span><span class="p">)</span>
<span class="c1"># CUBIN embedding is only supported on Unix systems</span>
<span class="k">if</span> <span class="n">embed_cubin</span> <span class="ow">and</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">NotImplementedError</span><span class="p">(</span><span class="s2">"CUBIN embedding is not yet supported on Windows"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">NotImplementedError</span><span class="p">(</span><span class="s2">"CUBIN embedding is not yet supported on Windows"</span><span class="p">)</span>
<span class="c1"># Write CUBIN files to build directory if needed (for Unix systems)</span>
<span class="c1"># These will be embedded using the embed_cubin utility during ninja build</span>
<span class="k">if</span> <span class="n">embed_cubin</span><span class="p">:</span>
<span class="k">for</span> <span class="n">cubin_name</span><span class="p">,</span> <span class="n">cubin_bytes</span> <span class="ow">in</span> <span class="n">embed_cubin</span><span class="o">.</span><span class="n">items</span><span class="p">():</span>
- <span class="n">cubin_path</span> <span class="o">=</span> <span class="n">build_dir</span> <span class="o">/</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">.cubin"</span>
+ <span class="n">cubin_path</span> <span class="o">=</span> <span class="n">build_dir</span> <span class="o">/</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">cubin_name</span><span class="si">}</span><span class="s2">.cubin"</span>
<span class="n">cubin_path</span><span class="o">.</span><span class="n">write_bytes</span><span class="p">(</span><span class="n">cubin_bytes</span><span class="p">)</span>
<span class="c1"># generate build.ninja</span>
@@ -1650,29 +1547,29 @@ window.runMermaid = runMermaid;</script>
<span class="p">)</span>
<span class="c1"># may not hold lock when build_directory is specified, prevent deadlock</span>
- <span class="k">with</span> <span class="n">FileLock</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"lock"</span><span class="p">))</span> <span class="k">if</span> <span class="n">need_lock</span> <span class="k">else</span> <span class="n">nullcontext</span><span class="p">():</span>
+ <span class="k">with</span> <span class="n">FileLock</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"lock"</span><span class="p">))</span> <span class="k">if</span> <span class="n">need_lock</span> <span class="k">else</span> <span class="n">nullcontext</span><span class="p">():</span>
<span class="c1"># write build.ninja if it does not already exist</span>
- <span class="n">_maybe_write</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"build.ninja"</span><span class="p">),</span> <span class="n">ninja_source</span><span class="p">)</span>
+ <span class="n">_maybe_write</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"build.ninja"</span><span class="p">),</span> <span class="n">ninja_source</span><span class="p">)</span>
<span class="c1"># build the module</span>
<span class="n">build_ninja</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span><span class="p">))</span>
<span class="c1"># Determine the output filename (mirrors _generate_ninja_build logic)</span>
<span class="k">if</span> <span class="n">output</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="n">out_ext</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">output</span><span class="p">)</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span>
- <span class="n">object_mode</span> <span class="o">=</span> <span class="n">out_ext</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">".o"</span><span class="p">,</span> <span class="s2">".obj"</span><span class="p">)</span>
+ <span class="n">object_mode</span> <span class="o">=</span> <span class="n">out_ext</span> <span class="ow">in</span> <span class="p">(</span><span class="s2">".o"</span><span class="p">,</span> <span class="s2">".obj"</span><span class="p">)</span>
<span class="n">output_name</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">output</span><span class="p">)</span><span class="o">.</span><span class="n">name</span>
<span class="k">else</span><span class="p">:</span>
<span class="n">object_mode</span> <span class="o">=</span> <span class="kc">False</span>
- <span class="n">output_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}{</span><span class="s1">'.dll'</span><span class="w"> </span><span class="k">if</span><span class="w"> </span><span class="n">IS_WINDOWS</span><span class="w"> </span><span class="k">else</span><span class="w"> </span><span class="s1">'.so'</span><span class="si">}</span><span [...]
+ <span class="n">output_name</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}{</span><span class="s1">'.dll'</span><span class="w"> </span><span class="k">if</span><span class="w"> </span><span class="n">IS_WINDOWS</span><span class="w"> </span><span class="k">else</span><span class="w"> </span><span class="s1">'.so'</span><span class="si">}</span><span class="s2">"</span>
<span class="k">if</span> <span class="n">object_mode</span> <span class="ow">and</span> <span class="n">IS_WINDOWS</span><span class="p">:</span>
<span class="c1"># Windows has no ld -r; the actual target is the first intermediate object.</span>
<span class="c1"># The name must match _generate_ninja_build: c_0.o / cpp_0.o / cuda_0.o.</span>
- <span class="n">first_ext</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="nb">sorted</span><span class="p">(</span><span class="n">source_path_list</span><span class="p">)[</span><span class="mi">0</span><span class="p">])</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="k">if</span> <span class="n">source_path_list</span> <span [...]
- <span class="k">if</span> <span class="n">first_ext</span> <span class="o">==</span> <span class="s2">".c"</span><span class="p">:</span>
- <span class="n">obj_name</span> <span class="o">=</span> <span class="s2">"c_0.o"</span>
- <span class="k">elif</span> <span class="n">first_ext</span> <span class="o">==</span> <span class="s2">".cu"</span><span class="p">:</span>
- <span class="n">obj_name</span> <span class="o">=</span> <span class="s2">"cuda_0.o"</span>
+ <span class="n">first_ext</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="nb">sorted</span><span class="p">(</span><span class="n">source_path_list</span><span class="p">)[</span><span class="mi">0</span><span class="p">])</span><span class="o">.</span><span class="n">suffix</span><span class="o">.</span><span class="n">lower</span><span class="p">()</span> <span class="k">if</span> <span class="n">source_path_list</span> <span [...]
+ <span class="k">if</span> <span class="n">first_ext</span> <span class="o">==</span> <span class="s2">".c"</span><span class="p">:</span>
+ <span class="n">obj_name</span> <span class="o">=</span> <span class="s2">"c_0.o"</span>
+ <span class="k">elif</span> <span class="n">first_ext</span> <span class="o">==</span> <span class="s2">".cu"</span><span class="p">:</span>
+ <span class="n">obj_name</span> <span class="o">=</span> <span class="s2">"cuda_0.o"</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">obj_name</span> <span class="o">=</span> <span class="s2">"cpp_0.o"</span>
+ <span class="n">obj_name</span> <span class="o">=</span> <span class="s2">"cpp_0.o"</span>
<span class="k">return</span> <span class="nb">str</span><span class="p">((</span><span class="n">build_dir</span> <span class="o">/</span> <span class="n">obj_name</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
<span class="k">return</span> <span class="nb">str</span><span class="p">((</span><span class="n">build_dir</span> <span class="o">/</span> <span class="n">output_name</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
@@ -1694,7 +1591,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="n">output</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Compile and build a C++/CUDA module from inline source code.</span>
+<span class="w"> </span><span class="sd">"""Compile and build a C++/CUDA module from inline source code.</span>
<span class="sd"> This function compiles the given C++ and/or CUDA source code into a shared library or object file.</span>
<span class="sd"> Both ``cpp_sources`` and ``cuda_sources`` are compiled to an object file. When ``output`` is</span>
@@ -1741,8 +1638,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The extra compiler flags for C++ compilation.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
-<span class="sd"> - On Windows: ['/std:c++17', '/O2']</span>
+<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
+<span class="sd"> - On Windows: ['/std:c++17', '/O2']</span>
<span class="sd"> extra_cuda_cflags</span>
<span class="sd"> The extra compiler flags for CUDA compilation.</span>
@@ -1751,8 +1648,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The extra linker flags.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-shared']</span>
-<span class="sd"> - On Windows: ['/DLL']</span>
+<span class="sd"> - On Linux/macOS: ['-shared']</span>
+<span class="sd"> - On Windows: ['/DLL']</span>
<span class="sd"> extra_include_paths</span>
<span class="sd"> The extra include paths.</span>
@@ -1770,13 +1667,13 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> See the `examples/cubin_launcher` directory for examples how to use cubin launcher to launch CUBIN kernels in TVM-FFI.</span>
<span class="sd"> backend</span>
-<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
+<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
<span class="sd"> If not specified, the backend will be automatically determined based on the available GPU and the provided source code.</span>
<span class="sd"> output</span>
<span class="sd"> Output filename that determines the build type from its extension. When ``None``</span>
<span class="sd"> (the default), builds a shared library (``.so`` on Unix, ``.dll`` on Windows).</span>
-<span class="sd"> Use an object-file extension (e.g., ``"hello.o"``) to skip linking and produce</span>
+<span class="sd"> Use an object-file extension (e.g., ``"hello.o"``) to skip linking and produce</span>
<span class="sd"> a relocatable object file. The file is placed in the build directory.</span>
<span class="sd"> Returns</span>
@@ -1794,26 +1691,26 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi.cpp</span>
<span class="sd"> # define the cpp source code</span>
-<span class="sd"> cpp_source = '''</span>
+<span class="sd"> cpp_source = '''</span>
<span class="sd"> void add_one_cpu(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {</span>
<span class="sd"> // implementation of a library function</span>
-<span class="sd"> TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
<span class="sd"> DLDataType f32_dtype{kDLFloat, 32, 1};</span>
-<span class="sd"> TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
-<span class="sd"> TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
-<span class="sd"> TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
-<span class="sd"> TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
+<span class="sd"> TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
<span class="sd"> for (int i = 0; i < x.size(0); ++i) {</span>
<span class="sd"> static_cast<float*>(y.data_ptr())[i] = static_cast<float*>(x.data_ptr())[i] + 1;</span>
<span class="sd"> }</span>
<span class="sd"> }</span>
-<span class="sd"> '''</span>
+<span class="sd"> '''</span>
<span class="sd"> # compile the cpp source code and load the module</span>
<span class="sd"> lib_path: str = tvm_ffi.cpp.build_inline(</span>
-<span class="sd"> name="hello",</span>
+<span class="sd"> name="hello",</span>
<span class="sd"> cpp_sources=cpp_source,</span>
-<span class="sd"> functions="add_one_cpu",</span>
+<span class="sd"> functions="add_one_cpu",</span>
<span class="sd"> )</span>
<span class="sd"> # load the module</span>
@@ -1825,14 +1722,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> mod.add_one_cpu(x, y)</span>
<span class="sd"> torch.testing.assert_close(x + 1, y)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="n">cpp_source_list</span> <span class="o">=</span> <span class="n">_str_seq2list</span><span class="p">(</span><span class="n">cpp_sources</span><span class="p">)</span>
- <span class="n">cpp_source</span> <span class="o">=</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cpp_source_list</span><span class="p">)</span>
+ <span class="n">cpp_source</span> <span class="o">=</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cpp_source_list</span><span class="p">)</span>
<span class="n">with_cpp</span> <span class="o">=</span> <span class="nb">bool</span><span class="p">(</span><span class="n">cpp_source_list</span><span class="p">)</span>
<span class="k">del</span> <span class="n">cpp_source_list</span>
<span class="n">cuda_source_list</span> <span class="o">=</span> <span class="n">_str_seq2list</span><span class="p">(</span><span class="n">cuda_sources</span><span class="p">)</span>
- <span class="n">cuda_source</span> <span class="o">=</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cuda_source_list</span><span class="p">)</span>
+ <span class="n">cuda_source</span> <span class="o">=</span> <span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="n">cuda_source_list</span><span class="p">)</span>
<span class="n">with_backend</span> <span class="o">=</span> <span class="nb">bool</span><span class="p">(</span><span class="n">cuda_source_list</span><span class="p">)</span>
<span class="k">del</span> <span class="n">cuda_source_list</span>
@@ -1845,11 +1742,11 @@ window.runMermaid = runMermaid;</script>
<span class="k">if</span> <span class="n">functions</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
<span class="n">function_map</span><span class="p">:</span> <span class="nb">dict</span><span class="p">[</span><span class="nb">str</span><span class="p">,</span> <span class="nb">str</span><span class="p">]</span> <span class="o">=</span> <span class="p">{}</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">functions</span><span class="p">,</span> <span class="nb">str</span><span class="p">):</span>
- <span class="n">function_map</span> <span class="o">=</span> <span class="p">{</span><span class="n">functions</span><span class="p">:</span> <span class="s2">""</span><span class="p">}</span>
+ <span class="n">function_map</span> <span class="o">=</span> <span class="p">{</span><span class="n">functions</span><span class="p">:</span> <span class="s2">""</span><span class="p">}</span>
<span class="k">elif</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">functions</span><span class="p">,</span> <span class="n">Mapping</span><span class="p">):</span>
<span class="n">function_map</span> <span class="o">=</span> <span class="nb">dict</span><span class="p">(</span><span class="n">functions</span><span class="p">)</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">function_map</span> <span class="o">=</span> <span class="p">{</span><span class="n">name</span><span class="p">:</span> <span class="s2">""</span> <span class="k">for</span> <span class="n">name</span> <span class="ow">in</span> <span class="n">functions</span><span class="p">}</span>
+ <span class="n">function_map</span> <span class="o">=</span> <span class="p">{</span><span class="n">name</span><span class="p">:</span> <span class="s2">""</span> <span class="k">for</span> <span class="n">name</span> <span class="ow">in</span> <span class="n">functions</span><span class="p">}</span>
<span class="k">if</span> <span class="n">with_cpp</span><span class="p">:</span>
<span class="n">cpp_source</span> <span class="o">=</span> <span class="n">_decorate_with_tvm_ffi</span><span class="p">(</span><span class="n">cpp_source</span><span class="p">,</span> <span class="n">function_map</span><span class="p">)</span>
@@ -1860,7 +1757,7 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># determine the cache dir for the built module</span>
<span class="n">build_dir</span><span class="p">:</span> <span class="n">Path</span>
<span class="k">if</span> <span class="n">build_directory</span> <span class="ow">is</span> <span class="kc">None</span><span class="p">:</span>
- <span class="n">cache_dir</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_CACHE_DIR"</span><span class="p">,</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="s2">"~/.cache/tvm-ffi"</span><span class="p">)</span><span class="o">. [...]
+ <span class="n">cache_dir</span> <span class="o">=</span> <span class="n">os</span><span class="o">.</span><span class="n">environ</span><span class="o">.</span><span class="n">get</span><span class="p">(</span><span class="s2">"TVM_FFI_CACHE_DIR"</span><span class="p">,</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="s2">"~/.cache/tvm-ffi"</span><span class="p">)</span><span class="o">.</span><span class=" [...]
<span class="n">source_hash</span><span class="p">:</span> <span class="nb">str</span> <span class="o">=</span> <span class="n">_hash_sources</span><span class="p">(</span>
<span class="n">cpp_source</span><span class="p">,</span>
<span class="n">cuda_source</span><span class="p">,</span>
@@ -1873,15 +1770,15 @@ window.runMermaid = runMermaid;</script>
<span class="n">extra_include_paths_list</span><span class="p">,</span>
<span class="n">embed_cubin</span><span class="p">,</span>
<span class="p">)</span>
- <span class="n">build_dir</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">cache_dir</span><span class="p">)</span><span class="o">.</span><span class="n">expanduser</span><span class="p">()</span> <span class="o">/</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}</span><span class="s2">_</span><span class="si">{</span><span class="n">source_hash</ [...]
+ <span class="n">build_dir</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">cache_dir</span><span class="p">)</span><span class="o">.</span><span class="n">expanduser</span><span class="p">()</span> <span class="o">/</span> <span class="sa">f</span><span class="s2">"</span><span class="si">{</span><span class="n">name</span><span class="si">}</span><span class="s2">_</span><span class="si">{</span><span class="n">source_hash</span> [...]
<span class="k">else</span><span class="p">:</span>
<span class="n">build_dir</span> <span class="o">=</span> <span class="n">Path</span><span class="p">(</span><span class="n">build_directory</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">()</span>
<span class="n">build_dir</span><span class="o">.</span><span class="n">mkdir</span><span class="p">(</span><span class="n">parents</span><span class="o">=</span><span class="kc">True</span><span class="p">,</span> <span class="n">exist_ok</span><span class="o">=</span><span class="kc">True</span><span class="p">)</span>
- <span class="n">cpp_file</span> <span class="o">=</span> <span class="nb">str</span><span class="p">((</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"main.cpp"</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
- <span class="n">cuda_file</span> <span class="o">=</span> <span class="nb">str</span><span class="p">((</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"cuda.cu"</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
+ <span class="n">cpp_file</span> <span class="o">=</span> <span class="nb">str</span><span class="p">((</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"main.cpp"</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
+ <span class="n">cuda_file</span> <span class="o">=</span> <span class="nb">str</span><span class="p">((</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"cuda.cu"</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
- <span class="k">with</span> <span class="n">FileLock</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"lock"</span><span class="p">)):</span>
+ <span class="k">with</span> <span class="n">FileLock</span><span class="p">(</span><span class="nb">str</span><span class="p">(</span><span class="n">build_dir</span> <span class="o">/</span> <span class="s2">"lock"</span><span class="p">)):</span>
<span class="c1"># write source files if they do not already exist</span>
<span class="n">_maybe_write</span><span class="p">(</span><span class="n">cpp_file</span><span class="p">,</span> <span class="n">cpp_source</span><span class="p">)</span>
<span class="k">if</span> <span class="n">with_backend</span><span class="p">:</span>
@@ -1925,10 +1822,10 @@ window.runMermaid = runMermaid;</script>
<span class="n">keep_module_alive</span><span class="p">:</span> <span class="nb">bool</span> <span class="o">=</span> <span class="kc">True</span><span class="p">,</span>
<span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">Module</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Compile, build and load a C++/CUDA module from inline source code.</span>
+<span class="w"> </span><span class="sd">"""Compile, build and load a C++/CUDA module from inline source code.</span>
<span class="sd"> This function compiles the given C++ and/or CUDA source code into a shared library. Both ``cpp_sources`` and</span>
-<span class="sd"> ``cuda_sources`` are compiled to an object file, and then linked together into a shared library. It's possible to only</span>
+<span class="sd"> ``cuda_sources`` are compiled to an object file, and then linked together into a shared library. It's possible to only</span>
<span class="sd"> provide cpp_sources or cuda_sources.</span>
<span class="sd"> The ``functions`` parameter is used to specify which functions in the source code should be exported to the tvm ffi</span>
@@ -1970,8 +1867,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The extra compiler flags for C++ compilation.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
-<span class="sd"> - On Windows: ['/std:c++17', '/O2']</span>
+<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
+<span class="sd"> - On Windows: ['/std:c++17', '/O2']</span>
<span class="sd"> extra_cuda_cflags</span>
<span class="sd"> The extra compiler flags for CUDA compilation.</span>
@@ -1980,8 +1877,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The extra linker flags.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-shared']</span>
-<span class="sd"> - On Windows: ['/DLL']</span>
+<span class="sd"> - On Linux/macOS: ['-shared']</span>
+<span class="sd"> - On Windows: ['/DLL']</span>
<span class="sd"> extra_include_paths</span>
<span class="sd"> The extra include paths.</span>
@@ -2001,7 +1898,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> for the duration of the program until libtvm_ffi.so is unloaded.</span>
<span class="sd"> backend</span>
-<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
+<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
<span class="sd"> If not specified, the backend will be automatically determined based on the available GPU and the provided source code.</span>
<span class="sd"> Returns</span>
@@ -2023,26 +1920,26 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> import tvm_ffi.cpp</span>
<span class="sd"> # define the cpp source code</span>
-<span class="sd"> cpp_source = '''</span>
+<span class="sd"> cpp_source = '''</span>
<span class="sd"> void add_one_cpu(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {</span>
<span class="sd"> // implementation of a library function</span>
-<span class="sd"> TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
<span class="sd"> DLDataType f32_dtype{kDLFloat, 32, 1};</span>
-<span class="sd"> TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
-<span class="sd"> TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
-<span class="sd"> TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
-<span class="sd"> TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
+<span class="sd"> TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
+<span class="sd"> TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
<span class="sd"> for (int i = 0; i < x.size(0); ++i) {</span>
<span class="sd"> static_cast<float*>(y.data_ptr())[i] = static_cast<float*>(x.data_ptr())[i] + 1;</span>
<span class="sd"> }</span>
<span class="sd"> }</span>
-<span class="sd"> '''</span>
+<span class="sd"> '''</span>
<span class="sd"> # compile the cpp source code and load the module</span>
<span class="sd"> mod: Module = tvm_ffi.cpp.load_inline(</span>
-<span class="sd"> name="hello",</span>
+<span class="sd"> name="hello",</span>
<span class="sd"> cpp_sources=cpp_source,</span>
-<span class="sd"> functions="add_one_cpu",</span>
+<span class="sd"> functions="add_one_cpu",</span>
<span class="sd"> )</span>
<span class="sd"> # use the function from the loaded module to perform</span>
@@ -2051,7 +1948,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> mod.add_one_cpu(x, y)</span>
<span class="sd"> torch.testing.assert_close(x + 1, y)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="n">load_module</span><span class="p">(</span>
<span class="n">build_inline</span><span class="p">(</span>
<span class="n">name</span><span class="o">=</span><span class="n">name</span><span class="p">,</span>
@@ -2087,7 +1984,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="n">output</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Compile and build a C/C++/CUDA module from source files.</span>
+<span class="w"> </span><span class="sd">"""Compile and build a C/C++/CUDA module from source files.</span>
<span class="sd"> This function compiles the given C, C++, and/or CUDA source files into a shared library or</span>
<span class="sd"> object file. The compiler is selected automatically based on file extension:</span>
@@ -2141,8 +2038,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> Extra compiler flags applied to both C and C++ compilation.</span>
<span class="sd"> The C++ default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
-<span class="sd"> - On Windows: ['/std:c++17', '/MD', '/O2']</span>
+<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
+<span class="sd"> - On Windows: ['/std:c++17', '/MD', '/O2']</span>
<span class="sd"> The C default flags omit ``-std=c++17`` and ``/EHsc``.</span>
@@ -2150,15 +2047,15 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The extra compiler flags for CUDA compilation.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - ['-Xcompiler', '-fPIC', '-std=c++17', '-O2'] (Linux/macOS)</span>
-<span class="sd"> - ['-Xcompiler', '/std:c++17', '/O2'] (Windows)</span>
+<span class="sd"> - ['-Xcompiler', '-fPIC', '-std=c++17', '-O2'] (Linux/macOS)</span>
+<span class="sd"> - ['-Xcompiler', '/std:c++17', '/O2'] (Windows)</span>
<span class="sd"> extra_ldflags</span>
<span class="sd"> The extra linker flags.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-shared', '-L<tvm_ffi_lib_path>', '-ltvm_ffi']</span>
-<span class="sd"> - On Windows: ['/DLL', '/LIBPATH:<tvm_ffi_lib_path>', '<tvm_ffi_lib_name>.lib']</span>
+<span class="sd"> - On Linux/macOS: ['-shared', '-L<tvm_ffi_lib_path>', '-ltvm_ffi']</span>
+<span class="sd"> - On Windows: ['/DLL', '/LIBPATH:<tvm_ffi_lib_path>', '<tvm_ffi_lib_name>.lib']</span>
<span class="sd"> extra_include_paths</span>
<span class="sd"> The extra include paths for header files. Both absolute and relative paths are supported.</span>
@@ -2169,13 +2066,13 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> specify the cache directory.</span>
<span class="sd"> backend</span>
-<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
+<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
<span class="sd"> If not specified, the backend will be automatically determined based on the available GPU and the provided source code.</span>
<span class="sd"> output</span>
<span class="sd"> Output filename that determines the build type from its extension. When ``None``</span>
<span class="sd"> (the default), builds a shared library (``.so`` on Unix, ``.dll`` on Windows).</span>
-<span class="sd"> Use an object-file extension (e.g., ``"my_ops.o"``) to skip linking and produce</span>
+<span class="sd"> Use an object-file extension (e.g., ``"my_ops.o"``) to skip linking and produce</span>
<span class="sd"> a relocatable object file. The file is placed in the build directory.</span>
<span class="sd"> Returns</span>
@@ -2192,7 +2089,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> from tvm_ffi import Module</span>
<span class="sd"> import tvm_ffi.cpp</span>
-<span class="sd"> # Assume we have a C++ source file "my_ops.cpp" with the following content:</span>
+<span class="sd"> # Assume we have a C++ source file "my_ops.cpp" with the following content:</span>
<span class="sd"> # ```cpp</span>
<span class="sd"> # #include <tvm/ffi/container/tensor.h></span>
<span class="sd"> # #include <tvm/ffi/dtype.h></span>
@@ -2201,12 +2098,12 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # #include <tvm/ffi/function.h></span>
<span class="sd"> #</span>
<span class="sd"> # void add_one_cpu(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {</span>
-<span class="sd"> # TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
<span class="sd"> # DLDataType f32_dtype{kDLFloat, 32, 1};</span>
-<span class="sd"> # TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
-<span class="sd"> # TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
-<span class="sd"> # TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
-<span class="sd"> # TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
+<span class="sd"> # TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
<span class="sd"> # for (int i = 0; i < x.size(0); ++i) {</span>
<span class="sd"> # static_cast<float*>(y.data_ptr())[i] = static_cast<float*>(x.data_ptr())[i] + 1;</span>
<span class="sd"> # }</span>
@@ -2217,8 +2114,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # compile the cpp source file and get the library path</span>
<span class="sd"> lib_path: str = tvm_ffi.cpp.build(</span>
-<span class="sd"> name="my_ops",</span>
-<span class="sd"> sources="my_ops.cpp",</span>
+<span class="sd"> name="my_ops",</span>
+<span class="sd"> sources="my_ops.cpp",</span>
<span class="sd"> )</span>
<span class="sd"> # load the module</span>
@@ -2230,7 +2127,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> mod.add_one_cpu(x, y)</span>
<span class="sd"> torch.testing.assert_close(x + 1, y)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="c1"># Merge sources, cpp_files, and cuda_files (backward compat aliases)</span>
<span class="n">merged</span> <span class="o">=</span> <span class="n">_str_seq2list</span><span class="p">(</span><span class="n">sources</span><span class="p">)</span> <span class="o">+</span> <span class="n">_str_seq2list</span><span class="p">(</span><span class="n">cpp_files</span><span class="p">)</span> <span class="o">+</span> <span class="n">_str_seq2list</span><span class="p">(</span><span class="n">cuda_files</span><span class="p">)</span>
<span class="k">return</span> <span class="n">_build_impl</span><span class="p">(</span>
@@ -2264,7 +2161,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">keep_module_alive</span><span class="p">:</span> <span class="nb">bool</span> <span class="o">=</span> <span class="kc">True</span><span class="p">,</span>
<span class="n">backend</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">Module</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Compile, build and load a C/C++/CUDA module from source files.</span>
+<span class="w"> </span><span class="sd">"""Compile, build and load a C/C++/CUDA module from source files.</span>
<span class="sd"> This function compiles the given source files into a shared library and loads it as a tvm ffi</span>
<span class="sd"> module. The compiler is selected automatically based on file extension.</span>
@@ -2302,22 +2199,22 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The extra compiler flags for C++ compilation.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
-<span class="sd"> - On Windows: ['/std:c++17', '/MD', '/O2']</span>
+<span class="sd"> - On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']</span>
+<span class="sd"> - On Windows: ['/std:c++17', '/MD', '/O2']</span>
<span class="sd"> extra_cuda_cflags</span>
<span class="sd"> The extra compiler flags for CUDA compilation.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - ['-Xcompiler', '-fPIC', '-std=c++17', '-O2'] (Linux/macOS)</span>
-<span class="sd"> - ['-Xcompiler', '/std:c++17', '/O2'] (Windows)</span>
+<span class="sd"> - ['-Xcompiler', '-fPIC', '-std=c++17', '-O2'] (Linux/macOS)</span>
+<span class="sd"> - ['-Xcompiler', '/std:c++17', '/O2'] (Windows)</span>
<span class="sd"> extra_ldflags</span>
<span class="sd"> The extra linker flags.</span>
<span class="sd"> The default flags are:</span>
-<span class="sd"> - On Linux/macOS: ['-shared', '-L<tvm_ffi_lib_path>', '-ltvm_ffi']</span>
-<span class="sd"> - On Windows: ['/DLL', '/LIBPATH:<tvm_ffi_lib_path>', '<tvm_ffi_lib_name>.lib']</span>
+<span class="sd"> - On Linux/macOS: ['-shared', '-L<tvm_ffi_lib_path>', '-ltvm_ffi']</span>
+<span class="sd"> - On Windows: ['/DLL', '/LIBPATH:<tvm_ffi_lib_path>', '<tvm_ffi_lib_name>.lib']</span>
<span class="sd"> extra_include_paths</span>
<span class="sd"> The extra include paths for header files. Both absolute and relative paths are supported.</span>
@@ -2332,7 +2229,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> for the duration of the program until libtvm_ffi.so is unloaded.</span>
<span class="sd"> backend</span>
-<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
+<span class="sd"> The GPU backend to use. It can be "cuda" or "hip".</span>
<span class="sd"> If not specified, the backend will be automatically determined based on the available GPU and the provided source code.</span>
<span class="sd"> Returns</span>
@@ -2353,7 +2250,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> from tvm_ffi import Module</span>
<span class="sd"> import tvm_ffi.cpp</span>
-<span class="sd"> # Assume we have a C++ source file "my_ops.cpp" with the following content:</span>
+<span class="sd"> # Assume we have a C++ source file "my_ops.cpp" with the following content:</span>
<span class="sd"> # ```cpp</span>
<span class="sd"> # #include <tvm/ffi/container/tensor.h></span>
<span class="sd"> # #include <tvm/ffi/dtype.h></span>
@@ -2362,12 +2259,12 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # #include <tvm/ffi/function.h></span>
<span class="sd"> #</span>
<span class="sd"> # void add_one_cpu(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {</span>
-<span class="sd"> # TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";</span>
<span class="sd"> # DLDataType f32_dtype{kDLFloat, 32, 1};</span>
-<span class="sd"> # TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
-<span class="sd"> # TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
-<span class="sd"> # TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
-<span class="sd"> # TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
+<span class="sd"> # TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";</span>
+<span class="sd"> # TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";</span>
<span class="sd"> # for (int i = 0; i < x.size(0); ++i) {</span>
<span class="sd"> # static_cast<float*>(y.data_ptr())[i] = static_cast<float*>(x.data_ptr())[i] + 1;</span>
<span class="sd"> # }</span>
@@ -2378,8 +2275,8 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # compile the cpp source file and load the module</span>
<span class="sd"> mod: Module = tvm_ffi.cpp.load(</span>
-<span class="sd"> name="my_ops",</span>
-<span class="sd"> sources="my_ops.cpp",</span>
+<span class="sd"> name="my_ops",</span>
+<span class="sd"> sources="my_ops.cpp",</span>
<span class="sd"> )</span>
<span class="sd"> # use the function from the loaded module</span>
@@ -2388,7 +2285,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> mod.add_one_cpu(x, y)</span>
<span class="sd"> torch.testing.assert_close(x + 1, y)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">return</span> <span class="n">load_module</span><span class="p">(</span>
<span class="n">build</span><span class="p">(</span>
<span class="n">name</span><span class="o">=</span><span class="n">name</span><span class="p">,</span>
@@ -2445,7 +2342,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -2493,8 +2390,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/cpp/nvrtc.html b/ffi/_modules/tvm_ffi/cpp/nvrtc.html
index 37e96e01462..7e519f03efd 100644
--- a/ffi/_modules/tvm_ffi/cpp/nvrtc.html
+++ b/ffi/_modules/tvm_ffi/cpp/nvrtc.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../../_static/downloads/require.min.js"></script>
- <script src="../../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/cpp/nvrtc';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../../genindex.html" />
<link rel="search" title="Search" href="../../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../../_static/searchtools.js"></script>
+ <script src="../../../_static/language_data.js"></script>
+ <script src="../../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,18 +777,18 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""NVRTC (NVIDIA Runtime Compilation) utilities for compiling CUDA source to CUBIN."""</span>
+<span class="sd">"""NVRTC (NVIDIA Runtime Compilation) utilities for compiling CUDA source to CUBIN."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -903,11 +800,11 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">nvrtc_compile</span><span class="p">(</span> <span class="c1"># noqa: PLR0912, PLR0915</span>
<span class="n">source</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span>
<span class="o">*</span><span class="p">,</span>
- <span class="n">name</span><span class="p">:</span> <span class="nb">str</span> <span class="o">=</span> <span class="s2">"kernel.cu"</span><span class="p">,</span>
+ <span class="n">name</span><span class="p">:</span> <span class="nb">str</span> <span class="o">=</span> <span class="s2">"kernel.cu"</span><span class="p">,</span>
<span class="n">arch</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="n">extra_opts</span><span class="p">:</span> <span class="n">Sequence</span><span class="p">[</span><span class="nb">str</span><span class="p">]</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="nb">bytes</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Compile CUDA source code to CUBIN using NVRTC.</span>
+<span class="w"> </span><span class="sd">"""Compile CUDA source code to CUBIN using NVRTC.</span>
<span class="sd"> This function uses the NVIDIA Runtime Compilation (NVRTC) library to compile</span>
<span class="sd"> CUDA C++ source code into a CUBIN binary that can be loaded and executed</span>
@@ -919,14 +816,14 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> The CUDA C++ source code to compile.</span>
<span class="sd"> name : str, optional</span>
-<span class="sd"> The name to use for the source file (for error messages). Default: "kernel.cu"</span>
+<span class="sd"> The name to use for the source file (for error messages). Default: "kernel.cu"</span>
<span class="sd"> arch : str, optional</span>
-<span class="sd"> The target GPU architecture (e.g., "sm_75", "sm_80", "sm_89"). If not specified,</span>
+<span class="sd"> The target GPU architecture (e.g., "sm_75", "sm_80", "sm_89"). If not specified,</span>
<span class="sd"> attempts to auto-detect from the current GPU.</span>
<span class="sd"> extra_opts : Sequence[str], optional</span>
-<span class="sd"> Additional compilation options to pass to NVRTC (e.g., ["-I/path/to/include", "-DDEFINE=1"]).</span>
+<span class="sd"> Additional compilation options to pass to NVRTC (e.g., ["-I/path/to/include", "-DDEFINE=1"]).</span>
<span class="sd"> Returns</span>
<span class="sd"> -------</span>
@@ -944,24 +841,24 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> from tvm_ffi.cpp import nvrtc</span>
-<span class="sd"> cuda_source = '''</span>
-<span class="sd"> extern "C" __global__ void add_one(float* x, float* y, int n) {</span>
+<span class="sd"> cuda_source = '''</span>
+<span class="sd"> extern "C" __global__ void add_one(float* x, float* y, int n) {</span>
<span class="sd"> int idx = blockIdx.x * blockDim.x + threadIdx.x;</span>
<span class="sd"> if (idx < n) {</span>
<span class="sd"> y[idx] = x[idx] + 1.0f;</span>
<span class="sd"> }</span>
<span class="sd"> }</span>
-<span class="sd"> '''</span>
+<span class="sd"> '''</span>
<span class="sd"> cubin_bytes = nvrtc.nvrtc_compile(cuda_source)</span>
<span class="sd"> # Use cubin_bytes with tvm_ffi.cpp.load_inline and embed_cubin parameter</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">try</span><span class="p">:</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">cuda.bindings</span><span class="w"> </span><span class="kn">import</span> <span class="n">driver</span><span class="p">,</span> <span class="n">nvrtc</span> <span class="c1"># noqa: PLC0415</span>
<span class="k">except</span> <span class="ne">ImportError</span> <span class="k">as</span> <span class="n">e</span><span class="p">:</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="s2">"CUDA bindings not available. Install with: pip install cuda-python"</span>
+ <span class="s2">"CUDA bindings not available. Install with: pip install cuda-python"</span>
<span class="p">)</span> <span class="kn">from</span><span class="w"> </span><span class="nn">e</span>
<span class="c1"># Auto-detect architecture if not specified</span>
@@ -970,7 +867,7 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># Initialize CUDA driver API</span>
<span class="p">(</span><span class="n">result</span><span class="p">,)</span> <span class="o">=</span> <span class="n">driver</span><span class="o">.</span><span class="n">cuInit</span><span class="p">(</span><span class="mi">0</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="o">!=</span> <span class="n">driver</span><span class="o">.</span><span class="n">CUresult</span><span class="o">.</span><span class="n">CUDA_SUCCESS</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to initialize CUDA driver: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to initialize CUDA driver: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="c1"># Get current device</span>
<span class="n">result</span><span class="p">,</span> <span class="n">device</span> <span class="o">=</span> <span class="n">driver</span><span class="o">.</span><span class="n">cuCtxGetDevice</span><span class="p">()</span>
@@ -983,31 +880,31 @@ window.runMermaid = runMermaid;</script>
<span class="n">driver</span><span class="o">.</span><span class="n">CUdevice_attribute</span><span class="o">.</span><span class="n">CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR</span><span class="p">,</span> <span class="n">device</span>
<span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="o">!=</span> <span class="n">driver</span><span class="o">.</span><span class="n">CUresult</span><span class="o">.</span><span class="n">CUDA_SUCCESS</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get compute capability major: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get compute capability major: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="n">result</span><span class="p">,</span> <span class="n">minor</span> <span class="o">=</span> <span class="n">driver</span><span class="o">.</span><span class="n">cuDeviceGetAttribute</span><span class="p">(</span>
<span class="n">driver</span><span class="o">.</span><span class="n">CUdevice_attribute</span><span class="o">.</span><span class="n">CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR</span><span class="p">,</span> <span class="n">device</span>
<span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="o">!=</span> <span class="n">driver</span><span class="o">.</span><span class="n">CUresult</span><span class="o">.</span><span class="n">CUDA_SUCCESS</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get compute capability minor: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get compute capability minor: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">arch</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">"</span>
+ <span class="n">arch</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"sm_</span><span class="si">{</span><span class="n">major</span><span class="si">}{</span><span class="n">minor</span><span class="si">}</span><span class="s2">"</span>
<span class="k">except</span> <span class="ne">Exception</span> <span class="k">as</span> <span class="n">e</span><span class="p">:</span>
<span class="c1"># Fallback to a reasonable default</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span>
- <span class="sa">f</span><span class="s2">"Failed to auto-detect GPU architecture: </span><span class="si">{</span><span class="n">e</span><span class="si">}</span><span class="s2">. "</span>
- <span class="s2">"Please specify 'arch' parameter explicitly."</span>
+ <span class="sa">f</span><span class="s2">"Failed to auto-detect GPU architecture: </span><span class="si">{</span><span class="n">e</span><span class="si">}</span><span class="s2">. "</span>
+ <span class="s2">"Please specify 'arch' parameter explicitly."</span>
<span class="p">)</span> <span class="kn">from</span><span class="w"> </span><span class="nn">e</span>
<span class="c1"># Create program</span>
<span class="n">result</span><span class="p">,</span> <span class="n">prog</span> <span class="o">=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcCreateProgram</span><span class="p">(</span><span class="nb">str</span><span class="o">.</span><span class="n">encode</span><span class="p">(</span><span class="n">source</span><span class="p">),</span> <span class="nb">str</span><span class="o">.</span><span class="n">encode</span><span class="p">(</span>< [...]
<span class="k">if</span> <span class="n">result</span> <span class="o">!=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcResult</span><span class="o">.</span><span class="n">NVRTC_SUCCESS</span><span class="p">:</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to create NVRTC program: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to create NVRTC program: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="c1"># Compile options</span>
<span class="n">opts</span> <span class="o">=</span> <span class="p">[</span>
- <span class="sa">b</span><span class="s2">"--gpu-architecture="</span> <span class="o">+</span> <span class="n">arch</span><span class="o">.</span><span class="n">encode</span><span class="p">(),</span>
- <span class="sa">b</span><span class="s2">"-default-device"</span><span class="p">,</span>
+ <span class="sa">b</span><span class="s2">"--gpu-architecture="</span> <span class="o">+</span> <span class="n">arch</span><span class="o">.</span><span class="n">encode</span><span class="p">(),</span>
+ <span class="sa">b</span><span class="s2">"-default-device"</span><span class="p">,</span>
<span class="p">]</span>
<span class="c1"># Add extra options if provided</span>
@@ -1020,14 +917,14 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># Get compilation log</span>
<span class="n">result_log</span><span class="p">,</span> <span class="n">log_size</span> <span class="o">=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcGetProgramLogSize</span><span class="p">(</span><span class="n">prog</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result_log</span> <span class="o">==</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcResult</span><span class="o">.</span><span class="n">NVRTC_SUCCESS</span> <span class="ow">and</span> <span class="n">log_size</span> <span class="o">></span> <span class="mi">0</span><span class="p">:</span>
- <span class="n">log_buf</span> <span class="o">=</span> <span class="sa">b</span><span class="s2">" "</span> <span class="o">*</span> <span class="n">log_size</span>
+ <span class="n">log_buf</span> <span class="o">=</span> <span class="sa">b</span><span class="s2">" "</span> <span class="o">*</span> <span class="n">log_size</span>
<span class="p">(</span><span class="n">result_log</span><span class="p">,)</span> <span class="o">=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcGetProgramLog</span><span class="p">(</span><span class="n">prog</span><span class="p">,</span> <span class="n">log_buf</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result_log</span> <span class="o">==</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcResult</span><span class="o">.</span><span class="n">NVRTC_SUCCESS</span><span class="p">:</span>
- <span class="n">error_msg</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"NVRTC compilation failed:</span><span class="se">\n</span><span class="si">{</span><span class="n">log_buf</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="s1">'utf-8'</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span>
+ <span class="n">error_msg</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"NVRTC compilation failed:</span><span class="se">\n</span><span class="si">{</span><span class="n">log_buf</span><span class="o">.</span><span class="n">decode</span><span class="p">(</span><span class="s1">'utf-8'</span><span class="p">)</span><span class="si">}</span><span class="s2">"</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">error_msg</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"NVRTC compilation failed (couldn't get log): </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span>
+ <span class="n">error_msg</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"NVRTC compilation failed (couldn't get log): </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span>
<span class="k">else</span><span class="p">:</span>
- <span class="n">error_msg</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"NVRTC compilation failed: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span>
+ <span class="n">error_msg</span> <span class="o">=</span> <span class="sa">f</span><span class="s2">"NVRTC compilation failed: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span>
<span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcDestroyProgram</span><span class="p">(</span><span class="n">prog</span><span class="p">)</span>
<span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="n">error_msg</span><span class="p">)</span>
@@ -1036,13 +933,13 @@ window.runMermaid = runMermaid;</script>
<span class="n">result</span><span class="p">,</span> <span class="n">cubin_size</span> <span class="o">=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcGetCUBINSize</span><span class="p">(</span><span class="n">prog</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="o">!=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcResult</span><span class="o">.</span><span class="n">NVRTC_SUCCESS</span><span class="p">:</span>
<span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcDestroyProgram</span><span class="p">(</span><span class="n">prog</span><span class="p">)</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get CUBIN size from NVRTC: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get CUBIN size from NVRTC: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
- <span class="n">cubin_buf</span> <span class="o">=</span> <span class="sa">b</span><span class="s2">" "</span> <span class="o">*</span> <span class="n">cubin_size</span>
+ <span class="n">cubin_buf</span> <span class="o">=</span> <span class="sa">b</span><span class="s2">" "</span> <span class="o">*</span> <span class="n">cubin_size</span>
<span class="p">(</span><span class="n">result</span><span class="p">,)</span> <span class="o">=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcGetCUBIN</span><span class="p">(</span><span class="n">prog</span><span class="p">,</span> <span class="n">cubin_buf</span><span class="p">)</span>
<span class="k">if</span> <span class="n">result</span> <span class="o">!=</span> <span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcResult</span><span class="o">.</span><span class="n">NVRTC_SUCCESS</span><span class="p">:</span>
<span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcDestroyProgram</span><span class="p">(</span><span class="n">prog</span><span class="p">)</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get CUBIN from NVRTC: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="sa">f</span><span class="s2">"Failed to get CUBIN from NVRTC: </span><span class="si">{</span><span class="n">result</span><span class="si">}</span><span class="s2">"</span><span class="p">)</span>
<span class="c1"># Clean up</span>
<span class="n">nvrtc</span><span class="o">.</span><span class="n">nvrtcDestroyProgram</span><span class="p">(</span><span class="n">prog</span><span class="p">)</span>
@@ -1089,7 +986,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1137,8 +1034,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/error.html b/ffi/_modules/tvm_ffi/error.html
index f0936f9bce1..5af3fd868df 100644
--- a/ffi/_modules/tvm_ffi/error.html
+++ b/ffi/_modules/tvm_ffi/error.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/error';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,19 +777,19 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
<span class="c1"># pylint: disable=invalid-name</span>
-<span class="sd">"""Error handling."""</span>
+<span class="sd">"""Error handling."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -906,7 +803,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">_parse_backtrace</span><span class="p">(</span><span class="n">backtrace</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="nb">list</span><span class="p">[</span><span class="nb">tuple</span><span class="p">[</span><span class="nb">str</span><span class="p">,</span> <span class="nb">int</span><span class="p">,</span> <span class="nb">str</span><spa [...]
-<span class="w"> </span><span class="sd">"""Parse the backtrace string into a list of (filename, lineno, func).</span>
+<span class="w"> </span><span class="sd">"""Parse the backtrace string into a list of (filename, lineno, func).</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -918,10 +815,10 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> result</span>
<span class="sd"> The list of (filename, lineno, func)</span>
-<span class="sd"> """</span>
- <span class="n">pattern</span> <span class="o">=</span> <span class="sa">r</span><span class="s1">'File "(.+?)", line (\d+), in (.+)'</span>
+<span class="sd"> """</span>
+ <span class="n">pattern</span> <span class="o">=</span> <span class="sa">r</span><span class="s1">'File "(.+?)", line (\d+), in (.+)'</span>
<span class="n">result</span> <span class="o">=</span> <span class="p">[]</span>
- <span class="k">for</span> <span class="n">line</span> <span class="ow">in</span> <span class="n">backtrace</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">):</span>
+ <span class="k">for</span> <span class="n">line</span> <span class="ow">in</span> <span class="n">backtrace</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="s2">"</span><span class="se">\n</span><span class="s2">"</span><span class="p">):</span>
<span class="n">match</span> <span class="o">=</span> <span class="n">re</span><span class="o">.</span><span class="n">match</span><span class="p">(</span><span class="n">pattern</span><span class="p">,</span> <span class="n">line</span><span class="o">.</span><span class="n">strip</span><span class="p">())</span>
<span class="k">if</span> <span class="n">match</span><span class="p">:</span>
<span class="k">try</span><span class="p">:</span>
@@ -935,10 +832,10 @@ window.runMermaid = runMermaid;</script>
<span class="k">class</span><span class="w"> </span><span class="nc">TracebackManager</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Helper to manage traceback generation."""</span>
+<span class="w"> </span><span class="sd">"""Helper to manage traceback generation."""</span>
<span class="k">def</span><span class="w"> </span><span class="fm">__init__</span><span class="p">(</span><span class="bp">self</span><span class="p">)</span> <span class="o">-></span> <span class="kc">None</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Initialize the traceback manager and its cache."""</span>
+<span class="w"> </span><span class="sd">"""Initialize the traceback manager and its cache."""</span>
<span class="bp">self</span><span class="o">.</span><span class="n">_code_cache</span><span class="p">:</span> <span class="nb">dict</span><span class="p">[</span><span class="nb">tuple</span><span class="p">[</span><span class="nb">str</span><span class="p">,</span> <span class="nb">int</span><span class="p">,</span> <span class="nb">str</span><span class="p">],</span> <span class="n">types</span><span class="o">.</span><span class="n">CodeType</span><span class="p">]</span> <sp [...]
<span class="k">def</span><span class="w"> </span><span class="nf">_get_cached_code_object</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">filename</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span> <span class="n">lineno</span><span class="p">:</span> <span class="nb">int</span><span class="p">,</span> <span class="n">func</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</sp [...]
@@ -950,25 +847,25 @@ window.runMermaid = runMermaid;</script>
<span class="k">return</span> <span class="bp">self</span><span class="o">.</span><span class="n">_code_cache</span><span class="p">[</span><span class="n">key</span><span class="p">]</span>
<span class="c1"># Parse to AST and zero out column info</span>
<span class="c1"># since column info are not accurate in original trace</span>
- <span class="n">tree</span> <span class="o">=</span> <span class="n">ast</span><span class="o">.</span><span class="n">parse</span><span class="p">(</span><span class="s2">"_getframe()"</span><span class="p">,</span> <span class="n">filename</span><span class="o">=</span><span class="n">filename</span><span class="p">,</span> <span class="n">mode</span><span class="o">=</span><span class="s2">"eval"</span><span class="p">)</span>
+ <span class="n">tree</span> <span class="o">=</span> <span class="n">ast</span><span class="o">.</span><span class="n">parse</span><span class="p">(</span><span class="s2">"_getframe()"</span><span class="p">,</span> <span class="n">filename</span><span class="o">=</span><span class="n">filename</span><span class="p">,</span> <span class="n">mode</span><span class="o">=</span><span class="s2">"eval"</span><span class="p">)</span>
<span class="k">for</span> <span class="n">node</span> <span class="ow">in</span> <span class="n">ast</span><span class="o">.</span><span class="n">walk</span><span class="p">(</span><span class="n">tree</span><span class="p">):</span>
- <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">node</span><span class="p">,</span> <span class="s2">"col_offset"</span><span class="p">):</span>
+ <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">node</span><span class="p">,</span> <span class="s2">"col_offset"</span><span class="p">):</span>
<span class="n">node</span><span class="o">.</span><span class="n">col_offset</span> <span class="o">=</span> <span class="mi">0</span> <span class="c1"># ty: ignore[invalid-assignment]</span>
- <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">node</span><span class="p">,</span> <span class="s2">"end_col_offset"</span><span class="p">):</span>
+ <span class="k">if</span> <span class="nb">hasattr</span><span class="p">(</span><span class="n">node</span><span class="p">,</span> <span class="s2">"end_col_offset"</span><span class="p">):</span>
<span class="n">node</span><span class="o">.</span><span class="n">end_col_offset</span> <span class="o">=</span> <span class="mi">0</span> <span class="c1"># ty: ignore[invalid-assignment]</span>
<span class="c1"># call into get frame, bt changes the context</span>
- <span class="n">code_object</span> <span class="o">=</span> <span class="nb">compile</span><span class="p">(</span><span class="n">tree</span><span class="p">,</span> <span class="n">filename</span><span class="p">,</span> <span class="s2">"eval"</span><span class="p">)</span>
+ <span class="n">code_object</span> <span class="o">=</span> <span class="nb">compile</span><span class="p">(</span><span class="n">tree</span><span class="p">,</span> <span class="n">filename</span><span class="p">,</span> <span class="s2">"eval"</span><span class="p">)</span>
<span class="c1"># replace the function name and line number</span>
<span class="n">code_object</span> <span class="o">=</span> <span class="n">code_object</span><span class="o">.</span><span class="n">replace</span><span class="p">(</span><span class="n">co_name</span><span class="o">=</span><span class="n">func</span><span class="p">,</span> <span class="n">co_firstlineno</span><span class="o">=</span><span class="n">lineno</span><span class="p">)</span>
<span class="bp">self</span><span class="o">.</span><span class="n">_code_cache</span><span class="p">[</span><span class="n">key</span><span class="p">]</span> <span class="o">=</span> <span class="n">code_object</span>
<span class="k">return</span> <span class="n">code_object</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_create_frame</span><span class="p">(</span><span class="bp">self</span><span class="p">,</span> <span class="n">filename</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span> <span class="n">lineno</span><span class="p">:</span> <span class="nb">int</span><span class="p">,</span> <span class="n">func</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span [...]
-<span class="w"> </span><span class="sd">"""Create a frame object from the filename, lineno, and func."""</span>
+<span class="w"> </span><span class="sd">"""Create a frame object from the filename, lineno, and func."""</span>
<span class="n">code_object</span> <span class="o">=</span> <span class="bp">self</span><span class="o">.</span><span class="n">_get_cached_code_object</span><span class="p">(</span><span class="n">filename</span><span class="p">,</span> <span class="n">lineno</span><span class="p">,</span> <span class="n">func</span><span class="p">)</span>
<span class="c1"># call into get frame, but changes the context so the code</span>
<span class="c1"># points to the correct frame</span>
- <span class="n">context</span> <span class="o">=</span> <span class="p">{</span><span class="s2">"_getframe"</span><span class="p">:</span> <span class="n">sys</span><span class="o">.</span><span class="n">_getframe</span><span class="p">}</span>
+ <span class="n">context</span> <span class="o">=</span> <span class="p">{</span><span class="s2">"_getframe"</span><span class="p">:</span> <span class="n">sys</span><span class="o">.</span><span class="n">_getframe</span><span class="p">}</span>
<span class="c1"># pylint: disable=eval-used</span>
<span class="k">return</span> <span class="nb">eval</span><span class="p">(</span><span class="n">code_object</span><span class="p">,</span> <span class="n">context</span><span class="p">,</span> <span class="n">context</span><span class="p">)</span>
@@ -979,7 +876,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">lineno</span><span class="p">:</span> <span class="nb">int</span><span class="p">,</span>
<span class="n">func</span><span class="p">:</span> <span class="nb">str</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">types</span><span class="o">.</span><span class="n">TracebackType</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Append a traceback to the given traceback.</span>
+<span class="w"> </span><span class="sd">"""Append a traceback to the given traceback.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -997,7 +894,7 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> new_tb</span>
<span class="sd"> The new traceback with the appended frame.</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="c1"># This approach avoids binding the created frame object to a local variable</span>
<span class="c1"># in `append_traceback`, which would create a reference cycle. By using a</span>
@@ -1016,7 +913,7 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">_with_append_backtrace</span><span class="p">(</span><span class="n">py_error</span><span class="p">:</span> <span class="ne">BaseException</span><span class="p">,</span> <span class="n">backtrace</span><span class="p">:</span> <span class="nb">str</span><span class="p">)</span> <span class="o">-></span> <span class="ne">BaseException</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Append the backtrace to the py_error and return it."""</span>
+<span class="w"> </span><span class="sd">"""Append the backtrace to the py_error and return it."""</span>
<span class="c1"># We manually delete py_error and tb to avoid reference cycle, making it faster to gc the locals inside the frame</span>
<span class="c1"># please see pull request #327 for more details</span>
<span class="c1">#</span>
@@ -1055,24 +952,24 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># We explicitly break the reference cycle here. The `finally` block is</span>
<span class="c1"># executed just before the function returns, after the `return` expression</span>
<span class="c1"># in the `try` block has been evaluated. Deleting `py_error` and `tb`</span>
- <span class="c1"># here ensures they are not held by this function's frame's locals,</span>
+ <span class="c1"># here ensures they are not held by this function's frame's locals,</span>
<span class="c1"># which resolves the cycle.</span>
<span class="k">del</span> <span class="n">py_error</span><span class="p">,</span> <span class="n">tb</span>
<span class="k">def</span><span class="w"> </span><span class="nf">_traceback_to_backtrace_str</span><span class="p">(</span><span class="n">tb</span><span class="p">:</span> <span class="n">types</span><span class="o">.</span><span class="n">TracebackType</span> <span class="o">|</span> <span class="kc">None</span><span class="p">)</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Convert the traceback to a string."""</span>
+<span class="w"> </span><span class="sd">"""Convert the traceback to a string."""</span>
<span class="n">lines</span> <span class="o">=</span> <span class="p">[]</span>
<span class="k">while</span> <span class="n">tb</span> <span class="ow">is</span> <span class="ow">not</span> <span class="kc">None</span><span class="p">:</span>
<span class="n">frame</span> <span class="o">=</span> <span class="n">tb</span><span class="o">.</span><span class="n">tb_frame</span>
<span class="n">lineno</span> <span class="o">=</span> <span class="n">tb</span><span class="o">.</span><span class="n">tb_lineno</span>
<span class="n">filename</span> <span class="o">=</span> <span class="n">frame</span><span class="o">.</span><span class="n">f_code</span><span class="o">.</span><span class="n">co_filename</span>
<span class="n">funcname</span> <span class="o">=</span> <span class="n">frame</span><span class="o">.</span><span class="n">f_code</span><span class="o">.</span><span class="n">co_name</span>
- <span class="n">lines</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s1">' File "</span><span class="si">{</span><span class="n">filename</span><span class="si">}</span><span class="s1">", line </span><span class="si">{</span><span class="n">lineno</span><span class="si">}</span><span class="s1">, in </span><span class="si">{</span><span class="n">funcname</span><span class="si">}</span><span [...]
+ <span class="n">lines</span><span class="o">.</span><span class="n">append</span><span class="p">(</span><span class="sa">f</span><span class="s1">' File "</span><span class="si">{</span><span class="n">filename</span><span class="si">}</span><span class="s1">", line </span><span class="si">{</span><span class="n">lineno</span><span class="si">}</span><span class="s1">, in </span><span class="si">{</span><span class="n">funcname</span><span class="si">}</span><span class="se">\n [...]
<span class="n">tb</span> <span class="o">=</span> <span class="n">tb</span><span class="o">.</span><span class="n">tb_next</span>
<span class="c1"># needs to reverse the order of the lines so backtrace stores in</span>
<span class="c1"># the reverse order of python traceback</span>
- <span class="k">return</span> <span class="s2">""</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="nb">reversed</span><span class="p">(</span><span class="n">lines</span><span class="p">))</span>
+ <span class="k">return</span> <span class="s2">""</span><span class="o">.</span><span class="n">join</span><span class="p">(</span><span class="nb">reversed</span><span class="p">(</span><span class="n">lines</span><span class="p">))</span>
<span class="n">core</span><span class="o">.</span><span class="n">_WITH_APPEND_BACKTRACE</span> <span class="o">=</span> <span class="n">_with_append_backtrace</span>
@@ -1085,7 +982,7 @@ window.runMermaid = runMermaid;</script>
<span class="n">name_or_cls</span><span class="p">:</span> <span class="nb">str</span> <span class="o">|</span> <span class="nb">type</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="bp">cls</span><span class="p">:</span> <span class="nb">type</span> <span class="o">|</span> <span class="kc">None</span> <span class="o">=</span> <span class="kc">None</span><span class="p">,</span>
<span class="p">)</span> <span class="o">-></span> <span class="n">Any</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Register an error class so it can be recognized by the ffi error handler.</span>
+<span class="w"> </span><span class="sd">"""Register an error class so it can be recognized by the ffi error handler.</span>
<span class="sd"> Parameters</span>
<span class="sd"> ----------</span>
@@ -1114,17 +1011,17 @@ window.runMermaid = runMermaid;</script>
<span class="sd"> # Convert a Python exception to an FFI Error and back</span>
-<span class="sd"> ffi_err = tvm_ffi.convert(MyError("boom"))</span>
+<span class="sd"> ffi_err = tvm_ffi.convert(MyError("boom"))</span>
<span class="sd"> py_err = ffi_err.py_error()</span>
<span class="sd"> assert isinstance(py_err, MyError)</span>
-<span class="sd"> """</span>
+<span class="sd"> """</span>
<span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">name_or_cls</span><span class="p">,</span> <span class="nb">type</span><span class="p">):</span>
<span class="bp">cls</span> <span class="o">=</span> <span class="n">name_or_cls</span>
<span class="n">name_or_cls</span> <span class="o">=</span> <span class="bp">cls</span><span class="o">.</span><span class="vm">__name__</span>
<span class="k">def</span><span class="w"> </span><span class="nf">register</span><span class="p">(</span><span class="n">mycls</span><span class="p">:</span> <span class="nb">type</span><span class="p">)</span> <span class="o">-></span> <span class="nb">type</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Register the error class name with the FFI core."""</span>
+<span class="w"> </span><span class="sd">"""Register the error class name with the FFI core."""</span>
<span class="n">err_name</span> <span class="o">=</span> <span class="n">name_or_cls</span> <span class="k">if</span> <span class="nb">isinstance</span><span class="p">(</span><span class="n">name_or_cls</span><span class="p">,</span> <span class="nb">str</span><span class="p">)</span> <span class="k">else</span> <span class="n">mycls</span><span class="o">.</span><span class="vm">__name__</span>
<span class="n">core</span><span class="o">.</span><span class="n">ERROR_NAME_TO_TYPE</span><span class="p">[</span><span class="n">err_name</span><span class="p">]</span> <span class="o">=</span> <span class="n">mycls</span>
<span class="n">core</span><span class="o">.</span><span class="n">ERROR_TYPE_TO_NAME</span><span class="p">[</span><span class="n">mycls</span><span class="p">]</span> <span class="o">=</span> <span class="n">err_name</span>
@@ -1136,14 +1033,14 @@ window.runMermaid = runMermaid;</script>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"RuntimeError"</span><span class="p">,</span> <span class="ne">RuntimeError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"ValueError"</span><span class="p">,</span> <span class="ne">ValueError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"TypeError"</span><span class="p">,</span> <span class="ne">TypeError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"AttributeError"</span><span class="p">,</span> <span class="ne">AttributeError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"KeyError"</span><span class="p">,</span> <span class="ne">KeyError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"IndexError"</span><span class="p">,</span> <span class="ne">IndexError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"AssertionError"</span><span class="p">,</span> <span class="ne">AssertionError</span><span class="p">)</span>
-<span class="n">register_error</span><span class="p">(</span><span class="s2">"MemoryError"</span><span class="p">,</span> <span class="ne">MemoryError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"RuntimeError"</span><span class="p">,</span> <span class="ne">RuntimeError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"ValueError"</span><span class="p">,</span> <span class="ne">ValueError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"TypeError"</span><span class="p">,</span> <span class="ne">TypeError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"AttributeError"</span><span class="p">,</span> <span class="ne">AttributeError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"KeyError"</span><span class="p">,</span> <span class="ne">KeyError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"IndexError"</span><span class="p">,</span> <span class="ne">IndexError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"AssertionError"</span><span class="p">,</span> <span class="ne">AssertionError</span><span class="p">)</span>
+<span class="n">register_error</span><span class="p">(</span><span class="s2">"MemoryError"</span><span class="p">,</span> <span class="ne">MemoryError</span><span class="p">)</span>
</pre></div>
</article>
@@ -1184,7 +1081,7 @@ By Apache TVM FFI contributors
<div class="footer-item">
<p class="last-updated">
- Last updated on Jul 30, 2026.
+ Last updated on Sep 15, 2026.
<br/>
</p>
</div>
@@ -1232,8 +1129,8 @@ By Apache TVM FFI contributors
</div>
<!-- Scripts loaded after <body> so the DOM is not blocked -->
- <script defer src="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf"></script>
-<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9"></script>
+<script defer src="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9"></script>
<footer class="bd-footer">
</footer>
diff --git a/ffi/_modules/tvm_ffi/libinfo.html b/ffi/_modules/tvm_ffi/libinfo.html
index 26d3dcb4ec6..f3b4830d985 100644
--- a/ffi/_modules/tvm_ffi/libinfo.html
+++ b/ffi/_modules/tvm_ffi/libinfo.html
@@ -26,23 +26,23 @@
</noscript>
<!-- Loaded before other Sphinx assets -->
- <link href="../../_static/styles/theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
-<link href="../../_static/styles/pydata-sphinx-theme.css?digest=8878045cc6db502f8baf" rel="stylesheet" />
+ <link href="../../_static/styles/theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
+<link href="../../_static/styles/pydata-sphinx-theme.css?digest=90905a2f556bf617f1a9" rel="stylesheet" />
<link rel="stylesheet" type="text/css" href="../../_static/pygments.css?v=03e43079" />
- <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=3c74b3bc" />
+ <link rel="stylesheet" type="text/css" href="../../_static/styles/sphinx-book-theme.css?v=4418a689" />
<link rel="stylesheet" type="text/css" href="../../_static/copybutton.css?v=76b2166b" />
<link rel="stylesheet" type="text/css" href="../../_static/sphinx-design.min.css?v=95c83b7e" />
<link rel="stylesheet" type="text/css" href="../../_static/collapsible-lists/css/tree_view.css?v=a885cde7" />
<link rel="stylesheet" type="text/css" href="../../_static/custom.css?v=38aa6179" />
<!-- So that users can add custom icons -->
- <script src="../../_static/scripts/fontawesome.js?digest=8878045cc6db502f8baf"></script>
+ <script defer src="../../_static/scripts/fontawesome.js?digest=90905a2f556bf617f1a9"></script>
<!-- Pre-loaded scripts that we'll load fully later -->
- <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=8878045cc6db502f8baf" />
-<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=8878045cc6db502f8baf" />
+ <link rel="preload" as="script" href="../../_static/scripts/bootstrap.js?digest=90905a2f556bf617f1a9" />
+<link rel="preload" as="script" href="../../_static/scripts/pydata-sphinx-theme.js?digest=90905a2f556bf617f1a9" />
- <script src="../../_static/documentation_options.js?v=4b6634f4"></script>
+ <script src="../../_static/documentation_options.js?v=bd121c97"></script>
<script src="../../_static/doctools.js?v=9bcbadda"></script>
<script src="../../_static/sphinx_highlight.js?v=dc90522c"></script>
<script src="../../_static/clipboard.min.js?v=a7894cd8"></script>
@@ -52,182 +52,17 @@
<script src="../../_static/collapsible-lists/js/CollapsibleLists.compressed.js?v=73120307"></script>
<script src="../../_static/collapsible-lists/js/apply-collapsible-lists.js?v=660e4f45"></script>
<script crossorigin="anonymous" integrity="sha256-Ae2Vz/4ePdIu6ZyI/5ZGsYnb+m0JlOmKPjt6XZ9JJkA=" src="../../_static/downloads/require.min.js"></script>
- <script src="../../_static/downloads/d3.min.js"></script>
<script type="module">import mermaid from "https://cdn.jsdelivr.net/npm/mermaid@11.12.1/dist/mermaid.esm.min.mjs";
-
-
const initStyles = () => {
const defaultStyle = document.createElement('style');
- defaultStyle.textContent = `pre.mermaid {
- /* Same as .mermaid-container > pre */
- display: block;
- width: 100%;
-}
-
-pre.mermaid > svg {
- /* Same as .mermaid-container > pre > svg */
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}`;
+ defaultStyle.textContent = "pre.mermaid {\n /* Same as .mermaid-container > pre */\n display: block;\n width: 100%;\n}\n\npre.mermaid > svg {\n /* Same as .mermaid-container > pre > svg */\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}";
document.head.appendChild(defaultStyle);
const fullscreenStyle = document.createElement('style');
- fullscreenStyle.textContent = `.mermaid-container {
- display: flex;
- flex-direction: row;
- width: 100%;
-}
-
-.mermaid-container > pre {
- display: block;
- width: 100%;
-}
-
-.mermaid-container > pre > svg {
- height: 500px;
- width: 100%;
- max-width: 100% !important;
-}
-
-.mermaid-fullscreen-btn {
- width: 28px;
- height: 28px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.3);
- border-radius: 4px;
- cursor: pointer;
- display: flex;
- align-items: center;
- justify-content: center;
- transition: all 0.2s;
- box-shadow: 0 2px 6px rgba(0, 0, 0, 0.2);
- font-size: 14px;
- line-height: 1;
- padding: 0;
- color: #333;
-}
-
-.mermaid-fullscreen-btn:hover {
- opacity: 100% !important;
- background: rgba(255, 255, 255, 1);
- box-shadow: 0 3px 10px rgba(0, 0, 0, 0.3);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-btn.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.3);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-btn.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 3px 10px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal {
- display: none;
- position: fixed !important;
- top: 0 !important;
- left: 0 !important;
- width: 95vw;
- height: 100vh;
- background: rgba(255, 255, 255, 0.98);
- z-index: 9999;
- padding: 20px;
- overflow: auto;
-}
-
-.mermaid-fullscreen-modal.dark-theme {
- background: rgba(0, 0, 0, 0.98);
-}
-
-.mermaid-fullscreen-modal.active {
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen {
- position: relative;
- width: 95vw;
- height: 90vh;
- max-width: 95vw;
- max-height: 90vh;
- background: white;
- border-radius: 8px;
- padding: 20px;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.3);
- overflow: auto;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen.dark-theme {
- background: #1a1a1a;
- box-shadow: 0 10px 40px rgba(0, 0, 0, 0.8);
-}
-
-.mermaid-container-fullscreen pre.mermaid {
- width: 100%;
- height: 100%;
- display: flex;
- align-items: center;
- justify-content: center;
-}
-
-.mermaid-container-fullscreen .mermaid svg {
- height: 100% !important;
- width: 100% !important;
- cursor: grab;
-}
-
-.mermaid-fullscreen-close {
- position: fixed !important;
- top: 20px !important;
- right: 20px !important;
- width: 40px;
- height: 40px;
- background: rgba(255, 255, 255, 0.95);
- border: 1px solid rgba(0, 0, 0, 0.2);
- border-radius: 50%;
- cursor: pointer;
- z-index: 10000;
- display: flex;
- align-items: center;
- justify-content: center;
- box-shadow: 0 4px 12px rgba(0, 0, 0, 0.3);
- transition: all 0.2s;
- font-size: 24px;
- line-height: 1;
- color: #333;
-}
-
-.mermaid-fullscreen-close:hover {
- background: white;
- box-shadow: 0 6px 16px rgba(0, 0, 0, 0.4);
- transform: scale(1.1);
-}
-
-.mermaid-fullscreen-close.dark-theme {
- background: rgba(50, 50, 50, 0.95);
- border: 1px solid rgba(255, 255, 255, 0.2);
- color: #e0e0e0;
-}
-
-.mermaid-fullscreen-close.dark-theme:hover {
- background: rgba(60, 60, 60, 1);
- box-shadow: 0 6px 16px rgba(255, 255, 255, 0.2);
-}
-
-.mermaid-fullscreen-modal .mermaid-fullscreen-btn {
- display: none !important;
-}`;
+ fullscreenStyle.textContent = ".mermaid-container {\n position: relative;\n display: flex;\n flex-direction: row;\n width: 100%;\n}\n\n.mermaid-container > pre {\n display: block;\n width: 100%;\n}\n\n.mermaid-container > pre > svg {\n height: 500px;\n width: 100%;\n max-width: 100% !important;\n}\n\n.mermaid-fullscreen-btn {\n position: absolute;\n width: 28px;\n height: 28px;\n background: rgba(255, 255, 255, 0.95);\n border: 1px solid rgba [...]
document.head.appendChild(fullscreenStyle);
}
@@ -274,6 +109,13 @@ let modalContent = null;
let previousScrollOffset = [window.scrollX, window.scrollY];
let _lazyObserver = null;
let _renderQueue = Promise.resolve();
+let closeModal = () => {};
+
+document.addEventListener('keydown', (e) => {
+ if (e.key === 'Escape' && modal?.classList.contains('active')) {
+ closeModal();
+ }
+});
// Apply d3 zoom to each SVG. Idempotent: an SVG already wrapped is skipped,
// so this is safe to call both for the initial batch and for diagrams that
@@ -370,8 +212,10 @@ const runMermaid = async (rerun) => {
el.removeAttribute('data-mermaid-deferred');
// Apply zoom to the now-rendered diagram, matching
// the decoration visible diagrams receive.
- if ("False" === "True") {
- addZoomToSvgs(d3.selectAll(""));
+ if (false) {
+ addZoomToSvgs(d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ }));
}
} catch (e) {
console.error("Mermaid deferred rendering failed:", e);
@@ -392,10 +236,12 @@ const runMermaid = async (rerun) => {
const mermaids_active = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])");
const mermaids_processed = document.querySelectorAll(".mermaid:not([data-mermaid-deferred]):not([data-mermaid-render-failed])[data-processed='true']");
- if ("False" === "True") {
- const mermaids_to_add_zoom = -1 === -1 ? mermaids_active.length : -1;
+ if (false) {
+ const mermaids_to_add_zoom = 0 === -1 ? mermaids_active.length : 0;
if(mermaids_to_add_zoom > 0) {
- var svgs = d3.selectAll("");
+ var svgs = d3.selectAll("").select(function() {
+ return this.querySelector("svg");
+ });
// Wait until every active (visible) diagram has rendered. Deferred
// diagrams are excluded, so a fixed d3_node_count that counts a
// hidden diagram no longer loops forever waiting for its SVG.
@@ -412,7 +258,7 @@ const runMermaid = async (rerun) => {
}
// Stop here if not adding fullscreen capability
- if ("True" !== "True") return;
+ if (!true) return;
if (modal !== null ) {
// Destroy existing modal
@@ -435,7 +281,7 @@ const runMermaid = async (rerun) => {
modalContent = modal.querySelector('.mermaid-container-fullscreen');
const closeBtn = modal.querySelector('.mermaid-fullscreen-close');
- const closeModal = () => {
+ closeModal = () => {
modal.classList.remove('active');
modalContent.innerHTML = '';
document.body.style.overflow = ''
@@ -446,12 +292,6 @@ const runMermaid = async (rerun) => {
modal.addEventListener('click', (e) => {
if (e.target === modal) closeModal();
});
- document.addEventListener('keydown', (e) => {
- if (e.key === 'Escape' && modal.classList.contains('active')) {
- closeModal();
- }
- });
-
document.querySelectorAll('.mermaid').forEach((mermaidDiv) => {
if (mermaidDiv.parentNode.classList.contains('mermaid-container') ||
mermaidDiv.closest('.mermaid-fullscreen-modal')) {
@@ -471,8 +311,8 @@ const runMermaid = async (rerun) => {
const fullscreenBtn = document.createElement('button');
fullscreenBtn.className = 'mermaid-fullscreen-btn' + (darkTheme ? ' dark-theme' : '');
fullscreenBtn.setAttribute('aria-label', 'View diagram in fullscreen');
- fullscreenBtn.textContent = '⛶';
- fullscreenBtn.style.opacity = '50%';
+ fullscreenBtn.textContent = "\u26f6";
+ fullscreenBtn.style.opacity = "50%";
// Calculate dynamic position based on diagram's margin and padding
const diagramStyle = window.getComputedStyle(mermaidDiv);
@@ -484,7 +324,7 @@ const runMermaid = async (rerun) => {
fullscreenBtn.style.right = `${marginRight + paddingRight + 4}px`;
fullscreenBtn.addEventListener('click', () => {
- previousScrollOffset = [window.scroll, window.scrollY];
+ previousScrollOffset = [window.scrollX, window.scrollY];
const clone = mermaidDiv.cloneNode(true);
modalContent.innerHTML = '';
modalContent.appendChild(clone);
@@ -496,9 +336,9 @@ const runMermaid = async (rerun) => {
svg.style.width = '100%';
svg.style.height = 'auto';
svg.style.maxWidth = '100%';
- svg.style.sdisplay = 'block';
+ svg.style.display = 'block';
- if ("False" === "True") {
+ if (false) {
setTimeout(() => {
const g = svg.querySelector('g');
if (g) {
@@ -521,9 +361,12 @@ const runMermaid = async (rerun) => {
});
};
+
+
const load = async () => {
initStyles();
+
await runMermaid(true);
const reRunIfThemeChanges = async () => {
@@ -532,10 +375,8 @@ const load = async () => {
darkTheme = newDarkTheme;
console.log("Theme change detected, re-running mermaid with", darkTheme ? "dark" : "default", "theme");
await mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
await runMermaid(true);
@@ -560,37 +401,42 @@ const load = async () => {
console.log("Initializing mermaid with", darkTheme ? "dark" : "default", "theme");
mermaid.initialize(
- {...JSON.parse(
- `{"startOnLoad": false}`
- ),
- ...{ darkMode: darkTheme, theme: darkTheme ? 'dark' : 'default' },
+ {...{"startOnLoad": false},
+ ...{ darkMode: darkTheme, theme: darkTheme ? "dark" : "default" },
}
);
window.addEventListener("load", load);
window.runMermaid = runMermaid;</script>
<script>DOCUMENTATION_OPTIONS.pagename = '_modules/tvm_ffi/libinfo';</script>
+ <script>DOCUMENTATION_OPTIONS.search_as_you_type = false;</script>
<link rel="icon" href="https://tvm.apache.org/images/logo/tvm-logo-square.png"/>
<link rel="index" title="Index" href="../../genindex.html" />
<link rel="search" title="Search" href="../../search.html" />
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<meta name="docsearch:language" content="en"/>
- <meta name="docsearch:version" content="0.1.dev1+g68932438d" />
- <meta name="docbuild:last-update" content="Jul 30, 2026"/>
- </head>
+ <meta name="docsearch:version" content="0.1.dev1+gdaf594da8" />
+ <meta name="docbuild:last-update" content="Sep 15, 2026"/>
+
+ <script src="../../_static/searchtools.js"></script>
+ <script src="../../_static/language_data.js"></script>
+ <script src="../../searchindex.js"></script>
- <body data-bs-spy="scroll" data-bs-target=".bd-toc-nav" data-offset="180" data-bs-root-margin="0px 0px -60%" data-default-mode="">
-
+ </head>
+ <body data-default-mode="">
<div id="pst-skip-link" class="skip-link d-print-none"><a href="#main-content">Skip to main content</a></div>
+
<div id="pst-scroll-pixel-helper"></div>
<button type="button" class="btn rounded-pill" id="pst-back-to-top">
<i class="fa-solid fa-arrow-up"></i>Back to top</button>
-
+
+
+
<dialog id="pst-search-dialog">
@@ -616,7 +462,47 @@ window.runMermaid = runMermaid;</script>
</div>
- <header class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+ <header id="pst-header" class="bd-header navbar navbar-expand-lg bd-navbar d-print-none">
+<div class="bd-header__inner bd-page-width">
+ <button class="pst-navbar-icon sidebar-toggle primary-toggle" aria-label="Site navigation">
+ <span class="fa-solid fa-bars"></span>
+ </button>
+
+
+ <div class="col-lg-9 navbar-header-items">
+
+
+ <div class="navbar-header-items__end">
+
+ <div class="navbar-item navbar-persistent--container">
+
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+ </div>
+
+ </div>
+
+
+ <div class="navbar-persistent--mobile">
+
+<button class="btn search-button-field search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
+ <i class="fa-solid fa-magnifying-glass"></i>
+ <span class="search-button__default-text">Search</span>
+ <span class="search-button__kbd-shortcut"><kbd class="kbd-shortcut__modifier">Ctrl</kbd>+<kbd class="kbd-shortcut__modifier">K</kbd></span>
+</button>
+ </div>
+
+
+
+</div>
+
</header>
@@ -690,7 +576,7 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/tensor.html">Tensor and DLPack</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/func_module.html">Function and Module</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../concepts/exception_handling.html">Exception Handling</a></li>
-<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, and Walk</a></li>
+<li class="toctree-l1"><a class="reference internal" href="../../concepts/structural_eq_hash.html">Structural Equality, Hashing, Walking, and Mapping</a></li>
</ul>
<p aria-level="2" class="caption" role="heading"><span class="caption-text">Packaging</span></p>
<ul class="nav bd-sidenav">
@@ -717,13 +603,17 @@ window.runMermaid = runMermaid;</script>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.Map.html">tvm_ffi.Map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralKey.html">tvm_ffi.StructuralKey</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralVisitor.html">tvm_ffi.StructuralVisitor</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.StructuralMutator.html">tvm_ffi.StructuralMutator</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.VisitInterrupt.html">tvm_ffi.VisitInterrupt</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.DefRegionKind.html">tvm_ffi.DefRegionKind</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkOrder.html">tvm_ffi.WalkOrder</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.WalkResult.html">tvm_ffi.WalkResult</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.get_first_structural_mismatch.html">tvm_ffi.get_first_structural_mismatch</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_equal.html">tvm_ffi.structural_equal</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_hash.html">tvm_ffi.structural_hash</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_walk.html">tvm_ffi.structural_walk</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_visit.html">tvm_ffi.structural_visit</a></li>
+<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.structural_map.html">tvm_ffi.structural_map</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_error.html">tvm_ffi.register_error</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_object.html">tvm_ffi.register_object</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../reference/python/generated/tvm_ffi.register_global_func.html">tvm_ffi.register_global_func</a></li>
@@ -839,11 +729,18 @@ window.runMermaid = runMermaid;</script>
-<button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button pst-js-only" aria-label="Color mode" data-bs-title="Color mode" data-bs-placement="bottom" data-bs-toggle="tooltip">
- <i class="theme-switch fa-solid fa-sun fa-lg" data-mode="light" title="Light"></i>
- <i class="theme-switch fa-solid fa-moon fa-lg" data-mode="dark" title="Dark"></i>
- <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg" data-mode="auto" title="System Settings"></i>
-</button>
+<div class="theme-switch-container dropdown pst-js-only" data-bs-toggle="tooltip" data-bs-placement="bottom" title="Color mode">
+ <button class="btn btn-sm nav-link pst-navbar-icon theme-switch-button dropdown-toggle" aria-label="Color mode" data-bs-toggle="dropdown">
+ <i class="theme-switch fa-solid fa-sun fa-lg fa-fw" data-mode="light" title="Light"></i>
+ <i class="theme-switch fa-solid fa-moon fa-lg fa-fw" data-mode="dark" title="Dark"></i>
+ <i class="theme-switch fa-solid fa-circle-half-stroke fa-lg fa-fw" data-mode="auto" title="System Settings"></i>
+ </button>
+ <ul class="dropdown-menu dropdown-menu-end">
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="auto"><i class="fa-solid fa-circle-half-stroke fa-lg fa-fw me-1"></i>System Settings</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="light"><i class="fa-solid fa-sun fa-lg fa-fw me-1"></i>Light</button></li>
+ <li><button class="dropdown-item d-flex align-items-center theme-change-button" data-mode="dark"><i class="fa-solid fa-moon fa-lg fa-fw me-1"></i>Dark</button></li>
+ </ul>
+</div>
<button class="btn btn-sm pst-navbar-icon search-button search-button__button pst-js-only" title="Search" aria-label="Search" data-bs-placement="bottom" data-bs-toggle="tooltip">
@@ -880,22 +777,22 @@ window.runMermaid = runMermaid;</script>
<span class="c1"># distributed with this work for additional information</span>
<span class="c1"># regarding copyright ownership. The ASF licenses this file</span>
<span class="c1"># to you under the Apache License, Version 2.0 (the</span>
-<span class="c1"># "License"); you may not use this file except in compliance</span>
+<span class="c1"># "License"); you may not use this file except in compliance</span>
<span class="c1"># with the License. You may obtain a copy of the License at</span>
<span class="c1">#</span>
<span class="c1"># http://www.apache.org/licenses/LICENSE-2.0</span>
<span class="c1">#</span>
<span class="c1"># Unless required by applicable law or agreed to in writing,</span>
<span class="c1"># software distributed under the License is distributed on an</span>
-<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
+<span class="c1"># "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span>
<span class="c1"># KIND, either express or implied. See the License for the</span>
<span class="c1"># specific language governing permissions and limitations</span>
<span class="c1"># under the License.</span>
-<span class="sd">"""Utilities to locate tvm_ffi libraries, headers, and helper include paths.</span>
+<span class="sd">"""Utilities to locate tvm_ffi libraries, headers, and helper include paths.</span>
<span class="sd">This module also provides helpers to locate and load platform-specific shared</span>
<span class="sd">libraries by a target_name (e.g., ``tvm_ffi`` -> ``libtvm_ffi.so`` on Linux).</span>
-<span class="sd">"""</span>
+<span class="sd">"""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">__future__</span><span class="w"> </span><span class="kn">import</span> <span class="n">annotations</span>
@@ -911,108 +808,108 @@ window.runMermaid = runMermaid;</script>
<span class="k">def</span><span class="w"> </span><span class="nf">find_libtvm_ffi</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find libtvm_ffi.</span>
+<span class="w"> </span><span class="sd">"""Find libtvm_ffi.</span>
<span class="sd"> Returns</span>
<span class="sd"> -------</span>
<span class="sd"> path</span>
<span class="sd"> The full path to the located library.</span>
-<span class="sd"> """</span>
- <span class="n">candidate</span> <span class="o">=</span> <span class="n">_find_library_by_basename</span><span class="p">(</span><span class="s2">"apache-tvm-ffi"</span><span class="p">,</span> <span class="s2">"tvm_ffi"</span><span class="p">)</span>
+<span class="sd"> """</span>
+ <span class="n">candidate</span> <span class="o">=</span> <span class="n">_find_library_by_basename</span><span class="p">(</span><span class="s2">"apache-tvm-ffi"</span><span class="p">,</span> <span class="s2">"tvm_ffi"</span><span class="p">)</span>
<span class="k">if</span> <span class="n">ret</span> <span class="o">:=</span> <span class="n">_resolve_and_validate</span><span class="p">([</span><span class="n">candidate</span><span class="p">],</span> <span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">_</span><span class="p">:</span> <span class="kc">True</span><span class="p">):</span>
<span class="k">return</span> <span class="n">ret</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find libtvm_ffi"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find libtvm_ffi"</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_windows_implib</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find and return the Windows import library path for tvm_ffi.lib."""</span>
- <span class="c1"># implib = _find_library_by_basename("apache-tvm-ffi", "tvm_ffi").parent / "tvm_ffi.lib"</span>
+<span class="w"> </span><span class="sd">"""Find and return the Windows import library path for tvm_ffi.lib."""</span>
+ <span class="c1"># implib = _find_library_by_basename("apache-tvm-ffi", "tvm_ffi").parent / "tvm_ffi.lib"</span>
<span class="c1"># ret = _resolve_to_str(implib)</span>
- <span class="n">candidate</span> <span class="o">=</span> <span class="n">_find_library_by_basename</span><span class="p">(</span><span class="s2">"apache-tvm-ffi"</span><span class="p">,</span> <span class="s2">"tvm_ffi"</span><span class="p">)</span><span class="o">.</span><span class="n">parent</span> <span class="o">/</span> <span class="s2">"tvm_ffi.lib"</span>
+ <span class="n">candidate</span> <span class="o">=</span> <span class="n">_find_library_by_basename</span><span class="p">(</span><span class="s2">"apache-tvm-ffi"</span><span class="p">,</span> <span class="s2">"tvm_ffi"</span><span class="p">)</span><span class="o">.</span><span class="n">parent</span> <span class="o">/</span> <span class="s2">"tvm_ffi.lib"</span>
<span class="k">if</span> <span class="n">ret</span> <span class="o">:=</span> <span class="n">_resolve_and_validate</span><span class="p">([</span><span class="n">candidate</span><span class="p">],</span> <span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">_</span><span class="p">:</span> <span class="kc">True</span><span class="p">):</span>
<span class="k">return</span> <span class="n">ret</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find implib tvm_ffi.lib"</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find implib tvm_ffi.lib"</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_source_path</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find packaged source home path."""</span>
+<span class="w"> </span><span class="sd">"""Find packaged source home path."""</span>
<span class="k">if</span> <span class="n">ret</span> <span class="o">:=</span> <span class="n">_resolve_and_validate</span><span class="p">(</span>
<span class="n">paths</span><span class="o">=</span><span class="p">[</span>
<span class="n">_rel_top_directory</span><span class="p">(),</span>
<span class="n">_dev_top_directory</span><span class="p">(),</span>
<span class="p">],</span>
- <span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">p</span><span class="p">:</span> <span class="p">(</span><span class="n">p</span> <span class="o">/</span> <span class="s2">"src"</span><span class="p">)</span><span class="o">.</span><span class="n">is_dir</span><span class="p">(),</span>
+ <span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">p</span><span class="p">:</span> <span class="p">(</span><span class="n">p</span> <span class="o">/</span> <span class="s2">"src"</span><span class="p">)</span><span class="o">.</span><span class="n">is_dir</span><span class="p">(),</span>
<span class="p">):</span>
<span class="k">return</span> <span class="n">ret</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find home path."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find home path."</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_cmake_path</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find the preferred cmake path."""</span>
+<span class="w"> </span><span class="sd">"""Find the preferred cmake path."""</span>
<span class="k">if</span> <span class="n">ret</span> <span class="o">:=</span> <span class="n">_resolve_and_validate</span><span class="p">(</span>
<span class="n">paths</span><span class="o">=</span><span class="p">[</span>
- <span class="n">_rel_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"share"</span> <span class="o">/</span> <span class="s2">"cmake"</span> <span class="o">/</span> <span class="s2">"tvm_ffi"</span><span class="p">,</span> <span class="c1"># Standard install</span>
- <span class="n">_dev_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"cmake"</span><span class="p">,</span> <span class="c1"># Development mode</span>
+ <span class="n">_rel_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"share"</span> <span class="o">/</span> <span class="s2">"cmake"</span> <span class="o">/</span> <span class="s2">"tvm_ffi"</span><span class="p">,</span> <span class="c1"># Standard install</span>
+ <span class="n">_dev_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"cmake"</span><span class="p">,</span> <span class="c1"># Development mode</span>
<span class="p">],</span>
<span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">p</span><span class="p">:</span> <span class="n">p</span><span class="o">.</span><span class="n">is_dir</span><span class="p">(),</span>
<span class="p">):</span>
<span class="k">return</span> <span class="n">ret</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find cmake path."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find cmake path."</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_include_path</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find header files for C compilation."""</span>
+<span class="w"> </span><span class="sd">"""Find header files for C compilation."""</span>
<span class="k">if</span> <span class="n">ret</span> <span class="o">:=</span> <span class="n">_resolve_and_validate</span><span class="p">(</span>
<span class="n">paths</span><span class="o">=</span><span class="p">[</span>
- <span class="n">_rel_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
- <span class="n">_dev_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
+ <span class="n">_rel_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
+ <span class="n">_dev_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
<span class="p">],</span>
<span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">p</span><span class="p">:</span> <span class="n">p</span><span class="o">.</span><span class="n">is_dir</span><span class="p">(),</span>
<span class="p">):</span>
<span class="k">return</span> <span class="n">ret</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find include path."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find include path."</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_dlpack_include_path</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find dlpack header files for C compilation."""</span>
+<span class="w"> </span><span class="sd">"""Find dlpack header files for C compilation."""</span>
<span class="k">if</span> <span class="n">ret</span> <span class="o">:=</span> <span class="n">_resolve_and_validate</span><span class="p">(</span>
<span class="n">paths</span><span class="o">=</span><span class="p">[</span>
- <span class="n">_rel_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
- <span class="n">_dev_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"3rdparty"</span> <span class="o">/</span> <span class="s2">"dlpack"</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
+ <span class="n">_rel_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
+ <span class="n">_dev_top_directory</span><span class="p">()</span> <span class="o">/</span> <span class="s2">"3rdparty"</span> <span class="o">/</span> <span class="s2">"dlpack"</span> <span class="o">/</span> <span class="s2">"include"</span><span class="p">,</span>
<span class="p">],</span>
- <span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">p</span><span class="p">:</span> <span class="p">(</span><span class="n">p</span> <span class="o">/</span> <span class="s2">"dlpack"</span><span class="p">)</span><span class="o">.</span><span class="n">is_dir</span><span class="p">(),</span>
+ <span class="n">cond</span><span class="o">=</span><span class="k">lambda</span> <span class="n">p</span><span class="p">:</span> <span class="p">(</span><span class="n">p</span> <span class="o">/</span> <span class="s2">"dlpack"</span><span class="p">)</span><span class="o">.</span><span class="n">is_dir</span><span class="p">(),</span>
<span class="p">):</span>
<span class="k">return</span> <span class="n">ret</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find dlpack include path."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find dlpack include path."</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_cython_lib</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
-<span class="w"> </span><span class="sd">"""Find the path to tvm cython."""</span>
+<span class="w"> </span><span class="sd">"""Find the path to tvm cython."""</span>
<span class="kn">from</span><span class="w"> </span><span class="nn">tvm_ffi</span><span class="w"> </span><span class="kn">import</span> <span class="n">core</span> <span class="c1"># noqa: PLC0415</span>
<span class="k">try</span><span class="p">:</span>
<span class="k">return</span> <span class="nb">str</span><span class="p">(</span><span class="n">Path</span><span class="p">(</span><span class="n">core</span><span class="o">.</span><span class="vm">__file__</span><span class="p">)</span><span class="o">.</span><span class="n">resolve</span><span class="p">())</span>
<span class="k">except</span> <span class="ne">OSError</span><span class="p">:</span>
<span class="k">pass</span>
- <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find tvm cython path."</span><span class="p">)</span>
+ <span class="k">raise</span> <span class="ne">RuntimeError</span><span class="p">(</span><span class="s2">"Cannot find tvm cython path."</span><span class="p">)</span>
<span class="k">def</span><span class="w"> </span><span class="nf">find_python_helper_include_path</span><span class="p">()</span> <span class="o">-></span> <span class="nb">str</span><span class="p">:</span>
... 208293 lines suppressed ...