tile-ai/tilelang

Tile Language Documentation · Installation · Examples · Releases · Contributing · TileLang (Ascend 950) 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. 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 2025-12-18 — CuTe DSL backend : added compilation through NVIDIA CUTLASS CuTe DSL; follow ongoing work in issue #1454 . 2025-12-17 — Z3 integration : integrated SMT-based symbolic reasoning into the TVM arithmetic analyzer. 2025-10-31 — Apache TVM FFI migration : moved the runtime interface to apache-tvm-ffi to reduce host-side overhead. 2025-10-30 — TileLang v0.1.6.post2 : published the final TileLang release compatible with Python 3.8. 2025-10-07 — Apple Metal backend : introduced Metal device support for Apple silicon.