tile-ai/tilelang

Domain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels

What it solves

TileLang is a domain-specific language (DSL) designed to simplify the creation of high-performance kernels for GPUs, CPUs, and NPUs. It removes the need for developers to write complex low-level code for operations like matrix multiplication (GEMM) and FlashAttention, allowing them to focus on productivity without losing the ability to perform state-of-the-art hardware optimizations.

How it works

Built on top of Apache TVM, TileLang uses a Pythonic syntax to define kernel operations. It employs a compiler infrastructure that can target multiple backends, including NVIDIA CUDA, AMD ROCm/HIP, Huawei Ascend, and Apple Metal. The language provides high-level primitives for memory allocation (shared and fragment), pipelining, and hardware-specific operations like tensor-core matrix multiplication (T.gemm).

Who it’s for

It is intended for kernel developers and AI engineers who need to implement highly optimized mathematical operations for deep learning workloads across various hardware architectures.

Highlights

  • Multi-Backend Support: Native code generation for NVIDIA (SM70-SM120), AMD (CDNA/RDNA), Huawei Ascend 950, and Apple Metal (including M5 cooperative tensors).
  • Pythonic Syntax: Allows for concise kernel definitions using a @tilelang.jit decorator.
  • Advanced Hardware Features: Support for TMA (Tensor Memory Accelerator), WGMMA, and block-scaled GEMM on Blackwell GPUs.
  • C-style Integration: Includes a Language Server Protocol (LSP) for better IDE support with inlay hints and diagnostics.
  • Extensive Tooling: Features a pass visualizer, IR dump configuration, and an autotuner for performance optimization.

Related

  • Project
    tile-ai/tilelang-ascend

    A domain-specific language and compiler based on TVM for developing high-performance AI compute kernels optimized for Huawei Ascend NPUs.

  • Project
    NVIDIA/tilus

    Tilus is a tile-level GPU kernel programming language that enables high-performance, low-precision computation with explicit control over shared memory and registers.

  • Project
    tile-ai/TileRT

    A tile-based runtime engine for ultra-low-latency LLM inference that optimizes hardware utilization to achieve millisecond-level time per output token.

  • Project
    NVIDIA/TileGym

    A CUDA Tile kernel library providing tutorials and examples for tile-based GPU programming to optimize deep learning operators and LLM inference.

  • Project
    pytorch/helion

    Helion is a Python‑embedded DSL that lets you write GPU kernels like regular PyTorch code. It compiles the inner loops to Triton, automatically chooses tiling, grid sizes, and memory‑access strategies, and runs an extensive autotuner to find the fastest configuration. You can freeze the best config for production, and experimental back‑ends (TileIR, CuTe) broaden hardware support.