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_blockscaled 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 _activate_cuda_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).
warp_reduce no 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_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.atomic_max/T.atomic_min no longer silently corrupt fp32 values (#2780).
return_prev supported for scalar atomic_min/atomic_max (#2672), atomic_addx2 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 (e0f0ac9) 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 e**x, 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 (#2649); vectorized Select constraint 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.bulk destination 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_offset guard (#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.transpose swaps only the final two axes (#2757); contracting shared-buffer layouts rejected in T.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_type in common.h for CPU codegen (#2768).

Race analysis & warp-specialization

Two-instance modeling fixed in ThreadSync cross-thread race checks (#2805).
Flat Bind modeling 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.gemm rejected instead of silently producing wrong results (#2724).
MFMA DataType args 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 fragment spelling corrected (#2695); BufferStore cast warning context improved (#2733).

Autotuning

Early stop to skip slow configs during benchmarking (#2723), including decorator mode with an example (#2729).
Per-config pass_configs supported 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_selector memory-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 #2553
[TIR][Codegen] Preserve decoupled cast buffer scope by @LeiWang1999 in #2545
[Doc] Update SKILL.md to support editable installs and clarify develo… by @erhsh in #2533
[Refactor] Extract shared Int64Promoter into common header by @penguin-wwy in #2558
[BugFix] Preserve guard identity in LoopUnswitching by @zyy3077 in #2585
[Feature] Support iket profiler for CUDA backend by @Rachmanino in #2515
[BugFix] Fix metal stream bridge by @penguin-wwy in #2639
[CUDA] Support fp32x2 ops as reducers by @Yongqi-Zhuo in #2637
[BugFix] Fix grouped reduce_sum over-counts on straddle layout by @hhy3 in #2424
[Reduce][Codegen] Fix thread-segment projection for packed layouts by @LeiWang1999 in #2647
fix: cast ptxas register usage level to int before building the nvcc command by @gvr13n in #2641
[Arith] Gate canonical-simplify LT Case 2 on extra scale == +1 by @LeiWang1999 in #2649
[BugFix] Reject invalid dtypes in T.dp4a by @Chennesxu in #2652
[BugFix] Reject non-positive thread extents in T.Kernel by @Lfan-ke in #2653
[BugFix] Fix 1D bulk TMA transfer alignment check by @UnsettingGalaxy in #2646
[Autotune] Support per-config pass_configs in autotuning by @Da1L8-X in #2496
[TIR][Runtime] Enforce host-evaluable assumptions at runtime by @LeiWang1999 in #2655
Bump transformers from 5.3.0 to 5.5.0 in /examples/bitnet-1.58b by @dependabot[bot] in #2658
[BugFix] Fix segfault when tunable params default to None by @penguin-wwy in #2657
[BugFix][CUDA][HIP] Correct packed shared memory allocation sizes by @LeiWang1999 in #2660
[CUDA][Scan] Enable pipelining for multi-segment scans by @LeiWang1999 in #2664
[TIR][Transform] Fix flat Bind modeling in parallel race checks by @LeiWang1999 in #2665
[BugFix] Support T.infinity for float8_e5m2 by @Hughshine in #2671
[BugFix] Zero-extend signed int32 lanes in 256-bit vector pack by @ColmaLiu in #2673
[BugFix] Return previous value for scalar atomic_min/atomic_max by @Hughshine in #2672
[BugFix] Reject invalid atomic load and store memory orders by @Chennesxu in #2666
[JIT][TVM-FFI] Reuse executable across kernel launches by @LeiWang1999 in #2686
[BugFix] Support runtime-dependent vector negative indices by @Lyscoria in #2654
[BugFix] Respect safe value in tile copy OOB by @zyy3077 in #2636
[Feature] Add compiler pass timing profiling by @penguin-wwy in #2622
[BugFix] Track GEMM accumulator writes in warp-specialization liveness collector by @zyy3077 in #2685
[BugFix] Preserve bf16 NaN and Inf during RNE packing by @xiaoyouPREG in #2690
[BugFix] Reject unsafe non-warp-multiple partial thread sync by @Lyscoria in #2679
[BugFix] Fix WGMMA C-store layout for multiple warpgroups along M by @ColmaLiu in #2663
[BugFix] reject float dtype in bitwise reduce with an actionable error by @Hughshine in #2676
[BugFix] Fix operator precedence in increase_descriptor_offset guard by @Hughshine in #2675
[Fix] T.__exp must compute ex, not 2x (docstring + CuTeDSL codegen) by @Hughshine in #2696
[BugFix] Map CUDA atomic add consume ordering to acquire PTX by @morluto in #2713
[BugFix] Support return_prev for HIP vector atomic add by @morluto in #2712
[BugFix] Implement atomic load and store for HIP by @morluto in #2711
[BugFix] Correct the spelling of fragment by @penguin-wwy in #2695
[BugFix] Decode FP8 E4M3 special encodings correctly by @morluto in #2710
[Enhancement] Add local buffer reduction lowering by @LJC00118 in #2693
[BugFix][WS] Fix pipeline replacement under persistent T.serial by @RuneFang in #2674
[BugFix] Prevent partial 1-D TMA stores from bypassing bounds checks by @UnsettingGalaxy in #2716
[TIR][Transform] Allow unused fragment buffers without layouts by @LeiWang1999 in #2717
[Bugfix] Emit st.bulk destination as a shared write to fix missing barrier and potential compilation-introduced races by @arxgy in #2700
[BugFix] Pack 32-lane 8-bit CUDA vectors correctly by @xiaoyouPREG in #2701
[BugFix] Fix T.pow/T.power for constant integer exponent y <= 0 by @Hughshine in #2677
[BugFix] Correct IEEE math intrinsic names for fp64/fp16/bf16 by @ColmaLiu in #2619
[BugFix] Gate stochastic FP4/FP8 casts on sm_100a by @xiaoyouPREG in #2691
[BugFix] Fix wrong offset when scanning a non-zero-offset buffer sub-region by @li-ruinan in #2680
[BugFix] Guard T.atomic_addx4 return type for sliced destinations by @Lyscoria in #2590
[BugFix] Fix HIP predicated dword copy zero fill by @cla7aye15I4nd in #2721
[Example][Opt] deepseek_v32 topk_selector kernel memory access optimization (~1.9× faster) by @mengmeexix in #2659
[Transform] Replace CPU fallback thread placeholder with a constant-zero logical thread index by @penguin-wwy in #2718
[BugFix] Fix FP4 dequant symbolic exponent clamp by @mygitljf in #2656
[BugFix] Use blockDim as workspace stride in batch AllReduce by @ColmaLiu in #2621
[Autotune] Add early stop to skip slow configs during benchmark by @Da1L8-X in #2723
[BugFix] Preserve buffer element offsets in access pointers by @cla7aye15I4nd in #2727
[Autotune] Support early_stop in decorator mode and add decorator example by @Da1L8-X in #2729
[Feature] Add lower-trace support for debugging & rebased by @erhsh in #2725
[BugFix] Fix MFMA DataType args causing compilation failure on ROCm by @jayzlee147 in #2726
[TIR][Analyzer] Fix vectorized Select constraint handling by @LeiWang1999 in #2731
[Cherry][TIRx] Improve BufferStore cast warning context by @LeiWang1999 in #2733
[BugFix][Transform] Never classify side-effecting binds as replayable (atomics re-executed at every use site since v0.1.11) by @zkyue in #2651
[Example][DeepSeek-V3.2] Adaptive threads for sparse MLA backward by @Butterfingrz in #2592
[BugFix] Isolate cross-compiler options per invocation by @cla7aye15I4nd in #2728
[BugFix] Honor explicit row strides in Metal GEMM by @cla7aye15I4nd in #2730
[Build] Raise apache-tvm-ffi lower bound to 0.1.11 by @hhy3 in #2736
[CUDA][Pipeline] Fix 1D TMA selection for versioned layouts by @LeiWang1999 in #2737
[BugFix] Reject contracting shared-buffer layouts in T.annotate_layout by @Chennesxu in #2719
[FEATURE] Add block-causal attention for dLLM example by @perkyfever in #2499
[Feature] Support stochastic FP32 to FP16/BF16 casts by @LJC00118 in #2735
[BugFix] Emit Metal barriers for dynamic shared memory by @cla7aye15I4nd in #2738
[TIR][Python] Add typing wrappers for DSL ops by @LeiWang1999 in #2739
[BugFix] Gate TMEM and TMA builtins by CUDA architecture by @xiaoyouPREG in #2743
[BugFix] Use callee global symbols for cross-target calls by @cla7aye15I4nd in #2740
[BugFix] Preserve loop steps during unswitching by @cla7aye15I4nd in #2741
[BugFix] Make T.transpose swap only the final two axes by @morluto in #2757
[BugFix] Support BufferRegion destinations in atomic_addx2 return_prev by @morluto in #2753
[Language][Backend] Language dialect for multi-backends by @LeiWang1999 in #2734
[BugFix] Preserve re-evaluation of mutable if conditions by @cla7aye15I4nd in #2744
[Refactor] Move example-only helpers out of tilelang package by @SiriusNEO in #2761
[BugFix] Preserve if condition evaluation during fan-out by @SiriusNEO in #2764
[Quality] Fixes typings in ast frontend by @ppppqp in #2520
[BugFix] Add pre-SM80 fallback for bf16 __hfma by @Chennesxu in #2769
[Enhancement] Fix nondeterministic CanProve by @LJC00118 in #2772
[CI]: Bump actions/setup-python from 6 to 7 by @dependabot[bot] in #2773
[BugFix] Preserve explicit loop steps when transforms rebuild For nodes by @morluto in #2752
[BugFix] Add threadgroup address space qualifier in Metal codegen for shared memory pointer arithmetic by @GY-Bai in #2770
[BugFix] Add arithmetic operators to vec_type in common.h for CPU codegen by @GY-Bai in #2768
[BugFix] Handle strided global buffers correctly in 1D TMA copies by @cla7aye15I4nd in #2746
[BugFix] Marshal NVRTC scalar parameters and dynamic strides by @morluto in #2756
[BugFix] Fix rng_init after language dialect refactor by @LJC00118 in #2776
[CUDA][Reduce] Fix packed AllReduce workspace pointer by @LeiWang1999 in #2778
[BugFix] Reject non-power-of-two AllReduce widths by @zyy3077 in #2611
[BugFix] Cast to the destination dtype in the scalar T.copy path by @Chennesxu in #2771
[BugFix] Fix ROCm intrinsic resolution after language dialect refactor by @SiriusNEO in #2779
[BugFix] Fix fp16/bf16 T.atomic_max/atomic_min silently corrupting fp32 values by @jjppp in #2780
[TileOP] Add SM70 GEMM FMA fallback by @cklxx in #2339
[Enhancement] More compile-time guards for architecture-specific CUDA intrinsics by @Rachmanino in #2781
[BugFix] Fallback non-16B cluster bulk copies by @UnsettingGalaxy in #2683
[BugFix] Preserve loop step when unrolling loops by @penguin-wwy in #2784
[CUDA][Transform] Fix PCWS index dtype handling by @LeiWang1999 in #2783
[TIR] Inject source spans into tirx IR and surface source locations in compiler errors by @penguin-wwy in #2751
[BugFix] Honor nan_propagate in reduce max/min/absmax clear=False write-back by @bhaochen in #2788
[Metal] Add 16-byte alignment padding to shared/threadgroup memory by @GY-Bai in #2786
[Metal] M5 Cooperative Tensor T.gemm by @oraluben in #2252
[TIR][Language] Add typed vector lane extraction API by @LeiWang1999 in #2789
[BugFix] Pass buffer row stride to 2D scan kernel to fix silent miscomputation by @ColmaLiu in #2620
[FFI] Support apache-tvm-ffi 0.1.12 by @SiriusNEO in #2795
[CUDA] Support arbitrary TMEM layouts by @Yongqi-Zhuo in #2785
[BugFix] Preserve alloc_var initializer dtype by @erhsh in #2801
[BugFix] Reject floating-point predicates in vote intrinsics by @erhsh in #2800
[Metal] Add line-level threadgroup qualifier scanning (pass 5) by @GY-Bai in #2796
[BugFix] Check shared-TMEM buffer pointer types before dereference by @morluto in #2794
[BugFix] Prevent autotuner cache reuse across different outputs and validation settings by @morluto in #2793
fix: refine TCGEN05 architecture guards by @Rachmanino in #2790
[BugFix] Reject mixed packed x2 operand dtypes by @erhsh in #2802
[BugFix] Fix warp_reduce truncating int64/uint64 to 32 bits on sm_80+ by @jjppp in #2782
[BugFix] Skip descriptor TMA for device-bound copy bases by @SiriusNEO in #2803
[Refactor] Remove unused tilelang.common package by @SiriusNEO in #2810
[CUDA] Extend the GEMM FMA fallback to SM75 by @Chennesxu in #2811
[BugFix] Reject uncovered warp partitions in T.gemm instead of producing silently wrong results by @li-ruinan in #2724
[BugFix] Resolve partial scalar reduce barrier participation by @KellyFrog in #2777
[BugFix] Reject unsupported fast-math input dtypes by @erhsh in #2804
[Refactor] Remove intrinsic compatibility facade by @SiriusNEO in #2812
[CUDA][Reduce] Simplify scalar AllReduce thread range analysis by @LeiWang1999 in #2814
[Carver] Remove unused shape inference module by @SiriusNEO in #2813
[CUDA] Add SM120 NVF4 block-scale MMA support by @qqq-tao in #2364
[BugFix] Fix two-instance modeling in ThreadSync cross-thread race checks by @LJC00118 in #2805
[JIT] Remove legacy DLPack execution backend by @SiriusNEO in #2816
[Enhancement] Aggregate VerifyParallelLoop race diagnostics with span by @penguin-wwy in #2806
[BugFix] Reject unsupported TMA atomic add dtypes by @morluto in #2830
Revert "[BugFix] Preserve loop step when unrolling loops" by @Yongqi-Zhuo in #2834
[BugFix] Correctly preserve loop step when unrolling loops by @Yongqi-Zhuo in #2835
[Release] Bump versin into 0.1.13 by @LeiWang1999 in #2826

New Contributors

@zyy3077 made their first contribution in #2585
@hhy3 made their first contribution in #2424
@gvr13n made their first contribution in #2641
@Lfan-ke made their first contribution in #2653
@UnsettingGalaxy made their first contribution in #2646
@Da1L8-X made their first contribution in #2496
@Lyscoria made their first contribution in #2654
@xiaoyouPREG made their first contribution in #2690
@morluto made their first contribution in #2713
@arxgy made their first contribution in #2700
@li-ruinan made their first contribution in #2680
@cla7aye15I4nd made their first contribution in #2721
@mygitljf made their first contribution in #2656
@jayzlee147 made their first contribution in #2726
@zkyue made their first contribution in #2651
@Butterfingrz made their first contribution in #2592
@perkyfever made their first contribution in #2499
@GY-Bai made their first contribution in #2770
@bhaochen made their first contribution in #2788
@KellyFrog made their first contribution in #2777

Full Changelog: v0.1.12...v0.1.13

Navigation Menu

Uh oh!

v0.1.13

Choose a tag to compare

Sorry, something went wrong.

Sorry, something went wrong.

Uh oh!

No results found

TileLang v0.1.13

Highlights

New Features

Backend, API & Refactors

Language dialect refactor (#2734)

Removals (breaking)

FFI / JIT / Build

Bug Fixes

Loop & control-flow preservation

Reductions & scans

Atomics & memory ordering

Numerics, vectors & dtypes

TMA / copy / memory layout

Metal backend

Race analysis & warp-specialization

GEMM / misc compiler fixes

Autotuning

Examples

Full commit list

What's Changed

New Contributors

Contributors

Uh oh!