Skip to main content

A tile level programming language to generate high performance code.

Project description

Tile Language

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.

Latest News

  • 03/03/2025 🚀: Added high-performance MLA Decoding support using only 80 lines of Python code, achieving performance on par with FlashMLA on H100 (see example_mla_decode.py)! We also provide documentation explaining how TileLang achieves this.
  • 02/15/2025 ✨: Added WebGPU Codegen support, see Pull Request #86!
  • 02/12/2025 ✨: Excited to announce the release of v0.1.0!
  • 02/10/2025 🚀: Added debug tools for TileLang—T.print for printing variables/buffers (docs) and a memory layout plotter (examples/plot_layout).
  • 01/20/2025 ✨: We are excited to announce that tile-lang, a dsl for high performance AI workloads, is now open source and available to the public!

Tested Devices

Although tile-lang aims to be portable across a range of Devices, it has been specifically tested and validated on the following devices: for NVIDIA GPUs, this includes the H100 (with Auto TMA/WGMMA support), A100, V100, RTX 4090, RTX 3090, and RTX A6000; for AMD GPUs, it includes the MI250 (with Auto MatrixCore support) and the MI300X (with Async Copy support).

OP Implementation Examples

tile-lang provides the building blocks to implement a wide variety of operators. Some examples include:

Within the examples directory, you will also find additional complex kernels—such as convolutions, forward/backward passes for FlashAttention, more operators will continuously be added.

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

Installation

Method 1: Install with Pip

The quickest way to get started is to install the latest release from PyPI:

pip install tilelang

Alternatively, you can install directly from the GitHub repository:

pip install git+https://github.com/tile-ai/tilelang

Or install locally:

# install required system dependencies
sudo apt-get update
sudo apt-get install -y python3-setuptools gcc libtinfo-dev zlib1g-dev build-essential cmake libedit-dev libxml2-dev

pip install .  # with -e option if you want to install in editable mode

Method 2: Build from Source

We currently provide three ways to install tile-lang from source:

Quick Start

In this section, you'll learn how to write and execute a straightforward GEMM (matrix multiplication) kernel using tile-lang, followed by techniques for layout optimizations, pipelining, and L2-cache–friendly swizzling.

GEMM Example with Annotations (Layout, L2 Cache Swizzling, and Pipelining, etc.)

Below is an example that demonstrates more advanced features: layout annotation, parallelized copy, and swizzle for improved L2 cache locality. This snippet shows how to adapt your kernel to maximize performance on complex hardware.

import tilelang
import tilelang.language as T
# `make_mma_swizzle_layout` is a python defined layout function
# specifically designed for for MMA operations
# which ensures the consistency with the nvidia CUTLASS Library.
# to avoid bank conflicts and maximize the performance.
from tilelang.intrinsics import (
    make_mma_swizzle_layout as make_swizzle_layout,)

def matmul(M, N, K, block_M, block_N, block_K, dtype="float16", accum_dtype="float"):
    # add decorator @tilelang.jit if you want to return a torch function
    @T.prim_func
    def main(
        A: T.Buffer((M, K), dtype),
        B: T.Buffer((K, N), dtype),
        C: T.Buffer((M, N), dtype),
    ):
        # Initialize Kernel Context
        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), dtype)
            B_shared = T.alloc_shared((block_K, block_N), dtype)
            C_local  = T.alloc_fragment((block_M, block_N), accum_dtype)

            # Apply layout optimizations or define your own layout (Optional)
            # If not specified, we will deduce the layout automatically
            # T.annotate_layout({
            #     A_shared: make_swizzle_layout(A_shared),
            #     B_shared: make_swizzle_layout(B_shared),
            # })

            # Enable rasterization for better L2 cache locality (Optional)
            # T.use_swizzle(panel_size=10, enable=True)

            # Clear local accumulation
            T.clear(C_local)

            for ko in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):
                # Copy tile of A
                # This is a sugar syntax for parallelized copy
                T.copy(A[by * block_M, ko * block_K], A_shared)

                # Demonstrate parallelized copy from global to shared for B
                for k, j in T.Parallel(block_K, block_N):
                    B_shared[k, j] = B[ko * block_K + k, bx * block_N + j]

                # Perform a tile-level GEMM on the shared buffers
                # Currently we dispatch to the cute/hip on Nvidia/AMD GPUs
                T.gemm(A_shared, B_shared, C_local)

            # Copy result back to global memory
            T.copy(C_local, C[by * block_M, bx * block_N])

    return main


# 1. Define the kernel (matmul) with the desired dimensions
func = matmul(1024, 1024, 1024, 128, 128, 32)

# 2. Compile the kernel into a torch function
# out_idx specifies the index of the output buffer in the argument list
# if out_idx is specified, the tensor will be created during runtime
# target currently can be "cuda" or "hip" or "cpu".
jit_kernel = tilelang.compile(func, out_idx=[2], target="cuda")

# 3. Test the kernel in Python with PyTorch data
import torch

# Create random input tensors on the GPU
a = torch.randn(1024, 1024, device="cuda", dtype=torch.float16)
b = torch.randn(1024, 1024, device="cuda", dtype=torch.float16)


# Run the kernel through the JIT-compiled function
c = jit_kernel(a, b)

# Reference multiplication using PyTorch
ref_c = a @ b

# Validate correctness
torch.testing.assert_close(c, ref_c, rtol=1e-2, atol=1e-2)
print("Kernel output matches PyTorch reference.")

# 4. Retrieve and inspect the generated CUDA source (optional)
cuda_source = jit_kernel.get_kernel_source()
print("Generated CUDA kernel:\n", cuda_source)

# 5.Pofile latency with the profiler
profiler = jit_kernel.get_profiler()

latency = profiler.do_bench()

print(f"Latency: {latency} ms")

Dive Deep into TileLang Beyond GEMM

In addition to GEMM, we provide a variety of examples to showcase the versatility and power of TileLang, including:

  • Dequantize GEMM: Achieve high-performance dequantization by fine-grained control over per-thread operations, with many features now adopted as default behaviors in BitBLAS, which utilizing magic layout transformation and intrins to accelerate dequantize gemm.
  • FlashAttention: Enable cross-operator fusion with simple and intuitive syntax, and we also provide an example of auto tuning.
  • LinearAttention: Examples include RetNet and Mamba implementations.
  • Convolution: Implementations of Convolution with IM2Col.

Upcoming Features

Check our tilelang v0.2.0 release plan for upcoming features.


TileLang has now been used in project BitBLAS and AttentionEngine.

Join the Discussion

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

Join our Discord

Acknowledgements

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.

Project details


Download files

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

Source Distributions

No source distribution files available for this release.See tutorial on generating distribution archives.

Built Distributions

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

tilelang-0.1.2-cp312-cp312-manylinux1_x86_64.whl (68.8 MB view details)

Uploaded CPython 3.12

tilelang-0.1.2-cp311-cp311-manylinux1_x86_64.whl (68.8 MB view details)

Uploaded CPython 3.11

tilelang-0.1.2-cp310-cp310-manylinux1_x86_64.whl (68.8 MB view details)

Uploaded CPython 3.10

tilelang-0.1.2-cp39-cp39-manylinux1_x86_64.whl (68.8 MB view details)

Uploaded CPython 3.9

tilelang-0.1.2-cp38-cp38-manylinux1_x86_64.whl (68.8 MB view details)

Uploaded CPython 3.8

File details

Details for the file tilelang-0.1.2-cp312-cp312-manylinux1_x86_64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.2-cp312-cp312-manylinux1_x86_64.whl
Algorithm Hash digest
SHA256 2a12342bc9f9001f95740cd670b2ec11a814cfa6013a3a68f4de7adef7cd5dd4
MD5 191a2883e534cbd47d1657343d8e974b
BLAKE2b-256 d989b872bc0009a6a24e579b2a7129f8f2cd3f08c0c8899efbb7b03d37c87953

See more details on using hashes here.

File details

Details for the file tilelang-0.1.2-cp311-cp311-manylinux1_x86_64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.2-cp311-cp311-manylinux1_x86_64.whl
Algorithm Hash digest
SHA256 210624781962bf2b0122c70538366d4c30d6b894b1806a27fe862c1a1c303990
MD5 e02bb69189f6e42037b62b2c4e94ead4
BLAKE2b-256 982887797cb6b2399c5a103e3fd21478024b2782162b2378c8f14ef6f9d887af

See more details on using hashes here.

File details

Details for the file tilelang-0.1.2-cp310-cp310-manylinux1_x86_64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.2-cp310-cp310-manylinux1_x86_64.whl
Algorithm Hash digest
SHA256 c5a68203e8dff2ed92c35bea3d05b5413d9e7902d0831f697d851858fd693e15
MD5 6e68c7465f228cf4795460840214ee82
BLAKE2b-256 52cd4ab30464ab95e251892120a7707e817073a848b969760b0c3e9ba105263c

See more details on using hashes here.

File details

Details for the file tilelang-0.1.2-cp39-cp39-manylinux1_x86_64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.2-cp39-cp39-manylinux1_x86_64.whl
Algorithm Hash digest
SHA256 5d7917ea352232c82111a0ef2bea90b81fb0b870b1d350441b7adf0a06a2f308
MD5 16698edcbdaf13fbaa920213036ddde2
BLAKE2b-256 53705e6988a851b52e4fb4d42ee8217b7734817c083b03d6a0550953561aab6b

See more details on using hashes here.

File details

Details for the file tilelang-0.1.2-cp38-cp38-manylinux1_x86_64.whl.

File metadata

File hashes

Hashes for tilelang-0.1.2-cp38-cp38-manylinux1_x86_64.whl
Algorithm Hash digest
SHA256 9adf29bf3241a7b97b783f21007819fb95fc8c4d93659d2be6dc7d6eb1bf0a14
MD5 a756207304a13b2c849f4e3840742e98
BLAKE2b-256 728a01992f168a11d4797c59212edcb1933190ceaa72e7444454a445573af5a2

See more details on using hashes here.

Supported by

AWS Cloud computing and Security Sponsor Datadog Monitoring Depot Continuous Integration Fastly CDN Google Download Analytics Pingdom Monitoring Sentry Error logging StatusPage Status page