23 Commits

Author SHA1 Message Date
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
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
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
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
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
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 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
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 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
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
Tianqi Chen 1b9da40ce8 [REFACTOR][FFI] Phase out old VisitAttrs mechanism (#18106) 2025-07-03 20:18:19 +08:00
kavin-mcw 9eb8b3004b Add support for bucketize (#18040)
* add support for bucketize

* fix lint issue

* Fix lint issue

* Add GPU code for bucketize

* Resolve merge conflict

* Fix lint issue
2025-06-30 23:04:14 +08:00
Tianqi Chen ccaa534b2c [REFACTOR] Phase out relay python components (#17656)
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>
2025-02-15 13:48:28 -05:00
Leandro Nunes 89061fafa5 [CI] Bump black version to 22.3.0 (#10960)
* Make all required adjusts in the code to comply with the new version
* Upadte ci-lint to v0.71, based on tlcpackstaging/ci_lint:20220411-060305-45f3d4a52
2022-04-11 10:24:05 -07:00
masahi d12e7259c2 [TOPI] Fix compiing batch_matmul and dense when two args are the same tensor (#9207)
* Add explicit copy stage for batch_matmul(x, x) case

* do copy in relay strategy to avoid dup

* add copy to dense op and schedules

* black

* add batch_matmul test

* add dense test

* fix cuda int8 dense test

* remove need_copy flag

* do not use tag to decide if tensors are same

* rename to copy_if_identical and add comment

* black

* one more fix missed

* add length check on input tensors

* one more length check

* fix variable name
2021-10-08 18:32:43 +09:00
Egor Churaev 719d2f67ed [OpenCL] Add vectorization to cuda conv2d_nhwc schedule (#8636)
* Add vectorization to cuda conv2d_nhwc schedule

Adding vectorization significantly improved performance. About 6-7x
boost.

* Apply comment

* Move schedule to topi/gpu dir

* Add vectorization to inner loop

* Update values of vectorization factor
2021-09-30 09:04:13 -07:00
Chenfan 6d1ced0803 [Matmul] Add matmul op (#8234)
* Add Matmul Op

* Recover DenseAttrs

* Add grad for matmul & some update

* Update matmul cuda default schedule

* Add blas support for matmul

* Lint fix add update doc strings
2021-06-30 22:29:43 +08:00
Lunderberg ae58f2c387 [Topi][Unittests] Parametrized tests in test_topi_dense.py, split out gpu-independent implementations (#8336)
* [Topi][UnitTests] Parametrized tests in test_topi_dense.py

Now, tests run for multiple data types, can be extended with
additional datatypes.

* [Topi] Separated generic-gpu nn.dense implementations into topi.gpu.dense

As a follow-up to the renaming of "gpu" to "cuda", separating
implementations that require CUDA (e.g. dense_cublas.cuda) from
implementations that require any GPU, but not necessarily a CUDA GPU
(e.g. dense_small_batch.gpu).

My intent is to pair this migration with the extension of unit tests
to cover additional GPU runtimes, migrating only implementations that
run correctly on non-CUDA GPU devices.

* [Vulkan][Codegen] Updated storage sync to avoid incorrect matmul results on some GPUs

- In ThreadAllreduceBuilder, separate out load/store so that they can
  have a memory barrier in-between.

- In Vulkan codegen, added Workgroup memory sync for subgroup thread
  sync, since the different subgroup threads can still access
  workgroup memory.  Longer-term, may need tir enhancements to
  separate out sync of control/memory.

Co-authored-by: Eric Lunderberg <elunderberg@octoml.ai>
2021-06-30 13:59:16 +08:00