| Name | Modified | Size | Downloads / Week |
|---|---|---|---|
| Parent folder | |||
| tilelang-0.1.13.tar.gz | 2026-08-03 | 94.6 MB | |
| tilelang-0.1.13-cp38-abi3-win_amd64.whl | 2026-08-03 | 34.4 MB | |
| tilelang-0.1.13-cp38-abi3-manylinux_2_34_aarch64.whl | 2026-08-03 | 46.7 MB | |
| tilelang-0.1.13-cp38-abi3-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl | 2026-08-03 | 51.0 MB | |
| tilelang-0.1.13-cp38-abi3-macosx_11_0_arm64.whl | 2026-08-03 | 38.7 MB | |
| README.md | 2026-08-03 | 42.5 kB | |
| v0.1.13 source code.tar.gz | 2026-08-03 | 12.3 MB | |
| v0.1.13 source code.zip | 2026-08-03 | 13.2 MB | |
| Totals: 8 Items | 290.8 MB | 0 | |
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.mma_gemm_blockscalednow 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 publicmicro_pipelinestrategy knob was removed from the API. - [Metal] M5 cooperative tensor
T.gemm(#2252) — TileLang-owned cooperative-tensor intrinsics, Metal 4 MPPmatmul2dshader 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 ontilelang.language.commonwith 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.gemmnow 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).
fp32x2ops 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_profilepass-config option (with a configurable threshold) (#2622). lower-tracesupport 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.assumeconditions 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
.pyistubs + generator,py.typed, theglobals()-based__all__scraping, and_activate_cuda_facade(). - Backend dialects (
cuda/metal/rocm) now build ontilelang.language.commonwith per-backend TIR overlays. - Import-time dtype defaults in
mma/wgmma/mfmamacro 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.commonpackage removed (#2810).- Carver shape-inference module removed (#2813).
- Example-only helpers moved out of the
tilelangpackage (#2761).
FFI / JIT / Build
- Support for
apache-tvm-ffi0.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).
ptxasregister-usage level is cast tointbefore building the nvcc command (#2641).- Cross-compiler options isolated per invocation (#2728).
- Shared
Int64Promoterextracted into a common header (#2558). - CI:
actions/setup-python6 → 7 (#2773);transformersbumped inexamples/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
Fornodes (#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]).
warp_reduceno longer truncates int64/uint64 to 32 bits on sm_80+ (#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_sumover-counts on straddle layouts fixed (#2424). nan_propagatehonored in reduce max/min/absmaxclear=Falsewrite-back (#2788).- Float dtypes rejected in bitwise reduce with an actionable error (#2676).
Atomics & memory ordering
- fp16/bf16
T.atomic_max/T.atomic_minno longer silently corrupt fp32 values (#2780). return_prevsupported for scalaratomic_min/atomic_max(#2672),atomic_addx2withBufferRegiondestinations (#2753), and HIP vector atomic add (#2712).T.atomic_addx4return 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.infinitysupported 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.powerfixed for constant integer exponenty <= 0(#2677).T.__expcomputese**x, not2**x(docstring + CuTeDSL codegen) (#2696).- IEEE math intrinsic names corrected for fp64/fp16/bf16 (#2619).
- Unsupported fast-math input dtypes rejected (#2804); mixed packed
x2operand dtypes rejected (#2802); floating-point predicates rejected in vote intrinsics (#2800);alloc_varinitializer dtype preserved (#2801); invalid dtypes rejected inT.dp4a(#2652). - Scalar
T.copypath casts to the destination dtype (#2771). - Canonical-simplify LT Case 2 gated on extra scale
== +1(#2649); vectorizedSelectconstraint handling fixed (#052e6741).
TMA / copy / memory layout
- Strided global buffers handled correctly in 1D TMA copies (#2746); descriptor TMA skipped for device-bound copy bases (#2803).
- Partial 1-D TMA stores no longer bypass bounds checks (#2716); 1D bulk TMA transfer alignment check fixed (#2646); 1D TMA selection fixed for versioned layouts (#2737); non-16B cluster bulk copies fall back (#2683).
st.bulkdestination emitted as a shared write to fix a missing barrier and compilation-introduced races (#2700).- Tile copy OOB respects the safe value (#2636); runtime-dependent vector negative indices supported (#2654).
- Operator precedence fixed in the
increase_descriptor_offsetguard (#2675). - Packed shared memory allocation sizes corrected for CUDA/HIP (#2660); HIP predicated dword copy zero fill fixed (#2721).
- Buffer element offsets preserved in access pointers (#2727); decoupled cast buffer scope preserved in codegen (#2545).
T.transposeswaps only the final two axes (#2757); contracting shared-buffer layouts rejected inT.annotate_layout(#2719); unused fragment buffers allowed without layouts (#2717); shared-TMEM buffer pointer types checked before dereference (#2794).
Metal backend
- Threadgroup address-space qualifier emitted for shared-memory pointer arithmetic (#2770).
- Barriers emitted for dynamic shared memory (#2738).
- Explicit row strides honored in Metal GEMM (#2730).
- Metal stream bridge fixed (#2639).
- Arithmetic operators added to
vec_typeincommon.hfor CPU codegen (#2768).
Race analysis & warp-specialization
- Two-instance modeling fixed in ThreadSync cross-thread race checks (#2805).
- Flat
Bindmodeling fixed in parallel race checks (#2665). - VerifyParallelLoop race diagnostics aggregated with source spans (#2806).
- Side-effecting binds no longer classified as replayable — fixes atomics being re-executed at every use site since v0.1.11 (#2651).
- GEMM accumulator writes tracked in the warp-specialization liveness collector (#2685).
- Pipeline replacement fixed under persistent
T.serial(#2674). - WGMMA C-store layout fixed for multiple warpgroups along M (#2663).
- Unsafe non-warp-multiple partial thread sync rejected (#2679).
GEMM / misc compiler fixes
- Uncovered warp partitions in
T.gemmrejected instead of silently producing wrong results (#2724). - MFMA
DataTypeargs no longer break compilation on ROCm (#2726). - PCWS index dtype handling fixed (#2783).
- CPU-fallback thread placeholder replaced with a constant-zero logical thread index (#2718).
- Non-positive thread extents rejected in
T.Kernel(#2653). - Callee global symbols used for cross-target calls (#2740).
- TMEM/TMA builtins gated by CUDA architecture (#2743).
- Typo
fragmentspelling corrected (#2695);BufferStorecast warning context improved (#2733).
Autotuning
- Early stop to skip slow configs during benchmarking (#2723), including decorator mode with an example (#2729).
- Per-config
pass_configssupported in autotuning (#2496). - Autotuner cache no longer reused across different outputs and validation settings (#2793).
- Segfault fixed when tunable parameters default to
None(#2657).
Examples
- dLLM: block-causal attention example, including a varlen variant (#2499).
- DeepSeek-V3.2: adaptive thread selection for the sparse MLA backward kernel (launch width derived from head-block size) (#2592);
topk_selectormemory-access optimization with thread coarsening — ~1.9× faster with identical results (#2659). - SM120 NVFP4 block-scale GEMM example reworked to a non-persistent TileLang kernel using
T.mma_gemm_blockscaled(#2364).
Full commit list
138 commits (click to expand)
e0f0ac90 [CUDA] Refactor TMA atomic add layout validation 6b81bb87 [BugFix] Correctly preserve loop step when unrolling loops (#2835) 09526a27 Revert "[BugFix] Preserve loop step when unrolling loops" (#2834) 56a0f729 [BugFix] Reject unsupported TMA atomic add dtypes (#2830) bdb769ae [Enhancement] Aggregate VerifyParallelLoop race diagnostics with span (#2806) e01c498b [JIT] Remove legacy DLPack execution backend (#2816) 2bb0def9 [BugFix] Fix two-instance modeling in ThreadSync cross-thread race checks (#2805) 8f34abf4 [CUDA] Add SM120 NVF4 block-scale MMA support (#2364) 3e4a0544 [Carver] Remove unused shape inference module (#2813) e18d9699 [CUDA][Reduce] Simplify scalar AllReduce thread range analysis (#2814) 50481cce [Refactor] Remove intrinsic compatibility facade (#2812) 21e8c064 [BugFix] Reject unsupported fast-math input dtypes (#2804) 0bc1913d [BugFix] Resolve partial scalar reduce barrier participation (#2777) 5b1f3218 [BugFix] Reject uncovered warp partitions in T.gemm instead of producing silently wrong results (#2724) 7fd95363 [CUDA] Extend the GEMM FMA fallback to SM75 (#2811) dd92b781 [Refactor] Remove unused tilelang.common package (#2810) 6c3dd971 [BugFix] Skip descriptor TMA for device-bound copy bases (#2803) e41fadbe [BugFix] Fix warp_reduce truncating int64/uint64 to 32 bits on sm_80+ (#2782) 2a06036f [BugFix] Reject mixed packed x2 operand dtypes (#2802) 32e02e6c [Fix] Refine architecture guards (#2790) bc9515fe [BugFix] Prevent autotuner cache reuse across different outputs and validation settings (#2793) 2c84f4f9 [BugFix] Check shared-TMEM buffer pointer types before dereference (#2794) 28f70338 [Metal] Add line-level threadgroup qualifier scanning (pass 5) (#2796) 51f88a88 [BugFix] Reject floating-point predicates in vote intrinsics (#2800) b1b605da [BugFix] Preserve alloc_var initializer dtype (#2801) 1cb4d4f3 [CUDA] Support arbitrary TMEM layouts (#2785) 4086ba8e [FFI] Support apache-tvm-ffi 0.1.12 (#2795) 7ec5adbe [BugFix] Pass buffer row stride to 2D scan kernel to fix silent miscomputation (#2620) 6171343c [TIR][Language] Add typed vector lane extraction API (#2789) 1545f006 [Metal] M5 Cooperative Tensor T.gemm (#2252) aaf68d2e [Metal] Add 16-byte alignment padding to shared/threadgroup memory (#2786) 940b1061 [BugFix] Honor nan_propagate in reduce max/min/absmax clear=False write-back (#2788) a42bbc3c [TIR] Inject source spans into tirx IR and surface source locations in compiler errors (#2751) 500c3686 [CUDA][Transform] Fix PCWS index dtype handling (#2783) eb31994a [BugFix] Preserve loop step when unrolling loops (#2784) 51fbfc7e [BugFix] Fallback non-16B cluster bulk copies (#2683) b5e3eb93 [Enhancement] More compile-time guards for architecture-specific CUDA intrinsics (#2781) 2a17fffd [TileOP] Add SM70 GEMM FMA fallback (#2339) aa7df867 [BugFix] Fix fp16/bf16 T.atomic_max/atomic_min silently corrupting fp32 values (#2780) 9fb75728 [BugFix] Fix ROCm intrinsic resolution after language dialect refactor (#2779) 92072ab2 [BugFix] Cast to the destination dtype in the scalar T.copy path (#2771) a69708c1 [BugFix] Reject non-power-of-two AllReduce widths (#2611) 30aac1ac [CUDA][Reduce] Fix packed AllReduce workspace pointer (#2778) fc517bdc [BugFix] Fix rng_init after language dialect refactor (#2776) eceb0e66 [BugFix] Marshal NVRTC scalar parameters and dynamic strides (#2756) 5ef1500e [BugFix] Handle strided global buffers correctly in 1D TMA copies (#2746) 1dc86d71 [BugFix] Add arithmetic operators to vec_type in common.h for CPU codegen (#2768) 22a2452a [BugFix] Add threadgroup address space qualifier in Metal codegen for shared memory pointer arithmetic (#2770) 9609d3a5 [BugFix] Preserve explicit loop steps when transforms rebuild For nodes (#2752) b049f87d [CI]: Bump actions/setup-python from 6 to 7 (#2773) 7cb4b1d9 [Enhancement] Fix nondeterministic CanProve (#2772) c6294f07 [BugFix] Add pre-SM80 fallback for bf16 __hfma (#2769) 8ad82fa0 [Quality] Fixes typings in ast frontend (#2520) 8bb3300d [BugFix] Preserve if condition evaluation during fan-out (#2764) e9240d68 [Refactor] Move example-only helpers out of tilelang package (#2761) 1591d368 [BugFix] Preserve re-evaluation of mutable if conditions (#2744) f862dc38 [Language][Backend] Language dialect for multi-backends (#2734) ab1d2df4 [BugFix] Support BufferRegion destinations in atomic_addx2 return_prev (#2753) 2b4dd803 [BugFix] Make T.transpose swap only the final two axes (#2757) 235077cb [BugFix] Preserve loop steps during unswitching (#2741) 390d208d [BugFix] Use callee global symbols for cross-target calls (#2740) 192ddea6 [BugFix] Gate TMEM and TMA builtins by CUDA architecture (#2743) aae97c0e [TIR][Python] Add typing wrappers for DSL ops (#2739) 0c88682f [BugFix] Emit Metal barriers for dynamic shared memory (#2738) bff1b9a3 [Feature] Support stochastic FP32 to FP16/BF16 casts (#2735) 22baf2e2 [FEATURE] Add block-causal attention for dLLM example (#2499) a443dde9 [BugFix] Reject contracting shared-buffer layouts in T.annotate_layout (#2719) dff136d4 [CUDA][Pipeline] Fix 1D TMA selection for versioned layouts (#2737) f84825db [Build] Raise apache-tvm-ffi lower bound to 0.1.11 (#2736) 512d51f5 [BugFix] Honor explicit row strides in Metal GEMM (#2730) 322a9cbb [BugFix] Isolate cross-compiler options per invocation (#2728) e4e110e5 [Example][DeepSeek-V3.2] Adaptive threads for sparse MLA backward (#2592) 24a023c6 [BugFix][Transform] Never classify side-effecting binds as replayable (atomics re-executed at every use site since v0.1.11) (#2651) 1ea7530f [Cherry][TIRx] Improve BufferStore cast warning context (#2733) 052e6741 [TIR][Analyzer] Fix vectorized Select constraint handling (#2731) 9d819c3f [BugFix] Fix MFMA DataType args causing compilation failure on ROCm (#2726) cc106fa2 [Feature] Add lower-trace support for debugging & rebased (#2725) 96900c7d [Autotune] Support early_stop in decorator mode and add decorator example (#2729) bd5ca2f0 [BugFix] Preserve buffer element offsets in access pointers (#2727) 923c8a7d [Autotune] Add early stop to skip slow configs during benchmark (#2723) f8d8cd4b [BugFix] Use blockDim as workspace stride in batch AllReduce (#2621) ac576c63 [BugFix] Fix FP4 dequant symbolic exponent clamp (#2656) 25c0a155 [Transform] Replace CPU fallback thread placeholder with a constant-zero logical thread index (#2718) c4c5ec59 [Example][Opt] deepseek_v32 topk_selector kernel memory access optimization (~1.9× faster) (#2659) 8cfc90e0 [BugFix] Fix HIP predicated dword copy zero fill (#2721) 88e007d4 [BugFix] Guard T.atomic_addx4 return type for sliced destinations (#2590) 656c287a [BugFix] Fix wrong offset when scanning a non-zero-offset buffer sub-region (#2680) ce9ff0c1 [BugFix] Gate stochastic FP4/FP8 casts on sm_100a (#2691) 39c5b4ef [BugFix] Correct IEEE math intrinsic names for fp64/fp16/bf16 (#2619) cb26539a [BugFix] Fix T.pow/T.power for constant integer exponent y <= 0 (#2677) ffeda9f3 [BugFix] Pack 32-lane 8-bit CUDA vectors correctly (#2701) 30221e20 [Bugfix] Emit st.bulk destination as a shared write to fix missing barrier and potential compilation-introduced races (#2700) 172f6fbf [TIR][Transform] Allow unused fragment buffers without layouts (#2717) 8cdd4d62 [BugFix] Prevent partial 1-D TMA stores from bypassing bounds checks (#2716) 46f3b31a [BugFix][WS] Fix pipeline replacement under persistent T.serial (#2674) 4433981c [Enhancement] Add local buffer reduction lowering (#2693) 335afcf8 [BugFix] Decode FP8 E4M3 special encodings correctly (#2710) e3c3048f [BugFix] Correct the spelling of fragment (#2695) 134f9c2e [BugFix] Implement atomic load and store for HIP (#2711) bfe126d1 [BugFix] Support return_prev for HIP vector atomic add (#2712) 1354b610 [BugFix] Map CUDA atomic add consume ordering to acquire PTX (#2713) c5c10d26 [Fix] T.__exp must compute e**x, not 2**x (docstring + CuTeDSL codegen) (#2696) a1e3aee7 [BugFix] Fix operator precedence in increase_descriptor_offset guard (#2675) 6f70e5c9 [BugFix] reject float dtype in bitwise reduce with an actionable error (#2676) 0c7f14c3 [BugFix] Fix WGMMA C-store layout for multiple warpgroups along M (#2663) 4aea435b [BugFix] Reject unsafe non-warp-multiple partial thread sync (#2679) 9754ac44 [BugFix] Preserve bf16 NaN and Inf during RNE packing (#2690) 61f968b6 [BugFix] Track GEMM accumulator writes in warp-specialization liveness collector (#2685) 4a5cf099 [Feature] Add compiler pass timing profiling (#2622) 9ff4ef8d [BugFix] Respect safe value in tile copy OOB (#2636) 31755e71 [BugFix] Support runtime-dependent vector negative indices (#2654) 66837430 [JIT][TVM-FFI] Reuse executable across kernel launches (#2686) 4f442a39 [BugFix] Reject invalid atomic load and store memory orders (#2666) c5e53076 [BugFix] Return previous value for scalar atomic_min/atomic_max (#2672) 51e6c69f [BugFix] Zero-extend signed int32 lanes in 256-bit vector pack (#2673) d36ec37c [BugFix] Support T.infinity for float8_e5m2 (#2671) 3af87b7f [TIR][Transform] Fix flat Bind modeling in parallel race checks (#2665) 8164c9a0 [CUDA][Scan] Enable pipelining for multi-segment scans (#2664) 70548a17 [BugFix][CUDA][HIP] Correct packed shared memory allocation sizes (#2660) effebb6b [BugFix] Fix segfault when tunable params default to None (#2657) 917cd2b9 Bump transformers from 5.3.0 to 5.5.0 in /examples/bitnet-1.58b (#2658) 1ac5a01a [TIR][Runtime] Enforce host-evaluable assumptions at runtime (#2655) 6ed02aee [Autotune] Support per-config pass_configs in autotuning (#2496) d1ccb925 [BugFix] Fix 1D bulk TMA transfer alignment check (#2646) 207f3a75 [BugFix] Reject non-positive thread extents in T.Kernel (#2653) c8c49d50 [BugFix] Reject invalid dtypes in T.dp4a (#2652) 28101de3 [Arith] Gate canonical-simplify LT Case 2 on extra scale == +1 (#2649) 6c09e889 fix: cast ptxas register usage level to int before building the nvcc command (#2641) 2047357e [Reduce][Codegen] Fix thread-segment projection for packed layouts (#2647) 5dfa0f5d [BugFix] Fix grouped reduce_sum over-counts on straddle layout (#2424) 1e075149 [CUDA] Support fp32x2 ops as reducers (#2637) 2f4d0fe7 [BugFix] Fix metal stream bridge (#2639) 228c7c04 [Feature] Support iket profiler for CUDA backend (#2515) bd764738 [BugFix] Preserve guard identity in LoopUnswitching (#2585) 3b37333c [Refactor] Extract shared Int64Promoter into common header (#2558) 6611eaec [Doc] Update SKILL.md to support editable installs and clarify development workflow (#2533) 250c1fc9 [TIR][Codegen] Preserve decoupled cast buffer scope (#2545) 8533d2a9 [Language][Scheduler] Expose scalar tile scheduler state (#2553)What's Changed
- [Language][Scheduler] Expose scalar tile scheduler state by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2553
- [TIR][Codegen] Preserve decoupled cast buffer scope by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2545
- [Doc] Update SKILL.md to support editable installs and clarify develo… by @erhsh in https://github.com/tile-ai/tilelang/pull/2533
- [Refactor] Extract shared Int64Promoter into common header by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2558
- [BugFix] Preserve guard identity in LoopUnswitching by @zyy3077 in https://github.com/tile-ai/tilelang/pull/2585
- [Feature] Support iket profiler for CUDA backend by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2515
- [BugFix] Fix metal stream bridge by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2639
- [CUDA] Support fp32x2 ops as reducers by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2637
- [BugFix] Fix grouped reduce_sum over-counts on straddle layout by @hhy3 in https://github.com/tile-ai/tilelang/pull/2424
- [Reduce][Codegen] Fix thread-segment projection for packed layouts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2647
- fix: cast ptxas register usage level to int before building the nvcc command by @gvr13n in https://github.com/tile-ai/tilelang/pull/2641
- [Arith] Gate canonical-simplify LT Case 2 on extra scale == +1 by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2649
- [BugFix] Reject invalid dtypes in T.dp4a by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2652
- [BugFix] Reject non-positive thread extents in T.Kernel by @Lfan-ke in https://github.com/tile-ai/tilelang/pull/2653
- [BugFix] Fix 1D bulk TMA transfer alignment check by @UnsettingGalaxy in https://github.com/tile-ai/tilelang/pull/2646
- [Autotune] Support per-config pass_configs in autotuning by @Da1L8-X in https://github.com/tile-ai/tilelang/pull/2496
- [TIR][Runtime] Enforce host-evaluable assumptions at runtime by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2655
- Bump transformers from 5.3.0 to 5.5.0 in /examples/bitnet-1.58b by @dependabot[bot] in https://github.com/tile-ai/tilelang/pull/2658
- [BugFix] Fix segfault when tunable params default to None by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2657
- [BugFix][CUDA][HIP] Correct packed shared memory allocation sizes by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2660
- [CUDA][Scan] Enable pipelining for multi-segment scans by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2664
- [TIR][Transform] Fix flat Bind modeling in parallel race checks by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2665
- [BugFix] Support T.infinity for float8_e5m2 by @Hughshine in https://github.com/tile-ai/tilelang/pull/2671
- [BugFix] Zero-extend signed int32 lanes in 256-bit vector pack by @ColmaLiu in https://github.com/tile-ai/tilelang/pull/2673
- [BugFix] Return previous value for scalar atomic_min/atomic_max by @Hughshine in https://github.com/tile-ai/tilelang/pull/2672
- [BugFix] Reject invalid atomic load and store memory orders by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2666
- [JIT][TVM-FFI] Reuse executable across kernel launches by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2686
- [BugFix] Support runtime-dependent vector negative indices by @Lyscoria in https://github.com/tile-ai/tilelang/pull/2654
- [BugFix] Respect safe value in tile copy OOB by @zyy3077 in https://github.com/tile-ai/tilelang/pull/2636
- [Feature] Add compiler pass timing profiling by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2622
- [BugFix] Track GEMM accumulator writes in warp-specialization liveness collector by @zyy3077 in https://github.com/tile-ai/tilelang/pull/2685
- [BugFix] Preserve bf16 NaN and Inf during RNE packing by @xiaoyouPREG in https://github.com/tile-ai/tilelang/pull/2690
- [BugFix] Reject unsafe non-warp-multiple partial thread sync by @Lyscoria in https://github.com/tile-ai/tilelang/pull/2679
- [BugFix] Fix WGMMA C-store layout for multiple warpgroups along M by @ColmaLiu in https://github.com/tile-ai/tilelang/pull/2663
- [BugFix] reject float dtype in bitwise reduce with an actionable error by @Hughshine in https://github.com/tile-ai/tilelang/pull/2676
- [BugFix] Fix operator precedence in increase_descriptor_offset guard by @Hughshine in https://github.com/tile-ai/tilelang/pull/2675
- [Fix]
T.__expmust compute ex, not 2x (docstring + CuTeDSL codegen) by @Hughshine in https://github.com/tile-ai/tilelang/pull/2696 - [BugFix] Map CUDA atomic add consume ordering to acquire PTX by @morluto in https://github.com/tile-ai/tilelang/pull/2713
- [BugFix] Support return_prev for HIP vector atomic add by @morluto in https://github.com/tile-ai/tilelang/pull/2712
- [BugFix] Implement atomic load and store for HIP by @morluto in https://github.com/tile-ai/tilelang/pull/2711
- [BugFix] Correct the spelling of fragment by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2695
- [BugFix] Decode FP8 E4M3 special encodings correctly by @morluto in https://github.com/tile-ai/tilelang/pull/2710
- [Enhancement] Add local buffer reduction lowering by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2693
- [BugFix][WS] Fix pipeline replacement under persistent T.serial by @RuneFang in https://github.com/tile-ai/tilelang/pull/2674
- [BugFix] Prevent partial 1-D TMA stores from bypassing bounds checks by @UnsettingGalaxy in https://github.com/tile-ai/tilelang/pull/2716
- [TIR][Transform] Allow unused fragment buffers without layouts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2717
- [Bugfix] Emit st.bulk destination as a shared write to fix missing barrier and potential compilation-introduced races by @arxgy in https://github.com/tile-ai/tilelang/pull/2700
- [BugFix] Pack 32-lane 8-bit CUDA vectors correctly by @xiaoyouPREG in https://github.com/tile-ai/tilelang/pull/2701
- [BugFix] Fix T.pow/T.power for constant integer exponent y <= 0 by @Hughshine in https://github.com/tile-ai/tilelang/pull/2677
- [BugFix] Correct IEEE math intrinsic names for fp64/fp16/bf16 by @ColmaLiu in https://github.com/tile-ai/tilelang/pull/2619
- [BugFix] Gate stochastic FP4/FP8 casts on sm_100a by @xiaoyouPREG in https://github.com/tile-ai/tilelang/pull/2691
- [BugFix] Fix wrong offset when scanning a non-zero-offset buffer sub-region by @li-ruinan in https://github.com/tile-ai/tilelang/pull/2680
- [BugFix] Guard T.atomic_addx4 return type for sliced destinations by @Lyscoria in https://github.com/tile-ai/tilelang/pull/2590
- [BugFix] Fix HIP predicated dword copy zero fill by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2721
- [Example][Opt] deepseek_v32 topk_selector kernel memory access optimization (~1.9× faster) by @mengmeexix in https://github.com/tile-ai/tilelang/pull/2659
- [Transform] Replace CPU fallback thread placeholder with a constant-zero logical thread index by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2718
- [BugFix] Fix FP4 dequant symbolic exponent clamp by @mygitljf in https://github.com/tile-ai/tilelang/pull/2656
- [BugFix] Use blockDim as workspace stride in batch AllReduce by @ColmaLiu in https://github.com/tile-ai/tilelang/pull/2621
- [Autotune] Add early stop to skip slow configs during benchmark by @Da1L8-X in https://github.com/tile-ai/tilelang/pull/2723
- [BugFix] Preserve buffer element offsets in access pointers by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2727
- [Autotune] Support early_stop in decorator mode and add decorator example by @Da1L8-X in https://github.com/tile-ai/tilelang/pull/2729
- [Feature] Add lower-trace support for debugging & rebased by @erhsh in https://github.com/tile-ai/tilelang/pull/2725
- [BugFix] Fix MFMA DataType args causing compilation failure on ROCm by @jayzlee147 in https://github.com/tile-ai/tilelang/pull/2726
- [TIR][Analyzer] Fix vectorized Select constraint handling by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2731
- [Cherry][TIRx] Improve BufferStore cast warning context by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2733
- [BugFix][Transform] Never classify side-effecting binds as replayable (atomics re-executed at every use site since v0.1.11) by @zkyue in https://github.com/tile-ai/tilelang/pull/2651
- [Example][DeepSeek-V3.2] Adaptive threads for sparse MLA backward by @Butterfingrz in https://github.com/tile-ai/tilelang/pull/2592
- [BugFix] Isolate cross-compiler options per invocation by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2728
- [BugFix] Honor explicit row strides in Metal GEMM by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2730
- [Build] Raise apache-tvm-ffi lower bound to 0.1.11 by @hhy3 in https://github.com/tile-ai/tilelang/pull/2736
- [CUDA][Pipeline] Fix 1D TMA selection for versioned layouts by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2737
- [BugFix] Reject contracting shared-buffer layouts in T.annotate_layout by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2719
- [FEATURE] Add block-causal attention for dLLM example by @perkyfever in https://github.com/tile-ai/tilelang/pull/2499
- [Feature] Support stochastic FP32 to FP16/BF16 casts by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2735
- [BugFix] Emit Metal barriers for dynamic shared memory by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2738
- [TIR][Python] Add typing wrappers for DSL ops by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2739
- [BugFix] Gate TMEM and TMA builtins by CUDA architecture by @xiaoyouPREG in https://github.com/tile-ai/tilelang/pull/2743
- [BugFix] Use callee global symbols for cross-target calls by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2740
- [BugFix] Preserve loop steps during unswitching by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2741
- [BugFix] Make T.transpose swap only the final two axes by @morluto in https://github.com/tile-ai/tilelang/pull/2757
- [BugFix] Support BufferRegion destinations in atomic_addx2 return_prev by @morluto in https://github.com/tile-ai/tilelang/pull/2753
- [Language][Backend] Language dialect for multi-backends by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2734
- [BugFix] Preserve re-evaluation of mutable if conditions by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2744
- [Refactor] Move example-only helpers out of tilelang package by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2761
- [BugFix] Preserve if condition evaluation during fan-out by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2764
- [Quality] Fixes typings in ast frontend by @ppppqp in https://github.com/tile-ai/tilelang/pull/2520
- [BugFix] Add pre-SM80 fallback for bf16 __hfma by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2769
- [Enhancement] Fix nondeterministic CanProve by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2772
- [CI]: Bump actions/setup-python from 6 to 7 by @dependabot[bot] in https://github.com/tile-ai/tilelang/pull/2773
- [BugFix] Preserve explicit loop steps when transforms rebuild For nodes by @morluto in https://github.com/tile-ai/tilelang/pull/2752
- [BugFix] Add threadgroup address space qualifier in Metal codegen for shared memory pointer arithmetic by @GY-Bai in https://github.com/tile-ai/tilelang/pull/2770
- [BugFix] Add arithmetic operators to vec_type in common.h for CPU codegen by @GY-Bai in https://github.com/tile-ai/tilelang/pull/2768
- [BugFix] Handle strided global buffers correctly in 1D TMA copies by @cla7aye15I4nd in https://github.com/tile-ai/tilelang/pull/2746
- [BugFix] Marshal NVRTC scalar parameters and dynamic strides by @morluto in https://github.com/tile-ai/tilelang/pull/2756
- [BugFix] Fix rng_init after language dialect refactor by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2776
- [CUDA][Reduce] Fix packed AllReduce workspace pointer by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2778
- [BugFix] Reject non-power-of-two AllReduce widths by @zyy3077 in https://github.com/tile-ai/tilelang/pull/2611
- [BugFix] Cast to the destination dtype in the scalar T.copy path by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2771
- [BugFix] Fix ROCm intrinsic resolution after language dialect refactor by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2779
- [BugFix] Fix fp16/bf16
T.atomic_max/atomic_minsilently corrupting fp32 values by @jjppp in https://github.com/tile-ai/tilelang/pull/2780 - [TileOP] Add SM70 GEMM FMA fallback by @cklxx in https://github.com/tile-ai/tilelang/pull/2339
- [Enhancement] More compile-time guards for architecture-specific CUDA intrinsics by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2781
- [BugFix] Fallback non-16B cluster bulk copies by @UnsettingGalaxy in https://github.com/tile-ai/tilelang/pull/2683
- [BugFix] Preserve loop step when unrolling loops by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2784
- [CUDA][Transform] Fix PCWS index dtype handling by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2783
- [TIR] Inject source spans into tirx IR and surface source locations in compiler errors by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2751
- [BugFix] Honor nan_propagate in reduce max/min/absmax clear=False write-back by @bhaochen in https://github.com/tile-ai/tilelang/pull/2788
- [Metal] Add 16-byte alignment padding to shared/threadgroup memory by @GY-Bai in https://github.com/tile-ai/tilelang/pull/2786
- [Metal] M5 Cooperative Tensor T.gemm by @oraluben in https://github.com/tile-ai/tilelang/pull/2252
- [TIR][Language] Add typed vector lane extraction API by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2789
- [BugFix] Pass buffer row stride to 2D scan kernel to fix silent miscomputation by @ColmaLiu in https://github.com/tile-ai/tilelang/pull/2620
- [FFI] Support apache-tvm-ffi 0.1.12 by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2795
- [CUDA] Support arbitrary TMEM layouts by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2785
- [BugFix] Preserve alloc_var initializer dtype by @erhsh in https://github.com/tile-ai/tilelang/pull/2801
- [BugFix] Reject floating-point predicates in vote intrinsics by @erhsh in https://github.com/tile-ai/tilelang/pull/2800
- [Metal] Add line-level threadgroup qualifier scanning (pass 5) by @GY-Bai in https://github.com/tile-ai/tilelang/pull/2796
- [BugFix] Check shared-TMEM buffer pointer types before dereference by @morluto in https://github.com/tile-ai/tilelang/pull/2794
- [BugFix] Prevent autotuner cache reuse across different outputs and validation settings by @morluto in https://github.com/tile-ai/tilelang/pull/2793
- fix: refine TCGEN05 architecture guards by @Rachmanino in https://github.com/tile-ai/tilelang/pull/2790
- [BugFix] Reject mixed packed x2 operand dtypes by @erhsh in https://github.com/tile-ai/tilelang/pull/2802
- [BugFix] Fix warp_reduce truncating int64/uint64 to 32 bits on sm_80+ by @jjppp in https://github.com/tile-ai/tilelang/pull/2782
- [BugFix] Skip descriptor TMA for device-bound copy bases by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2803
- [Refactor] Remove unused tilelang.common package by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2810
- [CUDA] Extend the GEMM FMA fallback to SM75 by @Chennesxu in https://github.com/tile-ai/tilelang/pull/2811
- [BugFix] Reject uncovered warp partitions in T.gemm instead of producing silently wrong results by @li-ruinan in https://github.com/tile-ai/tilelang/pull/2724
- [BugFix] Resolve partial scalar reduce barrier participation by @KellyFrog in https://github.com/tile-ai/tilelang/pull/2777
- [BugFix] Reject unsupported fast-math input dtypes by @erhsh in https://github.com/tile-ai/tilelang/pull/2804
- [Refactor] Remove intrinsic compatibility facade by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2812
- [CUDA][Reduce] Simplify scalar AllReduce thread range analysis by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2814
- [Carver] Remove unused shape inference module by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2813
- [CUDA] Add SM120 NVF4 block-scale MMA support by @qqq-tao in https://github.com/tile-ai/tilelang/pull/2364
- [BugFix] Fix two-instance modeling in ThreadSync cross-thread race checks by @LJC00118 in https://github.com/tile-ai/tilelang/pull/2805
- [JIT] Remove legacy DLPack execution backend by @SiriusNEO in https://github.com/tile-ai/tilelang/pull/2816
- [Enhancement] Aggregate VerifyParallelLoop race diagnostics with span by @penguin-wwy in https://github.com/tile-ai/tilelang/pull/2806
- [BugFix] Reject unsupported TMA atomic add dtypes by @morluto in https://github.com/tile-ai/tilelang/pull/2830
- Revert "[BugFix] Preserve loop step when unrolling loops" by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2834
- [BugFix] Correctly preserve loop step when unrolling loops by @Yongqi-Zhuo in https://github.com/tile-ai/tilelang/pull/2835
- [Release] Bump versin into 0.1.13 by @LeiWang1999 in https://github.com/tile-ai/tilelang/pull/2826
New Contributors
- @zyy3077 made their first contribution in https://github.com/tile-ai/tilelang/pull/2585
- @hhy3 made their first contribution in https://github.com/tile-ai/tilelang/pull/2424
- @gvr13n made their first contribution in https://github.com/tile-ai/tilelang/pull/2641
- @Lfan-ke made their first contribution in https://github.com/tile-ai/tilelang/pull/2653
- @UnsettingGalaxy made their first contribution in https://github.com/tile-ai/tilelang/pull/2646
- @Da1L8-X made their first contribution in https://github.com/tile-ai/tilelang/pull/2496
- @Lyscoria made their first contribution in https://github.com/tile-ai/tilelang/pull/2654
- @xiaoyouPREG made their first contribution in https://github.com/tile-ai/tilelang/pull/2690
- @morluto made their first contribution in https://github.com/tile-ai/tilelang/pull/2713
- @arxgy made their first contribution in https://github.com/tile-ai/tilelang/pull/2700
- @li-ruinan made their first contribution in https://github.com/tile-ai/tilelang/pull/2680
- @cla7aye15I4nd made their first contribution in https://github.com/tile-ai/tilelang/pull/2721
- @mygitljf made their first contribution in https://github.com/tile-ai/tilelang/pull/2656
- @jayzlee147 made their first contribution in https://github.com/tile-ai/tilelang/pull/2726
- @zkyue made their first contribution in https://github.com/tile-ai/tilelang/pull/2651
- @Butterfingrz made their first contribution in https://github.com/tile-ai/tilelang/pull/2592
- @perkyfever made their first contribution in https://github.com/tile-ai/tilelang/pull/2499
- @GY-Bai made their first contribution in https://github.com/tile-ai/tilelang/pull/2770
- @bhaochen made their first contribution in https://github.com/tile-ai/tilelang/pull/2788
- @KellyFrog made their first contribution in https://github.com/tile-ai/tilelang/pull/2777
Full Changelog: https://github.com/tile-ai/tilelang/compare/v0.1.12...v0.1.13