628 Commits

Author SHA1 Message Date
Guan-Ming Chiu 277ae41efa [Relax] Legalize grouped conv with symbolic channels (#20039)
- `LegalizeOps` skips grouped `conv1d/2d/3d` when channel size is
symbolic
- The only blocker is `topi.nn.conv`'s divisibility `assert`s, which
fail on symbolic `PrimExpr`; the grouped compute already handles
symbolic dims

## How

- Skip the divisibility check when the channel size is not a constant
int
- Remove the symbolic-channel guards from the conv legalize functions
2026-07-23 15:31:55 -04:00
Tianqi Chen c717c5b217 [IR][Relax][TIRx] Unify Var identity (#20004) 2026-07-15 17:12:40 +08:00
Tianqi Chen 45e1b8233a Refactor Tensor arithmetic dispatch away from tirx.generic (#19943)
## Summary

- Move whole-Tensor arithmetic and cast dispatch onto `te.Tensor` while
scalar TIRx smart constructors decline whole-Tensor operands.
- Remove the legacy `tirx.generic` module, TOPI import-time mutation
bridge, and obsolete aliases.
- Migrate scan and cast callers while preserving identity-gated Thrust
sum selection.

Whole-Tensor behavior now lives with TE, leaving scalar TIRx
construction independent of TOPI initialization.
2026-07-04 20:18:43 -04:00
Tianqi Chen 99869414de [TIRX] Remove SizeVar in favor of contextual constraints (#19930)
## 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.
2026-07-03 11:33:14 -04:00
Tianqi Chen 275114b327 [REFACTOR][IR] Unify PrimExpr with Expr typed view (#19910)
## 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.
2026-07-01 18:55:33 -04:00
Guan-Ming Chiu 6383c7fd7f [Relax][ONNX] Support 3D AffineGrid (#19863)
## Related Issue

closes #19689

## Why

The Relax AffineGrid op only handled 2D (4D theta/grid); 5D 3D inputs
from ONNX failed.

## How

- Generalize struct-info inference to 2D/3D via spatial =
size_sinfo->ndim.
- Branch TOPI affine_grid compute on 2D vs 3D.
- Add the 3D permute path in the frontend and a test_affine_grid_3d
case.

---------

Signed-off-by: Guan-Ming (Wesley) Chiu <105915352+guan404ming@users.noreply.github.com>
2026-06-30 15:35:22 -04:00
Tianqi Chen 1e1920bcbd [REFACTOR][IR] Unify PrimExpr type mechanism to PrimType instead of DataType (#19875)
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
2026-06-24 21:31:47 -04:00
Guan-Ming (Wesley) Chiu 4650887f6a [Relax][ONNX] Support align_corners in AffineGrid op (#19864)
## Related Issue

closes #19690

## Why

ONNX AffineGrid carries an align_corners attribute, but the Relax op
ignored it and always produced the align_corners=1 grid.

## How

- Add align_corners field to AffineGridAttrs (mirrors GridSampleAttrs).
- Thread the flag through the op, legalize pass, and TOPI compute.
- Pass the ONNX attribute through in the frontend instead of dropping
it.
2026-06-22 12:46:36 +03:00
Tianqi Chen 878903f574 [REFACTOR][IR] Delete class Bool and class Integer boxed-type wrappers (#19636)
## Background

`class Integer : public IntImm` and `class Bool : public IntImm` were
thin
wrappers sharing `IntImmNode` with no separate node class and no FFI
registration. They existed to provide implicit int→Integer constructors
and
a `.IntValue()` / `operator bool()` accessor, but the same functionality
is
available directly through `IntImm`.

## What this PR does

Migrates all call sites away from `Integer` / `Bool` and then deletes
the
class definitions.  The changes are split into four commits, each
independently buildable:

**Commit 1 – [REFACTOR][TIR]** Replace IR-position `Integer(N)` /
`Bool(b)`
constructors with `IntImm(DataType::Int(32), N)` /
`IntImm(DataType::Bool(), b)`
across ~62 source files (arith, relax analysis, s_tir schedule state,
transform
passes, codegen).

**Commit 2 – [REFACTOR][SCHEDULE]** Migrate `Schedule` and
`MetaSchedule`
trace-boxing code: `Integer(N)` attrs in `TracedSchedule` →
`IntImm(DataType::Int(32), N)`;
`ffi::Array<Integer>` schedule-rule parameters → `int64_t`; `Bool(b)`
attrs →
`IntImm(DataType::Bool(), b)`.

**Commit 3 – [REFACTOR][TOPI]** Migrate topi container signatures
(`ffi::Array<Integer>` → `ffi::Array<int64_t>`) and update all internal
usages (`.IntValue()` → plain int64_t, `.defined()` → removed,
`->value` → direct indexing).  Also handles stray `Integer` / `Bool`
variables in clml codegen, make_packed_api, infer_layout_utils, and
relax distributed code.

**Commit 4 – [REFACTOR][IR]** Delete `class Bool`, `class Integer`,
`TypeTraits<Bool>`, and `TypeTraits<Integer>` from
`include/tvm/ir/expr.h`.

## Canonical replacements

| Old | New |
|-----|-----|
| `Integer(N)` | `IntImm(DataType::Int(32), N)` |
| `Bool(b)` | `IntImm(DataType::Bool(), b)` |
| `x.IntValue()` | `x->value` |
| `x` as bool | `x->value != 0` |
| `ffi::Array<Integer>` | `ffi::Array<int64_t>` |

## Testing

- All 118 C++ unit tests pass (`./cpptest`)
- `tests/python/s_tir/` — 1251 passed (14 pre-existing failures
unrelated to this change, all in TIR transform tests with
annotation-mismatch errors)
- `tests/python/relax/` — passes (excluding pre-existing
torch/torchvision import failures in frontend tests)
2026-05-29 09:48:37 -04:00
Bohan Hou 859498dc01 [TIRx] Bringup TIRx Infrastructure (#19581)
## 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.
2026-05-18 16:44:43 -07:00
ConvolutedDog e7a7447929 [Fix][CI]: remove astral-sh/setup-uv from lint workflow (#19554)
This PR fixes https://github.com/apache/tvm/issues/19552.

astral-sh/setup-uv is not on the ASF GitHub Enterprise action allowlist,
causing the Lint workflow to fail with "Startup failure" before any
pre-commit checks run. See
https://github.com/apache/tvm/actions/runs/25743684906 for the failed
reason.

This PR removes the uv setup and sync steps entirely; pre-commit/action
will install and manage pre-commit and all hook dependencies on its own.
This PR also corrected previous lint errors.

After the fix, the CI lint succeeded:
https://github.com/apache/tvm/actions/runs/25775499703/job/75707088129
2026-05-13 12:28:31 +08:00
as4230 fde09d2052 [BugFix][Relax] Fix scatter_elements and scatter_nd CUDA compilation (#19497)
`topi.scatter_elements` and `topi.scatter_nd` emit bare `T.parallel`
loops in their te.extern IRBuilder bodies which trips `VerifyMemory` on
CUDA targets:

    RuntimeError: Memory verification failed
    ...
    Did you forget to bind?

CPU (LLVM) is unaffected.

This fix makes the IRBuilder body in both `topi/scatter_elements.py` and
`topi/scatter.py` target-aware. When `Target.current()` is a GPU target
it emits thread bindings instead of `T.parallel`.

Fixes #19451.
2026-05-04 16:30:00 +08:00
as4230 772857d34c [Relax][Frontend][TFLite] Add ATAN2 op and TFLite mapping (#19485)
This PR adds the ATAN2 operator to the Relax TFLite frontend.

Introduces relax.op.atan2 as a new binary elementwise primitive (TOPI
broadcast op, Relax registration, legalization to topi.atan2, script
parser support) and registers ATAN2 in the TFLite convert_map. It reuses
the TIR primitive tvm::atan2 so this PR is the higher-layer plumbing.

Validation:
    python -m pytest tests/python/relax/test_op_binary.py
python -m pytest tests/python/relax/test_frontend_tflite.py -k binary

Addresses the ATAN2 item under #19412.
2026-05-01 12:04:51 +08:00
HoYi b915cac0bf [Relax][Frontend][TFLite] Add soft-NMS support for TFLite NON_MAX_SUPPRESSION_V5 (#19426)
## Summary

This PR completes the TFLite `NON_MAX_SUPPRESSION_V5` implementation in
Relax by adding support for `soft_nms_sigma != 0`.

It extends `relax.vision.non_max_suppression` with soft-NMS attributes,
updates the TFLite frontend to consume the soft-NMS outputs correctly,
and aligns the TOPI implementation with LiteRT's reference behavior.

Relates to #19412.

## Changes

1. **Relax / TOPI soft-NMS support**
- Extend `NonMaximumSuppressionAttrs` with `soft_nms_sigma` and
`score_threshold`.
- Thread the new attributes through Relax op registration, Python
wrapper, and legalization.
- Add soft-NMS handling to TOPI classic NMS so
`relax.vision.non_max_suppression` can represent the
`NON_MAX_SUPPRESSION_V5` behavior.

2. **TFLite frontend support for `NON_MAX_SUPPRESSION_V5`**
- Remove the previous `soft_nms_sigma != 0` unsupported-path guard in
the TFLite frontend.
- Forward `soft_nms_sigma` and `score_threshold` into
`relax.vision.non_max_suppression`.
- Handle the soft-NMS return path explicitly so the frontend reads
decayed scores from the processed NMS output instead of re-reading the
original score tensor.

3. **Soft-NMS correctness fixes**
- Fix the soft-NMS path so boxes whose scores fall below the threshold
after decay are invalidated consistently.
- Keep returned indices and decayed scores aligned in both the TOPI TIR
implementation and the NumPy reference implementation.
- Update the soft-NMS candidate selection logic to re-pick the current
best candidate after each decay step, matching LiteRT's
     reference behavior.
   - Align the Gaussian decay formula with LiteRT.

4. **Test coverage**
- Add Relax tests for soft-NMS struct-info inference and legalization.
- Add Relax E2E tests covering reordered outputs after score decay and
other soft-NMS follow-up cases.
- Add TFLite frontend tests for `NON_MAX_SUPPRESSION_V5` with
`soft_nms_sigma != 0`.
- Add IR checks to verify that `soft_nms_sigma` and `score_threshold`
are forwarded correctly.

## Testing

```bash
python -m pytest -n 1 tests/python/relax/test_op_vision.py -k "all_class_non_max_suppression or get_valid_counts or nms" -v
python -m pytest tests/python/relax/test_frontend_tflite.py -k "nms_v5" -v
```

## Result:
- Relax vision tests passed locally
- TFLite `NON_MAX_SUPPRESSION_V5` coverage added for both hard-NMS and
soft-NMS paths
2026-04-27 16:42:34 -04:00
Tianqi Chen 9edd5bd958 [REFACTOR] Remove tvm.runtime.packed_func and container shims; route via tvm_ffi (#19442)
## Summary

- Delete the three Python shim modules that re-exported tvm-ffi types
under `tvm.runtime` / `tvm.ir`:
`python/tvm/runtime/packed_func.py`, `python/tvm/runtime/container.py`,
`python/tvm/ir/container.py`.
- Drop the matching re-exports from `tvm.runtime`, `tvm.ir`, and `tvm`
package init files, so
`tvm.runtime.PackedFunc`, `tvm.runtime.ShapeTuple`,
`tvm.runtime.String`, `tvm.ir.Array`,
  `tvm.ir.Map`, and `tvm.container.Array` no longer exist.
- Migrate every productive caller, test, and tutorial to the canonical
names: `tvm_ffi.Function`,
`tvm_ffi.Shape`, `tvm_ffi.core.String`, `tvm_ffi.Array`, and
`tvm_ffi.Map`.

## Test plan

- [x] `pytest tests/python/all-platform-minimal-test` (75 passed, 77
skipped)
- [x] `pytest tests/python/runtime/test_runtime_container.py
tests/python/all-platform-minimal-test/test_runtime_packed_func.py` (20
passed)
- [x] `pytest tests/python/ir/test_node_reflection.py
tests/python/ir/test_container_structural_equal.py` (32 passed)
- [x] `pytest tests/python/relax/test_vm_build.py
tests/python/relax/test_vm_execbuilder.py
tests/python/relax/test_vm_codegen_only.py` (125 passed, 2 xfailed)
- [x] `pytest tests/python/relax/test_runtime_builtin.py
tests/python/relax/test_op_misc.py` (19 passed)
- [x] `pytest tests/python/target/test_target_target.py` (37 passed, 3
skipped)
- [x] `pre-commit run` clean on touched files
2026-04-25 11:02:08 -04:00
Shushi Hong b6f67b06db [Docs] Add Python API reference for tvm submodule docs (#19379)
as per title
2026-04-10 14:52:49 -04:00
Shushi Hong 3bc61d1fab [BugFix][TOPI] Fix get_const_tuple hanging indefinitely when passed a te.Tensor (#19380)
This pr fixes #18765: `topi.get_const_tuple` hangs indefinitely when
passed a `te.Tensor` instead of a shape tuple and adds a type check to
raise a clear `TypeError` with a helpful message suggesting
`get_const_tuple(tensor.shape)` instead
2026-04-10 14:51:39 -04:00
Soowon Jeong 2e6ee08eaf [BugFix] Align tir.round to ties-to-even across all backends (#19368)
## Problem

`tir.round` constant-folds using `std::nearbyint` (IEEE 754
ties-to-even), but all backends lower it to platform `round()` which
uses ties-away-from-zero. This means compiled code can produce different
results from constant-folded code for midpoint values:

| Input | Constant-fold (ties-to-even) | Compiled (ties-away) |
|-------|-----|------|
| 0.5   | 0.0 | 1.0  |
| 2.5   | 2.0 | 3.0  |
| -0.5  | 0.0 | -1.0 |

This was identified as a follow-up to #19367 — see [this
comment](https://github.com/apache/tvm/pull/19367#issuecomment-4201800320).

## Fix

Align all backends to use ties-to-even intrinsics, matching the
constant-folding behavior:

| Backend | Before | After |
|---------|--------|-------|
| LLVM/ROCm/Hexagon | `llvm::Intrinsic::round` |
`llvm::Intrinsic::nearbyint` |
| NVPTX | `__nv_round[f]` | `__nv_nearbyint[f]` |
| CUDA | `round`/`roundf` | `nearbyint`/`nearbyintf` (f16/bf16 already
used `hrint`) |
| Metal/OpenCL | `round` | `rint` |
| Vulkan/SPIR-V | `GLSLstd450Round` | `GLSLstd450RoundEven` |

Also fixes OpenCL codegen where `tir.nearbyint` was incorrectly mapped
to OpenCL `round()` instead of `rint()`.

Updates `op.h` documentation to explicitly state ties-to-even semantics
for both `round()` and `nearbyint()`.

## Testing

```
python -m pytest tests/python/tirx-base/test_tir_intrin.py -xvs
```

New `test_round_ties_to_even` verifies midpoint inputs `[0.5, 1.5, 2.5,
3.5, -0.5, -1.5, -2.5, -3.5]` produce ties-to-even results on the LLVM
backend. All 12 tests pass (10 passed, 2 skipped for CUDA).

---------

Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
2026-04-08 14:35:22 -04:00
Soowon Jeong acfc483738 [BugFix][ONNX] Fix Round op to use ties-to-even (#19367)
## Problem

The ONNX `Round` operator specification requires **ties-to-even**
(banker's) rounding:

> "For cases where number is exactly halfway between two integers, it
rounds to the nearest even integer."
> — https://onnx.ai/onnx/operators/onnx__Round.html

However, the current TVM implementation produces **ties-away-from-zero**
results on midpoint values:

| Input | Expected (ties-to-even) | Actual (ties-away) |
|-------|------------------------|--------------------|
| 0.5   | 0.0                    | 1.0                |
| 1.5   | 2.0                    | 2.0                |
| 2.5   | 2.0                    | 3.0                |
| -0.5  | 0.0                    | -1.0               |
| -2.5  | -2.0                   | -3.0               |

This was reported in issue #18590.

## Root Cause

The lowering chain for `relax.op.round`:

```
relax.op.round -> (LegalizeOps) -> topi.round() -> te.round -> tir.round -> llvm::round
```

`llvm::round` is defined as ties-away-from-zero (C99 `round()`), while
`llvm::nearbyint` uses the IEEE 754 default rounding mode
(ties-to-even).

## Fix

**`python/tvm/topi/math.py`**: Switch `topi.round()` from `te.round` to
`te.nearbyint`. This lowers to `tir.nearbyint` -> `llvm::nearbyint`,
which respects IEEE 754 ties-to-even.

**`src/target/source/intrin_rule_webgpu.cc`**: Register `tir.nearbyint`
for the WebGPU backend. WGSL `round()` is already ties-to-even per the
WGSL spec, so `tir.nearbyint` -> `round` is the correct mapping.

**`tests/python/relax/test_frontend_onnx.py`**: Add
`test_round_ties_to_even()` with explicit midpoint inputs to prevent
regression.

## Testing

```
python -m pytest tests/python/relax/test_frontend_onnx.py::test_round_ties_to_even -xvs
python -m pytest "tests/python/relax/test_frontend_onnx.py::test_unary[Round]" -xvs
```

Both pass. The new test compares TVM output against onnxruntime (which
correctly implements ties-to-even) for inputs `[0.5, 1.5, 2.5, -0.5,
-1.5, -2.5]`.

Fixes #18590
2026-04-07 15:52:47 -04:00
YinHanke 4df6b1750b [Relax][ONNX] Add roi_pool op and MaxRoiPool frontend support (#18952)
## Summary

Add Relax `roi_pool` support and wire it through the ONNX frontend for
`MaxRoiPool`.

## Changes

- add `relax.vision.roi_pool`, including attrs, Python wrapper, struct
info inference, and legalization
- add TOPI `roi_pool` compute for NCHW layout
- support ONNX `MaxRoiPool` in the Relax ONNX frontend
- handle empty / out-of-bound pooled bins according to ONNX/reference
semantics, returning `0` instead of propagating invalid reductions
- add regression tests for Relax op inference, legalization, and ONNX
frontend import
- add out-of-bound ROI coverage to make sure fully invalid pooled bins
still match ONNX Runtime

## Validation

- `pytest tests/python/relax/test_op_vision.py -k roi_pool`
- `pytest tests/python/relax/test_frontend_onnx.py -k 'max_roi_pool'`


This PR completes the `MaxRoiPool` portion of the Relax ONNX frontend
operator work tracked in #18945.
2026-03-29 12:29:43 -04:00
HoYi 38eb79c63f [Relax][Vision] Add get_valid_counts and classic NMS (#18943)
## Summary

Add `relax.vision.get_valid_counts` and classic
`relax.vision.non_max_suppression` for object-detection post-processing
pipelines.

`get_valid_counts` performs score-based bounding box filtering and
compacts valid boxes to the front of each batch. Classic
`non_max_suppression` performs flexible IoU-based suppression on
filtered boxes, complementing existing `all_class_non_max_suppression`
for custom post-processing workflows.

This PR implements the Relax-level registration, legalization, TOPI
compute, and test coverage for both operators.

## Changes

**Relax op registration and legalization:**
- C++ op functions, FFI registration, and struct info inference for both
operators (`vision.h`, `vision.cc`)
- Python wrappers with Relax docstrings (`vision.py`)
- Legalization to `topi.vision.get_valid_counts` and
`topi.vision.non_max_suppression`
- Additional struct-info validation for `score_index`, `id_index`, and
`coord_start` when `elem_length` is statically known

**TOPI and testing:**
- Full TOPI implementation for `get_valid_counts`
- Reimplementation of classic `non_max_suppression` in TOPI
- NumPy reference implementations in `tvm.topi.testing` for both
operators
- Op-level tests for struct info inference, legalization, invalid
attribute ranges, and e2e numerical correctness
- Stronger legalization tests that verify both `relax.call_tir`
introduction and removal of the original Relax vision op

## Limitations

- Attribute range validation for `score_index`, `id_index`, and
`coord_start` is only enforced when the input `elem_length` is
statically known during struct-info inference.
- Classic `non_max_suppression` follows the existing Relax / TOPI API
shape and is intended for single-class or class-aware custom
post-processing flows, distinct from `all_class_non_max_suppression`.

## Validation

```bash
pytest tests/python/relax/test_op_vision.py -k "get_valid_counts" -v
pytest tests/python/relax/test_op_vision.py -k "test_nms_" -v
```
All related tests passed.
2026-03-28 13:24:27 -04:00
Dayuxiaoshui 3eb86f78ed [Relax][TOPI] Add relax.vision.multibox_transform_loc for SSD/TFLite box decode (#18942)
Introduce relax.vision.multibox_transform_loc with
MultiboxTransformLocAttrs: decode center-size offsets against ltrb
priors, softmax on class logits, and optional clip, threshold masking,
and background score zeroing. Register the C++ op with FInferStructInfo
checks for shapes and dtypes (including batch and 4*N consistency).
Legalize to topi.vision.multibox_transform_loc.

Add tests for struct inference, invalid inputs, Legalize+e2e on LLVM,
attribute branches, and TVMScript roundtrip. Add a standalone numpy
reference under topi/testing (not exported from tvm.topi.testing to
avoid pulling scipy).

Update TFLite frontend NotImplementedError text for
DETECTION_POSTPROCESS and NON_MAX_SUPPRESSION_V5 to note multibox is
available and link tracking issue #18928.
2026-03-28 00:20:22 -04:00
YinHanke 1e08eb2fa9 [Relax][ONNX][Torch] Add roi_align support and frontend integration (#18936)
## Summary

Add Relax `roi_align` support and wire it through the ONNX and PyTorch
frontends.

## Changes

- add `relax.vision.roi_align`, including attrs, Python wrapper, struct
info inference, and legalization
- add TOPI `roi_align` compute and keep both legacy and aligned ROIAlign
semantics
- support ONNX `RoiAlign`, including `coordinate_transformation_mode`
handling for `output_half_pixel` and `half_pixel`
- support PyTorch `torchvision.ops.roi_align` in the exported-program
frontend, including the `aligned` flag
- add regression tests for Relax op inference, legalization, TVMScript
parsing, ONNX frontend import, and PyTorch frontend import
- add aligned ROIAlign test coverage to make sure sub-pixel RoIs no
longer use the legacy `min=1.0` clamp

## Validation

- `pytest tests/python/relax/test_op_vision.py -k roi_align`
- `pytest tests/python/relax/test_tvmscript_parser_op_vision.py -k
roi_align`
- `pytest tests/python/relax/test_frontend_onnx.py -k roi_align`
- `pytest tests/python/relax/test_frontend_from_exported_program.py -k
roi_align`

This PR completes the Relax/ONNX/Torch roi_align work tracked in #18928.
2026-03-26 10:41:54 -04:00
Tianqi Chen 141c22fd8a [Refactor] Bring up tirx namespace (#18913)
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
2026-03-19 21:27:54 -07:00
YinHanke a89b9f2880 [TOPI] Reject non-float inputs for inverse unary math ops (#18880)
## Summary

Reject non-float inputs for inverse trigonometric and hyperbolic unary
ops in TOPI.

## Changes

- add a shared floating-point dtype check for inverse unary math ops in
TOPI
- apply the check to `topi.acos`, `topi.acosh`, `topi.asin`,
`topi.asinh`, and `topi.atanh`
- add TE tests covering integer-input rejection for these ops
- add regression tests covering successful LLVM build for both `float32`
and `bfloat16`

## Validation

- `tests/python/te/test_te_create_primfunc.py -k 'topi_float_unary'`
- local repro now fails early with a clear `TypeError` for integer
inputs
- local regression check confirms the valid `float32` and `bfloat16`
paths still compile with LLVM

## Issue

Fixes #18729
2026-03-08 18:39:21 -04:00
Tianqi Chen 689d2b51b2 [REFACTOR][TIR] Remove body from AllocBuffer and DeclBuffer (#18876)
## 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++)
2026-03-06 06:47:20 -05:00
YinHanke 14e41c681a [Relax][ONNX] Support dynamic repeats for Tile (#18878)
## Summary

Support dynamic `repeats` for ONNX Tile in the Relax frontend.

## Changes

- add a dynamic Tile conversion path for ONNX when `repeats` is a graph
input
- expose `topi.dyn_tile` to the Python/packed TOPI interface
- add frontend tests for dynamic `repeats`

## Validation

- `tests/python/relax/test_frontend_onnx.py -k test_tile_dynamic_repeats
-q`
- local end-to-end repro matches ONNX Runtime

## Issue
Fixes #18752
2026-03-05 23:38:32 -05:00
Tianqi Chen 7d5c46e236 [TIR][FEAT] Require DeclBuffer before use in verify_well_formed (#18843) 2026-02-28 10:40:34 -05:00
Tianqi Chen 9a8320acbd [LINT][PYTHON] Modernize annotations with ruff UP rules (#18830)
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.
2026-02-27 21:29:47 -05:00
Tianqi Chen c1d32438a0 [BugFix][TOPI] Fix resize accuracy issue with non-floor rounding (#18838)
## Summary

The int_div optimization in `topi.image.resize` was applied
unconditionally
for `nearest_neighbor` + `asymmetric` mode, regardless of rounding
method.
This caused accuracy issues when `rounding_method` is not `"floor"`
(e.g.,
`"round"`, `"round_prefer_ceil"`), because integer division truncates
toward
zero rather than rounding.

**Fix**: Gate the int_div optimization on `rounding_method == "floor"`
or
`rounding_method == ""` (the default, which gets resolved to `"floor"`
for
non-align_corners modes).

- Updates `_resize_2d` in `python/tvm/topi/image/resize.py`
- Updates reference implementation in
`python/tvm/topi/testing/resize_python.py`
- Updates legalize test expected output to reflect the new behavior
2026-02-27 11:21:16 -05:00
Tianqi Chen 33dcea1686 [REFACTOR][LINT] Modernize ruff config (#18810)
This PR removes the extra lint violations from the codebase so lint
aligns with the latest style
2026-02-23 07:29:21 -05:00
Tianqi Chen aa2e609136 [LINT] Modernize lint to use pre-commit hooks (#18807)
This PR migrates existing lint to use pre-commit hooks
2026-02-22 11:03:21 -05:00
Qingchao Shen 9752557bf1 support integer types in fast_tanh and fast_exp (#18768)
Fix https://github.com/apache/tvm/issues/18767.

This PR fixes the issue by adding explicit type casting in `fast_tanh`
and `fast_exp` to convert integer inputs to float32.
2026-02-20 11:52:11 -05:00
Ruslan Baratov 52e45477de [DOC] Unify CUDA naming (#18797)
Fix CUDA naming in documentation and comments

- Cuda -> CUDA
- cuda -> CUDA
2026-02-19 08:04:00 -05:00
Tianqi Chen 283fd19683 [REFACTOR][TARGET] Further cleanup target python api (#18793)
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
2026-02-17 21:39:27 -05:00
Tianqi Chen c8140643d3 [REFACTOR][S-TIR] Move remaining data structures to s_tir (#18743)
This PR moves remaining related data structures to s_tir.
- Moves sblock_dependency_info and sblock_scope.
- Moves related analyssis.
- Hides the data_type_rewriter to private functions.
2026-02-10 18:24:46 -05:00
Tianqi Chen 87c1e471b0 [REFACTOR] Migrate old tir.ir_builder to tvmscript or builder (#18716)
This PR migrates legacy tir.ir_builder infavor of tvmscript or builder.
2026-02-06 10:34:13 -05:00
Nguyen Duy Loc 038b327cff [Relax][Onnx][PReLU] Handle slope and axis argument with different slope shapes (#18658)
This PR support handle slope and axis argument of PReLU op with
different slope shapes: (1xCx1x1) or (S,) or (1,1) etc.

### Description
- Handle slope and axis argument of PReLu op (to pass into
relax.op.nn.prelu function)
- If slope shape = (1xCx1x1), get axis = 1 and reshape slope to (C,)
- else if slope shape = (S,) or (1, 1), get axis = len(x_shape) - 1
(take the last axis of the input x)
(https://onnx.ai/onnx/repo-docs/Broadcasting.html)
 - else raise error
 
### Resolved
- Fixed 1: #18596
- Fixed 2: #18598
- Fixed 3: #18606
- Fixed 4: #18607
2026-01-12 20:55:54 +08:00
Asuka 5dc4e785d9 [Relax] Fix batch normalization computation logic (#18609)
Dear reviewers,

**Why**
The previous implementation of batch_norm incorrectly conflated the
computation of mean and variance between training and evaluation modes.
Additionally, for '_native_batch_norm_legit.no_stats‘, using
instance_norm to handle normalization ignored the batch dimension,
leading to incorrect behavior.

**How**
This PR includes the following fixes:
1. Corrects the computation logic to properly distinguish between
training and evaluation modes.
2. Ensures the batch dimension is properly accounted for in
`_batch_norm_legit_no_stats`.

**Environment**
GPU: NVIDIA A100-SXM4-80GB
2025-12-26 20:57:48 +08:00
Guan-Ming (Wesley) Chiu ec0026e0bc [Relax][PyTorch] Fix index_put with broadcast indices (#18533)
## Related Issue

closes https://github.com/apache/tvm/issues/18355

## Why

Converting PyTorch operations like M[:, rows, cols] = x failed because:
1. The TOPI index_put implementation called len() on TVM Tensor objects
(unsupported)
2. Index tensors with different shapes (e.g., (2,) and (10,)) couldn't
broadcast together

## How

- Added broadcasting support following NumPy rules to handle
multi-dimensional index tensors
- add tests for batched indexing pattern M[:, rows, cols] = x
2025-12-01 15:21:32 -05:00
Siva 4be951d710 [RELAX][PASS] Annotate Custom Scope layout pass for Adreno GPU (#17599)
This PR adds custom scope layout passes for Andreno GPU

https://discuss.tvm.apache.org/t/rfc-annotate-custom-scope-layout-relax-pass-for-adreno-gpu/18052/6

for details about texture scope handling.
2025-11-24 08:33:36 -05:00
Guan-Ming (Wesley) Chiu 0225d67d30 [Relax][PyTorch] Fix MultiheadAttention complie (#18459)
## Related Issus

closes #18440

## Why

- PyTorch `masked_fill` / `full_like` accept inf or nan and TVM couldn’t
handle these values when the tensor dtype was not float, which caused
wrong behavior or errors.

## How

- If `fill_value` is inf or nan and the tensor dtype is not float →
convert the fill to float32.
- For masked_fill → Create a float values tensor with full_like.
- Cast input to float if needed.
- In TOPI → Reject creating full with inf/nan on non-float dtypes.
2025-11-16 01:08:37 -05:00
Qingchao Shen d013dad06d [TOPI] Support integer type input for log and log2 (#18426)
Adds support for integer inputs in `topi.log` and `topi.log2` by
automatically converting them to float32, aligning with NumPy's
implicit float promotion behavior.
2025-11-12 11:11:03 -05:00
akaashrp 462eeb72b8 [WebLLM] Replace int64s with int32s in WebGPU kernels (#18361)
This PR replaces int64s with int32s in the argsort and parallel_sampling_from_prob
kernels when the target is WebGPU (since WGSL does not currently support i64)
2025-10-21 23:13:11 -04:00
Shushi Hong c00c66259a [Relax][ONNX] Support AllClassNMS Operator for ONNX Frontend (#18321)
Follow #18175 , this PR supports AllClassNMS Operator for ONNX Frontend
2025-10-01 16:09:59 -04:00
Tianqi Chen 543e64dbb1 [FFI][REFACTOR] Cleanup tvm_ffi python API and types (#18277)
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
2025-09-07 10:38:50 -04:00
Tianqi Chen 3c36ce2ec6 [FFI][REFACTOR][ABI] Rename NDArray to Tensor (#18275)
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.
2025-09-06 14:33:59 -07:00
Tianqi Chen a7a0168be5 [FFI][REFACTOR] Establish tvm_ffi python module (#18226)
* [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
2025-08-24 15:46:20 -07:00
Youngsik Yang 79368cea1a [Relax][ONNX][Transform] Add mode choice, new mode, and warning for take() (#18061)
[Relax][Transform] Add mode choice, NaN mode, and warning for take()

- Add a `mode` parameter to Relax’s `take()`
- Add `NaN` mode to `take()`
- Add unit tests covering all `take()` modes
- Add a warning log for `fast` mode
- Unify default modes in lower layers to `fast` for consistency with Relax
2025-07-08 19:58:35 +08:00
Tianqi Chen 1b9da40ce8 [REFACTOR][FFI] Phase out old VisitAttrs mechanism (#18106) 2025-07-03 20:18:19 +08:00