## Motivation and context
A TIRx pointer carries two pieces of information that later lowering
needs: the pointee element type and the storage scope. Both must survive
when a pointer-producing expression is assigned to a Python name and
then used as the backing storage of a buffer.
A concrete example is accessing an mbarrier in another CTA through
distributed shared memory:
```python
ptr_ty = PointerType(PrimType("uint64"), "shared")
remote_ptr = T.reinterpret(
ptr_ty,
T.ptx.map_shared_rank(mbar.ptr_to([0]), T.int32(0)),
)
remote_mbar = T.decl_buffer(
[1], "uint64", data=remote_ptr, scope="shared"
)
```
`map_shared_rank` returns the raw `uint64` address produced by PTX
`mapa`, and `reinterpret` gives that address the intended
`PointerType(uint64, shared)`. Because `decl_buffer(data=...)` requires
a pointer `Var`, assigning the expression to `remote_ptr` should create
an immutable typed pointer binding.
Before this PR, an unannotated assignment such as `remote_ptr = <pointer
expression>` followed the same parser path as a numeric assignment. That
path allocates a mutable local scalar and therefore cannot represent a
`PointerType`. The pointer expression could not be carried as a
correctly typed `Var` into `decl_buffer` and CUDA lowering.
This PR makes an unannotated pointer-valued assignment emit a TIRx
`Bind`. The bound `Var` has exactly the type of the right-hand side,
including its element type and storage scope. Pointer bindings are
immutable, so reassignment in the same scope is diagnosed; shadowing a
name supplied through `extra_vars` remains valid. Numeric assignments
keep their existing mutable-local behavior.
## Type propagation fixes
The parser fix exposed several other boundaries where pointer type
information must remain consistent:
| Boundary | Previous behavior | Behavior after this PR |
| --- | --- | --- |
| Unannotated pointer assignment | Tried to materialize the value as a
local scalar | Emits an immutable `Bind` with the RHS `PointerType` |
| `address_of(buffer)` / `buffer.ptr_to(...)` | Reused the raw backing
pointer type | Returns a pointer to `buffer.dtype` while preserving the
backing pointer storage scope |
| `tvm_access_ptr` / `ptr_byte_offset` | Accepted strings or annotation
expressions, but not a `PrimType` object directly | Accepts `PrimType`
and produces the corresponding typed pointer |
| Printed `T.ptx.mapa` call | The printer emits all intrinsic attributes
positionally, but the Python helper required keyword-only arguments |
Accepts the canonical printed form so pointer code round-trips through
TVMScript |
The `address_of` distinction matters for typed views over byte-addressed
storage. For example, if a `float32` buffer is backed by a `uint8*`
allocation in `shared.dyn`, the address of a buffer element must be
`PointerType(float32, shared.dyn)`, not `PointerType(uint8,
shared.dyn)`.
With these changes, the DSMEM example above round-trips through
TVMScript and CUDA codegen declares the remote buffer pointer as
`uint64_t*`.
## TMA dtype normalization
This PR also contains a small, separate type-representation fix in TMA
descriptor construction. `TmaPlan.elem_dtype` is a string consumed by
the host-side `runtime.cuTensorMapEncodeTiled` packed call, but
`_assemble_plan` stored `g_buf.dtype`, which is a `PrimType`. Converting
it with `str(g_buf.dtype)` ensures that the generated packed-call
argument is `StringImm("float16")` rather than an IR type object. This
does not change the pointer-binding semantics described above.
## Testing
- Verify that an unannotated pointer expression creates a `Bind` whose
`Var` type matches the RHS type.
- Verify that pointer reassignment is rejected while shadowing an
`extra_vars` name is allowed.
- Verify parser/printer structural round-tripping for the pointer
binding and canonical `T.ptx.mapa` call.
- Verify that `address_of` uses the logical buffer element type and
preserves the storage scope for byte-backed buffer views.
- Verify that `tvm_access_ptr` and `ptr_byte_offset` accept `PrimType`
inputs.
- Compile the DSMEM `map_shared_rank` example through the CUDA TIRx
pipeline and check for a typed `uint64_t*` remote buffer pointer.
- Verify that the TMA host initialization passes the descriptor dtype as
a `StringImm`.
## Motivation and context
`MBarrier.remote_view(rank)` represents an mbarrier owned by another CTA
in the same cluster. The existing view kept only a buffer whose pointer
had been mapped to the remote CTA with PTX `mapa`. Calling
`remote_bar.arrive(...)` then followed the inherited local-arrive path
and emitted the local `mbarrier.arrive.shared.b64` form against that
mapped remote address.
A remote arrival must instead use the `shared::cluster` instruction
form. TIRx models that form with the owner-local barrier pointer plus
the destination CTA rank and predicate. Using the local instruction with
a remote address is not equivalent and is reported by synccheck as a
local arrival on a remote mbarrier address.
## Changes
- Keep the owner-local buffer and target CTA rank when constructing a
remote mbarrier view.
- Route `MBarrier` remote arrivals through the cluster helper using the
local barrier pointer and stored CTA rank.
- Apply the same routing to `TMABar`, including
`mbarrier.arrive.expect_tx`.
- Keep a typed `PointerType(uint64, shared)` mapped buffer on the view
so `ptr_to` remains available to operations that explicitly consume a
remote shared-memory pointer.
- Reject operations with ambiguous or invalid ownership semantics:
- initializing or waiting on a remote view,
- supplying another `cta_id` to a view that already fixes its target,
- creating a remote view from another remote view.
- Preserve the existing local-CTA behavior for ordinary barriers.
## Testing
- Verify the typed `mapa` binding and remote buffer in TIRx IR.
- Verify CUDA codegen for plain and counted
`mbarrier.arrive.shared::cluster.b64` forms.
- Verify CUDA codegen for remote
`mbarrier.arrive.expect_tx.shared::cluster.b64`.
- Verify that the corresponding local instruction forms are not emitted
for remote views.
- Verify diagnostics for remote init, wait, nested views, and
conflicting `cta_id` arguments.
- Run changed-files pre-commit checks.
Focused result: 3 tests passed.
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.
## 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
This PR removes the suite-wide TIRx compute-capability gate and
localizes hardware skips to CUDA codegen and tile-primitive tests. It
keeps the original test parameterization unchanged, allowing parser,
printer, IR, transform, and other non-hardware TIRx tests to run in
regular CI while device-dependent cases are skipped when SM100 hardware
is unavailable. CUDA codegen helpers use explicit target architectures
where needed, and the run-only benchmark utility tests retain a local
SM100 gate.
This intentionally avoids adding separate compile/run parameter cases. A
follow-up PR can audit slow frontend execution tests and define a
focused runtime regression budget for real TIRx kernels.
Local validation:
- `pytest -n auto -m "not gpu" tests/python/tirx`: 486 passed, 79
skipped.
- `pytest -n auto -m gpu tests/python/tirx` without matching hardware:
1509 skipped.
- Pre-commit passed on all changed files.
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.
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`
## 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
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.
## 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.
Hongyi Jin