tile-ai/tilelang

Tile Language (tile-lang) is a concise domain-specific language designed to streamline the development of high-performance GPU/CPU/NPU kernels (e.g., GEMM, Dequant GEMM, FlashAttention, LinearAttention). By employing a Pythonic syntax with an underlying compiler infrastructure on top of TVM, tile-lang allows developers to focus on productivity without sacrificing the low-level optimizations necessary for state-of-the-art performance.

Latest News

• 2026-09-30 — Ascend 950 backend: TileLang now officially supports Huawei Ascend 950 NPUs with native code generation, automatic scheduling and synchronization, SIMD/SIMT vector programming, etc. Explore the Ascend examples for GEMM, FlashAttention, and more.
• 2026-08-04 — TileLang LSP open sourced: published a Language Server Protocol implementation for TileLang with inlay hints for buffer shapes, dtypes, scopes, and inferred layouts, plus hover details and precise diagnostics.
• 2026-08-03 — TileLang v0.1.13: shipped the multi-backend language dialect, source locations in compiler diagnostics, new CUDA and Metal hardware paths, and a broad set of correctness fixes. This release removes several legacy APIs; read the compatibility notes before upgrading.
• 2026-07-30 — SM120 NVF4 block-scaled MMA: added an optimized Blackwell path for T.mma_gemm_blockscaled and a corresponding SM120 example.
• 2026-07-28 — Metal 4 cooperative-tensor GEMM: added cooperative-tensor T.gemm support for Apple M5, while retaining the simdgroup fallback for unsupported shapes and systems.

See all releases for complete changelogs and compatibility notes.

Platform and Backend Support

TileLang is evolving into a multi-backend compiler (TileLang-X) built around a modular backend abstraction. See the backend architecture for the design, or ask a coding agent to use the backend integration skill when porting TileLang to a new backend.

The currently supported backends are listed below. Primary identifies TileLang’s core backend, while Supported and Experimental backends are implemented in the main repository. Ecosystem adapters live in separate repositories, are not included in TileLang release wheels, and may follow independent compatibility schedules. Prebuilt wheels are available for Linux x86-64/AArch64, Windows x86-64, and macOS arm64.

TileLang uses Target objects to represent compilation targets. The default auto target detects CUDA, HIP, Metal, and Ascend devices; select an explicit target when compiling for another backend or architecture. See the target guide for target syntax, architecture options, and backend-specific notes, or the corresponding adapter repository for installation and tested-device details.

Installation

Install the latest stable release from PyPI:

pip install tilelang

Verify the installation:

python -c "import tilelang; print(tilelang.__version__)"

Nightly wheels provide recent features and fixes before the next stable release:

pip install tilelang --find-links https://tile-ai.github.io/whl/nightly

On AMD GPUs the same Linux wheels work out of the box: install a ROCm build of PyTorch first (e.g. pip install torch --index-url https://download.pytorch.org/whl/rocm7.0), then pip install tilelang. A host ROCm installation is required at runtime; see the ROCm notes in the installation guide.

Nightly builds may be less stable than official releases. For source builds, editable installs, Docker, ROCm setup, pip-provided CUDA toolchains, or a custom TVM checkout, follow the complete installation guide.

Quick Start

The following example defines, compiles, runs, and verifies an FP16 GEMM kernel with FP32 accumulation and a fused ReLU epilogue. It uses PyTorch CUDA tensors; PyTorch uses the same cuda device name on ROCm systems. TileLang selects the target automatically from the current environment.

import torch
import tilelang
import tilelang.language as T

@tilelang.jit
def matmul_relu(A, B, block_M: int = 128, block_N: int = 128, block_K: int = 32):
    M, N, K = T.const("M, N, K")
    A: T.Tensor((M, K), T.float16)
    B: T.Tensor((K, N), T.float16)
    C = T.empty((M, N), T.float16)

    with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), T.float16)
        B_shared = T.alloc_shared((block_K, block_N), T.float16)
        C_local = T.alloc_fragment((block_M, block_N), T.float32)

        T.clear(C_local)
        for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):
            T.copy(A[by * block_M, k * block_K], A_shared)
            T.copy(B[k * block_K, bx * block_N], B_shared)
            T.gemm(A_shared, B_shared, C_local)

        for i, j in T.Parallel(block_M, block_N):
            C_local[i, j] = T.max(C_local[i, j], 0)

        T.copy(C_local, C[by * block_M, bx * block_N])

    return C

M = N = K = 1024
a = torch.randn((M, K), device="cuda", dtype=torch.float16)
b = torch.randn((K, N), device="cuda", dtype=torch.float16)
c = matmul_relu(a, b)
torch.testing.assert_close(c, torch.relu(a @ b), rtol=1e-2, atol=1e-2)
print("GEMM + ReLU passed.")

@tilelang.jit specializes the kernel for the input shape and compile-time arguments on first use. T.Pipelined stages global-to-shared transfers, T.gemm maps the tile operation to the target backend, and T.Parallel expresses the elementwise ReLU epilogue. Continue with the language basics, then explore the GEMM examples for layouts, autotuning, and architecture-specific optimizations.

Examples

• Start here: quickstart and elementwise kernels
• GEMM and quantization: GEMM, grouped GEMM, FP8 GEMM, dequantization GEMM, and block-scaled GEMM
• Attention and sequence models: FlashAttention, Flash Decoding, block-sparse attention, linear attention, and GDN
• Model workloads: DeepSeek MLA, DeepSeek V3.2, DeepSeek V4, and DeepSeek mHC
• Architecture-specific kernels: AMD, TCGEN05, and SM120
• Compiler and debugging tools: analyzer, layout visualization, AutoDD, and IKET

Browse the complete examples directory for additional operators, tests, and architecture-specific implementations.

Benchmark Summary

TileLang achieves exceptional performance across a variety of computational patterns. Comprehensive benchmark scripts and settings are available at tilelang-benchmark. Below are selected results showcasing its capabilities:

• MLA Decoding Performance on H100

• Flash Attention Performance on H100

• Matmul Performance on GPUs (RTX 4090, A100, H100, MI300X)

• Dequantize Matmul Performance on A100

---

Join the Discussion

Welcome to join our Discord community for discussions, support, and collaboration!

Acknowledgments

We would like to express our gratitude to the TVM community for their invaluable contributions. The initial version of this project was mainly developed by LeiWang1999, chengyupku and nox-410 with supervision from Prof. Zhi Yang at Peking University. Part of this work was carried out during an internship at Microsoft Research, where Dr. Lingxiao Ma, Dr. Yuqing Xia, Dr. Jilong Xue, and Dr. Fan Yang offered valuable advice and support. We deeply appreciate their mentorship and contributions.