tile-ai/tilelang

What it solves

TileLang simplifies the creation of high-performance GPU and CPU kernels (such as GEMM, FlashAttention, and LinearAttention) by providing a concise, Pythonic domain-specific language. It removes the need for developers to manually handle low-level optimizations while maintaining the state-of-the-art performance typically required for AI workloads.

How it works

TileLang acts as a high-level interface built on top of the Apache TVM compiler infrastructure. It allows developers to define kernels using a tiled approach—allocating shared memory and fragments, and using primitives like T.gemm and T.Pipelined to manage data movement and computation. The compiler then lowers this high-level description into optimized code for various backends, including NVIDIA CUDA, AMD ROCm/HIP, Apple Metal, and LLVM for CPUs.

Who it’s for

Kernel developers and AI engineers who need to implement custom, highly optimized operators for machine learning models on diverse hardware architectures without writing raw low-level GPU code.

Highlights

• Multi-Backend Support: Native support for NVIDIA (SM70 to SM120), AMD (CDNA/RDNA), Apple Silicon (Metal), and CPUs (LLVM).
• Hardware-Specific Optimizations: Includes specialized paths for Blackwell MXFP8 block-scaled GEMM, Metal 4 cooperative tensors, and structured-sparse matrix multiplication.
• Productivity Tools: Features a dedicated Language Server Protocol (LSP) for IDE support, an IR trace tool for debugging compiler passes, and a pass visualizer.
• Tiled Programming Model: Uses a Pythonic syntax to express pipelined transfers and parallel operations, reducing the complexity of manual memory management.