TileLang v0.1.15 发布：新增华为 Ascend 950 原生后端与 CUDA 自动 warp specialization

Official release notes (https://github.com/tile-ai/tilelang/releases/tag/v0.1.15) ## Highlights - **Native Huawei Ascend 950 support** (#3308): an end-to-end NPU backend with native code generation, automatic Cube/Vector scheduling and synchronization, and mixed SIMD/SIMT programming. - **Automatic CUDA warp specialization** (#3059, #3185): an opt-in role-based scheduler that assigns TMA loads, MMA computation, TMA stores, and worker operations to specialized warp groups. - **Unified block-scaled GEMM** (#3237, #3257, #3284): common `T.gemm_blockscaled` semantics with dedicated backend dispatch, improved SM100 instruction selection, and expanded SM120 fragment support. - **More expressive Python frontend** (#3230): compile-time iteration over Python iterables, `enumerate`, `zip`, comprehensions, and generator expressions. ## Ascend 950 - Add `tilelang.ascend.language` and `target="ascend"` for Huawei Ascend 950 (`dav-3510`). - Combine Cube GEMM and Vector computation in one kernel, with `T.SimdVF` and `T.SimtVF` regions. - Support explicit UB/L1/L0 storage, tiled copies, cross-core transfers, and MXFP8/MXFP4 block-scaled GEMM. - Add automatic scheduling, pipelining, multi-buffering, layout inference, and synchronization insertion. - Integrate Bisheng compilation, `tvm_ffi` and Cython execution, PyTorch NPU tensors and streams, and NPU profiling. - Include GEMM, DeepGEMM-style kernels, FlashAttention forward/backward, RMSNorm, and FP8 quantization examples. See the [Ascend 950 guide](https://github.com/tile-ai/tilelang/blob/main/tilelang/ascend/README.md) for installation and usage. Ascend A2/A3 support remains in the community-maintained TileLang-Ascend projects. ## CUDA - Enable role-based automatic warp specialization with `TL_ENABLE_AUTO_WARP_SPECIALIZATION: "role_based"`; add GEMM and FlashAttention examples (#3059, #3185). - Extend SM120 block-scaled GEMM to fragment-resident A operands, row-major scale fragments, and odd per-warp atom grids; diagnose conflicting scale-fragment layouts (#3257, #3284). - Add round-to-nearest `T.fma` and `T.fmul`, and complete additional FP16/BF16 math bridges (#3134, #3132, #3163). - Preserve packed FP8 vector copies and fuse exact FP4-to-FP8 conversions through FP32 (#3276, #3204). - Fix WGMMA/UMMA K-panel strides for operand layouts and insert async-proxy fences for sparse MMA (#2965, #3139). - Correct atomic vectorization for invariant or non-contiguous destinations; keep shared FP32 atomics scalar on SM90 (#3129, #3219, #3238). - Fix NVRTC warp reductions, kernel-body assertions, boolean bitwise negation, and non-constant int4/uint4 broadcasts; restrict 256-bit global loads/stores to SM100+ (#3260, #3206, #3228, #3114, #3248). - Add an FP8 sparse MLA forward example for DeepSeek V3.2 on Hopper (#3224). ## Language and Compiler - Make kernel launch encoding target-neutral, with launch options and operation hints owned by their backend dialects (#3186, #3203). - Support Python compile-time iteration and comprehensions while preserving device-loop semantics for `for ... in range(...)` (#3230). - Support unary plus on symbolic expressions, fix loop-variable binding, remove warnings for immutable rebinding, and reject unsupported loop `else` clauses (#3141, #3232, #3262, #3142). - Treat `T.copy` / `T.async_copy` `coalesced_width` as a hint and clamp it to the achievable vector width (#3246). - Preserve dynamic reduction tail guards, fix parallel-loop lowering with let inlining disabled, and restore async-copy lowering with partitioned layouts (#3294, #3269, #3278). - Improve reducer layout planning and symbolic layout validation; reject unsupported AllReduce thread strides and reduction NaN-propagation dtypes (#3171, #3233, #3266, #3273). - Fix NaN-propagating clamp lowering and eliminate unused bindings during simplification (#3205, #3293). ## ROCm, CPU, and Metal - ROCm: support the DeepSeek V3.2 Top-K selector, wave64 Hadamard transforms, and GLM-5.3 k-pool examples, including Top-K index transformation (#3147, #3154, #3254). - Expand ROCm validation for attention kernels and portable examples, and document pip installation (#3146, #3148, #3165, #3168). - CPU: compute scalar GEMM products in the accumulator dtype and legalize BF16 arithmetic (#2917, #3201). - Metal: support 32-bit integer atomic add and respect GEMM buffer-region offsets (#3211, #3209). ## Runtime, Build, and Tooling - Store cached kernel parameters as JSON instead of cloudpickle (#3143). - Publish CUDA binaries and metadata atomically in immutable cache directories, preventing readers from observing partially published entries (#3177). - Fix dynamic-output allocation when the sizing input follows the output, and add missing `uint64` argument mappings (#3207, #3229). - Add wall-clock benchmarking and MPS-compatible timing helpers (#3234). - Improve CUDA/ROCm compiler discovery and library loading for symlinked installations (#2839, #3166, #3227). - Enable optimization for default single-config native builds and restore effective Windows wheel-build caching (#3191, #3133, #3305). - Validate built wheels on GPU runners and make performance regressions fail CI (#3167, #3182). ## Compatibility Notes - Remove the unused `sync` and `group` parameters from `T.Pipelined`, and `k_pack` from `T.gemm_sp` (#3202). - ROCm kernels using `T.gemm(k_pack=...)` should import `tilelang.rocm.language` (#3203). - Kernels querying thread extents during tracing must specify `threads=` explicitly in `T.Kernel` (#3186). - `T.symbolic` remains available as a deprecated alias; use `T.dynamic` for new code (#3216). - Remove `TILELANG_CACHE_VERIFY_HASH`; binary artifact hash verification is now mandatory. Legacy cache formats are rebuilt automatically (#3143, #3177). **Full Changelog**: https://github.com/tile-ai/tilelang/compare/v0.1.14...v0.1.15 ## What's Changed * [CI] Cap torch<2.14 for CUDA tests until flash-attn supports torch 2.14 by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3136 * Bump transformers from 5.5.0 to 5.10.1 in /examples/bitnet-1.58b by @dependabot[bot] in https://github.com/tile-ai/tilelang/pull/3131 * [CI] Fix slow Windows wheel builds by making ccache effective by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3133 * [Language][CUDA] Add T.fma and T.fmul round-to-nearest intrinsics by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3134 * [CUDA] Complete 16-bit bridges for CUDA-lowered unary math by @Chennesxu in https://github.com/tile-ai/tilelang/pull/3132 * [Refactor] Restore LoopUnswitching test altered by #3121 by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/3124 * [BugFix][CUDA] Support non-constant int4/uint4 broadcast by @jjppp in https://github.com/tile-ai/tilelang/pull/3114 * [BugFix][CPU] Compute scalar GEMM products in accum dtype instead of input dtype by @Dino1844 in https://github.com/tile-ai/tilelang/pull/2917 * [BugFix][Hopper][Blackwell] Read the WGMMA/UMMA K-panel stride from the layout instead of the operand extent by @bigSheep123 in https://github.com/tile-ai/tilelang/pull/2965 * [BugFix][Carver] Gate _legalize_info fallback on sm_version by @mocusez in https://github.com/tile-ai/tilelang/pull/3145 * [ROCm] Validate block-causal attention kernels by @andyluo7 in https://github.com/tile-ai/tilelang/pull/3146 * [ROCm] Validate attention sink kernels by @andyluo7 in https://github.com/tile-ai/tilelang/pull/3148 * [ROCm] Support DeepSeek-V3.2 Top-K selector by @andyluo7 in https://github.com/tile-ai/tilelang/pull/3147 * [Bugfix][Language] Reject loop else clauses in the eager frontend by @rishabhsinha17 in https://github.com/tile-ai/tilelang/pull/3142 * [Example] Fix hadamard reference precision on TF32 and add pytest entry for CI by @Guan-jeans in https://github.com/tile-ai/tilelang/pull/3126 * [BugFix][Carver] Compare sm_version numerically in plan_rasterization by @mocusez in https://github.com/tile-ai/tilelang/pull/3144 * [Misc] Expose .agents skills to Claude Code via .claude/skills symlinks by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3161 * [Language] Cache dtype.as_torch and demote storage-dtype fallback logs to debug by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3160 * [Refactor][CUDA] Replace per-lane-count vector constructors with variadic packers by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3158 * [CUDA] Add 16-bit hpow and hfmod bridges by @Chennesxu in https://github.com/tile-ai/tilelang/pull/3163 * [BugFix][ROCm] Resolve hipcc robustly instead of trusting bare PATH lookup by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3166 * [CI] Validate built wheels on GPU runners in the Dist workflow by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3167 * [Doc] Document pip installation on AMD GPUs (ROCm) by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3168 * [Bugfix][Language] Support unary plus on PrimExpr in the eager frontend by @rishabhsinha17 in https://github.com/tile-ai/tilelang/pull/3141 * [CUDA] Add async-proxy fences for sparse MMA intrinsics by @ZenAlexa in https://github.com/tile-ai/tilelang/pull/3139 * [Feature] [CUDA] Role-based automatic warp specialization by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/3059 * [CUDA] Resolve symlinked nvcc before deriving CUDA_HOME by @morluto in https://github.com/tile-ai/tilelang/pull/2839 * [Language] Remove deprecated T.symbolic alias by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/3122 * [ROCm] Support wave64 Hadamard transforms by @andyluo7 in https://github.com/tile-ai/tilelang/pull/3154 * [Cache] Store cached kernel params as JSON instead of cloudpickle by @rishabhsinha17 in https://github.com/tile-ai/tilelang/pull/3143 * [BugFix][Transform] Keep shared FP32 atomics scalar on SM90 by @ZenAlexa in https://github.com/tile-ai/tilelang/pull/3129 * [Layout] Consider scalar reducer plans in register-count search by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3171 * [Transform] Add snapshot mode to FreshenMutableReads by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3174 * [CUDA][Cache] Publish immutable binary cache directories by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3177 * [CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in https://github.com/tile-ai/tilelang/pull/3181 * [CI] Fail the perf bot on a regression instead of only reporting it by @cklxx in https://github.com/tile-ai/tilelang/pull/3182 * [Refactor] [CUDA] Rename `AutoSchedule` to `AutoWarpSpecialization` by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/3185 * [Refactor][Language] Make T.Kernel target-neutral with dialect-owned launch annotations by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3186 * [Cleanup][Language] Remove dead T.Pipelined(sync/group) and T.gemm_sp(k_pack) parameters by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3202 * [Refactor][Language] Move backend-specific op hints into their owning dialects by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3203 * [BugFix][CPU] Legalize BF16 arithmetic in the CPU pipeline by @anerli in https://github.com/tile-ai/tilelang/pull/3201 * [Language] Restore deprecated T.symbolic alias by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3216 * [Refactor][JIT] Replace callee-allocated-output target sniffing with a backend capability flag by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3218 * [Transform] Restore bind-before-use order when merging constraint sets by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3221 * [BugFix][Metal] Respect GEMM buffer region offsets by @anerli in https://github.com/tile-ai/tilelang/pull/3209 * [Runtime] Fix library loading for symlink installs by @sepcnt in https://github.com/tile-ai/tilelang/pull/3227 * [CUDA] Fix boolean bitwise negation codegen by @sepcnt in https://github.com/tile-ai/tilelang/pull/3228 * [Example] FP8 sparse MLA forward for DeepSeek V3.2 on Hopper by @xuebozhang525-alt in https://github.com/tile-ai/tilelang/pull/3224 * [Profiler] Split timing helpers and add wall-clock benchmarking by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/3234 * [CPU] Import the CPU dialect in CPU tests by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/3236 * [BugFix][Vectorize] Keep atomic_add scalar for invariant/non-contiguous destinations by @Dino1844 in https://github.com/tile-ai/tilelang/pull/3219 * [Transform][CUDA] Plan atomic vector widths from destination addresses by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3238 * [Do not review][Op][Language][CUDA] Add common block-scaled GEMM semantics and backend dispatch by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3237 * [JIT] Add missing uint64 argument type mappings by @sepcnt in https://github.com/tile-ai/tilelang/pull/3229 * [BugFix] Restore symbolic loop-layout injectivity proof; reject layouts on symbolic shared tiles by @sepcnt in https://github.com/tile-ai/tilelang/pull/3233 * [ROCm] Run portable example validation in CI by @andyluo7 in https://github.com/tile-ai/tilelang/pull/3165 * [BugFix][CUDA] Only emit 256-bit global load/store on SM100+ targets by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/3248 * [Metal] Support 32-bit integer atomic add by @anerli in https://github.com/tile-ai/tilelang/pull/3211 * [Testing] Drop duplicated codegen-smoke tests; fix fastmath self-comparison assertions by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/3252 * [ROCm] Add GLM-5.3 k-pool Top-K transform by @andyluo7 in https://github.com/tile-ai/tilelang/pull/3254 * [Frontend] Support Python iterables and comprehensions by @sepcnt in https://github.com/tile-ai/tilelang/pull/3230 * [Frontend] Remove warnings for immutable variable rebinding by @LJC00118 in https://github.com/tile-ai/tilelang/pull/3262 * [BugFix] Reject non-power-of-two AllReduce thread strides by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3266 * [NVRTC] Fix warp reduction compilation by @sepcnt in https://github.com/tile-ai/tilelang/pull/3260 * [BugFix][CUDA] Lower a kernel-body assert to a device-legal check by @xy200303 in https://github.com/tile-ai/tilelang/pull/3206 * [BugFix][JIT] Allocate a dynamic-shape output that precedes its sizing input by @xy200303 in https://github.com/tile-ai/tilelang/pull/3207 * [CUDA] Fuse exact FP4 to FP8 conversion through FP32 by @ZenAlexa in https://github.com/tile-ai/tilelang/pull/3204 * [Fix]Clamp T.copy/T.async_copy coalesced_width to achievable vector size instead of LOG(FATAL) by @edragain2nd in https://github.com/tile-ai/tilelang/pull/3246 * [BugFix] Bind loop targets independently of mutable scalar variables by @sepcnt in https://github.com/tile-ai/tilelang/pull/3232 * [BugFix][Language] Lower NaN-propagating clamp through device templates by @xy200303 in https://github.com/tile-ai/tilelang/pull/3205 * [Fix] Reject unsupported reduction NaN propagation dtypes by @ZenAlexa in https://github.com/tile-ai/tilelang/pull/3273 * [BugFix] Fix parallel loop lowering with let inlining disabled by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/3269 * [Transform][CUDA] Fix async copy lowering with partitioned layouts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3278 * [CUDA] Keep FP8 vector copies packed by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3276 * [CUDA] Support SM120 block-scaled GEMM fragments and odd warp atom grids by @sepcnt in https://github.com/tile-ai/tilelang/pull/3257 * [CUDA] Reject conflicting SM120 scale fragment layouts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3284 * [Build] Optimize default single-config builds by @KellyFrog in https://github.com/tile-ai/tilelang/pull/3191 * [Fix][TVM] Update TVM to preserve dynamic reduction tail guards by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3294 * [BugFix] Eliminate unused bindings in Simplify by @LJC00118 in https://github.com/tile-ai/tilelang/pull/3293 * [Build][Windows] Restore ccache hits for clang-cl wheel builds by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3305 * [Public Release 9/30] Introduce Ascend 950 backend by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/3308 * [Release] Bump version into 0.1.15 by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/3309 ## New Contributors * @Dino1844 made their first contribution in https://github.com/tile-ai/tilelang/pull/2917 * @mocusez made their first contribution in https://github.com/tile-ai/tilelang/pull/3145 * @rishabhsinha17 made their first contribution in https://github.com/tile-ai/tilelang/pull/3142 * @Guan-jeans made their first contribution in https://github.com/tile-ai/tilelang/pull/3126 * @anerli made their first contribution in https://github.com/tile-ai/tilelang/pull/3201 * @xuebozhang525-alt made their first contribution in https://github.com/tile-ai/tilelang/pull/3224 * @xy200303 made their first contribution in https://github.com/tile-ai/tilelang/pull/3206 **Full Changelog**: https://github.com/tile-ai/tilelang/compare/v0.1.14...v0.1.15 Repository README (current repository documentation, not release changes) # Tile Language [](https://badge.fury.io/py/tilelang) [](https://deepwiki.com/tile-ai/tilelang) [](https://discord.gg/TUrHyJnKPG) [](https://github.com/tile-ai/tilelang-puzzles) [](https://github.com/tile-ai/tilelang-lsp) [Documentation](https://tilelang.com/) · [Installation](https://tilelang.com/get_started/Installation.html) · [Examples](https://github.com/tile-ai/tilelang/tree/main/examples) · [Releases](https://github.com/tile-ai/tilelang/releases) · [Contributing](https://github.com/tile-ai/tilelang/blob/main/CONTRIBUTING.md) · [TileLang (Ascend 950)](./tilelang/ascend/README.md) 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](https://tvm.apache.org/), 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/ascend/README.md):** 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](https://github.com/tile-ai/tilelang-lsp):** 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](https://github.com/tile-ai/tilelang/releases/tag/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](https://github.com/tile-ai/tilelang/pull/2364):** added an optimized Blackwell path for `T.mma_gemm_blockscaled` and a corresponding SM120 example. - **2026-07-28 — [Metal 4 cooperative-tensor GEMM](https://github.com/tile-ai/tilelang/pull/2252):** 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](https://github.com/tile-ai/tilelang/pull/2751):** carried Python source locations into TIRX and surfaced them in compiler errors. - **2026-07-24 — [Multi-backend language dialect](https://github.com/tile-ai/tilelang/pull/2734):** reorganized the language layer around shared semantics with static CUDA, ROCm, and Metal dialects. - **2026-07-23 — [Block-causal attention for dLLM](https://github.com/tile-ai/tilelang/pull/2499):** added fixed-length and variable-length block-causal attention examples for diffusion language models. - **2026-07-22 — [IR Lower Trace](https://github.com/tile-ai/tilelang/pull/2725):** 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](https://github.com/tile-ai/tilelang/pull/2592):** selected the launch width adaptively from the head-block size. - **2026-07-21 — [DeepSeek V3.2 top-k optimization](https://github.com/tile-ai/tilelang/pull/2659):** 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](https://github.com/tile-ai/tilelang/pull/2622):** added profiling for compiler passes with a configurable reporting threshold. - **2026-07-12 — [IKET profiler integration](https://github.com/tile-ai/tilelang/pull/2515):** added CUDA timeline instrumentation and profiling support. - **2026-07-08 — [TileLang v0.1.12](https://github.com/tile-ai/tilelang/releases/tag/v0.1.12):** added the LLVM backend, tile scheduler, backend registry, pass visualizer, and expanded Blackwell support. - **2026-07-06 — [Pass Visualizer](https://github.com/tile-ai/tilelang/pull/2449):** introduced a structure-tree browser for inspecting compiler transformations. - **2026-06-26 — [Cross-host CUDA binary cache](https://github.com/tile-ai/tilelang/pull/2459):** enabled compiled CUDA binaries to be reused across compatible hosts. - **2026-06-24 — [Tile scheduler](https://github.com/tile-ai/tilelang/pull/2441):** introduced persistent tile-scheduling primitives for kernel authors. - **2026-06-24 — [Backend CodeGen registry](https://github.com/tile-ai/tilelang/pull/2442):** moved device and [host CodeGen](https://github.com/tile-ai/tilelang/pull/2446) dispatch behind a backend registry. - **2026-06-18 — [LLVM backend](https://github.com/tile-ai/tilelang/pull/2409):** added CPU lowering and execution through LLVM. - **2026-06-18 — [Arbitrary-layout TMA lowering](https://github.com/tile-ai/tilelang/pull/2380):** enabled TMA transfers for swizzled shared-memory layouts. - **2026-06-16 — [Pass Diff](https://github.com/tile-ai/tilelang/pull/2375):** added compiler-pass IR comparison for debugging lowering changes. - **2026-06-08 — [TileLang v0.1.11](https://github.com/tile-ai/tilelang/releases/tag/v0.1.11):** expanded scan, pipeline, backend, CUDA, ROCm, and Metal functionality. - **2026-05-25 — [Scan operators](https://github.com/tile-ai/tilelang/pull/2262):** introduced tile-level scan primitives. - **2026-05-25 — [TileLang v0.1.10](https://github.com/tile-ai/tilelang/releases/tag/v0.1.10):** broadened AMD and Blackwell support, added initial Metal GEMM, improved Windows packaging, and expanded autotuning. - **2026-05-24 — [CDNA4 MXFP4](https://github.com/tile-ai/tilelang/pull/2132):** added FP4 E2M1 matrix-core support for AMD gfx950. - **2026-05-22 — [Metal simdgroup GEMM](https://github.com/tile-ai/tilelang/pull/1869):** added the first Metal `T.gemm` path using `simdgroup_matrix` MMA. - **2026-05-20 — [Cluster copies](https://github.com/tile-ai/tilelang/pull/1908):** introduced `T.copy_cluster` for TMA multicast and SM-to-SM cluster transfers. - **2026-05-20 — [TMA gather/scatter](https://github.com/tile-ai/tilelang/pull/2129):** added `tile::gather4` and `tile::scatter4` support. - **2026-05-20 — [Native SM75 MMA GEMM](https://github.com/tile-ai/tilelang/pull/2198):** added FP16, INT8, and INT4 tensor-core paths for Turing GPUs. - **2026-05-20 — [TIRX migration](https://github.com/tile-ai/tilelang/pull/2216):** moved TileLang IR usage to TVM's TIRX representation. - **2026-05-11 — [Parallel autotuning](https://github.com/tile-ai/tilelang/pull/2159):** added pipelined compilation, grouped compilation, and multi-GPU benchmarking. - **2026-05-07 — [DeepSeek V4 operators](https://github.com/tile-ai/tilelang/pull/2148):** added TileLang examples for DeepSeek V4 workloads. - **2026-05-06 — [Windows support](https://github.com/tile-ai/tilelang/pull/2093):** added complete Windows build and runtime support with cross-platform fixes. - **2026-04-28 — [MXFP8 grouped GEMM](https://github.com/tile-ai/tilelang/pull/2098):** added block-scaled grouped GEMM examples with transposed-B support on Blackwell. - **2026-04-25 — [HISA sparse-attention indexer](https://github.com/tile-ai/tilelang/pull/2069):** added hierarchical sparse-attention indexing examples. - **2026-04-24 — [Blackwell MXFP8 block-scaled GEMM](https://github.com/tile-ai/tilelang/pull/1945):** added MXFP8 block-scaled matrix multiplication on SM100. - **2026-04-22 — [TileLang v0.1.9](https://github.com/tile-ai/tilelang/releases/tag/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](https://github.com/tile-ai/tilelang/pull/2044):** added WMMA lowering for AMD gfx11 GPUs. - **2026-04-20 — [INT4 `T.gemm`](https://github.com/tile-ai/tilelang/pull/2063):** added INT4 matrix multiplication to the CUDA GEMM path. - **2026-04-17 — [CUDA source kernels](https://github.com/tile-ai/tilelang/pull/1970):** introduced `T.CUDASourceCodeKernel` for embedding custom CUDA source. - **2026-04-15 — [AutoDD frozen regions](https://github.com/tile-ai/tilelang/pull/2045):** added `__freeze__` annotations to preserve selected code during automatic delta debugging. - **2026-03-27 — [TMA stores](https://github.com/tile-ai/tilelang/pull/1981):** added store support to `T.tma_copy`. - **2026-03-24 — [Two-SM Blackwell kernels](https://github.com/tile-ai/tilelang/pull/1882):** added two-SM TMA, TMEM, and TCGEN5 MMA support. - **2026-03-23 — [AMD RDNA4](https://github.com/tile-ai/tilelang/pull/1951):** upgraded the ROCm path and added RDNA4 GPU support. - **2026-03-22 — [FlashAttention on SM100](https://github.com/tile-ai/tilelang/pull/1910):** added Blackwell FlashAttention examples. - **2026-03-18 — [Producer-consumer warp specialization](https://github.com/tile-ai/tilelang/pull/1909):** added automatic warp-specialized pipelines and the `T.tma_copy` API. - **2026-03-12 — [Eager-mode autotuning](https://github.com/tile-ai/tilelang/pull/1906):** enabled the autotuner with eager JIT kernels. - **2026-03-10 — [CPU `T.gemm`](https://github.com/tile-ai/tilelang/pull/1904):** added matrix multiplication support for the CPU target. - **2026-03-05 — [IR dump configuration](https://github.com/tile-ai/tilelang/pull/1903):** added a TileLang pass configuration for dumping intermediate IR. - **2026-02-28 — [CUDA cluster primitives](https://github.com/tile-ai/tilelang/pull/1874):** added cluster launch, query, synchronization, and barrier operations. - **2026-02-28 — [TCGEN5 MMA tensor-shared path](https://github.com/tile-ai/tilelang/pull/1866):** added the tensor-memory/shared-memory Blackwell GEMM path. - **2026-02-24 — [Host-toolchain-free builds](https://github.com/tile-ai/tilelang/pull/1833):** enabled installation without a host C/C++ toolchain when supported artifacts are available. - **2026-02-23 — [CuTe DSL GEMM V2](https://github.com/tile-ai/tilelang/pull/1855):** added SM90 and SM100 GEMM V2 support to the CuTe DSL backend. - **2026-02-16 — [TileLang v0.1.8](https://github.com/tile-ai/tilelang/releases/tag/v0.1.8):** shipped dynamic pipeline improvements, logging documentation, richer layout representations, and AMD fixes. - **2026-02-14 — [Cross-CUDA release wheels](https://github.com/tile-ai/tilelang/pull/1826):** unified multiple CUDA versions behind a single wheel. - **2026-02-14 — [Hierarchical reductions](https://github.com/tile-ai/tilelang/pull/1762):** added hierarchical and warp-level reduction intrinsics. - **2026-02-09 — [CUDA runtime stubs](https://github.com/tile-ai/tilelang/pull/1821):** added lazy-loading CUDART and NVRTC stubs for CUDA 11, 12, and 13 compatible wheels. - **2026-02-08 — [Layout visualization improvements](https://github.com/tile-ai/tilelang/pull/1811):** improved rendering and inspection of TileLang layouts. - **2026-02-02 — [TileLang Puzzles](https://github.com/tile-ai/tilelang-puzzles):** published ten progressively harder exercises for learning TileLang interactively. ### 2025 - **2025-12-18 — [CuTe DSL backend](https://github.com/tile-ai/tilelang/pull/1421):** added compilation through NVIDIA CUTLASS CuTe DSL; follow ongoing work in [issue #1454](https://github.com/tile-ai/tilelang/issues/1454). - **2025-12-17 — [Z3 integration](https://github.com/tile-ai/tilelang/pull/1367):** integrated SMT-based symbolic reasoning into the TVM arithmetic analyzer. - **2025-10-31 — [Apache TVM FFI migration](https://github.com/tile-ai/tilelang/pull/1108):** moved the runtime interface to `apache-tvm-ffi` to reduce host-side overhead. - **2025-10-30 — [TileLang v0.1.6.post2](https://github.com/tile-ai/tilelang/releases/tag/v0.1.6.post2):** published the final TileLang release compatible with Python 3.8. - **2025-10-07 — [Apple Metal backend](https://github.com/tile-ai/tilelang/pull/799):** introduced Metal device support for Apple silicon. - **2025-09-29 — [Huawei Ascend adapters](https://github.com/tile-ai/tilelang-ascend):** published AscendC and Ascend NPU IR backend work in the external TileLang Ascend project. - **2025-07-04 — [2:4 sparse tensor cores](https://github.com/tile-ai/tilelang/pull/526):** introduced `T.gemm_sp` for structured-sparse matrix multiplication. - **2025-06-05 — [NVRTC execution backend](https://github.com/tile-ai/tilelang/pull/461):** added an NVRTC path to reduce compilation time for generated CUDA templates. - **2025-04-14 — [FlashMLA on AMD MI300X](https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_mla/amd):** published the optimized AMD implementation and accompanying documentation. - **2025-03-03 — [MLA decoding on H100](https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_mla):** published the compact TileLang implementation, benchmarks, and optimization walkthrough. - **2025-02-15 — [WebGPU code generation](https://github.com/tile-ai/tilelang/pull/86):** added the initial WebGPU backend. - **2025-02-12 — [TileLang v0.1.0](https://github.com/tile-ai/tilelang/releases/tag/v0.1.0):** published the first v0.1 public release. - **2025-02-10 — [Debugging and layout tools](https://tilelang.com/tutorials/debug_tools_for_tilelang.html):** added `T.print` and fragment-layout visualization workflows. - **2025-01-20 — [TileLang open sourced](https://github.com/tile-ai/tilelang):** made the project publicly available. See [all releases](https://github.com/tile-ai/tilelang/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](tilelang/backend/README.md) for the design, or ask a coding agent to use the [backend integration skill](.agents/skills/tilelang-backend/SKILL.md) 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](https://tilelang.com/get_started/targets.html) 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. | | Huawei Ascend 950 | `ascend` | Linux; Ascend 950 NPUs | Supported | Build from source with `USE_ASCEND=ON`; requires CANN and `torch_npu`. See the [Ascend guide](./tilelang/ascend/README.md). | | 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 A2/A3 | `ascendc` / `pto` / `npuir` | Ascend A2 and A3 | Ecosystem | Developed in [tilelang-ascend](https://github.com/tile-ai/tilelang-ascend) and the MLIR-based [tilelang-mlir-ascend](https://github.com/tile-ai/tilelang-mlir-ascend). | | MetaX MACA | `maca` | MetaX C500 and C600 | Ecosystem | Developed in [tilelang-metax](https://github.com/tile-ai/tilelang-metax); requires the MACA software stack. | | Moore Threads MUSA | `musa` | S5000, S4000, and M1000 | Ecosystem | Developed in [tilelang-musa](https://github.com/tile-ai/tilelang-musa) and released independently. | | HYGON | `hcu` | Linux; BW1000, BW1100, BW150 and K100_AI | Ecosystem | Developed in [tilelang-hygon](https://github.com/tile-ai/tilelang-hygon); requires the DTK software stack. | | Sunrise-AI TANG | `tang` | Sunrise S2 and S3 | Ecosystem | Developed in [tilelang-sunrise](https://github.com/tile-ai/tilelang-sunrise). The TANG software stack is required. | ## Installation Install the latest stable release from PyPI: ```bash pip install tilelang ``` Verify the installation: ```bash python -c "import tilelang; print(tilelang.__version__)" ``` Nightly wheels provide recent features and fixes before the next stable release: ```bash 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](https://tilelang.com/get_started/Installation.html#installing-on-amd-gpus-rocm) 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](https://tilelang.com/get_started/Installation.html). ## 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. ```python 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](https://tilelang.com/programming_guides/language_basics.html), then explore the [GEMM examples](https://github.com/tile-ai/tilelang/tree/main/examples/gemm) for layouts, autotuning, and architecture-specific optimizations. ## Examples - **Start here:** [quickstart](https://github.com/tile-ai/tilelang/blob/main/examples/quickstart.py) and [elementwise kernels](https://github.com/tile-ai/tilelang/tree/main/examples/elementwise) - **GEMM and quantization:** [GEMM](https://github.com/tile-ai/tilelang/tree/main/examples/gemm), [grouped GEMM](https://github.com/tile-ai/tilelang/tree/main/examples/grouped_gemm), [FP8 GEMM](https://github.com/tile-ai/tilelang/tree/main/examples/gemm_fp8), [dequantization GEMM](https://github.com/tile-ai/tilelang/tree/main/examples/dequantize_gemm), and [block-scaled GEMM](https://github.com/tile-ai/tilelang/tree/main/examples/blockscaled_gemm_sm100) - **Attention and sequence models:** [FlashAttention](https://github.com/tile-ai/tilelang/tree/main/examples/flash_attention), [Flash Decoding](https://github.com/tile-ai/tilelang/tree/main/examples/flash_decoding), [block-sparse attention](https://github.com/tile-ai/tilelang/tree/main/examples/blocksparse_attention), [linear attention](https://github.com/tile-ai/tilelang/tree/main/examples/linear_attention), and [GDN](https://github.com/tile-ai/tilelang/tree/main/examples/gdn) - **Model workloads:** [DeepSeek MLA](https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_mla), [DeepSeek V3.2](https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_v32), [DeepSeek V4](https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_v4), and [DeepSeek mHC](https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_mhc) - **Architecture-specific kernels:** [AMD](https://github.com/tile-ai/tilelang/tree/main/examples/amd), [TCGEN05](https://github.com/tile-ai/tilelang/tree/main/examples/gemm_tcgen05), and [SM120](https://github.com/tile-ai/tilelang/tree/main/examples/gemm_sm120) - **Compiler and debugging tools:** [analyzer](https://github.com/tile-ai/tilelang/tree/main/examples/analyze), [layout visualization](https://github.com/tile-ai/tilelang/tree/main/examples/plot_layout), [AutoDD](https://github.com/tile-ai/tilelang/tree/main/examples/autodd), and [IKET](https://github.com/tile-ai/tilelang/tree/main/examples/iket) Browse the [complete examples directory](https://github.com/tile-ai/tilelang/tree/main/examples) 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](https://github.com/tile-ai/tilelang-benchmark). Below are selected results showcasing its capabilities: - MLA Decoding Performance on H100 - Flash Attention Performance on H100