TileLang Beats Triton by 5x Writing GPU Kernels in 80 Lines
TileLang is a Python DSL that lets you write GPU kernels with Triton-like ease but CUTLASS-level control — and DeepSeek just used it in production.
PRO- TileLang is a Python DSL for writing high-performance GPU kernels that sits between Triton (easy, less control) and CUTLASS (full control, complex C++ templates).
- Benchmarks show up to 5× speedup over Triton on H100 and 6× on AMD GPUs, with 90% less code than manual CUDA implementations.
- DeepSeek used TileLang in production for DeepSeek-V4, replacing hundreds of Torch ATen operators with fused kernels for training and inference.
- Microsoft's BitBLAS and AttentionEngine are both built on TileLang; Qwen's FlashQLA linear attention library also uses it.
- Supports NVIDIA (H100, A100, RTX 4090), AMD (MI300X, MI250), Apple Metal, WebGPU, and Huawei Ascend NPUs.
- Free and open-source, installable via
pip install tilelang, with a GitHub repo at ~6.9k stars and growing.
TileLang has quietly become one of the most interesting tools in the GPU kernel programming space. Triton is fast to write but surrenders too much control to the compiler. CUTLASS gives you full control at the cost of hundreds of lines of C++ template machinery. TileLang offers a third path: write Python, explicitly decide what lives in shared memory and how the pipeline is staged, and let a layout inference pass handle the rest.
The problem it solves
Modern AI workloads depend on optimized kernels for training and inference. These kernels follow clear data-flow patterns: move tiles between DRAM and SRAM, compute on those tiles, repeat. The patterns are simple; the implementation is not. Achieving peak performance requires hardware-centric decisions that most compilers either hide from you or force you to express in verbose, brittle code.
The original author of TileLang described writing 500 lines of scheduling primitives per operator during his Microsoft Research internship, code that only he could understand. Fused kernels are increasingly necessary for performance, but Triton lacks the fine-grained control to implement them well, and hand-writing CUDA or CUTLASS produces code that is hard to read and harder to maintain.
How it works
TileLang is a domain-specific language built on TVM. It targets GPU and CPU kernels, including GEMM, dequantized GEMM, FlashAttention, and linear attention variants. The programming model rests on three ideas:
- Tiles are first-class objects. A shaped chunk of data, say
block_M × block_K, is owned by a thread block, warp, or thread. You reason at the tile level, not the individual thread level. - You place buffers in the memory hierarchy yourself.
T.alloc_sharedputs data in shared memory;T.alloc_fragmentputs it in registers. Triton hides shared-memory allocation inside the compiler. TileLang does not. - The compiler infers the thread mapping. Once you declare where a tile lives and what operation runs on it, a layout inference pass parallelizes it across threads and works out the register and shared-memory layouts automatically, with the option to override.
This story is for Pro members
You've reached the end of the free preview. Upgrade to AlphaSignal Pro to read the full article - and everything else behind the paywall.
Comments
