TVM 0.25's TIRx reveals what Triton deliberately conceals

Apache TVM just shipped a kernel DSL that does the opposite of what most compilers brag about: instead of automating away hardware detail, TIRx hands it back to the programmer on purpose.

What Exactly Is TIRx, and Why Did It Land in TVM 0.25.0?

TIRx is Apache TVM's next-generation kernel-level compiler structure, defined in the project's official documentation as "an open-source, hardware-native DSL and compiler for machine-learning kernels" built on existing TVM infrastructure . It entered the tree through PR #19581, "[TIRx] Bringup TIRx Infrastructure," merged into apache/tvm:main by TVM creator Tianqi Chen on May 18, 2026 , and was packaged in the v0.25.0 release dated June 19, 2026 .

Quick Answer: TIRx is Apache TVM's hardware-native kernel DSL and compiler, merged via PR #19581 on May 18, 2026 and shipped in v0.25.0 on June 19, 2026 . Unlike TensorIR's automated scheduling, it keeps loops, synchronization, and buffer layout in author code.

The architectural shift is the point. Where TensorIR leaned on automated scheduling passes to derive performant code, TIRx keeps loop structure, synchronization, buffer layout, and pipeline state in the author's source. The compiler reasons only over the recurring tile-level structure — the official overview frames it as orchestration staying "in hardware-native source code, while the recurring tile-level structure becomes visible to the compiler" . The CUDA-native docs are blunt about the tradeoff: there is "no automatic scheduling" — what you write is what gets emitted.

The initial target is NVIDIA Blackwell-class GPUs, with the PR describing support for "low-level programming of Blackwell-class GPU architectures" and exposing intrinsics like WGMMA, tcgen05, and TMA . Bringup CI reported 1,723 passed / 47 skipped for tests/python/tirx , a signal that the surface is functional even if still early.

Installing it is a one-liner: pip install apache-tvm==0.25.0, which requires Python >=3.10 and ships built wheels for Windows x86-64, manylinux x86-64, manylinux ARM64, and macOS 11.0+ ARM64 . The package is Apache-2.0 licensed, so there is no pricing tier to weigh — only the question of whether the explicit-control model fits how your team writes kernels.

Where TIRx and Triton Deliberately Diverge

TIRx and Triton solve the same problem — writing fast GPU kernels — but draw the programmer-compiler boundary in opposite places. Triton sets a high boundary: block-level abstractions hide warp coordination, shared-memory management, and prefetch scheduling from the author, letting the compiler orchestrate them. TIRx sets a low, explicit one: warp IDs, synchronization calls, buffer layouts, and pipeline stages all stay in author-written hardware-native source. The official overview states the split plainly — "orchestration stays in hardware-native source code, while the recurring tile-level structure becomes visible to the compiler" .

That choice has a direct consequence at the native layer: there is no autoscheduler. What you write is what gets emitted — no loop reordering, no compiler-inserted prefetch hints, no implicit thread predication. The CUDA-native basics page is blunt that "no automatic scheduling" applies, and that the native layer is where you go when a hardware feature has no primitive yet . For engineers used to Triton inferring a schedule from a block program, this is a deliberate trade: you lose the convenience of automation in exchange for predictability on hardware that moves faster than any autoscheduler can keep up with.

The design explicitly targets frontier kernels — cases where new instructions, memory spaces, cooperation patterns, and algorithms arrive before a compiler can reliably automate them . On a Blackwell-class GPU shipping instructions a stable compiler hasn't yet learned to schedule, an explicit boundary lets you reach the silicon today rather than waiting for backend support to mature.

The second divergence is tooling. TIRx records reusable operations as inspectable TilePrimitiveCall IR nodes that are later resolved by dispatch, and it exposes that IR and the compiler's utilities through TVM FFI across Python, C++, and Rust . The compiler surface is therefore open to static analyzers, validators, and programmatic writers rather than sealed behind a single Python frontend. This is the angle TVM creator Tianqi Chen leans on:

"Agents should have access to a predictable DSL that offers maximum expressiveness, paired with a minimal compiler they can directly open up, build toolings, and improve for," — Tianqi Chen, Apache TVM creator (source: @tqchenml).

Whether that bet pays off is unproven — Triton's higher boundary is exactly what makes it productive for most kernels, and TIRx enters a contested space amid heavy 2025–2026 work on Triton for Blackwell, including NVIDIA's CUDA Tile IR backend . The practical question is not which boundary is correct, but which matches how your team writes kernels.

Scope IDs, Tensor Layout, and Tile Dispatch: How TIRx Structures Its IR

TIRx organizes its IR around three documented constructs — Execution Scope, Tensor Layout, and Tile Primitive Dispatch — that together let authors write hardware-native orchestration while the compiler still reasons over reusable tile structure . The split is deliberate: orchestration stays in source code you control, and only the recurring tile-level structure becomes visible to the compiler. That is the mechanism behind the lower boundary discussed above.

Execution Scope selects which hardware participants are active using ordinary control flow rather than implicit thread predication. You gate work with scope-id intrinsics such as if wg_id == 0 or warp_id == ..., which map directly to CUDA workgroup and warp roles . Nothing is auto-broadcast across threads; the participants you name in the branch are the ones that execute, and the cleanup stage later lowers those scope ids into concrete thread axes .

Tensor Layout is a storage-first interface: it describes how a logical tensor maps to a physical resource — global memory, shared memory, registers, or accelerator SRAM — rather than leaving placement to a scheduler . You declare buffers with T.match_buffer and request scratch with explicit allocations; during LowerTIRxCleanup, TileLayout buffer access resolves into physical address arithmetic . The layout you write is the layout that ships.

Tile Primitive Dispatch is the reusable middle ground. Operations like copy or gemm are first recorded as unresolved TilePrimitiveCall IR nodes that carry the primitive type, execution scope, operand layouts, target backend, and optional hints . The LowerTIRx pass then runs TilePrimitiveDispatch, replacing each node with a concrete backend dispatch body chosen from those attributes . This is where one written call can become TMA copies, vectorized loads, WGMMA, or tcgen05 depending on the resolved backend.

On the authoring surface, you import the dialect with from tvm.script import tirx as T, annotate kernels with @T.prim_func or @T.jit, mark the device boundary with T.device_entry(), and write ordinary loops, branches, and buffers. Compilation runs through tvm.compile(mod, target, tir_pipeline="tirx"), which binds the target and dispatches into the TIRx pipeline . The three constructs map cleanly onto that flow: scopes and layouts are what you express, and dispatch is what the compiler resolves on the way down.

Inside TIRx's 18-Pass Lowering Chain

Once tvm.compile(mod, target, tir_pipeline="tirx") is invoked, the program runs through an ordered chain of 18 module-level passes that turn explicit, hardware-native source into CUDA device code. The pipeline is deterministic and documented end to end, which matters for the agent-visible workflow TIRx is built around: every transformation is a named pass you can inspect rather than a hidden heuristic . The full sequence is LowerTIRx, UnifyThreadBinding, StmtSimplify, LowerTIRxOpaque, FlattenBuffer, BF16ComputeLegalize, NarrowDataType(32), VectorizeLoop, UnrollLoop, a second StmtSimplify, CommonSubexprElim, FP8ComputeLegalize, VerifyMemory, AnnotateEntryFunc, SplitHostDevice, MakePackedAPI, FP8StorageLegalize, and BF16StorageLegalize .

The chain front-loads the TIRx-specific work. LowerTIRx is where the abstract tile layer collapses into concrete code: its TilePrimitiveDispatch step replaces unresolved TilePrimitiveCall IR nodes with the selected backend dispatch bodies, and LowerTIRxCleanup resolves TileLayout buffer access into physical address arithmetic while lowering scope ids into thread axes . After that, the passes are largely standard TIR machinery — buffer flattening, vectorization, loop unrolling, common-subexpression elimination — operating on code that is now backend-specific.

Two details are worth reading closely. First, BF16 and FP8 legalization bracket the chain: compute legalization (BF16ComputeLegalize, FP8ComputeLegalize) runs early, before the arithmetic-oriented passes, while storage legalization (FP8StorageLegalize, BF16StorageLegalize) runs last, after the host/device split. That ordering ensures low-precision types reach CUDA PTX codegen in a form the backend accepts — arithmetic is legalized while the IR still reasons about values, and storage representation is fixed only once the device function is finalized . Second, the tail of the pipeline handles the runtime boundary: VerifyMemory catches illegal memory access patterns before codegen, SplitHostDevice separates host and device functions, and MakePackedAPI wraps device functions in TVM's runtime calling convention .

The practical takeaway for kernel engineers: because the boundary between programmer and compiler sits low, what you write survives mostly intact through these passes, and the points where automation does intervene — dispatch, legalization, memory verification — are explicit, ordered, and open to inspection .

The 29 Tile Primitives and How Each Backend Lowers Them

TIRx ships 29 tile primitives, each registered as tirx.tile.<name> in src/tirx/op/tirx.cc, with Python wrappers and TVMScript builders layered on top . These are the reusable middle ground between fully hand-written native code and a higher-level scheduling DSL: an author calls copy, copy_async, gemm, gemm_async, reductions, casts, elementwise math, or fused forms, and the compiler resolves the concrete instruction later. The catalog is explicit that these are not the whole vocabulary of a kernel — the native layer is still where you go when a hardware feature has no primitive.

The key design decision is late dispatch. When you call a tile primitive, it is recorded as an unresolved TilePrimitiveCall IR node rather than being lowered immediately . Those nodes survive in the IR until LowerTIRx, where TilePrimitiveDispatch replaces each one with a backend-specific body chosen by primitive type, execution scope, operand layouts, target backend, and optional hints . The practical payoff: the same kernel source retargets to a different backend by swapping dispatch rules, with no change to author code.

How a single primitive maps to instructions depends entirely on the resolved target:

So a copy in your source might become a TMA bulk transfer on one target and a plain vectorized load on another, and a gemm might lower to WGMMA, to Blackwell's tcgen05 path, or to a systolic-array instruction — without you rewriting the call . This is the mechanism behind TIRx's "bring up new hardware as a staged process" claim: a new accelerator needs new dispatch rules, not a new kernel.

One caveat worth respecting before you build on this surface: the official tile-primitives documentation warns that signatures and variants "may change" . Combined with the fact that the TIRx docs are labeled 0.26.dev0 while the shipped wheel is 0.25.0, treat the 29-primitive catalog as a pre-stable API in v0.25.0 — usable for experimentation, but not yet a frozen contract you should pin production kernels against.

Megakernels, Accelerator Bringup, and Programmatic IR Extension

TIRx is built for three workloads where explicit hardware control pays for itself: megakernels, new-accelerator bringup, and programmatic IR extension by tools and agents. The official overview describes megakernels as fusing tasks with "fine-grained dependencies and in-kernel scheduling" . In practice that means an author can fuse, say, an attention computation with a normalization step inside one kernel, expressing the dependency in explicit in-kernel control flow instead of writing an intermediate tensor back to global memory and paying runtime dispatch overhead between two separate launches.

Accelerator bringup is the second target. Because kernel logic and backend dispatch are separated — tile primitives are recorded as unresolved nodes and resolved later by TilePrimitiveDispatch — adding a new hardware target becomes, per the docs, "a staged process rather than a redesign" . You register new dispatch rules for an instruction or memory space without rewriting the kernel that uses them and without forking the compiler. TIRx is explicit that this is the layer you reach for when a hardware feature "arrives before a compiler can reliably automate" it .

The third workload — agentic kernel programming — is where the project's framing is most opinionated. TVM creator Tianqi Chen, who is faculty at Carnegie Mellon University, framed the June 2026 launch around the role of AI compilers in the age of agents, arguing that agents should have:

"access to a predictable DSL that offers maximum expressiveness, paired with a minimal compiler they can directly open up, build toolings, and improve for," — Tianqi Chen, TVM creator and CMU faculty (source: @tqchenml, 2026-06).

The concrete mechanism behind that claim is the FFI surface: TIRx exposes "IR and compiler utilities through TVM FFI across Python, C++, and Rust" , so a tool or LLM agent can inspect and extend the IR rather than treat it as a black box.

What makes agent-driven kernel search practical is a pre-benchmark static feedback loop. The docs describe a structured search space with feedback such as well-formedness, synchronization validity, race-freedom, and value simulation that runs before any GPU execution . An agent can therefore reject malformed or racy candidates at search time without burning accelerator cycles. Secondary reporting references search-space levels L1–L4, but that detail does not appear on the official documentation pages, so treat it as provisional .

TIRx Readiness Indicators: Where the Packaged Wheel and Documentation Diverge

Treat TIRx in v0.25.0 as a substantial bringup, not a frozen product. The packaged wheel — apache-tvm==0.25.0 on PyPI, released June 19, 2026 — contains the initial TIRx infrastructure plus follow-up fixes: a removed T.block API in TIRx script tests, dialect redirect fixes, and S-TIR transform-test repairs after the bringup . The release page is not titled a dedicated TIRx release, so the accurate reading is that 0.25.0 ships major TIRx bringup and documentation rather than a stable standalone compiler .

The most concrete divergence is versioning. The TIRx documentation pages are labeled tvm 0.26.dev0, while the latest stable wheel is 0.25.0 . Some documented surfaces — primitive signatures, builder APIs, pipeline pass names — may reflect post-0.25 development rather than what the 0.25.0 wheel actually exposes. Before depending on any TIRx API in production, verify it against tagged source at the v0.25.0 release, not against the docs alone.

Backend coverage carries the second caveat. A PR review comment references AWS Trainium, but the official author summary and documentation focus exclusively on CUDA and Blackwell-class codegen . Non-CUDA dispatch should be treated as not-yet-confirmed; the maturity you can verify today is the CUDA path.

Finally, remember what TIRx removes by design. The CUDA-native layer states plainly that there is no automatic scheduling — what the author writes is what gets emitted . There is no performance-portability guarantee: every property of the emitted kernel — vectorization, synchronization, occupancy, memory layout — is the author's responsibility, not the compiler's.

The takeaway: install apache-tvm==0.25.0 to explore TIRx today, pin to the tagged source for anything you ship, scope your expectations to CUDA/Blackwell, and budget for owning performance end to end. That is the trade TIRx makes explicit where Triton keeps it hidden.

Frequently asked questions

How does TIRx differ from Triton at the author level?

Both are tile-based DSLs, but they draw the programmer/compiler boundary in different places. Triton hides hardware orchestration behind block-level abstractions; TIRx keeps loop structure, warp IDs, synchronization, and pipeline state in the author's source code, while the compiler still reasons about reusable tile structure. As the official overview puts it, "orchestration stays in hardware-native source code, while the recurring tile-level structure becomes visible to the compiler" . In practice, no autoscheduler rewrites your iteration order — you express loops, branches, buffers, and backend intrinsics directly, then reach for tile primitives (copy, gemm, reductions) where reuse is worth it.

What does "no automatic scheduling" mean for performance?

It means what you write is exactly what gets emitted. The CUDA-native basics page states plainly that there is "no automatic scheduling" — TIRx does not reorder loops, insert prefetch hints, or add implicit thread predication on your behalf . The native layer is where you go when a hardware feature lacks a primitive. The result is full control paired with full responsibility: vectorization, occupancy, synchronization, and memory layout are encoded by the author, so performance is whatever those choices produce — not what an optimizer infers.

Which GPU architectures does TIRx support in v0.25.0?

NVIDIA Blackwell-class GPUs are the primary documented target, with dispatch examples lowering matrix multiply to WGMMA and tcgen05 instructions and copies to TMA . PR #19581, merged May 18, 2026, describes the work as low-level programming of "Blackwell-class GPU architectures" . AWS Trainium appears in a PR review comment but is not formally documented, and CUDA is the only confirmed codegen path — treat backend maturity beyond CUDA/Blackwell as unverified.

Is TIRx stable enough to depend on in production today?

Not as a stable standalone product yet. The v0.25.0 wheel, released June 19, 2026, passes 1,723 tests in tests/python/tirx (47 skipped) , which is a meaningful signal of basic soundness. But the TIRx documentation is labeled 0.26.dev0 while the latest stable wheel is 0.25.0 , and the tile-primitive page warns that signatures and variants may change . Pin to apache-tvm==0.25.0 and verify any API against the tagged source before shipping.

How does TIRx make its IR accessible to programmatic tooling or LLM-based writers?

TIRx exposes IR nodes and compiler utilities through TVM FFI across Python, C++, and Rust, so tools and agents can open up the IR rather than treat it as a black box . Tile primitive calls are recorded as unresolved TilePrimitiveCall nodes and remain inspectable before dispatch . The docs also describe static feedback — well-formedness, synchronization validity, race-freedom, and value simulation — that runs without GPU execution, enabling search-time validation for agent-generated kernels .