main
6 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
 Tianqi Chen
|
adf8d6a463 |
[TIRx] Phase out duplicate Var type_annotation (#19944)
## Rationale
TIRx variables use inherited `ExprNode::ty` as their single semantic
type. Retaining a primitive handle surrogate erases the distinction
between scalar values, typed pointers, and true opaque pointers, then
forces later passes and code generators to reconstruct information that
the IR already owns.
## Changes
- Remove the duplicate reflected `Var::type_annotation` state and
preserve exact `PrimType` or `PointerType` through construction,
visitors, transforms, specialization, builders, printers, and code
generation.
- Keep scalar-only boundaries explicit through `PrimExpr`, `PrimVar`,
and `PrimType`; pointer-capable values remain general `Expr` or `Var`.
- Keep helper boundaries no broader than their contracts: TE tensor
variable indices use `PrimVar`, while expression deep equality recurses
through general `Expr` only where pointer-bearing `Call` arguments
require it and does not generalize private arithmetic subclasses.
- Keep core statement reflection typed as `Expr`, name general
reinterpret targets as `target_ty`, and preserve exact pointer calls in
the general vectorization path with explicit scalarization behavior.
- Delete `PrimType::Handle()` and `PrimType::IsHandle()`. True opaque
pointers use `PointerType::VoidPointerTy()`; TVMScript renders the
canonical global type as `T.handle`, standalone values as `T.handle()`,
and scoped void pointers with a keyword-only storage scope.
- Make `CodeGenSourceBase::SSAGetID` a single `Type` boundary across
source backends, without a separate primitive-type or runtime-dtype
variant.
- Keep WebGPU semantic argument classification type-aware: storage
buffers are identified from `PointerType`, POD arguments from
`PrimType`, and only the final `FunctionInfo` launch ABI is serialized
to `DLDataType`.
- Preserve exact pointer semantics at runtime boundaries, including
access pointers, packed calls and returns, external calls, storage
rewrites, and target-specific lowering.
## Migration guide
- **Variable types:** In C++, replace `var->type_annotation` with
`var->ty`; in Python, replace `var.type_annotation` with `var.ty`. The
result is the exact `Type`: scalar variables carry `PrimType`, while
pointer variables carry `PointerType`.
- **Scalar boundaries:** Use `PrimVar` and `PrimExpr` for variables and
expressions that are semantically scalar. When starting from a general
view, narrow explicitly with `var.as_or_throw<PrimVar>()` or
`expr.as_or_throw<PrimExpr>()`. Keep pointer-capable fields and call
arguments as `Var` or `Expr`. A default-constructed `PrimVar` is
nullable, so construct local scalar variables explicitly, for example
`PrimVar i("i")`.
- **Opaque pointers:** Replace `PrimType::Handle()` with
`PointerType::VoidPointerTy()`. Replace `IsHandle()` tests with explicit
`PointerType` inspection; use `PointerType(element_type, storage_scope)`
when the pointee type is known instead of erasing it to a runtime handle
dtype.
- **TVMScript handles:** Use `arg: T.handle` for a global void-pointer
annotation and `arg = T.handle()` for a standalone value. Use
`T.handle(storage_scope="shared")` for a scoped void pointer. Typed
pointers use forms such as `T.handle("float32")`, `T.handle("float32",
"global")`, or `T.handle("float32", "shared")`. Legacy
`T.handle("void")` input remains parse-compatible, but the printer
canonicalizes it to `T.handle` (or the keyword-only scoped form).
- The separate `tirx.type_annotation` intrinsic used by access-pointer
APIs is unchanged; this migration removes only the duplicate variable
field.
## Validation
- Complete native C++ test executable: 122/122 passed, including
`IRF.CountVar`.
- Relax binding-rewrite suite: 12/12 passed, including transferred-user
bookkeeping.
- Canonical typed/void/scoped TVMScript handle printer and round-trip
checks: 5/5 passed.
|
||
|
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. |
||
|
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 |
||
|
Bohan Hou
|
16d0a7edae | [TIRX][CUDA] Framework support for FA4, CLC intrinsics, and nvfp4 tcgen05 GEMM (#19785) | ||
|
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. |
||
|
Tianqi Chen
|
7504e3ed1a |
[REFACTOR][SCRIPT] TVMScript dialect-friendly refactor: per-dialect restructure + dialect registry (#19479)
## Summary Restructure TVMScript to be dialect-agnostic at the script-core layer while letting each extension dialect (TIRX, Relax) own its own per-dialect script subtree. IR is below script in the dependency stack and is NOT a peer dialect — its script handlers stay in the shared core. This PR folds together two coupled refactors that were initially opened as separate PRs (#19478 and the original #19479); they share rename / relocation surface so they ship as one cohesive change. ## What this PR does ### Per-dialect script subtree (originally #19479) - Moves per-dialect printer + builder from `src/script/{printer,ir_builder}/{tirx,relax}/` to `src/{tirx,relax}/script/{printer,builder}/`. - Tightens `src/script/*.cc` CMake glob to the dialect-free core. - Refactors `IRBuilder::DeclFunction` to dispatch via FFI registry (`script.ir_builder.decl_function.<type-key>`); removes cross-dialect includes from the shared core. - Adds `tvm.script.register_dialect` API + `__getattr__` + a `sys.meta_path` finder for Python-side dialect discovery. In-tree dialects (tirx, relax) registered centrally in `python/tvm/__init__.py`. - Drops the obsolete static re-export shims at `python/tvm/script/{parser,ir_builder}/{tirx,relax}/`. ### Dialect-agnostic printer config (originally #19478) - Relocates `include/tvm/ir/script_printer.h` → `include/tvm/script/printer/config.h` next to the rest of the printer's public surface. The header is not IR-specific. - Renames `TVM_SCRIPT_REPR` → `TVM_REGISTER_SCRIPT_AS_REPR` for clarity (the macro registers Script as the kRepr callback + per-type vtable dispatch). Aligns with the `TVM_REGISTER_*` family. - Drops dialect-hardcoded `PrinterConfig` fields (`tir_prefix`, `relax_prefix`, `show_all_struct_info`, `buffer_dtype`) in favor of a generic `ffi::Map<String, Any> extra_config` keyed by `"<dialect>.<knob>"`. Each call site reads via the templated accessor `config->GetExtraConfig<T>("...", default)`. - Promotes `std::string` config fields to `ffi::String`. After this lands, the script-printer core knows nothing specific about any dialect — new dialects plug in via the registry pattern with zero core edits. Public Python API surface unchanged. |