v0.1.12

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

New Contributors

• @Triang-jyed-driung made their first contribution in #2348
• @jjjxia made their first contribution in #2353
• @ppppqp made their first contribution in #2362
• @LaiQuan-conquer made their first contribution in #2386
• @Federicorao made their first contribution in #2390
• @Yongqi-Zhuo made their first contribution in #2411
• @Witherstrike made their first contribution in #2409
• @jjppp made their first contribution in #2437
• @RuneFang made their first contribution in #2455
• @ring00 made their first contribution in #2494
• @mengmeexix made their first contribution in #2513
• @net-snix made their first contribution in #2517
• @shuyilinn made their first contribution in #2449

Full Changelog: v0.1.11...v0.1.12