TileLang v0.1.11 → v0.1.12 Changes

Summary of the main changes between v0.1.11 and v0.1.12 (91 commits).

New Features

• LLVM backend support (#2409), with follow-up fixes for auto backend resolution (#2519) and module export (#2467)
• Tile scheduler introduced (#2441)
• Backend registry architecture: host and device CodeGen are now dispatched through a backend registry (#2442, [#2446]), target detection/normalization is registration-based, and ExecutionBackend was merged into the backend module (#2323); kernel launch is materialized per backend (#2387)
• Developer tooling: pass_visualizer structure-tree pass browser (#2449) and pass-diff display for debugging (#2375)
• New CUDA intrinsics exposed (#2473), st.bulk shared-memory zero fill on SM100+ (#2403), stmatrix m16n8 on Blackwell (#2417), SM75 MMA dispatchers for FP16 accumulation and UINT8 (#2392)

CUDA / Codegen Improvements

• Optimized fp8↔half/bf16 casts: vectorized and scalar cast codegen (#2511, [#2475]), plus a fix for vectorized fp16↔bf16 cast compilation (#2407)
• TMA lowering for arbitrary/swizzled SMEM layouts (#2380) and GMMA/UMMA lowering for sliced (arbitrary-layout) SMEM (#2452)
• Reduced CUDA template include overhead (#2474), swizzled TMA buffer alignment (#2391), RNG state kept in kernel scope (#2540)
• Warp-specialization fixes: register over-subscription (#2406), register reallocation for 1P1C (#2440)

JIT / Caching / Build

• Cross-host CUDA binary cache (#2459); compile options now included in the cache key (#2532)
• PyTorch extensions and perf wheels are cached (#2509); frontend disk cache removed (#2363); lazy kernel lookup caching improved (#2357)
• JIT diagnostics and configurable NVCC timeout (#2350), -ccbin support for choosing the C++ compiler (#2348), TILELANG_VERBOSE env var to control compile output (#2453)

Notable Bug Fixes

• Pipeline: fixed physical async wait counts (#2505) and simplified async copy lowering (#2444)
• Layout/Transform: divide-by-zero in LayoutInference on non-power-of-two broadcast (#2469), avoided thread-indexed replicated fragment readback (#2514), kept all-rep reducers from scalarizing vector plans (#2507), reducer workspace only allocated for cross-warp AllReduce (#2494)
• Correctness: T.Persistent dropping tiles when the last dim isn't a multiple of group_size (#2455), sign-extension bugs in packed uint32 decode (#2500) and make_int negative int8 lanes (#2438), PTX v4 atomics for fp16/bf16 atomic_addx4 (#2492), vectorized atomic_add dtype mismatch (#2414), bf16 exp self-recursion (#2402) and rsqrt overload (#2386)
• DeepSeek V3.2 topk threshold on exact-boundary inputs (#2513); flash attention bwd varlen NaN fix (#2461); SM100 CLC GEMM schedule-state lifetime (#2423)
• Autotuning benchmarking stabilized across devices (#2370); do_bench gained a cache_size option (#2531)

Other

• T.view / T.reshape enhancements (#2450), better T.assume/loop-bound handling to eliminate redundant boundary checks (#2502), improved diagnostics for T.serial fragment access (#2462)
• C++ style guide added and API naming/namespace normalization across the C++ codebase (#2430, [#2434]–2436)
• Auto target arch now detected from the current device instead of device 0 (#2517)

Overall

This release centers on the new LLVM backend and backend-registry refactor, major TMA/GMMA layout flexibility on CUDA, fp8/fp16/bf16 cast performance, and a large batch of correctness fixes across layout inference, atomics, and pipelining.

What's Changed

• [NVCC] add -ccbin arguments to specify C++ compiler by @Triang-jyed-driung in https://github.com/tile-ai/tilelang/pull/2348
• [Backend] Add target detector/normalizer registeration and merge ExecutionBackend into backend module by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2323
• [Feature]Add JIT diagnostics and configurable NVCC timeout by @TerminusAkivili in https://github.com/tile-ai/tilelang/pull/2350
• [JIT] Improve lazy kernel lookup caching by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2357
• [Backend] Cleanup Metal Codegen, split AsyncCopy lowering and common target utils by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2361
• [JIT] Remove frontend disk cache by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2363
• Fix storage rewrite source codegen test by @jjjxia in https://github.com/tile-ai/tilelang/pull/2353
• [BugFix][CuTeDSL] Complete TileKernels benchmark support by @JayceSu98 in https://github.com/tile-ai/tilelang/pull/2319
• [Build] Pin apache-tvm-ffi to compatible versions by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2373
• [Build] Fix sdist version metadata by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2374
• [TIR][CUDA] Remove unused instruction annotation plumbing by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2389
• [Transform] Materialize kernel launch per backend by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2387
• [CI][Cache] Enable local ccache for self-hosted runners by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2388
• [Fix] Bind statement missing role marker by @ppppqp in https://github.com/tile-ai/tilelang/pull/2362
• Fix bf16 CUDA rsqrt overload by @LaiQuan-conquer in https://github.com/tile-ai/tilelang/pull/2386
• [CUDA] Add SM75 MMA dispatchers for FP16 accumulation and UINT8 by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2392
• Add assert that block_K is a multiple of micro_size_k in CUDA MMA GEMM backends to prevent silent miscompilation. by @Federicorao in https://github.com/tile-ai/tilelang/pull/2390
• [CUDA] Align swizzled TMA shared buffers by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2391
• [CI]: Bump pypa/cibuildwheel from 4.0 to 4.1 by @dependabot[bot] in https://github.com/tile-ai/tilelang/pull/2400
• [Language] Fix TMA 1D load test lowering by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2401
• [BugFix] Fix bf16 CUDA exp self-recursion by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2402
• [Feature] Support st.bulk for shared zero fill on SM100+ by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2403
• [Autotune] Stabilize autotune benchmarking across devices and fix the unconsistency in the autotuning process. by @Wazrrr in https://github.com/tile-ai/tilelang/pull/2370
• [feat] add pass diff show for debugging by @erhsh in https://github.com/tile-ai/tilelang/pull/2375
• [BugFix] Fix warp-specialized register over-subscription by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2406
• Try to preload z3 to resolve tvm deps by @oraluben in https://github.com/tile-ai/tilelang/pull/2405
• [BugFix][CuTeDSL] Fix TileKernels scan, optional-shape, and e5m6 paths by @JayceSu98 in https://github.com/tile-ai/tilelang/pull/2369
• [BugFix] Fix vectorized fp16<->bf16 cast compilation by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2407
• [CI] Trigger CI when a PR is marked ready for review by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2411
• [CI] Exclude agent files from markdown fixes by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2412
• [Feature] Add LLVM backend support by @Witherstrike in https://github.com/tile-ai/tilelang/pull/2409
• [Backend] Support TMA lowering for arbitrary (swizzled) SMEM layout by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2380
• [BugFix] Fix vectorized atomic_add dtype mismatch reinterpret by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2414
• fix: remove abandoned TILELANG_CLEAR_CACHE env var and dead clear_cache() by @erhsh in https://github.com/tile-ai/tilelang/pull/2425
• Fix SM100 CLC GEMM schedule-state lifetime by @VitalyAnkh in https://github.com/tile-ai/tilelang/pull/2423
• [Docs][Dev] Add C++ style guide by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2430
• [Docs][C++] Clarify namespace style policy by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2432
• [CI]: Bump actions/checkout from 6 to 7 by @dependabot[bot] in https://github.com/tile-ai/tilelang/pull/2428
• [C++][Headers] Clarify namespace boundaries by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2434
• [C++][Style] Normalize API context argument names by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2435
• [C++][Style] Normalize API helper names by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2436
• [BugFix] Fix packed vector type printing for bfloat16 with metal codegen by @jjppp in https://github.com/tile-ai/tilelang/pull/2437
• [Transform] Warn on vectorized loop serial fallback by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2439
• [Feature] Support stmatrix m16n8 on Blackwell by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2417
• [BugFix] Support warpgroup register reallocation for 1P1C by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2440
• [BugFix] Fix make_int sign-extending negative int8 lanes by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2438
• [Backend] Dispatch device CodeGen through backend registry by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2442
• [CUDA] Remove __shfl_sync from tl_shuffle_elect by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2445
• [Backend] Dispatch host CodeGen through backend registry by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2446
• [Pipeline] Simplify async copy lowering by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2444
• [Feature] Introduce tile scheduler by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2441
• Enhance T.view and T.reshape by @bucket-xv in https://github.com/tile-ai/tilelang/pull/2450
• [Docs] Optimize TileLang C++ coding style by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2447
• [CUDA] Add target code attribute support by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2454
• Use TILELANG_VERBOSE environment var to control the compile output info by @bucket-xv in https://github.com/tile-ai/tilelang/pull/2453
• [CUDA] Increase MMA descriptor without touching high bits by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2460
• [BugFix] Fix T.Persistent dropping tiles when last dim is not a multiple of group_size by @RuneFang in https://github.com/tile-ai/tilelang/pull/2455
• [CI][BugFix] Flash bwd varlen: zero-init lse/Delta padding to avoid NaN in Dk by @RuneFang in https://github.com/tile-ai/tilelang/pull/2461
• [CUDA][JIT][Cache] Add cross-host CUDA binary cache by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2459
• [BugFix] Improve diagnostic for T.serial fragment access by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2462
• [BugFix] Ignore flat Bind nodes in ForBodyContainsSeqStmt by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2464
• [Enhancement] Add vectorized fp8x2 <-> fp16/bf16 cast codegen by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2475
• [Env] Require JSON for default target config by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2491
• [Testing][CUDA][CI] Improve regression workflow and CUDA selection by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2495
• [Testing][CI] Isolate perf regression runner imports by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2498
• [BugFix] Only allocate reducer workspace for cross-warp AllReduce by @ring00 in https://github.com/tile-ai/tilelang/pull/2494
• [CUDA] Reduce template include overhead by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2474
• [Enhancement] Fix T.assume and loop bounds to eliminate redundant boundary checks by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2502
• [Backend] [CUDA] Support GMMA/UMMA lowering for sliced SMEM layout (actually arbitrary layout) by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2452
• [BugFix] Sign-extend packed uint32 signed decode by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2500
• [Pipeline] Fix physical async wait counts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2505
• [BugFix][CUDA] Use PTX v4 atomics for fp16/bf16 atomic_addx4 by @JayceSu98 in https://github.com/tile-ai/tilelang/pull/2492
• [BugFix] Skip source-compilation options when exporting LLVM module by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2467
• [Transform] Keep all-rep reducers from scalarizing vector plans by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2507
• [Feature] Expose multiple CUDA intrinsics by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2473
• Bump transformers from 5.0.0rc3 to 5.3.0 in /examples/bitnet-1.58b by @dependabot[bot] in https://github.com/tile-ai/tilelang/pull/2512
• [JIT][Cache] Cache PyTorch extensions and perf wheels by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2509
• [BugFix] Fix LayoutInference divide-by-zero on non-power-of-two broadcast by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2469
• [BugFix] Fix DeepSeek V3.2 topk threshold on exact-boundary inputs by @mengmeexix in https://github.com/tile-ai/tilelang/pull/2513
• [Transform][Layout] Avoid thread-indexed replicated fragment readback by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2514
• fix(cuda): detect auto target arch from current device, not device 0 by @net-snix in https://github.com/tile-ai/tilelang/pull/2517
• [BugFix] Fix llvm auto backend resolution by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2519
• [Feature][Tool] Add pass_visualizer: structure-tree pass browser by @shuyilinn in https://github.com/tile-ai/tilelang/pull/2449
• [CUDA][Cache] Include compile options in binary cache key by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2532
• [CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in https://github.com/tile-ai/tilelang/pull/2535
• [Enhancement] Add optimized fp8↔half/bf16 vectorized and scalar cast codegen by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2511
• [Enhancement] Add cache_size option to do_bench by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2531
• [CUDA][Codegen] Keep RNG state in kernel scope by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2540
• [Feature] Clean up CPU pass pipeline by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2534
• [CUDA][ROCm] Rename GPU stub library artifacts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2541
• [Release] Bump version to 0.1.12 by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2544

New Contributors

• @Triang-jyed-driung made their first contribution in https://github.com/tile-ai/tilelang/pull/2348
• @jjjxia made their first contribution in https://github.com/tile-ai/tilelang/pull/2353
• @ppppqp made their first contribution in https://github.com/tile-ai/tilelang/pull/2362
• @LaiQuan-conquer made their first contribution in https://github.com/tile-ai/tilelang/pull/2386
• @Federicorao made their first contribution in https://github.com/tile-ai/tilelang/pull/2390
• @Yongqi-Zhuo made their first contribution in https://github.com/tile-ai/tilelang/pull/2411
• @Witherstrike made their first contribution in https://github.com/tile-ai/tilelang/pull/2409
• @jjppp made their first contribution in https://github.com/tile-ai/tilelang/pull/2437
• @RuneFang made their first contribution in https://github.com/tile-ai/tilelang/pull/2455
• @ring00 made their first contribution in https://github.com/tile-ai/tilelang/pull/2494
• @mengmeexix made their first contribution in https://github.com/tile-ai/tilelang/pull/2513
• @net-snix made their first contribution in https://github.com/tile-ai/tilelang/pull/2517
• @shuyilinn made their first contribution in https://github.com/tile-ai/tilelang/pull/2449

Full Changelog: https://github.com/tile-ai/tilelang/compare/v0.1.11...v0.1.12