v0.1.12

[LeiWang1999]

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 ([Feature] Add LLVM backend support #2409), with follow-up fixes for auto backend resolution ([BugFix] Fix llvm auto backend resolution #2519) and module export ([BugFix] Skip source-compilation options when exporting LLVM module #2467)
• Tile scheduler introduced ([Feature] Introduce tile scheduler #2441)
• Backend registry architecture: host and device CodeGen are now dispatched through a backend registry ([Backend] Dispatch device CodeGen through backend registry #2442, [Backend] Dispatch host CodeGen through backend registry #2446), target detection/normalization is registration-based, and ExecutionBackend was merged into the backend module ([Backend] Add target detector/normalizer registeration and merge ExecutionBackend into backend module #2323); kernel launch is materialized per backend ([Transform] Materialize kernel launch per backend #2387)
• Developer tooling: pass_visualizer structure-tree pass browser ([Feature][Tool] Add pass_visualizer: structure-tree pass browser #2449) and pass-diff display for debugging ([feat] add pass diff show for debugging #2375)
• New CUDA intrinsics exposed ([Feature] Expose multiple CUDA intrinsics #2473), st.bulk shared-memory zero fill on SM100+ ([Feature] Support st.bulk for shared zero fill on SM100+ #2403), stmatrix m16n8 on Blackwell ([Feature] Support stmatrix m16n8 on Blackwell #2417), SM75 MMA dispatchers for FP16 accumulation and UINT8 ([CUDA] Add SM75 MMA dispatchers for FP16 accumulation and UINT8 #2392)

CUDA / Codegen Improvements

• Optimized fp8↔half/bf16 casts: vectorized and scalar cast codegen ([Enhancement] Add optimized fp8↔half/bf16 vectorized and scalar cast codegen #2511, [Enhancement] Add vectorized fp8x2 <-> fp16/bf16 cast codegen #2475), plus a fix for vectorized fp16↔bf16 cast compilation ([BugFix] Fix vectorized fp16<->bf16 cast compilation #2407)
• TMA lowering for arbitrary/swizzled SMEM layouts ([Backend] Support TMA lowering for arbitrary (swizzled) SMEM layout #2380) and GMMA/UMMA lowering for sliced (arbitrary-layout) SMEM ([Backend] [CUDA] Support GMMA/UMMA lowering for sliced SMEM layout (actually arbitrary layout) #2452)
• Reduced CUDA template include overhead ([CUDA] Reduce template include overhead #2474), swizzled TMA buffer alignment ([CUDA] Align swizzled TMA shared buffers #2391), RNG state kept in kernel scope ([CUDA][Codegen] Keep RNG state in kernel scope #2540)
• Warp-specialization fixes: register over-subscription ([BugFix] Fix warp-specialized register over-subscription #2406), register reallocation for 1P1C ([BugFix] Support warpgroup register reallocation for 1P1C #2440)

JIT / Caching / Build

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

Notable Bug Fixes

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

Other

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

New Contributors

• @Triang-jyed-driung made their first contribution in [NVCC] add -ccbin arguments to specify C++ compiler #2348
• @jjjxia made their first contribution in Fix storage rewrite source codegen test #2353
• @ppppqp made their first contribution in [Fix] Bind statement missing role marker #2362
• @LaiQuan-conquer made their first contribution in Fix bf16 CUDA rsqrt overload #2386
• @Federicorao made their first contribution in Add assert that block_K is a multiple of micro_size_k in CUDA MMA GEMM backends to prevent silent miscompilation. #2390
• @Yongqi-Zhuo made their first contribution in [CI] Trigger CI when a PR is marked ready for review #2411
• @Witherstrike made their first contribution in [Feature] Add LLVM backend support #2409
• @jjppp made their first contribution in [BugFix] Fix packed vector type printing for bfloat16 with metal codegen #2437
• @RuneFang made their first contribution in [BugFix] Fix T.Persistent dropping tiles when last dim is not a multiple of group_size #2455
• @ring00 made their first contribution in [BugFix] Only allocate reducer workspace for cross-warp AllReduce #2494
• @mengmeexix made their first contribution in [BugFix] Fix DeepSeek V3.2 topk threshold on exact-boundary inputs #2513
• @net-snix made their first contribution in fix(cuda): detect auto target arch from current device, not device 0 #2517
• @shuyilinn made their first contribution in [Feature][Tool] Add pass_visualizer: structure-tree pass browser #2449

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

---

This discussion was created from the release v0.1.12.