## 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.
## Summary
Follow-up work on top of the TIRx infrastructure bring-up (#19581). It
extends the TIRx operator-dispatch, codegen, and TVMScript surfaces with
the next batch of low-level programming features for Blackwell-class
GPUs, while keeping `s_tir` script support intact.
## Main Changes
- **op-dispatch**: warp `ldmatrix`/`stmatrix` copy dispatch; split CUDA
copy into register / gmem-smem / `ldgsts` paths; `tcgen05.ld/st`
`.16x{64,128,256}b` dispatch with a factory and M=128 layout;
element-wise broadcast at the layout level with a copy vec-alignment
fix.
- **gemm**: CUDA synchronous `mma.sync` tensor-core dispatch; accept a
Layout F C operand for M=64 MMAs.
- **op**: add the `permute_layout` primitive (replaces `permute_dims`).
- **tvmscript**: add the `Tx.jit` decorator, `Tx.constexpr` compile-time
params, and `Tx.wg_reg_tile`.
- **lower-tirx**: introduce the `Tx.device_entry()` marker (replacing
`ScopeKind::kKernel`); canonical thread filters that drop the
`Tx.filter` wrapper.
- **codegen**: add a typed-pointer byte-offset intrinsic; remove the
`entry_cluster_sync` codegen attribute.
## Validation
- `pre-commit run` (changed files) — clean
- `ninja -C build -j$(nproc)` — builds
- `pytest tests/python/tirx/ -n 16`
- `1997 passed, 39 skipped, 3 xpassed`
- `python -m pytest tests/python/all-platform-minimal-test`
- `37 passed, 105 skipped`
- `TVM_TEST_TARGETS=llvm pytest tests/python/tirx-analysis
tests/python/tirx-base tests/python/tirx-transform -n 16`
- `630 passed, 25 skipped, 8 xfailed, 1 xpassed`
## Local CI Notes
Several full CI-equivalent jobs are not locally reproducible because
this machine is missing parts of the Apache TVM CI environment (e.g.,
specific `llvm-config` versions, Vulkan, ROCm, ARM/QEMU cross-toolchain,
and web/wasm components). The Blackwell/Trainium kernel tests are
maintained downstream and are intentionally not part of this PR.
## 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.
Tianqi Chen