This PR fixes invalid pointer arithmetic emitted by C-family codegen for
vector-typed `tvm_access_ptr`.
A vector access pointer is lowered to `address_of(BufferLoad(...))` with
a `Ramp` index describing its lane indices. For example, `Ramp(4, 1, 2)`
represents scalar elements `[4, 5]`, so its address should be the
address of the first lane, `&A[4]`.
LLVM codegen already extracts `Ramp::base` in this case. However,
`CodeGenC` previously passed the complete ramp to pointer arithmetic,
which could generate invalid CUDA code such as:
```cpp
(float*)A + make_int2(4, 5)
```
This PR makes `CodeGenC` use `Ramp::base` when generating the address of
a vector `BufferLoad`. The normalized index is applied to both the
direct pointer-offset path and the general `GetBufferRef` path.
The existing scalar-buffer plus `Ramp` lowering is preserved. This
avoids regressions for padded vector types such as `float32x3` and
packed vector types such as `int4x4`, while making C-family codegen
consistent with LLVM codegen.
Regression tests cover:
- `float32x2` C codegen.
- Padded `float32x3` LLVM codegen.
- Packed `int4x4` CUDA codegen.
Return is control flow, but TIRx currently represents it as an
Evaluate-wrapped intrinsic call. This prevents return values from
participating naturally in statement traversal and requires special-case
handling across the pipeline.
This change introduces a reflected tirx.Return statement carrying an
Expr, wires it through TVMScript, statement visitors and mutators,
lowering, storage planning, and C/LLVM code generation, and removes the
legacy tirx.ret and T.ret surfaces.
This PR reduces the runtime of several slow Python test groups:
- Parameterize LLVM division and CUDA vectorized-cast cases so
pytest-xdist can schedule them independently.
- Replace exhaustive ONNX execution with structural importer checks plus
representative numerical cases, and avoid registering unsupported
backend cases.
- Reuse compiled paged-attention kernels across compatible test cases.
Targeted measurements showed:
- LLVM division: 25.34s → 19.87s
- CUDA vectorized casts: 142.85s → 108.40s
- Paged-attention CPU: 316.32s → 192.50s
- ONNX Conv: 25.84s → 1.81s
- ONNX Reduce: 20.66s → 4.04s
This PR also fixes a latent CUDA Graph cleanup bug that could leave
`cudaErrorStreamCaptureInvalidated` in the worker thread and cause
unrelated subsequent GPU tests to fail.
## Summary
Compiler Targets can carry device-type semantics that runtime
device-name parsing does not preserve.
- add `tvm.device_from_target` for canonical Target-to-Device
translation
- use explicit runtime constructors where the device kind is fixed
- update target-derived utilities, tests, and documentation to use the
explicit boundary
A tvm_callback_metal_compile used purely for debugging/inspection may
return the MSL source unchanged. Previously, merely registering the
callback forced the module format to "metallib", so the runtime tried to
load the text source as a binary metallib and failed with "Invalid
library file" (issue #18798).
The callback may now return a (payload, format) pair to declare the
payload format. A bare str/bytes return keeps the legacy metallib
behavior. All kernels of a module share a single declared format, so a
callback that mixes formats across kernels (including a legacy metallib
return alongside a (payload, "metal") return) is rejected at codegen
time instead of producing a module that fails to load.
## Summary
- Keep the Python test launcher close to plain `pytest -n auto`, move
nightly tests under `tests/nightly/python`, remove obsolete launchers
and collection bookkeeping, and partition CPU/GPU jobs with explicit
`gpu` marker expressions.
- Repair exact-pointer regressions at their owning boundaries: packed
raw-string ABI values, CUDA/Metal matrix intrinsic pointers, internal TE
extern offsets, MetaSchedule scalar annotations, localized
auto-tensorization scope matching, and typed DLTensor fixture fields.
- Preserve typed workspace calls in TIR and cast pointer-returning
external calls in CodeGenC, covered by a plain-TIRx 1024-byte global
workspace that is compiled as C++.
- Finish phasing out value-bearing Relax `R.Prim` annotations by
requiring an explicit dtype, removing obsolete value-based contracts,
and expressing the DISCO rank-dependent slices as explicit scalar
`call_tir` inputs.
- Gate the distributed callback on the optional DISCO runtime, NCCL, and
at least two GPUs so capability-limited jobs skip instead of failing.
- Remove the non-demonstrating pointer probe, use direct TVMScript
comparison for packed strings, and remove the four designated legacy
testing modules.
The seven repaired CPU categories cover packed raw strings (7 failures),
CUDA/Metal matrix access-pointer types (7), internal TE extern offsets
(1), a typed DLTensor fixture (1), MetaSchedule scalar annotations (1),
CodeGenC workspace return casts (12), and localized auto-tensorization
storage-scope matching (19).
## Validation
- Base: `ded6ad8dd212869c881efb5590f8a33fc972728e`
- Head: `a7277e86dbcfe0638c8c252d36760859c4ab4297`
- All 35 locally available original failing node IDs pass across the
focused runs.
- The full focused TE, TIR builtin-lowering, and CodeGenC files pass: 61
tests.
- The complete touched Relax/TVMScript set plus
PlanAndUpdateBufferAllocationLocation passes with 784 passed, 20
skipped, and 1 expected failure.
- The DISCO callback collects and skips when its runtime or two-GPU
environment is unavailable.
- Six direct mapping tests, twelve tensor-core sketches, and the dp4a
sketch pass unchanged.
- The compiler rebuild, branch-wide pre-commit hooks, and full-range
whitespace checks pass.
- The 13 broad CBLAS/TFLite nodes remain dependency-gated; their owning
TE and generated-C regressions compile.
No merge is included in this change.
Add tvm.testing.run_with_gpu_lock backed by the existing
tvm_ffi.utils.FileLock. Migrate live local GPU tests to acquire the
machine-local lock around device execution, synchronization, host
transfer, and checks while leaving target construction and compilation
outside the critical section.
Replace the custom xdist scheduler with standard xdist_group placement
for the order-dependent test family. RPC tests retain dynamic port
allocation and per-test process isolation rather than gaining a broad
category lock.
## Rationale
`SizeVar` encodes nonnegativity in runtime subtype identity, which is
fragile under cloning and remapping. Symbolic integer values should use
one `Var` representation, with nonnegative facts recorded in the
analyzer at the use sites that establish them.
## Changes
- Remove `SizeVar` from the C++, Python, TE, TVMScript, FFI, visitor,
and serialization surfaces, and migrate callers to `Var`.
- Preserve the existing Relax constraint ownership model and use
`MarkGlobalNonNegValue` as the canonical path for global nonnegative
facts.
- Preserve `T.handle()` as the normal opaque-handle form. An optional
dtype constructs a typed pointer, with `T.handle("void")` reserved for
an explicit pointer-to-void.
## Summary
- Make `PrimExpr` a typed C++ view over `Expr` values whose
`ExprNode::ty` is `PrimType`, instead of using a separate runtime node
class as the proof of primitive-ness.
- Use the shared `ir::Call` node for Relax, TIRX, and primitive-valued
calls, while keeping primitive-only APIs explicit at their semantic
boundaries.
- Keep Python on the general `Expr` surface for primitive-typed values
so `isinstance` behavior does not imply a nominal primitive-expression
subclass.
## Design Rationale
The main advantage of this change is that common expression nodes such
as `Call` can be unified without specializing each one to `PrimType`. A
single `ir::Call` can represent a Relax tensor call, a Relax scalar
call, or a primitive-valued intrinsic call; the result type stored in
`ExprNode::ty` determines whether that particular value can be viewed as
`PrimExpr`.
This keeps the IR node hierarchy focused on expression structure rather
than result-type categories. Nodes that are intrinsically primitive,
such as integer and floating-point literals or TIRX primitive operators,
still have strongly typed C++ APIs and data structures. General nodes
whose result type may vary, such as `Call`, remain general `Expr` nodes
and are narrowed to `PrimExpr` only where primitive-only semantics are
required.
The PR also keeps the compatibility surface practical: C++
primitive-only APIs continue to accept `PrimExpr`, Python exposes a
compatibility predicate for checking the primitive typed category, and
visitors/printers use one natural `Call` path rather than duplicating
Relax and primitive call handling. Missing expression types are
represented explicitly with `Type::Missing()` so constructors can leave
type inference to later analysis without relying on nullable `Type`
values.
The existing Ramp lowering path in CodeGenLLVM constructs fixed-width
vectors by inserting each lane explicitly. This does not work for
scalable vectors, whose runtime lane count is not known at compile time.
Previously, CodeGenLLVM rejected scalable-vector Ramp expressions. This
prevents vectorized TIR/TIRx programs from lowering induction
expressions to RVV/SVE-style scalable vectors.
This patch adds a separate lowering path for scalable integer Ramp
expressions using LLVM stepvector. A scalable Ramp expression:
Ramp(base, stride, lanes)
is lowered as:
splat(base) + stepvector() * splat(stride)
For LLVM >= 20, this uses llvm.stepvector. For older LLVM versions, this
uses llvm.experimental.stepvector.
A RISC-V RVV codegen test is added to verify that a vectorized induction
expression lowers to RVV lane-id and arithmetic instructions, such as
vid.v, vmul.v*, and vadd.v*.
In the past we have been using `DataType` in PrimExpr.dtype field to
check type information for PrimExpr while still having BaseExpr.ty for
richer type information. DataType is also used both in runtime and
compiler. This PR streamlines the boundary:
- PrimExpr.ty now carries PrimType that replaces original use of
`DataType`
- Runtime use will now favor DLPack DLDataType, removing one layer of
indirection.
- Constants attributes where values are usually runtime values, will use
`DLDataType`
- DataType will be phased out after this PR
We also brings up helper functions in PrimType, but also limits them to
a more concise set so the functions do not grow with the data type codes
in DLPack.
This is a major refactor that changes the IR primitive. It helps to
bring possible future benefits:
- Unified type mechanism through Expr.ty
- Possibility of carry future Type nodes
Migration Guide:
- Use `PrimType` when code reasons about compiler expression types,
tensor element compiler types, or constructs a `PrimExpr`/compiler type.
- Use existing source types such as `expr.ty()`, `ExprOp.expr_ty()`, or
TE tensor element `dtype` where possible instead of rebuilding a type
from dtype text.
- Use raw `DLDataType` for runtime constants, ABI paths, dtype-valued
attrs, and storage/runtime helper logic.
- Prefer direct `PrimType` equality, `MatchesCode(...)`,
`MatchesElementType(...)`, and `WithCode(...)` over local wrappers or
string dtype checks.
Performance:
Using Object type instead of DLDataType would indeed bring some
performance impact to the IR. We have done the following performance
optimizations:
- Make sure most of the outputs reuse one of the PrimType from inputs
- Cache a thread local PrimType based on input so we don't repeatly
realloc
We did benchmarks show that rewrite simplify operation stays within
+-10% overhead of original one. Which merits the refactor given the
benefit the unfication brings
## Summary
- unify Relax's former StructInfo surface into the Type vocabulary and
Expr.ty storage path
- remove leftover DependentTypeNode and legacy OpNode::op_type storage
- keep base Type nullable while concrete Relax/DTensor type refs are
non-nullable
- clean stale StructInfo/TensorStructInfo/sinfo vocabulary in
Python/docs and distributed-op macros
- address Gemini follow-ups for parser annotations, BlockBuilder
docstring, and Adreno TensorType cast audit
This pr phases out the custom `tvm.testing.parameters()` helper in favor
of native `pytest.mark.parametrize`. `parameters()` itself is left in
place for now and removed in a follow-up, together with updating the
framework self-test
(`tests/python/testing/test_tvm_testing_features.py`) that exercises it.
Migration rules
- A group consumed only by test functions becomes
`pytest.mark.parametrize`.
- Single-name groups are unwrapped from 1-tuples to bare values.
- A group shared by multiple tests uses a module-level named list; a
test that uses only a subset of a group's names is parametrized only on
the names in its signature.
- `pytest.mark.parametrize` is stacked above the existing, unrelated
`tvm.testing.parametrize_targets(...)`, which is kept as-is.
Per-file pytest collection case counts are unchanged, except the two
intentional changes below.
Behavior changes (intentional)
- tests/python/relax/test_training_optimizer_numeric.py: the names `lr`
and `weight_decay` were rebound across three `parameters()` groups, so
`test_sgd` and `test_momentum_sgd` silently used the *adam* group's
`lr`/`weight_decay` (and `test_momentum_sgd` cross-producted with it:
2/6/2 = 10 cases). Native parametrize gives each test its own co-located
group: 2/3/2 = 7 cases. This fixes that latent rebinding bug; the case
count drops 10 -> 7 and `test_momentum_sgd` now exercises its own
`weight_decay` values.
- tests/python/target/test_arm_target.py: its `parameters()` group was
orphaned (no test consumed those names) — removed the dead definition.
Note: for tests that also use `tvm.testing.parametrize_targets`, the
generated test ids reorder the target (e.g. `test_unary[abs-True-llvm]`
-> `test_unary[llvm-abs-True]`); values and case counts are unchanged.
## Why
ASF INFRA enforces that external GitHub Actions must be pinned to a
commit SHA on the approved allowlist, failing the workflow with "not
allowed in apache/tvm". See the
[policy](https://infra.apache.org/github-actions-policy.html) and the
[approved
allowlist](https://github.com/apache/infrastructure-actions/blob/main/approved_patterns.yml).
## How
- Pin `pre-commit/action` to `2c7b3805fd2a0fd8c1884dcaebf91fc102a13ecd`
(v3.0.1)
- Pin `pypa/cibuildwheel` to `294735312765b09d24a2fbec22660ce817587d55`
(v4.1.0)
- Pin `pypa/gh-action-pypi-publish` to
`ed0c53931b1dc9bd32cbe73a98c7f6766f8a527e` (v1.13.0)
- Leave GitHub-owned `actions/*` and the allowlisted
`conda-incubator/setup-miniconda@*` pattern untouched
---------
Signed-off-by: Guan-Ming (Wesley) Chiu <105915352+guan404ming@users.noreply.github.com>
This PR improves TIRx vectorization for RISC-V RVV targets.
Fixed-width `T.vectorized` loops can be lowered to fixed LLVM vectors
such as `<16 x float>`, which LLVM/RVV may scalarize into repeated
scalar `flw/fsub.s/fsw` instructions. This PR rewrites fixed-width
vectorized loops on RVV targets into scalable `T.vscale() * 4` chunks
with lane masks, allowing LLVM to generate RVV load/store instructions
instead.
The change is limited to RISC-V RVV and does not enable the same
automatic rewrite for Arm SVE.
Tested on a RISC-V K3 board:
Before: flw/fsub.s/fsw = 16/16/16, vle32/vse32 = 0/0
After: flw/fsub.s/fsw = 0/0/0, vle32/vse32 = 1/1
Also added a RISC-V LLVM codegen regression test.
This pr modernizes test gating. It replaces the heavy
`tvm.testing.Feature` machinery with a thin `tvm.testing.env` module of
`has_*()` capability probes, used via standard pytest.mark + skipif. And
markers move to `pyproject.toml`
`tests/python/codegen/test_target_codegen_aarch64.py` cross-compiles
AArch64 SVE kernels and regex-matches the generated assembly for
specific instruction forms and counts. Several of those assertions
encode the exact code shape produced by the LLVM versions used in CI
(15-17). On a TVM built against LLVM 20 the tests fail, even though the
emitted IR is correct (+sve target-features and vscale_range(1,16) are
present) -- the difference is entirely inside LLVM's loop vectorizer /
cost model, not in TVM's codegen.
Since LLVM 19, SVE/SVE2 are optional extensions of Armv9.0-A
(llvm/llvm-project#96007), so "+v9a" no longer implies "+sve" and
CodeGenAArch64 does not add the vscale_range() function attribute. Gate
the expectation in test_vscale_range_function_attribute on
llvm_version_major() < 19 so the test passes on both older LLVM (15-17,
as used in CI) and LLVM 19+.
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.
The generic-target tests in test_target_codegen_vulkan.py are
auto-parametrized over all enabled targets, which may include nvptx. For
nvptx, TVM produces PTX codegen output but not a directly launchable
runtime module, so executing the compiled function fails with errors
such as cuModuleGetFunction CUDA_ERROR_NOT_FOUND.
Add a small helper that skips runtime execution for nvptx after a
successful compile, so codegen is still exercised while the invalid
runtime launch is avoided. Vulkan-only tests are unchanged.
Newer LLVM versions (observed with LLVM 20) print a scalable broadcast
store as a splat constant, e.g.
`store <vscale x 4 x float> splat (float 1.000000e+00)`, instead of the
older `shufflevector (<vscale x 4 x float> insertelement (...` form.
Accept either representation in test_scalable_broadcast so the test
passes across LLVM versions while still verifying scalable vector
codegen.
---------
Co-authored-by: tqchen <tianqi.tchen@gmail.com>
Validating the apache-tvm wheel in a minimal environment (no torch,
scipy, cloudpickle, or tornado installed) produced 33 pytest collection
errors from module-level imports of optional packages. Add
pytest.importorskip guards so these modules are reported as skipped
instead of erroring during collection.
Indirect import chains guarded:
- tvm.topi.testing imports scipy
- tvm.s_tir.meta_schedule.testing.local_rpc (tvm.rpc.tracker) requires
tornado
- tvm.s_tir.dlight.benchmark imports cloudpickle
Also remove a stray pre-license-header `import pytest` in
test_runtime_builtin_paged_attention_kv_cache_flashinfer.py.
## 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 updates CUDA fast math intrinsic lowering to use a PassContext
option instead of a CUDA Target attribute.
The new option is:
```python
with tvm.transform.PassContext(config={"tirx.enable_fast_math": True}):
...
```
When unset or false, CUDA math intrinsics continue to lower to the
precise CUDA math functions such as expf. When true, tirx.LowerIntrin
prioritizes the cuda.fastmath.* lowering rules, producing fast math
intrinsics such as __expf.
Fix CUDA lowering of standard TIR math intrinsics so they use precise
CUDA math functions by default instead of fast-math `__*f` functions.
This fixes the default behavior reported in #19546, where operators such
as `tirx.exp` could lower to `__expf` even though fast math was not
explicitly requested.
This change adds a CUDA target attribute, `enable_fast_math`, which
defaults to `false`. When the attribute is unset or false, standard math
intrinsics lower through the normal CUDA math rule, for example `expf`,
`logf`, `sinf`, `cosf`, `powf`, and `rsqrtf` for `float32`. When users
explicitly enable the attribute on the target, the lowering pass also
checks the `cuda.fastmath.FLowerIntrinsic` rules before the normal CUDA
lowering rules.
Users can opt in to fast math by constructing a CUDA target with the
attribute:
```py
tvm.target.Target({"kind": "cuda", "enable_fast_math": True})
target = tvm.target.Target({
"tag": "nvidia/nvidia-a100",
"enable_fast_math": True,
})
```
The fast-math lowering path currently covers the CUDA math operators
registered with `cuda.fastmath.FLowerIntrinsic`: `tirx.exp`,
`tirx.exp10`, `tirx.log`, `tirx.log2`, `tirx.log10`, `tirx.tan`,
`tirx.cos`, `tirx.sin`, `tirx.tanh`, and `tirx.pow`.
`tirx.rsqrt` is also registered for CUDA lowering so it maps to the CUDA
reciprocal-square-root intrinsic instead of being legalized as `1 /
sqrt(x)`.
Add CUDA codegen tests
`tests/python/codegen/test_target_codegen_cuda_fastmath.py` that check
the lowered IR, generated CUDA source, and runtime results for the
supported math intrinsics across floating point dtypes and both default
and fast-math targets.
## Summary
This PR adds the initial TIRx support needed for low-level programming
of Blackwell-class GPU architectures. As part of the ongoing TIRx
refactor, it introduces TVMScript support for directly scripting
advanced hardware features without relying on scheduling as the primary
programming interface.
The change keeps existing `s_tir` script support intact while making
direct scripting a first-class path for TIRx programs.
## Main Changes
- Add TIRx operator dispatch and layout infrastructure.
- Add TVMScript support for new low-level TIRx operations.
- Add analysis, transform, and lowering support for TIRx IR nodes.
- Add CUDA/Blackwell-oriented codegen and intrinsic coverage.
- Add Python and C++ integration points for TIRx scripting and runtime
support.
## Validation
- `pre-commit run --all-files`
- `ninja -C build -j32`
- `CUDA_VISIBLE_DEVICES=2 pytest tests/python/tirx/ -n 16`
- `1723 passed, 47 skipped, 32 warnings`
- `CUDA_VISIBLE_DEVICES=2 python -m pytest -v
tests/python/all-platform-minimal-test`
- `37 passed, 105 skipped`
- `TVM_TEST_TARGETS=llvm python -m pytest -v tests/python/tirx-analysis
tests/python/tirx-base tests/python/tirx-transform -n 16`
- `664 passed, 25 skipped, 9 xfailed, 1 xpassed`
## Local CI Notes
Some full CI-equivalent jobs were not locally reproducible because this
machine is missing parts of the Apache TVM CI environment, including
`llvm-config-15/17`, Vulkan, ROCm, Maven, Sphinx, Doxygen, Emscripten,
and ARM/QEMU cross-toolchain components. Metal-specific tests were
skipped locally because no Metal runtime is available.
## Why
This refactor reshapes each backend into a self-contained
`src/target/<X>/`
cluster (with optional `src/target/<X>/llvm/` for LLVM-dependent
codegen) and
introduces a per-backend fallback module that absorbs the cross-compile
role
cleanly — without `target/opt/` stubs, without `DeviceSourceModuleNode`,
and
without leaking a synthetic `kind()` to consumers.
## High-level principles
- **One directory per backend.** All codegen-side files for backend
`<X>` live
under `src/target/<X>/`. Optional `src/target/<X>/llvm/` subdir for
files
that require `USE_LLVM` at build time. Backend grouping wins over
build-dependency grouping (the latter being upstream's `target/llvm/`).
- **Plugin-only runtime modules.** `src/runtime/<X>/<X>_module.h` is
deleted.
The runtime's real `<X>ModuleNode` is reachable only via the FFI
registry
(`ffi.Module.create.<kind>`, `ffi.Module.load_from_bytes.<kind>`). No
C++ API surface other than the static registrations.
- **Per-backend fallback module for cross-compile.** Each `<X>` gets a
`<X>FallbackModuleNode` in `src/target/<X>/<X>_fallback_module.{h,cc}`.
Same `kind()` as the real module. Codegen-time only — never reachable
via
load. `GetFunction` errors with a backend-specific "runtime not linked"
message; `InspectSource` works.
- **Codegen-side wrapper does the fallback selection.** Codegen calls
`<X>ModuleCreateWithFallback(...)`, which tries
`ffi.Module.create.<kind>`
via the registry; on miss, falls through to `<X>FallbackModuleCreate`
(plain C++; reachable directly from the fallback header). When
`USE_<X>=ON`
is in effect, the registry hit returns the real module; when
`USE_<X>=OFF`, the fallback is what codegen gets. No CMake `if/else`
gating; fallback always compiled.
## Specific changes
### New per-backend directories (codegen + fallback)
- `src/target/cuda/` — `codegen_cuda.cc` + `intrin_rule_cuda.cc` +
fallback module pair + `llvm/codegen_nvptx.cc`
- `src/target/rocm/` — fallback module pair + `llvm/codegen_amdgpu.cc` +
`llvm/intrin_rule_rocm.cc`
- `src/target/hexagon/` — fallback module pair +
`llvm/codegen_hexagon.cc` + `llvm/intrin_rule_hexagon.cc`
- `src/target/metal/` — `codegen_metal.cc` + `intrin_rule_metal.cc` +
fallback module pair
- `src/target/vulkan/` — `build_vulkan.cc` + the rest of `target/spirv/`
absorbed + fallback module pair
- `src/target/opencl/` — `codegen_opencl.cc` + `intrin_rule_opencl.cc` +
fallback module pair
- `src/target/webgpu/` — `codegen_webgpu.cc` + fallback module pair (was
`WebGPUSourceModuleNode`, renamed)
### New fallback classes
`<X>FallbackModuleNode` for X in {`CUDA`, `ROCm`, `Hexagon`, `Metal`,
`Vulkan`, `OpenCL`, `WebGPU`}. Each:
- `kind()` matches the real backend
- Stores `(code or smap, fmt, fmap, source)` — no driver/runtime calls
- `GetFunction` errors with backend-specific "runtime not linked"
message
- `InspectSource` works fully
- `SaveToBytes` byte-identical to real
SPIR-V codegen currently emits `ArrayStride` and `Offset` decorations
for non-interface allocations in `GetStructArrayType()`. That is correct
for descriptor-backed interface blocks, but not for static workgroup
allocations.
I hit this while bringing up tilelang vulkan shared memory allocation
path: the vulkan validation rejected shaders that used shared memory
lowered through this path:
```
tvm.error.InternalError: Check failed: res == SPV_SUCCESS (-10 vs. 0) :
index=44 error:[VUID-StandaloneSpirv-None-10684] Invalid explicit layout decorations on type for operand '25[%_ptr_Workgroup__struct_24]'
%A_shared = OpVariable %_ptr_Workgroup__struct_24 Workgroup
```
FIX: This PR keeps layout decoration for interface blocks, and skips for
non-interface allocations such as static shared/workgroup memory. A new
compile-only test is added for this.
One possible concern is that there's already a pre-existing test using
`fetch_to_shared`.
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
## Summary
- Remove `body` field from `AllocBufferNode` and `DeclBufferNode`,
making them flat statements consistent with `Bind`
- Buffer scope extends to end of enclosing scope via flat `SeqStmt`
semantics
- 60 files changed across core IR, codegen backends, transforms, script
IR builder, and tests
## Test plan
- All existing test suites pass (tir-transform, tir-base, tvmscript,
s_tir, codegen, C++)
## Summary
This PR introduces `AllocBufferNode`/`AllocBuffer` as a single TIR
statement that both allocates memory and declares a buffer into scope.
This replaces the previous pattern of `Allocate(var, dtype, shape, cond,
DeclBuffer(buf, body))` with the simpler `AllocBuffer(buf, body)`.
### Main changes
- **New IR node** `AllocBufferNode` with fields `{buffer, annotations,
body}` — same semantics as `DeclBuffer` but also allocates memory
- **TVMScript**: `T.alloc_buffer(shape, dtype, scope)` now emits
`AllocBuffer` directly (statement-level allocation).
`T.sblock_alloc_buffer(...)` for SBlock-level buffer allocation (full
parameter set)
- **All codegen backends** (C, CUDA, Metal, OpenCL, WebGPU, LLVM, NVPTX,
AMDGPU, SPIR-V) updated to handle `AllocBufferNode`
- **All TIR transforms** (storage_rewrite, flatten_buffer,
vectorize_loop, lower_warp_memory, etc.) updated
- **All S-TIR transforms** (compact_buffer_region, merge_shared_memory,
inject_double_buffer, etc.) updated
- **Removed `AllocateNode`** entirely — `AllocBuffer` is now the sole
allocation primitive
- **Removed `AllocDescriptor`** from merge_shared_memory_allocations —
uses `Buffer` objects directly
- **Added `AllocBuffer::ConstantAllocationSize()`** inline helper method
### Design rationale
The old `Allocate + DeclBuffer` pair was a historical artifact:
`AllocateNode` stored raw fields (`buffer_var`, `dtype`, `extents`,
`condition`) separate from the `Buffer` object, requiring pattern
matching (`IsAllocateDeclBufferPattern`) to reconstruct the buffer
association. `AllocBuffer` unifies this into a single node with a proper
`Buffer` reference, simplifying codegen backends and transform passes.
225 files changed, ~3500 insertions/deletions (net near-zero, mostly
mechanical migration).
## Test plan
- [x] All TIR base tests pass
- [x] All TIR transform tests pass
- [x] TVMScript roundtrip tests pass
- [x] S-TIR transform tests pass
- [x] Codegen tests pass
- [x] All-platform minimal tests pass
- [x] C++ functor tests pass
- [x] Pre-commit clean (clang-format, ruff, etc.)
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 removes unused `from tvm import te` imports from 25 test files
across the codebase, continuing the ongoing TE → TVMScript migration
cleanup.
Changes:
- Remove unused `from tvm import te` from 24 test files in
tir-transform/, s_tir/transform/, codegen/, arith/,
all-platform-minimal-test/, and testing/
- Replace `te.var("x")` with `tvm.tir.Var("x", "int32")` in
`test_s_tir_transform_decorate_device_scope.py` (the only file where te
was actually used)
- Clean up stale `tvm.tir.ir_builder` comment references in
`test_tir_transform_convert_ssa.py`
This PR cleans up the target python api.
- Removes the indirections of attribute exposure
- Move tag registry to python so it is easily configurable
- Remove legacy constructors in favor of tags
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 migrates the s_tir related transform passes into s_tir namespace
instead. This set of changes can minimize the overall tir namespace to
make it more focused.
This PR migrates all the codegen tests to explicitly using tvmscript
instead of indirectly via s_tir.Schedule. They makes the test surface
more unit, contains less dep and more maintainable.
Shushi Hong