TileLang: Writing Near-Hardware-Limit GPU Kernels in 70 Lines of Python

Opening Scenario: A Typical Day for a CUDA Programmer

Imagine you are a kernel engineer at DeepSeek, tasked with writing a FlashAttention kernel for the attention layer of the V4 model. You open VS Code and start writing CUDA. Three days later, you have 3,000 lines of code: TMA copies, warp specialization, asynchronous pipelines, Blackwell MXFP8 block-scaled matrix multiplies—every line fighting the GPU memory hierarchy.

Then the hardware team tells you: they also need an AMD MI300X version, a Huawei Ascend NPU version, and an Apple Metal version.

Your three days become three months.

This is exactly the problem TileLang aims to solve: turning high-performance GPU kernel development from "rewrite for every hardware platform" into "write once, compile to five backends."

What It Is

TileLang is a domain-specific language (DSL) built on Apache TVM, using Pythonic syntax to write high-performance GPU/CPU/NPU kernels. It is not yet another CUDA wrapper—it is a layered programming model that lets you switch freely between abstraction levels:

• High level: one line of T.gemm(), automatically selecting the optimal tensor core path
• Mid level: manually managing tile layout, pipeline stages, warp specialization
• Low level: embedding T.CUDASourceCodeKernel and writing CUDA source directly

This "piercable abstraction" is TileLang's core design philosophy—abstraction is a tool, not a cage. When the high-level API falls short, you don't need to start over; you simply "drop down" to a lower abstraction at that spot.

Analogy: The History of Compiler Optimization

TileLang's position resembles compiler optimization in the 1970s. Before GCC, every programmer wrote assembly by hand because compiler-generated code "wasn't fast enough." GCC took 20 years to catch up with handwritten assembly. LLVM then spent another decade turning optimization into a modular pass pipeline.

GPU kernel programming is following the same trajectory:

• Handwritten CUDA = handwritten assembly (best performance, but not portable)
• Triton = early GCC (easier than CUDA, but limited backends)
• TileLang = LLVM (multi-backend, modular passes, extensible)

The key difference: TileLang was designed from the start with a multi-backend IR (intermediate representation). The same TileLang program, through different lowering passes, can compile to CUDA, ROCm, Metal, Ascend, and WebGPU. The TIRX migration in May 2026 moved this IR onto TVM's new representation layer, paving the way for future backend expansion.

The Data: Five Months of New Progress

This forum covered TileLang in late April 2026 (then called "TileLang + TileKernels"). Over the past five months, the project has undergone explosive expansion:

Backend Coverage

| Backend | Supported Hardware | Key PR | |:--------|:-------------------|:-------| | CUDA | SM70-SM100 (Volta-Blackwell) | INT4/FP16/MXFP8 tensor cores | | ROCm | RDNA3/RDNA4/CDNA | WMMA lowering (2026-04-22) | | Metal | Apple M1-M5 | Metal 4 cooperative-tensor GEMM | | Ascend | 910B/950 | Dual path: AscendC + NPU IR | | WebGPU | Browser/WASM | Supported since 2025-02 |

Five backends, one codebase. This is not "portable in theory"—DeepSeek V4's MLA (Multi-head Latent Attention) kernels are already running at production-grade performance on CUDA, ROCm, and Ascend.

Key Technical Milestones

• 2026-05-07: DeepSeek V4 operator support, including sparse MLA backward pass
• 2026-05-20: TIRX migration, unifying TileLang IR onto TVM's new representation layer
• 2026-05-11: Parallel autotuning with multi-GPU benchmarking and grouped compilation
• 2026-04-28: Blackwell MXFP8 grouped GEMM with block scaling + transposed B matrix
• 2026-04-25: HISA sparse attention indexer—hierarchical sparse attention
• 2026-03-24: Dual-SM Blackwell kernels with TMA + TMEM + TCGEN5 MMA
• 2026-03-18: Producer-consumer warp specialization with automatic pipelining
• 2026-02-02: TileLang Puzzles, 10 progressive exercises

Ecosystem Tools

TileLang is no longer just a DSL—it is becoming a language ecosystem:

• TileLang Puzzles: 10 interactive exercises, from simple GEMM to FlashAttention, similar to CUDA Mode puzzles but in TileLang syntax
• TileLang LSP: Language Server Protocol support—completion, go-to-definition, and type checking in VS Code
• DeepWiki integration: ask AI questions about TileLang directly
• Discord community: active developer discussion

Core Insight: The Right to Choose Your Abstraction Level

TileLang's most interesting design decision is not multi-backend—multi-backend is a natural TVM capability. What's truly interesting is the right to choose your abstraction level.

Traditional GPU programming frameworks give you two options:

1. High-level frameworks (PyTorch, JAX): write attention in one line, but no control over tile sizes, pipeline stages, or warp specialization 2. Low-level CUDA: full control, but rewrite for every hardware platform

TileLang offers a third option: mix different abstraction levels within the same codebase. You can write matrix multiplication with T.gemm() (high level), manually manage FlashAttention's pipeline (mid level), and embed CUDA source where needed (low level).

This resembles Python's design philosophy: give you choices without forcing them. You can write scripts or large systems; use OOP or functional style. TileLang brings this flexibility to GPU kernel programming.

Why It Matters Now

AI systems engineering in 2026 faces a structural contradiction:

• Accelerating model iteration: DeepSeek V3 → V3.2 → V4, with attention mechanisms changing every generation (MLA, sparse MLA, HISA)
• Intensifying hardware fragmentation: NVIDIA Blackwell, AMD MI300X, Huawei Ascend 950, Apple M5—each with different tensor core instruction sets
• Scarcity of kernel engineers: perhaps no more than a few hundred people worldwide can write production-grade FlashAttention

TileLang's value proposition: let an algorithm-savvy engineer write a kernel once and deploy it to five hardware platforms. This is not a "performance-for-portability" trade-off—in the DeepSeek MLA case, TileLang-generated kernels approach handwritten CUDA performance.

Cross-Domain Analogy: From CUDA to TileLang

This evolution also parallels the database field:

• Handwritten CUDA = handwritten SQL execution plans (1970s)
• Triton = early SQL optimizers (1980s—works, but not smart enough)
• TileLang = modern cost-based optimizers (2000s—automatically picking the best path)

Each stage transfers "expert knowledge" from human brains into tools. TileLang's autotuner (parallel autotuning introduced 2026-05-11) does what a CBO does: given a kernel description, automatically search for optimal tile sizes, pipeline stage counts, and warp specialization strategies.

Limitations

TileLang is not a silver bullet:

• Learning curve: easier than CUDA, but still requires understanding GPU memory hierarchy (shared memory, registers, TMA, TMEM)
• Performance ceiling: in extreme scenarios (custom sparse patterns, non-standard data layouts), handwritten CUDA still has an edge
• Backend maturity gaps: the CUDA backend is most mature; Ascend and Metal backends are still iterating rapidly
• Ecosystem size: compared to Triton (backed by OpenAI) and CUDA (NVIDIA native), TileLang's community is still small

Conclusion

TileLang represents an engineering philosophy: finding a balance between abstraction and performance, not either/or. It does not aim to "eliminate CUDA"—it aims to let 80% of kernel engineers avoid writing CUDA, and let the remaining 20% of experts iterate more efficiently.

When DeepSeek V4's MLA kernel can be written in 70 lines of TileLang instead of 3,000 lines of CUDA, that is not a "performance compromise"—that is a productivity revolution. Just as GCC didn't eliminate assembly but made most programmers no longer need to write it, TileLang is turning GPU kernel programming from a "craft" into "engineering." Eighteen months into this path, it has only just begun.

---

Project: github.com/tile-ai/tilelang Docs: tilelang.com Paper: arXiv:2504.17577 Puzzles: github.com/tile-ai/tilelang-puzzles