TileLang logo

Tile Language (tile-lang) is a concise domain-specific language designed to streamline the development of high-performance GPU/CPU 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.

TileLang tiled matrix multiplication example

Latest News

  • 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.
Earlier news (2025–2026)

2026

  • 2026-07-28 — Source-aware compiler diagnostics: carried Python source locations into TIRX and surfaced them in compiler errors.
  • 2026-07-24 — Multi-backend language dialect: reorganized the language layer around shared semantics with static CUDA, ROCm, and Metal dialects.
  • 2026-07-23 — Block-causal attention for dLLM: added fixed-length and variable-length block-causal attention examples for diffusion language models.
  • 2026-07-22 — IR Lower Trace: introduced a debugging tool for inspecting IR changes across every compiler pass and the final code-generation step.
  • 2026-07-22 — DeepSeek V3.2 sparse MLA backward: selected the launch width adaptively from the head-block size.
  • 2026-07-21 — DeepSeek V3.2 top-k optimization: improved the sparse-attention top-k selector's memory access pattern, delivering approximately 1.9× higher performance in the reported benchmark.
  • 2026-07-16 — Compiler pass timing: added profiling for compiler passes with a configurable reporting threshold.
  • 2026-07-12 — IKET profiler integration: added CUDA timeline instrumentation and profiling support.
  • 2026-07-08 — TileLang v0.1.12: added the LLVM backend, tile scheduler, backend registry, pass visualizer, and expanded Blackwell support.
  • 2026-07-06 — Pass Visualizer: introduced a structure-tree browser for inspecting compiler transformations.
  • 2026-06-26 — Cross-host CUDA binary cache: enabled compiled CUDA binaries to be reused across compatible hosts.
  • 2026-06-24 — Tile scheduler: introduced persistent tile-scheduling primitives for kernel authors.
  • 2026-06-24 — Backend CodeGen registry: moved device and host CodeGen dispatch behind a backend registry.
  • 2026-06-18 — LLVM backend: added CPU lowering and execution through LLVM.
  • 2026-06-18 — Arbitrary-layout TMA lowering: enabled TMA transfers for swizzled shared-memory layouts.
  • 2026-06-16 — Pass Diff: added compiler-pass IR comparison for debugging lowering changes.
  • 2026-06-08 — TileLang v0.1.11: expanded scan, pipeline, backend, CUDA, ROCm, and Metal functionality.
  • 2026-05-25 — Scan operators: introduced tile-level scan primitives.
  • 2026-05-25 — TileLang v0.1.10: broadened AMD and Blackwell support, added initial Metal GEMM, improved Windows packaging, and expanded autotuning.
  • 2026-05-24 — CDNA4 MXFP4: added FP4 E2M1 matrix-core support for AMD gfx950.
  • 2026-05-22 — Metal simdgroup GEMM: added the first Metal T.gemm path using simdgroup_matrix MMA.
  • 2026-05-20 — Cluster copies: introduced T.copy_cluster for TMA multicast and SM-to-SM cluster transfers.
  • 2026-05-20 — TMA gather/scatter: added tile::gather4 and tile::scatter4 support.
  • 2026-05-20 — Native SM75 MMA GEMM: added FP16, INT8, and INT4 tensor-core paths for Turing GPUs.
  • 2026-05-20 — TIRX migration: moved TileLang IR usage to TVM's TIRX representation.
  • 2026-05-11 — Parallel autotuning: added pipelined compilation, grouped compilation, and multi-GPU benchmarking.
  • 2026-05-07 — DeepSeek V4 operators: added TileLang examples for DeepSeek V4 workloads.
  • 2026-05-06 — Windows support: added complete Windows build and runtime support with cross-platform fixes.
  • 2026-04-28 — MXFP8 grouped GEMM: added block-scaled grouped GEMM examples with transposed-B support on Blackwell.
  • 2026-04-25 — HISA sparse-attention indexer: added hierarchical sparse-attention indexing examples.
  • 2026-04-24 — Blackwell MXFP8 block-scaled GEMM: added MXFP8 block-scaled matrix multiplication on SM100.
  • 2026-04-22 — TileLang v0.1.9: delivered CuTe DSL GEMM V2, Metal code generation improvements, and build-without-host-toolchain support.
  • 2026-04-22 — RDNA3/RDNA3.5 WMMA: added WMMA lowering for AMD gfx11 GPUs.
  • 2026-04-20 — INT4 T.gemm: added INT4 matrix multiplication to the CUDA GEMM path.
  • 2026-04-17 — CUDA source kernels: introduced T.CUDASourceCodeKernel for embedding custom CUDA source.
  • 2026-04-15 — AutoDD frozen regions: added __freeze__ annotations to preserve selected code during automatic delta debugging.
  • 2026-03-27 — TMA stores: added store support to T.tma_copy.
  • 2026-03-24 — Two-SM Blackwell kernels: added two-SM TMA, TMEM, and TCGEN5 MMA support.
  • 2026-03-23 — AMD RDNA4: upgraded the ROCm path and added RDNA4 GPU support.
  • 2026-03-22 — FlashAttention on SM100: added Blackwell FlashAttention examples.
  • 2026-03-18 — Producer-consumer warp specialization: added automatic warp-specialized pipelines and the T.tma_copy API.
  • 2026-03-12 — Eager-mode autotuning: enabled the autotuner with eager JIT kernels.
  • 2026-03-10 — CPU T.gemm: added matrix multiplication support for the CPU target.
  • 2026-03-05 — IR dump configuration: added a TileLang pass configuration for dumping intermediate IR.
  • 2026-02-28 — CUDA cluster primitives: added cluster launch, query, synchronization, and barrier operations.
  • 2026-02-28 — TCGEN5 MMA tensor-shared path: added the tensor-memory/shared-memory Blackwell GEMM path.
  • 2026-02-24 — Host-toolchain-free builds: enabled installation without a host C/C++ toolchain when supported artifacts are available.
  • 2026-02-23 — CuTe DSL GEMM V2: added SM90 and SM100 GEMM V2 support to the CuTe DSL backend.
  • 2026-02-16 — TileLang v0.1.8: shipped dynamic pipeline improvements, logging documentation, richer layout representations, and AMD fixes.
  • 2026-02-14 — Cross-CUDA release wheels: unified multiple CUDA versions behind a single wheel.
  • 2026-02-14 — Hierarchical reductions: added hierarchical and warp-level reduction intrinsics.
  • 2026-02-09 — CUDA runtime stubs: added lazy-loading CUDART and NVRTC stubs for CUDA 11, 12, and 13 compatible wheels.
  • 2026-02-08 — Layout visualization improvements: improved rendering and inspection of TileLang layouts.
  • 2026-02-02 — TileLang Puzzles: published ten progressively harder exercises for learning TileLang interactively.

2025

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, and Metal 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.

Backend Target Platforms and hardware Support level Notes
NVIDIA CUDA cuda Linux x86-64/AArch64, Windows x86-64; code paths from SM70 through SM120 Primary Release wheels and CI coverage; TMA, WGMMA, and TMEM features require the corresponding GPU architecture.
AMD ROCm/HIP hip Linux; CDNA and RDNA GPUs, including gfx942/gfx950 paths Supported Included in Linux wheels; a ROCm runtime is required. CI runs on a self-hosted gfx942 (MI300X) runner; gfx950 is not yet covered.
Apple Metal metal macOS on Apple silicon Supported Release wheels and CI coverage; Metal 4 cooperative tensors are available on supported M5 systems.
LLVM CPU llvm Host CPUs Experimental Build from source with USE_LLVM=ON; LLVM 15 or newer is required.
NVIDIA CuTe DSL cutedsl NVIDIA GPUs Experimental Requires nvidia-cutlass-dsl.
WebGPU webgpu WebGPU runtimes Experimental Code generation and runtime integration are still evolving.
Huawei Ascend Ascend C / NPU IR Ascend A2 and A3 Ecosystem Developed in tilelang-ascend and the MLIR-based tilelang-mlir-ascend.
MetaX MACA maca MetaX C500 and C600 Ecosystem Developed in tilelang-metax; requires the MACA software stack.
Moore Threads MUSA musa S5000, S4000, and M1000 Ecosystem Developed in tilelang-musa and released independently.
HYGON hcu Linux; BW1000, BW1100, BW150 and K100_AI Ecosystem Developed in tilelang-hygon; requires the DTK software stack.
Sunrise-AI TANG tang Sunrise S2 and S3 Ecosystem Developed in tilelang-sunrise. The TANG software stack is required.

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

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

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

    mla decode performance bs64 on H100
    mla decode performance bs128 on H100

  • Flash Attention Performance on H100

    operator performance on H100

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

    gemm fp16 performance on Gpus

  • Dequantize Matmul Performance on A100

    dequantize gemv performance on A100

Join the Discussion

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

Join our Discord

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.

Key dates

PyPI data
Data sourced directly from PyPI's database.
  • Released:
Latest release

1 maintainer

PyPI data
Data sourced directly from PyPI's database.
Avatar for tilelang from gravatar.com tilelang

Credits

Author: TileLang Contributors, Tile-AI

Maintainer: Lei Wang

License expression

MIT
View SPDX License List

Requires

Python >=3.10

Provides Extra

fp4 vis nvcc

Tags

BLAS CUDA HIP Code Generation TVM

Classifiers

Report project as malware

Download files

Download the file for your platform. If you're not sure which to choose, learn more about installing packages.

Source Distribution

tilelang-0.1.14.tar.gz (91.7 MB view details)

Uploaded Source

Built Distributions

If you're not sure about the file name format, learn more about wheel file names.

tilelang-0.1.14-cp39-abi3-win_amd64.whl (28.5 MB view details)

Uploaded CPython 3.9+Windows x86-64

tilelang-0.1.14-cp39-abi3-manylinux_2_34_aarch64.whl (41.0 MB view details)

Uploaded CPython 3.9+manylinux: glibc 2.34+ ARM64

tilelang-0.1.14-cp39-abi3-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl (45.4 MB view details)

Uploaded CPython 3.9+manylinux: glibc 2.27+ x86-64manylinux: glibc 2.28+ x86-64

tilelang-0.1.14-cp39-abi3-macosx_11_0_arm64.whl (32.7 MB view details)

Uploaded CPython 3.9+macOS 11.0+ ARM64

File details

Details for the file tilelang-0.1.14.tar.gz.

File metadata

  • Download URL: tilelang-0.1.14.tar.gz
  • Upload date:
  • Size: 91.7 MB
  • Tags: Source
  • Uploaded using Trusted Publishing? Yes
  • Uploaded via: twine/7.0.0 CPython/3.13.14

File hashes

Hashes for tilelang-0.1.14.tar.gz
Algorithm Hash digest
SHA256 7ae0fba6c63dae80b3c9c5c11a70bb8d7165794994f290f36e2fe87bfb143cfc
MD5 0466736f510650678c33268b2246f2e7
BLAKE2b-256 98b6b2a79ca220c63f0ec43ece86c0a37f92af9e22261b13951e58700b008a2b

See more details on using hashes here.

Provenance

The following attestation bundles were made for tilelang-0.1.14.tar.gz:

Publisher: dist.yml on tile-ai/tilelang

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file tilelang-0.1.14-cp39-abi3-win_amd64.whl.

File metadata

  • Download URL: tilelang-0.1.14-cp39-abi3-win_amd64.whl
  • Upload date:
  • Size: 28.5 MB
  • Tags: CPython 3.9+, Windows x86-64
  • Uploaded using Trusted Publishing? Yes
  • Uploaded via: twine/7.0.0 CPython/3.13.14

File hashes

Hashes for tilelang-0.1.14-cp39-abi3-win_amd64.whl
Algorithm Hash digest
SHA256 3eff7e03e8e86d69ba2afd6b9a7df00c5d44323b0323e8cb06cb489601ecd52c
MD5 fe6970cd4b7685eecb98c1fe65a8dd0c
BLAKE2b-256 aa1d97fb5b21437a9fae43f16523b7eae5f98f27a87e56c6026c6e5d2ae03f6d

See more details on using hashes here.

Provenance

The following attestation bundles were made for tilelang-0.1.14-cp39-abi3-win_amd64.whl:

Publisher: dist.yml on tile-ai/tilelang

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file tilelang-0.1.14-cp39-abi3-manylinux_2_34_aarch64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.14-cp39-abi3-manylinux_2_34_aarch64.whl
Algorithm Hash digest
SHA256 99308791d356732aa30525d28ce686e32c9b3961852ab380be0d688fff448e74
MD5 308bb77469e93ade836a364ad865168e
BLAKE2b-256 6aea1463c3c736f7ace24169caeaa14189fa4aa486206179d92a47283339883a

See more details on using hashes here.

Provenance

The following attestation bundles were made for tilelang-0.1.14-cp39-abi3-manylinux_2_34_aarch64.whl:

Publisher: dist.yml on tile-ai/tilelang

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file tilelang-0.1.14-cp39-abi3-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.14-cp39-abi3-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl
Algorithm Hash digest
SHA256 fcc5d47ac1baeb7a8555ff5384af9ce9f3efef46dd926bc23d2c1cd425f68600
MD5 bdd6ce3ad5c6647ece6f8550f3097439
BLAKE2b-256 4069e69a18eaf239efaa1718922a26d0a4605a0525a2263d3224f506354c02ae

See more details on using hashes here.

Provenance

The following attestation bundles were made for tilelang-0.1.14-cp39-abi3-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl:

Publisher: dist.yml on tile-ai/tilelang

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file tilelang-0.1.14-cp39-abi3-macosx_11_0_arm64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.14-cp39-abi3-macosx_11_0_arm64.whl
Algorithm Hash digest
SHA256 a78003b810b63092de380accc1e73c0139a1f11b0377e1335774b72845ba0571
MD5 1a921f6573e50d9ce19fcb3e88b90fd1
BLAKE2b-256 9113aa13f149ea8c83c1115f715c2e611cec9a09dfdec1d0871551dda3979c16

See more details on using hashes here.

Provenance

The following attestation bundles were made for tilelang-0.1.14-cp39-abi3-macosx_11_0_arm64.whl:

Publisher: dist.yml on tile-ai/tilelang

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

Release history Release notifications | RSS feed

This release

0.1.14 This release

5 files

0.1.13

5 files

0.1.12

5 files

0.1.11

5 files

0.1.10

5 files

0.1.9

4 files

0.1.8

4 files

0.1.7.post3

4 files

0.1.7.post2

3 files

0.1.7.post1

3 files

0.1.7

4 files

0.1.6.post2

2 files

0.1.6.post1

5 files

Yanked

0.1.6

5 files

Yanked reason: static link g++

0.1.5

5 files

0.1.4

5 files

0.1.3

5 files

0.1.2.post1

5 files

0.1.2

5 files

0.1.1

5 files

0.1.0

5 files

0.0.1

1 file

Pre-release

0.0.1.dev0

1 file

Anthropic, PBC Visionary sponsor Bloomberg Visionary sponsor Hudson River Trading Visionary sponsor Meta Visionary sponsor NVIDIA Visionary sponsor Microsoft Sustainability sponsor Depot Continuous Integration AWS Cloud computing and Security Sponsor Datadog Monitoring Fastly CDN Google Download Analytics Sentry Error logging StatusPage Status page