v0.1.152026-09-30
TileLang 这版的重头戏是华为昇腾 950 原生支持（#3308）：新增 tilelang.ascend.language 与 target="ascend"，Cube GEMM 和向量计算可在同一 kernel 内混排，自动完成流水线、多缓冲与同步插入，并接入毕昇编译链、PyTorch NPU 张量和性能剖析，FlashAttention、RMSNorm、FP8
v0.1.142026-09-02
Highlights
- Reducer v2 (#2940, #3093, #3043, #3044, #3079, #3100): T.alloc_reducer reworked into first-class deferred reduction epochs whose physical lowering is planned by layout inference (via a first-class PartialFragment layout). Adds loop-scoped epochs, conditional reducer finalization, and automatic vectorization of contiguous reducer updates.
- Warp specialization schedules (#2892): new scheduling and materialization mechanism for warp-specialized kernels.
- Layout inference cost models (#2960, #3055, #3061): new IO-aware cost model for free-mode layout selection; register-count restored as the default, with an environment override to switch models.
- Unified backend resolution policy (#2318) plus backend split-up (#2855, #2870, #2850): backend selection is now resolved through a single policy, and builtin ops / Python op proxies are split per backend (CUDA/ROCm/Metal).
- Compilation speed: up to ~4x faster cold parallel/AOT compilation (#2809); Z3 solvers materialized lazily (#3105) and analyzer contexts isolated per kernel compilation (#2890).
- TMA rework: TMA copy lowering unified on CuTe algebra (#3106); TMA layouts made region-aware to keep slices contiguous (#3089).
Language
- Recycle T.unroll(explicit=True) for early explicit unrolling (#2859)
- Make the region bridge a builtin intrinsic (#2983)
- Expose clustermask on T.tmacopy (#2932)
- Unify contiguous stride construction under a single implementation (#3016); honor declared strides in pointer helpers (#3073)
- Stricter validation: reject symbolic T.gemm tile dimensions with a clear message (#3113), validate T.gemm kpack arguments (#3094), reject non-positive arrivecount in allocbarrier/alloccluster_barrier (#3112), reject T.Parallel indexing of local buffers (#3041), reject break in fully expanded loops (#3078)
CUDA
- tcgen05: pack logical TMEM buffers into shared tcgen05.alloc arenas (#2831); support half-subpartition (M=64) TMEM tiles in tcgen05.ld/st (#2880); fix ld/st segment pointer advancement in b32 columns (#2952)
- Select the widest legal WGMMA N instead of gcd (#2931)
- FP32x2 accumulation for reductions: per-reduce control (#3057) and a global PassConfig (#3128)
- Pre-SM80 fallback for bf16 atomic add (#2938); int4x2/uint4x2 codegen (#3036); 16-bit CUTLASS type overloads for fast-math, __ldg, and htan intrinsics (#3097, #3077, #3028, #2894)
- Fixes: warp shuffle for half/bfloat16/FP8 (#3056), FP8 min/max codegen (#3047), UB in packed 8-bit vector stores (#3092), logical not for vectorized bool (#3117, also HIP), vectorized Select codegen (#2843), ldmatrix source offsets wrapped within shared-memory regions (#3110), NVRTC kernel handles isolated per adapter (#2950), flat CUDA include discovery for NVRTC (#2829), masked warpsync in in-warp allreduce (#2865)
- Remove T.{reads,writes} for T.tma_{gather4,scatter4} (#3053); separate TMA atomic-add dtype support from layout encoding (#2846)
ROCm and other backends
- Remove the Composable Kernel dependency (#3111); ROCm CI re-enabled on a gfx942 runner (#2874, #2910)
- Fixes: preserve FP8 bits in warp shuffles (#3104), lower vector Select conditions lane-wise (#2889), emit a compiler barrier for tl.sync_warp on HIP (#2872), reject sub-wavefront block sizes instead of crashing (#2918), resolve versioned device properties in the HIP stub (#2919); emit #line directives for the HIP target (#3058)
- CPU backend: support atomic ops (#2941) and reduce ops (#2893)
- Metal: preserve pointer address spaces for byte offsets (#2925); resolve auto backend to torch and skip disk cache for torch (#2856)
- CuTeDSL: port backend intrinsics to CUTLASS DSL primitives (#2871)
Compiler / Transform
- Refactor the loop vectorization plan with ConstraintKind (#2935); always vectorize T.Parallel loops (#3121); scalarize Select in automatic vectorization (#3060)
- Add VerifyBufferInit, a general buffer-initialization check (#2956)
- Debug info: preserve source spans across lowering passes (#2966); emit #line directives from TIR spans (#3048)
- Fixes: don't drop syncs from the other if branch (#3085), fix wait parity for explicit mbarriers in pipelined loops (#3087), avoid int32 overflow in vector analysis (#3066), fix non-divisible nested modulo simplification (#3065), fix ties-away-from-zero round compile on bfloat16/float8 (#2873), fix absmax/abssum for uint dtypes (#2845), carry memoryorder through vectorized atomicadd (#2924), fix unsigned zero-point decode underflow (#3118), keep cp.async operands in their address spaces (#2869), handle grid barriers and unbounded pointer ranges (#3050), deduplicate replicated reducer updates (#2881), bind symbolic coordinate ranges in FragmentThreadIndexProbe (#3096), reject non-round-tripping inferred layout inverses (#3090)
- Cleanup: remove obsolete compiler and runtime paths (#3086); remove the obsolete disable-fast-math pass config (#3098)
Runtime / JIT / Build
- Kernel cache: detect and repair corrupted cache entries (#3074), export libraries after disk-cache hits (#3116), remove the separate cache temporary directory (#3069)
- Allocate kernel outputs through the packed API (#2937); fix getparentlocals frame self-reference leak (#2934)
- Support Cython 3.3 with the Python 3.9 limited API (#3068); fix CMake reconfigure aborting in FindPipCUDAToolkit before project() (#3102)
Tooling / Ecosystem
- Official compile-only CLI (#3045)
- Unified pass instrumentation per compilation (#2923); Pass Visualizer driven by PassInstrument (#2866); show kernel name in the pass timing report (#2905)
- Data race check disabled by default, opt-in via env var (#2851)
- Open-source TileLang LSP announced (#2862); new agent skills: simplification (#3080), semantic validation (#3054), PR submission (#3082), backend architecture (#2900)
- Examples: generalize TCGEN05 GEMMs for Thor (sm110a) and add Stream-K scheduling (#2902); adopt multi-staged buffers in examples (#2836)
- Docs: add Sunrise-AI TANG (#3107), HYGON (#3101), and MetaX MACA (#3095) to supported platforms; link the multi-backend architecture design (#3123)
-
What's Changed
* [CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2831
* [Enhancement] Speed up cold parallel/AOT compilation up to ~4x by @cklxx in https://github.com/tile-ai/tilelang/pull/2809
* [Docs] Refresh README news and onboarding by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2849
* [Enhancement] Disable data race check by default, opt-in via env var by @KellyFrog in https://github.com/tile-ai/tilelang/pull/2851
* [BugFix] Fix absmax and abssum for uint dtypes by @jjppp in https://github.com/tile-ai/tilelang/pull/2845
* [CUDA][TMA] Separate atomic-add dtype support from layout encoding by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2846
* [TIR][Python] Trim redundant op proxy wrappers by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2850
* [CUDA][ROCm][Metal] Split backend-specific builtin ops by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2855
* [CUDA] Adopt multi-staged buffers in examples by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2836
* [CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in https://github.com/tile-ai/tilelang/pull/2861
* [Docs][LSP] Announce the open-source TileLang LSP by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2862
* [BugFix][Metal] Resolve auto backend to torch and skip disk cache for torch by @oraluben in https://github.com/tile-ai/tilelang/pull/2856
* [Typo] Correct source spelling errors by @morluto in https://github.com/tile-ai/tilelang/pull/2858
* [Refactor] Recycle T.unroll(explicit=True) for early explicit unrolling by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2859
* [BugFix] Use masked warpsync in in-w
v0.1.132026-08-03
TileLang v0.1.13
This release contains 138 commits (79 bug fixes plus features, refactors, and examples) accumulated since v0.1.12 (2026-07-08 → 2026-08-02).
The headline work is a multi-backend language-dialect refactor that replaces the runtime-activated language facade with static per-backend re-exports, alongside two major new hardware paths: SM120 NVF4 block-scale MMA for Blackwell and Metal 4 (M5) cooperative-tensor GEMM. On top of that, a large batch of correctness fixes landed across reductions/scans, atomics, TMA/copy lowering, loop-step preservation, and FP encoding edge cases.
> Breaking changes: this release removes several legacy APIs and packages. See Backend, API & Refactors before upgrading.
---
Highlights
- [CUDA] SM120 (Blackwell) NVF4 block-scale MMA support (#2364) — T.mmagemmblockscaled now routes the packed-scale SM120 path internally through package-pingpong lowering, with an optimized non-persistent example (8192³ measured at ~1527 TFLOPS on SM120). The public micro_pipeline strategy knob was removed from the API.
- [Metal] M5 cooperative tensor T.gemm (#2252) — TileLang-owned cooperative-tensor intrinsics, Metal 4 MPP matmul2d shader emission, and a shape-aware instruction selector that keeps the simdgroup fallback for fragment accumulators and unsupported tiles.
- [CUDA] Arbitrary TMEM layouts (#2785) — TMEM buffers are no longer restricted to a fixed set of layouts.
- [Language/Backend] Language dialect for multi-backends (#2734) — the runtime-activated language facade was replaced with a static from tilelang.cuda.language import * re-export; CUDA/Metal/ROCm dialects now build on tilelang.language.common with per-backend TIR overlays (details below).
- [TIR] Source-span injection (#2751) — source locations are now carried into the TIRX IR and surfaced in compiler error messages.
New Features
- CUDA
- SM70 GEMM FMA fallback (#2339) and SM75 extension of the GEMM FMA fallback (#2811) — T.gemm now works on older architectures instead of erroring out.
- Pre-SM80 fallback for bf16 __hfma (#2769).
- Stochastic FP32 → FP16/BF16 casts (#2735), with stochastic FP4/FP8 casts gated on sm_100a (#2691).
- Arbitrary TMEM layout support (#2785).
- Pipelining for multi-segment scans (#2664).
- fp32x2 ops usable as reducers (#2637).
- IKET profiler support for the CUDA backend (#2515).
- Metal
- M5 cooperative-tensor GEMM (#2252).
- Line-level threadgroup qualifier scanning for shared memory (pass 5) (#2796).
- 16-byte alignment padding for shared/threadgroup memory (#2786).
- Compiler / IR / Runtime
- Compiler pass timing profiling via the pass_profile pass-config option (with a configurable threshold) (#2622).
- lower-trace support for debugging (new doc: docs/tools/lower_trace.md) (#2725).
- Local buffer reduction lowering (#2693).
- Typed vector lane extraction API (#2789) and typing wrappers for DSL ops (#2739).
- Scalar tile scheduler state exposed (#2553).
- Host-evaluable T.assume conditions are now enforced at runtime (#2655).
- Deterministic CanProve (#2772).
Backend, API & Refactors
Language dialect refactor (#2734)
The runtime-activated language facade has been replaced by a static re-export architecture:
- Dropped the .pyi stubs + generator, py.typed, the globals()-based all scraping, and activatecuda_facade().
- Backend dialects (cuda / metal / rocm) now build on tilelang.language.common with per-backend TIR overlays.
- Import-time dtype defaults in mma/wgmma/mfma macro generators are pinned to the dtypes leaf and no longer touch the half-initialized facade during bootstrap.
- 2:4 sparsity layout metadata extracted into tilelang/cuda/intrinsics/sparse_layout.py (a dtypes-only leaf).
Follow-up fixes: ROCm intrinsic resolution (#2779), rng_init (#2776), and shared-intrinsic resolution across backends.
Removals (breaking)
- Legacy DLPack execution backend removed (#2816).
- Intrinsic compatibility facade removed (#2812).
- tilelang.common package removed (#2810).
- Carver shape-inference module removed (#2813).
- Example-only helpers moved out of the tilelang package (#2761).
FFI / JIT / Build
- Support for apache-tvm-ffi 0.1.12, while keeping 0.1.11 compatibility (#2795); lower bound raised to >=0.1.11 (#2736).
- JIT now reuses the compiled executable across kernel launches (#2686).
- NVRTC scalar parameters and dynamic strides are marshaled correctly (#2756).
- ptxas register-usage level is cast to int before building the nvcc command (#2641).
- Cross-compiler options isolated per invocation (#2728).
- Shared Int64Promoter extracted into a common header (#2558).
- CI: actions/setup-python 6 → 7 (#2773); transformers bumped in examples/bitnet-1.58b (#2658).
- Docs: SKILL.md updated for editable installs and clarified development workflow (#2533).
Bug Fixes
Loop & control-flow preservation
- Loop steps preserved when unrolling loops — a fix (#2784) was reverted (#2834) and then correctly re-landed (#2835).
- Explicit loop steps preserved when transforms rebuild For nodes (#2752).
- Loop steps preserved during unswitching (#2741) and guard identity preserved in LoopUnswitching (#2585).
- If-condition evaluation preserved during fan-out (#2764) and re-evaluation of mutable if conditions (#2744).
Reductions & scans
- Scalar AllReduce thread-range analysis simplified; partial scalar reduce barrier participation fixed (#2777, #2814).
- warpreduce no longer truncates int64/uint64 to 32 bits on sm80+ (#2782).
- Non-power-of-two AllReduce widths rejected (#2611); packed AllReduce workspace pointer fixed (#2778); blockDim used as workspace stride in batch AllReduce (#2621).
- 2D scan kernel now receives the buffer row stride, fixing silent miscomputation (#2620); wrong offset when scanning a non-zero-offset buffer sub-region fixed (#2680).
- Thread-segment projection for packed layouts fixed (#2647); grouped reduce_sum over-counts on straddle layouts fixed (#2424).
- nan_propagate honored in reduce max/min/absmax clear=False write-back (#2788).
- Float dtypes rejected in bitwise reduce with an actionable error (#2676).
Atomics & memory ordering
- fp16/bf16 T.atomicmax/T.atomicmin no longer silently corrupt fp32 values (#2780).
- returnprev supported for scalar atomicmin/atomicmax (#2672), atomicaddx2 with BufferRegion destinations (#2753), and HIP vector atomic add (#2712).
- T.atomic_addx4 return type guarded for sliced destinations (#2590).
- Atomic load/store implemented for HIP (#2711); invalid atomic memory orders rejected (#2666); CUDA consume ordering mapped to acquire PTX (#2713).
- TMA atomic-add layout validation refactored (e0f0ac90) and unsupported dtypes rejected (#2830).
Numerics, vectors & dtypes
- FP8 E4M3 special encodings decoded correctly (#2710); T.infinity supported for float8_e5m2 (#2671).
- bf16 NaN/Inf preserved during RNE packing (#2690).
- Signed int32 lanes zero-extended in 256-bit vector pack (#2673); 32-lane 8-bit CUDA vectors packed correctly (#2701).
- FP4 dequant symbolic exponent clamp fixed (#2656).
- T.pow/T.power fixed for constant integer exponent y <= 0 (#2677).
- T.exp computes ex, not 2**x (docstring + CuTeDSL codegen) (#2696).
- IEEE math intrinsic names corrected for fp64/fp16/bf16 (#2619).
- Unsupported fast-math input dtypes rejected (#2804); mixed packed x2 operand dtypes rejected (#2802); floating-point predicates rejected in vote intrinsics (#2800); alloc_var initializer dtype preserved (#2801); invalid dtypes rejected in T.dp4a (#2652).
- Scalar T.copy path casts to the destination dtype (#2771).
- Canonical-simplify LT Case 2 gated on extra scale == +1 (#26
v0.1.122026-07-08
TileLang v0.1.12 共 91 个提交，主线是新增 LLVM 后端与后端注册表重构：host/device CodeGen 改由 registry 分发，目标检测与规范化走注册机制，ExecutionBackend 并入后端模块。CUDA 侧放开 TMA/GMMA 对任意及 swizzled 共享内存布局的支持，fp8 与 half/bf16 转换改为向量化代码生成，性能更好。另有 tile scheduler、passvisualizer 结构树浏览器与 pass 差异显示，以及 SM100+ 的 st.bulk 清零、Blackwell 的 stmatrix m16n8 等新内置函数。JIT 引入跨主机 CUDA 二进制缓存、可配置 NVCC 超时与 TILELANGVERBOSE 输出开关，编译选项纳入缓存键。修复相当密集：流水线异步等待计数、布局推断除零、T.Persistent 尾块丢失、原子操作类型不匹配、bf16 exp 自递归与 rsqrt 重载、DeepSeek V3.2 topk 边界等；autotuning 跨设备测量更稳定，dobench 新增 cachesize 选项。
v0.1.112026-06-08
TileLang 本次更新集中在算子能力与编译管线两方面。功能上，T.tmacopy 新增 threadextent 参数，TMA 拷贝更灵活；TileOP 补上 scan 算子，CUDA 侧支持 ffs 位操作、命名 barrier arrive，以及 f32 到 FP8/FP4 的随机舍入转换；新增 float4e2m1_unpacked 数据类型，tfloat32 在 CuTeDSL 中按 Float32 表示。管线方面，软件流程（software pipeline）变换做了整体重构，共享内存分配合并与 SIMT 循环切分的分析也一并修正，同时把 PassPipeline 移入后端目录，为多后端铺路。修复涉及动态索引原子加载、混合精度 reduce 的类型转换、FP32 MMA 操作数按 TF32 下降低、缓存源码恢复与 Metal 回退，以及 Windows 下 libtvm 符号链接的兼容问题。另新增 CLC 流水线 2-CTA GEMM（sm100）示例。
v0.1.102026-05-25
TileLang 这版把重心放在后端铺开、GPU 指令补齐和编译管线打磨上。AMD 阵营收获最大：RDNA3/RDNA3.5 的 WMMA、gfx950 的 160K LDS 与 copy.async、LDS 转置读取、INT8 MFMA、MXFP4 以及 gfx1151 目标都已就位；CUDA 侧新增 Blackwell 的 MXFP8 块缩放 GEMM、FP4 TensorMap 拷贝、gather4/scatter4 和 T.copy_cluster 多播通信，另外补上了 SM75 上 FP16/INT8/INT4 的原生 MMA，Metal 也首次跑通 GEMM。
编译器这边，自动调优支持流水线化与分组编译、多卡基准测试和 donotspecialize；CUDA、ROCm、Metal、CPU、WebGPU 的降级与代码生成被拆成独立后端模块，IR 正逐步迁往 tirx。软件流水线的标量绑定、TMA 守卫和归约向量化都有改进，另加 TLDISABLESHAREDMEMORYREUSE 配置。
修掉的坑不少：TMA 的对齐与布局推断、SM90 WGMMA 类型笔误、SM75 上 T.gemm()、ROCm 的 FP4 缓冲键及若干 HIP 代码生成问题。
注意两处破坏性变更：不再支持 Python 3.9（需 >=3.10），apache-tvm-ffi 要求升到 >=0.1.10。PyPI 发布流程和 Windows 下的 DLL 拆分也已上线。
v0.1.92026-04-22
TileLang 这版更新面较广，核心围绕编译优化与新硬件特性展开。布局系统新增 Layout.repeat 与 Layout.expand，支持把 tiling atom 布局扩展到更高维度，并补齐了高维 GEMM 布局能力。算子层面引入 cluster 的启动、查询、同步与 barrier 操作，CUDA 侧支持 tcgen5mma 的 gemm ts，SM90/SM100 的 CuTeDSL 后端也接上了 GEMM V2。
编译器方面，InjectFenceProxy 经过重构并放宽了 async proxy 判定条件，cp.async 下降逻辑同步改进并新增 async_copy 算子；循环向量化的独立变量识别、AMD gfx950 的 16x16x32 bf16/fp16 指令支持也已跟进。
修复了几个容易踩坑的问题：Hopper 无 warp specialization 时的 TMA 下降、flatten buffer 的 elem_offset 重复计数、AtomicAddx2 模板类型不匹配，以及 FP4 在 CUDA 13.0 以下误用 cvt 指令。另外新增 DumpIR 配置项，便于排查 IR。
v0.1.82026-02-16
本次更新集中在编译器正确性与性能优化：修复 GQA 解码切分 kernel 布局、T.maketensor 分配位置不当、analyzer 绑定冲突等缺陷；重构向量化不变量检查，并实现 FP8 到 FP32 的向量转换；inflight 计算改为支持动态流水线 extent，T.cumsum 新增区域输入，新增 T.ldg 只读访问与 CUDA 代码生成的只读参数标注，InjectAssumes 在 SplitHostDevice 之后也能生效。AMD 方面，MI300X 已启用 FlashAttention-2 前向，MI3x 系列 Docker 构建的三处问题也已修复。工程侧，依赖升级至 apache-tvm-ffi >=0.1.2 并新增 torch-c-dlpack-ext，Yapf 退役、改用 ruff format，primitives 目录并入 tileop，测试引入 pytest.mark.parameterize 以加快并行执行；另修正 CMake 不再向 sys.path 注入项目根目录，以及 compileflags 经 passconfigs 传入 FFI 编译路径的问题。
v0.1.7.post32026-01-18
本次更新以修复与重构为主：修复变长注意力前向在 Sq != Skv 时的错误、sm120 上归约导致的挂起，以及 gemmspv2 缺失 clearaccum、向量化加载覆盖不全、local.var 缓冲区被误判与 T.Fill 失效等问题。功能上，DeepSeek 稀疏 MLA 反向通过拆分 H 支持更大的 H 规模，gemm v2 接入 cutedsl 后端，z3 超时改为 rlimit，证明行为更稳定。KernelCache 改为继承式设计，为并行与拷贝节点引入布局注解，补齐缓冲区操作类型标注，宿主代码新增后处理钩子。构建方面已按 CUDA 13.1 编译并对 nvfp8e8m0 类型做版本检查，隔离构建本地版本的报错也已修复。
v0.1.7.post22025-12-31
本轮更新集中在几条线上。编译器侧，Inject Pipeline Pass 重构了缓冲区分配，为 ParallelOPNode 与 CopyNode 引入布局标注，VectorizeLoop 的标量化处理得到改进，向量化加载补齐常用数据类型。KernelCache 改为继承式设计，LazyJIT 将类型标注移入函数体，Z3 超时改用 rlimit，证明行为更可复现。功能上，cutedsl 后端适配 gemm v2，DeepSeek 稀疏 MLA 反向通过拆分 H 支持更大的 H，gemmspv2 补上了遗漏的 clearaccum 逻辑。修复包括 sm120 上归约卡死、varlen attention 在 Sq != S_kv 时的示例错误、local.var 缓冲被误判及 T.Fill 的问题。CI 已切换到 CUDA 13.1 构建，并新增本地回归测试脚本。
