TIRx: A Unified Tile Primitive ML Compiler Across Heterogeneous Backends
Organizations: Carnegie Mellon University · Shanghai Jiao Tong University · NVIDIA · Princeton University ·
Abstract
Modern ML kernels evolve rapidly with new optimizations and hardware features. ML compiler abstractions must accommodate richer memory hierarchies, more forms of thread cooperation, and increasingly heterogeneous backends while retaining the control required for peak performance. We present TIRx, an ML compiler for GPUs and AI accelerators that exposes these ingredients through explicit interfaces. TIRx introduces an explicit execution-scope abstraction that captures participating hardware teams at multiple granularities. It also uses tensor layouts to map logical tensors to heterogeneous hardware spaces through named axes and shard/replica/offset semantics. Finally, it introduces tile-primitive dispatch, which lowers high-level operations to efficient implementations according to scope, layout, and backend. This design preserves native-level control while automating recurring address calculation, operand preparation, and instruction-loop generation. Our evaluation shows vendor-level performance across the latest GPUs and AI accelerators, while native-code case studies show substantial reductions in low-level implementation code.
Figures & tables
| Axis | Backend | Extent | Use |
|---|---|---|---|
| m | all | buffer-defined | linear memory coordinate |
| laneid | NVIDIA | 32 | lane within a warp |
| warpid | NVIDIA | launch-defined | warp within a CTA |
| wid_in_wg | NVIDIA | 4 | warp within a warpgroup |
| tid_in_wg | NVIDIA | 128 | thread within a warpgroup |
| cbx/cby/cbz | NVIDIA | launch-defined | CTA coordinate in a cluster |
| Primitive | Specialized family | Fallback |
|---|---|---|
| copy | ldmatrix/stmatrix , NKI copy | vector/scalar load/store |
| copy_async | TMA, tcgen05.cp/ld/st , cp.async | none |
| gemm | mma.sync.aligned , nki.matmul | scalar FMA |
| gemm_async | tcgen05.mma | none |
| reduce | warp shuffle, NKI tensor reduce | vector/scalar arithmetic |
| pointwise | NKI/vector arithmetic | scalar arithmetic |
GPU Kernel Generation
GPU (graphics processing unit) kernel generation automates writing and optimizing device code for faster computation across accelerators.
See all 70 papers