## Summary
Lifts 10 host-toolchain / CLI / process / utility modules from
`python/tvm/contrib/` to a new `python/tvm/support/` package, and
deletes two dead contrib shims.
`tvm.support` is the home for Python helpers that integrate TVM with
external CLIs and host-side tools — compilers, archivers, subprocess
pools, and build-info queries. These are load-bearing internal pieces
that TVM's compile/link/run paths depend on. `tvm.contrib` is reserved
for optional vendor SDK integrations and experimental features. The
distinction is documented in the `tvm.support` package docstring.
Moved (one commit each):
- `tvm.contrib.cc` → `tvm.support.cc`
- `tvm.contrib.nvcc` → `tvm.support.nvcc`
- `tvm.contrib.rocm` → `tvm.support.rocm`
- `tvm.contrib.ndk` → `tvm.support.ndk`
- `tvm.contrib.xcode` → `tvm.support.xcode`
- `tvm.contrib.clang` → `tvm.support.clang`
- `tvm.contrib.emcc` → `tvm.support.emcc`
- `tvm.contrib.popen_pool` → `tvm.support.popen_pool`
- `tvm.contrib.utils` → `tvm.support.utils`
- `tvm.contrib.tar` → `tvm.support.tar`
Deleted:
- `tvm.contrib.spirv` — single `optimize()` wrapping `spirv-opt`; zero
importers.
- `tvm.contrib.rpc` — self-deprecation shim with "removed in 0.5"
banner; honoring it.
Package conversion:
- `python/tvm/support.py` → `python/tvm/support/__init__.py` with
inclusion-rule docstring.
- `libinfo()` extracted into `python/tvm/support/libinfo.py`.
- `FrontendTestModule` dropped (audit confirmed zero callers outside its
own definition).
## Compatibility
Hard break — no `tvm.contrib.<mod>` re-export shims. All callers updated
in this PR.
C++-side FFI registry keys (`tvm.contrib.nvcc.*`, etc.) are unchanged —
only the Python module path moves. Renaming the FFI keys is a separate
follow-up.
This PR brings up the tirx namespace. We have been spliting out the
original tir namespace to include high-level component s_tir and this PR
updates the remaining low-level part as tirx namespace
This PR enables ruff pyupgrade (UP) rules with py310 target, auto-fixing
~5600 annotation modernizations (PEP 585 generics, PEP 604 unions,
deprecated typing imports).
Also removes from __future__ import annotations from ir/module.py and
rmsnorm.py, bumps requires-python to >=3.10, and removes absolute_import
aliases from topi/contrib files.
This PR phases out legacy target string format in favor of the json
style format that is more well formed. It also simplfies our overall
code in handling multiple formats.
This PR initalizes the s_tir for scheduable TensorIR. The change mainly
starts from python side, the we will gradually move towards the c++ side
in followup PRs. The python main change:
tir.Schedule => s_tir.Schedule
This PR renames tir.Block to SBlock. This clearly indicate the
scheduable property of the block and is a prereq for followup stir
passes refactor.
Main changes:
- Data structure change from Block to SBlock
- Syntax change from T.block to T.sblock
This PR fixes a few places where the python import of runtime tensor
is incorrect. The error wasn't revealed in the previous
NDArray->Tensor rename PR since these imports are not at the top
level.
This PR cleans up the python API to make things more consistent
with existing python array api and torch.
Device update
- device_id => index, to be consistent with torch
- device_type => dlpack_device_type() returns int
- added type property same as torch.device
API updates:
- Move the convenient method like cpu() out into tvm runtime to keep device minimal
- tvm_ffi._init_api => tvm_ffi.init_ffi_api
- tvm_ffi.register_func => tvm_ffi.register_global_func
This PR Updates the NDArray => Tensor.
Both tensor and ndarray are commonly used terms.
Because the term Tensor is getting more common in the context of ML,
we do the rename to stay more aligned with torch.Tensor and DLTensor.
* [FFI][REFACTOR] Establish tvm_ffi as a standalone python module
This PR establishes tvm_ffi as a standalone python module.
The ffi is structured as a minimal pip module that can be
directly install by path or url.
examples/get_started provided a minimal example.
This is a major change as we are decoupling tvm_ffi as a
separate package, users need to install tvm_ffi separately.
Thanks to its minimal dependency, tvm_ffi can be easily installed
even just from the source by pip install ./ffi
This change would enable future improvement for library plugins
to have lightweight dependencies by just working on top of
the tvm_ffi, while the main compiler toolchain and runtime
can be layered on top.
* [FFI] Improve traceback setups
This PR improves traceback related setups
This PR hides StringObj/BytesObj into details and bring
implementations to directly focus on the String/Bytes.
This change will prepare us for future changes such as SmallStr support.
Also moves more ObjectRef into Any in RPC.
This PR phases out tvm._ffi redirections in favor of new FFI
new functions are now called via tvm.ffi.
We also enabled limited API support for python 3.12+
so the compiled binary can be forward compatible to future
python versions.
* [REFACTOR] Phase out te.schedule python components
This PR phases out te.schedule python components.
te.compute is kept around for future usages.
tir.Schedule is a more modern version of the scheduling that we can use onwards.
Doing so also helps us to cleanup the testcases that relies on
explicit full build and execution. As we move future unit testcases
towards structural equality based unit tests.
* Simplify CI to focus on UT
The main rationale is that we should only have very few target
dependent UT in tests/python/codegen and possible
a new category in future for op-level integration if needed.
* Re-enable wasm
* fix lint
* remove hybrid,sparse autodoc and remove tests
---------
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
This PR starts the step 0 to phase out relay from the current
development main branch. This PR focuses on the python
components of relay, autotvm, auto_scheduler. To make the change
manageable, we will also do followup steps on te.Schedule and
c++ components in followup PRs.
To continue support community members who depends on
legacy flows, the [v0.19.0](https://github.com/apache/tvm/tree/v0.19.0)
branch will continue contain these components.
As noted in [discussion on phasing out legacy components](https://discuss.tvm.apache.org/t/phasing-out-legacy-components/17703/30),
this would help us to do two purposes:
- By removing outdated or redundant elements, we can significantly
reduce complexity and improve maintainability.
- Unify our focus: Concentrating our efforts on the new unity flow
will allow for more efficient development and innovation.
It is also a good opportunity for us to revisit and reduce CI time.
The past relay legacy flow contains a lot of end to end tests that
requires hardware resources to run and causing long CI time.
Moving onwards, we can focus more on unit-tests that focuses
on structural equality and runs within seconds, while be mindful
about tests that requires hardware resources (by restricting them
to specific folders and CI nightly in some cases).
---
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
In addition to the PackedFunc `"runtime.disco.worker_id"`, which
returns the worker ID wrapped in a `ShapeTuple`, this commit adds
`"runtime.disco.worker_rank"`, which returns the worker ID without
wrapping, and `"runtime.disco.device"`, which returns the device for
each worker.
The unit test added in this commit simulates loading of model weights
through a parameter transformation function.
Disco worker originally automatically import `tvm.testing.disco` for
convenient unittesting. However, `tvm.testing` is a special subpackage
that introduces many unnecessary dependencies, for example, pytest. This
PR removes such dependencies by directly moving the testing function
registration logic to the entry file.
This PR introduces `ProcessSession`, a new session implementation based
on multi-processing.
`ProcessSession` shares exactly the same communication protocol with
`ThreadedSession`, but all workers except for worker 0 are launched in a
separate process than thread. Workers communicate with the controller
via pipe provided by the OS, rather than SPSC message queue between
threads.
In our implementation, Python's `subproces.popen` is used to create
subprocesses, and the Python executable, or more specifically,
`sys.executable` calls into `tvm.exec.disco_worker` as the entrypoint.
Besides the launching logic that is only executed once in the very
beginning, the rest of the implementation resides in a C++-only
environment, including reads/writes to pipe file descriptors,
serialization and deserialization of messages, worker interpretation of
each message, etc.
Detailed engineering elements included in this PR:
- Refactors the MinRPC-based communication protocol out to be shared by
`ProcessSession` and `ThreadedSession` as `protocol.h`;
- Refactors a controller-side worker thread into `DiscoWorkerThread`,
which is shared by both session implementation to launch worker-0;
- Added two instructions `kDebugGetFromRemote` and `kDebugSetRegister`,
which are used to communicate with workers other than worker-0 in
debug mode;
- Introduces multi-processing infra including: `tvm.exec.disco_worker`
serving as the entrypoint that launches workers, and
`tvm/runtime/disco/process_pool.py` that exposes APIs to launch worker
processes. `tvm.exec.disco_worker` calls into a global function
`runtime.disco.WorkerProcess` that executes the worker main loop in
pure C++;
- Introduces `src/support/process_id.h` that provides cross-platform pid
and tid printing utilities;
- Refactors Disco's NCCL integration that get rids of initialized-once
global NCCL context, and switches to broadcasting `ncclUniqueId` from
controller to all workers, and then create NCCL communicators in each
worker thread/process accordingly. This is a thread/process-agnostic
way of using NCCL.
This PR updates the specific builtins for LM
and move them to lm_support.cc
The kv_create now takes an initial data and copies it instead of consumes it.
This will enable us to create kv within a VM more easily.
This PR makes various improvements web codegen in relax web runtime.
Correct support of shift operators.
Update relax vm to make most use of internal allocators.
Update the webgpu API to the latest spec.
* Finish support for list-of-targets
This finishes the work started in https://github.com/apache/tvm/pull/11173 to support
'external codegen' targets in the N build-like API surfaces.
- It turns out it's ok if a build is given only a single 'external codegen' target, so remove that check
in CompilationConfig::Init. When Collage builds a 'candidate partition' it does so for a single target.
As far as Collage is concerned it does not care whether the target is regular (eg Target("cuda")), or
for a specific external codegen (eg Target("cutlass")), it just passes the target into the build.
- Add CompilationConfig::FindPrimitiveTargetForKind which I'll later need to retrieve
the external codegen Target instance corresponding to a "Compiler" attribute value.
- Target.update_target_host_consist was supporting three API styles:
- single target
- map from device type to target
- map from target to IRModule (for the ir_to_runtime API)
I replaced all those calls with a more specialized 'canonicalize' call:
- Target.canonicalize_target_and_host
- Target.canonicalize_multi_targets_and_host
- Target.canonicalize_target_map_and_host
In particular, all the tuning interfaces (task extraction, tuning, tuning records) all explicitly
*do not* support multiple targets since the underlying code just doesn't support that.
* - Lints
- Revert unintended changes
* - more lints
* - Fix model_library_format handling of target.
- Improve comments in compilation_config.h
* - Lints
- Update target/target_host params documentation
* - Fix micro library format tests
- Rev micro library format from 5 to 6
- Use Target.current() in a few places
* - eta contract comprehension
* - Woops, one more device: target map left
- Handle host already being in Target
* - lint
* - lint
* - Bug with append
- Take device type from target
* - Fix hexagon
* fix bugs in the auto scheduler record:
* reformat the code
* reformat the code
* use the os.path.abspath
* change error to warning
* reformat the warning code
* fix some typos
* fix some typos
* fix some typos
* fix the port number typo
* Rename references to 0.0.0.0 to localhost. Also change references to 127.0.0.1 to localhost so that all references are consistent. 0.0.0.0 is not the same as localhost.
Previously the rpc server relies multiprocessing to start a new process and does not work under jupyter.
It also have a popen mode that does ensure the socket start listening before returning the port number.
This PR switches the implementations use PopenWorker. The port number is returned after the socket
get binded, which resolves some of the RPC flaky issues(need sleep to wait the server to start).
It also makes the RPC server jupyter friendly.
* Fix legacy code on target host
* Modify legacy code for target host change
* Add tests and fix merge issue
* Add condition for same host
* Modify all files for new target host api compatibility
* Add newline
* Change import format
* Optimize test file
* Add match error info for unit tests
* Fix for heterogeneous targets
* Fix format for dict iteration
* Fix target host type error
* Skip one testcase for tvm infinite loop bug
* Fixed bug for target map compatibility
* Fix another TargetsMap issue
* Fix typo and infinite loop error
* Temporary fix for handle issue
* Fix vm target
* Add condition support for str case
* Add GetHost function and fix previous bugs
* Fix measure_record.cc
* Fix search_task.cc
* Fix compiler.cc, memory_alloc.cc
* Fix driver_api.cc
* Fix format
* Fix bugs and GetHost function usage
* Fix clang format
* Fix bug
* Modify python tests
* Change python unit tests to new target api
* Fi test_runtime_heterogeneous.py
* Modify tutorials & remove extra print
* Update more tests to new api
* Refine the tutorial target usage
* change argument name for Target constructor function
* Fix target export function
* Fix and validate all tutorial usage
* Remove unused argument
* Fix format
* Fix bug in driver/build_module.py for heterogeneous target
* Fix bug in driver/build_module.py for heterogeneous target more
* Fix target host type error
* Fix cudnn target host bug
* Fix according to reviews, add helper function in python
* Refactor code as helper function
* Expand helper function
* Fix bug add and update python helper function
* Update target hosts
* Fix format & refresh function
* Fix unit test bug
* Fix bug in refreshing host
* Fix bug
* Add SetHost function
* Update export function
* Fix format
* Fix export bug in target
* Fix bug on host referencing
* Addtional tests
* Address review issues
* Fix format target.py
* Fix issues and format
* Add some 3rd party dependencies
* Merge main branch
* Fix target.h format
* Remove redundent import
* Fix function name
* Add parameter name
* Fix new code bug
* Fix bug in lowering
PopenPoolExecutor implements a ProcessPoolExecutor backed by popen.
- Only handles invoking functions in tvm namespace.
- Unlike multiprocessing, does not require __main__ block,
which means it can directly run on jupyter notebook.
- Come with timeout and fault tolerant support to timeout
long running jobs, and restart the process when an error happens.
Recommended usage: it is recommended to create a pool and reuse
it in a long running job(e.g. autotuning) so that the process
are reused when possible.
* [BUGFIX] Respect infinite-timed session start timeouts.
* When debugging, the intended behavior is to set the session start
timeout to infinite to allow the user to configure the debugger.
* At present, if a session start retry timeout is defined, the
current logic will bail after the retry timeout expires.
* This change makes the session start logic retry forever, once per
retry timeout.
* Document RPCEndpoint::Create.
* Add stm32f746xx to tvm.target.micro() call; fix parameter name.
* This API is expected to just be used with positional args, not
kwargs, so this change isn't expected to cause any breakage.
* model is more inline with the rest of the file, given TVM Target
Specification RFC.
* [BUGFIX] If session start fails, exit transport context manager.
* If an error occurred during session setup, then complex transports
e.g. DebugWrapperTransport would not de-initialize.
* Align transport writes/reads in TransportLogger
* fix syntax errors which were not exercised in previous PR
* Remove microTVM logic from standard RPC server, add debug shell.
* microTVM uses the host RPC server as a way to launch a debugger in
a dedicated, separate terminal window. microTVM needs to be able to
launch the debugger itself, because its model of the device
flash/debug flow separates these two things into distinct
operations implemented by shell commands (for maximum portability
across frameworks).
* microTVM can be configured to launch the debugger (e.g. GDB) in the
same terminal as is used for flashing, but this is sub-optimal
because then it hides any logs emitted by the device.
* Using the standard RPC server was hard because GDB expects the user
to issue SIGINT to interrupt program flow, but due to the RPC
server's necessary use of multiprocessing, multiple signal handlers
needed to be SIG_IGN'd, and further, because libtvm.so is
intentionally frontend-agnostic, it's difficult to include signal
handling directly in that binary (Python expects you to call
PyErr_CheckSignals, but we don't require and don't want to require
python-dev to compile libtvm.so, and this is the only such case
where libtvm.so is expected to block the main thread for a long
period of time).
* Here we implement a separate microTVM debug shell python script
using the non-blocking server implementation.
* Add serial transport, parameterize test_zephyr to work on real hardware
* add pytest test fixture, missed from previous change.
* this test fixture helps to parameterize the test case
* address leandron@ comment from #6703
* Split transport classes into transport package.
* Introduce transport timeouts.
* black format
* Add metadata-only artifacts
* Simplify utvm rpc server API and ease handling of short packets.
* add zephyr test against qemu
* Add qemu build config
* fix typo
* cleanup zephyr main
* fix nonblocking piping on some linux kernels
* don't double-open transport
* validate FD are in non-blocking mode
* gitignore test debug files
* cleanup zephyr compiler
* re-comment serial until added
* remove logging
* add zephyr exclusions to check_file_type
* add asf header
* lint
* black format
* more pylint
* kill utvm rpc_server bindings, which don't work anymore and fail pylint
* fix compiler warning
* fixes related to pylint
* clang-format again
* more black format
* add qemu regression
* Fix paths for qemu/ dir
* fix typo
* fix SETFL logic
* export SessionTerminatedError and update except after moving
* fix test_micro_artifact
* retrigger staging CI
* fix jenkins syntax hopefully
* one last syntax error
* Add ci_qemu to Jenkinsfile
* build in qemu
* address liangfu comments
* fix new bug with list passing
* retrigger CI
* [RPC] Lazily import micro when starting an RPC server
Since #6334 the RPC server cannot be started unless USE_MICRO is enabled. I've tracked this down to an import in `python/tvn/exec/rpc_server.py`: `from tvm import micro` in the top level list of imports. This will mean that we try to import micro when it's not been built. Fix this by lazily importing micro when initializing an rpc server with micro enabled.
Change-Id: I8f22d81e215cfe4ac0662b0a99bdf02a3e91f90c
* fix lint
Change-Id: I8b78b678374bc82b3b66a7b3595ed4f1684e7d90
This PR introduces WebGPU support to tvm.
The WebGPU runtime is directly built in javascript(as WebGPU uses JS as the first class citizen API)
and exposes back to the tvm's runtime via PackedFuncs.
One important note is that `ctx.sync` is not async.
This is due to the fact that WebGPU is a purely async API and we cannot block in the web environment.
So the current best way to use the js api is to wrap things in an async function.
When copy a GPU array to CPU, `await ctx.sync()` need to be called to wait for copy completion.
We use a AsyncIO rpc server to serve the async functions to the clients.
Tianqi Chen