$ show NVIDIA/tilus
Tilus is a tile-level GPU kernel programming language and compiler framework, primarily targeting NVIDIA architectures (SM80, SM90, SM100). It provides a Python DSL (`tilus.Script`) for writing high-performance CUDA kernels with explicit management of shared/register memory, tensor cores (WGMMA, tcgen05), and advanced synchronization primitives (mbarrier, cluster sync, TMA). The compiler stack includes a custom IR with layout inference (register/shared/tensor-memory), a pass pipeline (dead code elimination, bound-aware simplification, sub-byte type lowering), and backend emitters that generate PTX/CUDA code. Autotuning over tile sizes, warp counts, and pipeline stages is built into the framework. The architecture separates concerns into three layers: the user-facing `lang` package (Script, autotune, instructions), the `ir` package (tensors, layouts, instructions, inference rules), and the `backends` package (emitters, contexts, code generation). The `hidet` subpackage integrates a vendored copy of the Hidet IR and transforms, which provides the low-level expression/statement IR and a comprehensive set of optimization passes. Examples cover progressive matmul implementations (V0–V6) for Ampere, Hopper, and Blackwell, as well as flash attention, layer norm, and quantization kernels. This is a research project (ASPLOS '26 paper) with a strong emphasis on correctness and performance, evidenced by CUDA-level validation against PyTorch and cuBLAS. While the codebase is well-structured and demonstrates advanced compiler design patterns, it remains in beta (version 0.2.0) and carries some duplication between vendored Hidet code and native Tilus code, as well as hardware-specific constraints (NVGPU only). The layout inference system and autotuning infrastructure are particularly notable as reusable components.
The LLM identified 3 files worth highlighting for reuse:
A compiler pass that reorders and simplifies addition chains, including cancellation of div/mod pairs. It uses a `DepthAnalyzer` to determine loop nesting depth and reorders terms to exploit bound-derived simplifications. This is a self-contained, well-documented optimization that could be extracted into any compiler pipeline dealing with arithmetic-heavy code.
Implements fast integer division and modulo using precomputed magic multipliers (reciprocal approximation). Generates efficient PTX sequences via multiplication by magic constant and shift. This technique is hardware-independent (works on any GPU with multiply-high) and the implementation is cleanly parameterized, making it reusable for any compiler needing to lower integer div/mod operations.
Implements multi-function composition for tensor layouts, representing a layout as a pair of spatial and local functions. Provides operations like `compose`, `divide`, `permute`, `reshape`, and canonicalization. This abstraction is inspired by CuTe and can be reused in any tensor compiler or linear algebra library that needs to manipulate data layouts across memory hierarchies.