7 Commits

Author SHA1 Message Date
Bohan Hou 4224d51090 [TIRx] Bundle CUDA tile primitive and op dispatch updates (#19896)
## Summary

This bundles the 18 commits currently carried in `spectrometerHBH/tvm`
on top of `apache/tvm:main`.

Major areas:

- Extend CUDA TIRx tile primitives and op dispatch paths, including
vector PTX ld/st, shared-memory copy paths, TMA/tcgen05 descriptor
handling, dense FP8/TF32 `gemm_async`, and CUDA elementwise tile
dispatch.
- Add support utilities for benchmark timing, CUDA ptxas option
plumbing, and TMA/TFLOAT32 descriptors.
- Fix unsigned integer floormod/floordiv simplification rewrites without
overflow and update the corresponding TIRx constant-folding tests.
- Update TIRx dtype handling for upstream `PrimType` compatibility.
- Add and update TIRx CUDA/operator tests for copy, elementwise, permute
layout, and `gemm_async` behavior.

## Validation

- `git diff --check apache/main..HEAD`
- `python -m tirx_kernels.bench_suite --check-imports`
- `python -m tirx_kernels.registry --cc 10 --strict`
- `python -m pytest tests/python/tirx/ -n 16`
  - `2033 passed, 39 skipped, 3 xpassed`
- `python -m pytest tests/python/tirx-base/test_tir_imm_values.py -q`
  - `44 passed, 6 warnings`
- `pre-commit run --files tests/python/tirx-base/test_tir_imm_values.py`
- Focused TIRx regression tests after formatting:
  - `test_cast_vec2_packed_dispatch`
  - `test_cast_warpgroup_src_layout_to_flat_uses_vec2_intrinsic`
  - `test_gemm_tcgen05_cta_group_1[task0]`
- Full `bench_suite --impls all` sweep: 256/256 workloads completed
successfully.
- Apache PR CI on `928a0605d0`: all required GitHub Actions and Jenkins
checks passed.
2026-06-29 00:11:23 -04:00
Bohan Hou 16d0a7edae [TIRX][CUDA] Framework support for FA4, CLC intrinsics, and nvfp4 tcgen05 GEMM (#19785) 2026-06-16 06:52:14 -04:00
Shushi Hong 4899fa192d [CUDA] Narrow the cuda extra from cuda-python to cuda-bindings (#19784)
TVM's shipped code only uses cuda.bindings — cuda.bindings.nvrtc for the
NVRTC JIT path and cuda.bindings.driver for the NVSHMEM link path, both
in python/tvm/support/nvcc.py; it never uses cuda.core. cuda-python is
now a metapackage that pulls in cuda-bindings + cuda-core (and
cuda-pathfinder), so depending on it drags in cuda-core that TVM does
not need.
Depend directly on cuda-bindings, which provides exactly the nvrtc and
driver submodules TVM imports, and update the user-facing 'pip install
cuda-python' hints to match. A plain cuda-bindings install pulls no
nvidia-* toolkit wheels (those live behind the [all] extra); libnvrtc is
loaded from the system / TVM's CUDA install as before.
2026-06-15 18:47:33 -04:00
Tianqi Chen 96cba60464 [PYTHON] Autoload backends; simplify library loading; remove TVMError for native errors (#19727)
This PR adds an autoload mechanism for out-of-tree backends, simplifies
TVM's Python library loading, and removes `TVMError` in favor of native
Python errors.

## Autoload out-of-tree backends

Out-of-tree packages can register an autoload callable under the
`tvm.backends` entry-point group (mirroring torch's device-backend
autoload). At `import tvm` startup each entry point is discovered and
its callable invoked once, after the core runtime and the `tvm`
namespace are fully initialized, so an extension can register
ops/targets/funcs or load extra libraries.

```toml
[project.entry-points."tvm.backends"]
tvm_foo = "tvm_foo:_autoload"
```

A failing extension is caught and surfaced via `warnings.warn` so it
cannot break `import tvm`. Autoload can be disabled with
`TVM_DEVICE_BACKEND_AUTOLOAD=0`.

## Simplify library loading

The library-loading path in `base.py` is consolidated around a single
`_LOADED_LIBS` dict (basename to ctypes handle) so downstream and
autoloaded extensions can skip already-loaded libraries; the per-backend
runtime DSO list is folded into `load_backend_libs`. Accumulated cruft
is removed: the Python-3.9 check, the readline shim, the `_FFI_MODE`
ctypes check, the `base.__version__` re-export, and `py_str` (call sites
inline `.decode("utf-8")`).

## Remove TVMError in favor of native Python errors

`TVMError` added a layer atop `RuntimeError` that downstream code had to
import and learn. It is removed; the registered FFI error kinds
(`InternalError`, `RPCError`, `OpError`, `DiagnosticError`,
`ScheduleError`) now subclass `RuntimeError` directly while staying
registered, so the FFI keeps throwing the right kinds. All `TVMError`
imports, `except`/`raise`/`isinstance` uses, and
`pytest.raises(tvm.TVMError)` sites move to the `RuntimeError` builtin.
2026-06-11 13:50:38 -04:00
Tianqi Chen ebc9348bd7 [REFACTOR][PYTHON] Slim tvm.libinfo to info-only helpers (#19719)
This PR slims `tvm.libinfo` into a thin *info* layer that delegates path
discovery to the `tvm_ffi.libinfo` primitives and never loads libraries.
Loading responsibilities move to `tvm.base`, and the various ad-hoc
path-finding helpers are phased out in favor of the tvm-ffi resolvers.

## Changes

- **libinfo**: add `find_libtvm_runtime()` (resolves `libtvm_runtime`
via
  `_find_library_by_basename` + `_resolve_and_validate`) and
`find_tvm_include_path()` (TVM's own `include/`). `find_include_path()`
now
returns `[find_tvm_include_path(), *tvm_ffi.libinfo.include_paths()]`,
folding
in the FFI + dlpack + python-helper include dirs. Remove
`find_lib_path`,
  `get_dll_directories`, `use_runtime_lib`, `split_env_var`, and
  `load_backend_libs`.
- **base**: receive `load_backend_libs` and the backend DSO list; the
runtime-only switch becomes a strict `TVM_USE_RUNTIME_LIB == "1"` check.
- **rpc**: `with_minrpc` uses `find_libtvm_runtime()` (the `runtime`
kwarg is
retained as an inert back-compat parameter); the rpc server
`load_library`
resolves the literal library name against the current working directory.
- **wasm**: move the `web/dist` asset search into `emcc.find_wasm_lib`,
used by
  `emcc.create_tvmjs_wasm` and the tvmjs asset lookup.
- **hexagon**: fix a latent bug where `_get_hexagon_rpc_lib_dir` called
a
non-existent `tvm_ffi.libinfo.find_lib_path`; it now relies solely on
the
  `HEXAGON_RPC_LIB_DIR` environment variable.
2026-06-10 13:41:37 -04:00
Balint Cristian 349225ae23 [REFACTOR][PYTHON] Revisit lifted support modules from tvm.contrib (#19653)
In continuation of #19624 this catches some unlifted entries.

Hope there is no more left, for consistency it now covers comments and
perhaps non-active (hotpath) parts.
2026-06-01 08:21:07 -04:00
Tianqi Chen ffea531107 [REFACTOR][PYTHON] Lift compiler/CLI/process modules from tvm.contrib to tvm.support (#19624)
## 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.
2026-05-27 15:31:12 -04:00