v0.1.13

[LeiWang1999]

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 ([CUDA] Add SM120 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 ([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 ([CUDA] Support arbitrary TMEM layouts #2785) — TMEM buffers are no longer restricted to a fixed set of layouts.
• [Language/Backend] Language dialect for multi-backends ([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 ([TIR] Inject source spans into tirx IR and surface source locations in compiler errors #2751) — source locations are now carried into the TIRX IR and surfaced in compiler error messages.

New Features

• CUDA
  • SM70 GEMM FMA fallback ([TileOP] Add SM70 GEMM FMA fallback #2339) and SM75 extension of the GEMM FMA fallback ([CUDA] Extend the GEMM FMA fallback to SM75 #2811) — T.gemm now works on older architectures instead of erroring out.
  • Pre-SM80 fallback for bf16 __hfma ([BugFix] Add pre-SM80 fallback for bf16 __hfma #2769).
  • Stochastic FP32 → FP16/BF16 casts ([Feature] Support stochastic FP32 to FP16/BF16 casts #2735), with stochastic FP4/FP8 casts gated on sm_100a ([BugFix] Gate stochastic FP4/FP8 casts on sm_100a #2691).
  • Arbitrary TMEM layout support ([CUDA] Support arbitrary TMEM layouts #2785).
  • Pipelining for multi-segment scans ([CUDA][Scan] Enable pipelining for multi-segment scans #2664).
  • fp32x2 ops usable as reducers ([CUDA] Support fp32x2 ops as reducers #2637).
  • IKET profiler support for the CUDA backend ([Feature] Support iket profiler for CUDA backend #2515).
• Metal
  • M5 cooperative-tensor GEMM ([Metal] M5 Cooperative Tensor T.gemm #2252).
  • Line-level threadgroup qualifier scanning for shared memory (pass 5) ([Metal] Add line-level threadgroup qualifier scanning (pass 5) #2796).
  • 16-byte alignment padding for shared/threadgroup memory ([Metal] Add 16-byte alignment padding to shared/threadgroup memory #2786).
• Compiler / IR / Runtime
  • Compiler pass timing profiling via the pass_profile pass-config option (with a configurable threshold) ([Feature] Add compiler pass timing profiling #2622).
  • lower-trace support for debugging (new doc: docs/tools/lower_trace.md) ([Feature] Add lower-trace support for debugging & rebased #2725).
  • Local buffer reduction lowering ([Enhancement] Add local buffer reduction lowering #2693).
  • Typed vector lane extraction API ([TIR][Language] Add typed vector lane extraction API #2789) and typing wrappers for DSL ops ([TIR][Python] Add typing wrappers for DSL ops #2739).
  • Scalar tile scheduler state exposed ([Language][Scheduler] Expose scalar tile scheduler state #2553).
  • Host-evaluable T.assume conditions are now enforced at runtime ([TIR][Runtime] Enforce host-evaluable assumptions at runtime #2655).
  • Deterministic CanProve ([Enhancement] Fix nondeterministic 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 ([JIT] Remove legacy DLPack execution backend #2816).
• Intrinsic compatibility facade removed ([Refactor] Remove intrinsic compatibility facade #2812).
• tilelang.common package removed ([Refactor] Remove unused tilelang.common package #2810).
• Carver shape-inference module removed ([Carver] Remove unused shape inference module #2813).
• Example-only helpers moved out of the tilelang package ([Refactor] Move example-only helpers out of tilelang package #2761).

FFI / JIT / Build

• Support for apache-tvm-ffi 0.1.12, while keeping 0.1.11 compatibility ([FFI] Support apache-tvm-ffi 0.1.12 #2795); lower bound raised to >=0.1.11 ([Build] Raise apache-tvm-ffi lower bound to 0.1.11 #2736).
• JIT now reuses the compiled executable across kernel launches ([JIT][TVM-FFI] Reuse executable across kernel launches #2686).
• NVRTC scalar parameters and dynamic strides are marshaled correctly ([BugFix] Marshal NVRTC scalar parameters and dynamic strides #2756).
• ptxas register-usage level is cast to int before building the nvcc command (fix: cast ptxas register usage level to int before building the nvcc command #2641).
• Cross-compiler options isolated per invocation ([BugFix] Isolate cross-compiler options per invocation #2728).
• Shared Int64Promoter extracted into a common header ([Refactor] Extract shared Int64Promoter into common header #2558).
• CI: actions/setup-python 6 → 7 ([CI]: Bump actions/setup-python from 6 to 7 #2773); transformers bumped in examples/bitnet-1.58b (Bump transformers from 5.3.0 to 5.5.0 in /examples/bitnet-1.58b #2658).
• Docs: SKILL.md updated for editable installs and clarified development workflow ([Doc] Update SKILL.md to support editable installs and clarify develo… #2533).

Bug Fixes

Loop & control-flow preservation

• Loop steps preserved when unrolling loops — a fix ([BugFix] Preserve loop step when unrolling loops #2784) was reverted (Revert "[BugFix] Preserve loop step when unrolling loops" #2834) and then correctly re-landed ([BugFix] Correctly preserve loop step when unrolling loops #2835).
• Explicit loop steps preserved when transforms rebuild For nodes ([BugFix] Preserve explicit loop steps when transforms rebuild For nodes #2752).
• Loop steps preserved during unswitching ([BugFix] Preserve loop steps during unswitching #2741) and guard identity preserved in LoopUnswitching ([BugFix] Preserve guard identity in LoopUnswitching #2585).
• If-condition evaluation preserved during fan-out ([BugFix] Preserve if condition evaluation during fan-out #2764) and re-evaluation of mutable if conditions ([BugFix] Preserve re-evaluation of mutable if conditions #2744).

Reductions & scans

• Scalar AllReduce thread-range analysis simplified; partial scalar reduce barrier participation fixed ([BugFix] Resolve partial scalar reduce barrier participation #2777, [CUDA][Reduce] Simplify scalar AllReduce thread range analysis #2814).
• warp_reduce no longer truncates int64/uint64 to 32 bits on sm_80+ ([BugFix] Fix warp_reduce truncating int64/uint64 to 32 bits on sm_80+ #2782).
• Non-power-of-two AllReduce widths rejected ([BugFix] Reject non-power-of-two AllReduce widths #2611); packed AllReduce workspace pointer fixed ([CUDA][Reduce] Fix packed AllReduce workspace pointer #2778); blockDim used as workspace stride in batch AllReduce ([BugFix] Use blockDim as workspace stride in batch AllReduce #2621).
• 2D scan kernel now receives the buffer row stride, fixing silent miscomputation ([BugFix] Pass buffer row stride to 2D scan kernel to fix silent miscomputation #2620); wrong offset when scanning a non-zero-offset buffer sub-region fixed ([BugFix] Fix wrong offset when scanning a non-zero-offset buffer sub-region #2680).
• Thread-segment projection for packed layouts fixed ([Reduce][Codegen] Fix thread-segment projection for packed layouts #2647); grouped reduce_sum over-counts on straddle layouts fixed ([BugFix] Fix grouped reduce_sum over-counts on straddle layout #2424).
• nan_propagate honored in reduce max/min/absmax clear=False write-back ([BugFix] Honor nan_propagate in reduce max/min/absmax clear=False write-back #2788).
• Float dtypes rejected in bitwise reduce with an actionable error ([BugFix] reject float dtype 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 ([BugFix] Fix fp16/bf16 T.atomic_max/atomic_min silently corrupting fp32 values #2780).
• return_prev supported for scalar atomic_min/atomic_max ([BugFix] Return previous value for scalar atomic_min/atomic_max #2672), atomic_addx2 with BufferRegion destinations ([BugFix] Support BufferRegion destinations in atomic_addx2 return_prev #2753), and HIP vector atomic add ([BugFix] Support return_prev for HIP vector atomic add #2712).
• T.atomic_addx4 return type guarded for sliced destinations ([BugFix] Guard T.atomic_addx4 return type for sliced destinations #2590).
• Atomic load/store implemented for HIP ([BugFix] Implement atomic load and store for HIP #2711); invalid atomic memory orders rejected ([BugFix] Reject invalid atomic load and store memory orders #2666); CUDA consume ordering mapped to acquire PTX ([BugFix] Map CUDA atomic add consume ordering to acquire PTX #2713).
• TMA atomic-add layout validation refactored (e0f0ac9) and unsupported dtypes rejected ([BugFix] Reject unsupported TMA atomic add dtypes #2830).

Numerics, vectors & dtypes

• FP8 E4M3 special encodings decoded correctly ([BugFix] Decode FP8 E4M3 special encodings correctly #2710); T.infinity supported for float8_e5m2 ([BugFix] Support T.infinity for float8_e5m2 #2671).
• bf16 NaN/Inf preserved during RNE packing ([BugFix] Preserve bf16 NaN and Inf during RNE packing #2690).
• Signed int32 lanes zero-extended in 256-bit vector pack ([BugFix] Zero-extend signed int32 lanes in 256-bit vector pack #2673); 32-lane 8-bit CUDA vectors packed correctly ([BugFix] Pack 32-lane 8-bit CUDA vectors correctly #2701).
• FP4 dequant symbolic exponent clamp fixed ([BugFix] Fix FP4 dequant symbolic exponent clamp #2656).
• T.pow/T.power fixed for constant integer exponent y <= 0 ([BugFix] Fix T.pow/T.power for constant integer exponent y <= 0 #2677).
• T.__exp computes e**x, not 2**x (docstring + CuTeDSL codegen) ([Fix] T.__exp must compute e**x, not 2**x (docstring + CuTeDSL codegen) #2696).
• IEEE math intrinsic names corrected for fp64/fp16/bf16 ([BugFix] Correct IEEE math intrinsic names for fp64/fp16/bf16 #2619).
• Unsupported fast-math input dtypes rejected ([BugFix] Reject unsupported fast-math input dtypes #2804); mixed packed x2 operand dtypes rejected ([BugFix] Reject mixed packed x2 operand dtypes #2802); floating-point predicates rejected in vote intrinsics ([BugFix] Reject floating-point predicates in vote intrinsics #2800); alloc_var initializer dtype preserved ([BugFix] Preserve alloc_var initializer dtype #2801); invalid dtypes rejected in T.dp4a ([BugFix] Reject invalid dtypes in T.dp4a #2652).
• Scalar T.copy path casts to the destination dtype ([BugFix] Cast to the destination dtype in the scalar T.copy path #2771).
• Canonical-simplify LT Case 2 gated on extra scale == +1 ([Arith] Gate canonical-simplify LT Case 2 on extra scale == +1 #2649); vectorized Select constraint handling fixed (#052e6741).

TMA / copy / memory layout

• Strided global buffers handled correctly in 1D TMA copies ([BugFix] Handle strided global buffers correctly in 1D TMA copies #2746); descriptor TMA skipped for device-bound copy bases ([BugFix] Skip descriptor TMA for device-bound copy bases #2803).
• Partial 1-D TMA stores no longer bypass bounds checks ([BugFix] Prevent partial 1-D TMA stores from bypassing bounds checks #2716); 1D bulk TMA transfer alignment check fixed ([BugFix] Fix 1D bulk TMA transfer alignment check #2646); 1D TMA selection fixed for versioned layouts ([CUDA][Pipeline] Fix 1D TMA selection for versioned layouts #2737); non-16B cluster bulk copies fall back ([BugFix] Fallback non-16B cluster bulk copies #2683).
• st.bulk destination emitted as a shared write to fix a missing barrier and compilation-introduced races ([Bugfix] Emit st.bulk destination as a shared write to fix missing barrier and potential compilation-introduced races #2700).
• Tile copy OOB respects the safe value ([BugFix] Respect safe value in tile copy OOB #2636); runtime-dependent vector negative indices supported ([BugFix] Support runtime-dependent vector negative indices #2654).
• Operator precedence fixed in the increase_descriptor_offset guard ([BugFix] Fix operator precedence in increase_descriptor_offset guard #2675).
• Packed shared memory allocation sizes corrected for CUDA/HIP ([BugFix][CUDA][HIP] Correct packed shared memory allocation sizes #2660); HIP predicated dword copy zero fill fixed ([BugFix] Fix HIP predicated dword copy zero fill #2721).
• Buffer element offsets preserved in access pointers ([BugFix] Preserve buffer element offsets in access pointers #2727); decoupled cast buffer scope preserved in codegen ([TIR][Codegen] Preserve decoupled cast buffer scope #2545).
• T.transpose swaps only the final two axes ([BugFix] Make T.transpose swap only the final two axes #2757); contracting shared-buffer layouts rejected in T.annotate_layout ([BugFix] Reject contracting shared-buffer layouts in T.annotate_layout #2719); unused fragment buffers allowed without layouts ([TIR][Transform] Allow unused fragment buffers without layouts #2717); shared-TMEM buffer pointer types checked before dereference ([BugFix] Check shared-TMEM buffer pointer types before dereference #2794).

Metal backend

• Threadgroup address-space qualifier emitted for shared-memory pointer arithmetic ([BugFix] Add threadgroup address space qualifier in Metal codegen for shared memory pointer arithmetic #2770).
• Barriers emitted for dynamic shared memory ([BugFix] Emit Metal barriers for dynamic shared memory #2738).
• Explicit row strides honored in Metal GEMM ([BugFix] Honor explicit row strides in Metal GEMM #2730).
• Metal stream bridge fixed ([BugFix] Fix metal stream bridge #2639).
• Arithmetic operators added to vec_type in common.h for CPU codegen ([BugFix] Add arithmetic operators to vec_type in common.h for CPU codegen #2768).

Race analysis & warp-specialization

• Two-instance modeling fixed in ThreadSync cross-thread race checks ([BugFix] Fix two-instance modeling in ThreadSync cross-thread race checks #2805).
• Flat Bind modeling fixed in parallel race checks ([TIR][Transform] Fix flat Bind modeling in parallel race checks #2665).
• VerifyParallelLoop race diagnostics aggregated with source spans ([Enhancement] Aggregate VerifyParallelLoop race diagnostics with span #2806).
• Side-effecting binds no longer classified as replayable — fixes atomics being re-executed at every use site since v0.1.11 ([BugFix][Transform] Never classify side-effecting binds as replayable (atomics re-executed at every use site since v0.1.11) #2651).
• GEMM accumulator writes tracked in the warp-specialization liveness collector ([BugFix] Track GEMM accumulator writes in warp-specialization liveness collector #2685).
• Pipeline replacement fixed under persistent T.serial ([BugFix][WS] Fix pipeline replacement under persistent T.serial #2674).
• WGMMA C-store layout fixed for multiple warpgroups along M ([BugFix] Fix WGMMA C-store layout for multiple warpgroups along M #2663).
• Unsafe non-warp-multiple partial thread sync rejected ([BugFix] Reject unsafe non-warp-multiple partial thread sync #2679).

GEMM / misc compiler fixes

• Uncovered warp partitions in T.gemm rejected instead of silently producing wrong results ([BugFix] Reject uncovered warp partitions in T.gemm instead of producing silently wrong results #2724).
• MFMA DataType args no longer break compilation on ROCm ([BugFix] Fix MFMA DataType args causing compilation failure on ROCm #2726).
• PCWS index dtype handling fixed ([CUDA][Transform] Fix PCWS index dtype handling #2783).
• CPU-fallback thread placeholder replaced with a constant-zero logical thread index ([Transform] Replace CPU fallback thread placeholder with a constant-zero logical thread index #2718).
• Non-positive thread extents rejected in T.Kernel ([BugFix] Reject non-positive thread extents in T.Kernel #2653).
• Callee global symbols used for cross-target calls ([BugFix] Use callee global symbols for cross-target calls #2740).
• TMEM/TMA builtins gated by CUDA architecture ([BugFix] Gate TMEM and TMA builtins by CUDA architecture #2743).
• Typo fragment spelling corrected ([BugFix] Correct the spelling of fragment #2695); BufferStore cast warning context improved ([Cherry][TIRx] Improve BufferStore cast warning context #2733).

Autotuning

• Early stop to skip slow configs during benchmarking ([Autotune] Add early stop to skip slow configs during benchmark #2723), including decorator mode with an example ([Autotune] Support early_stop in decorator mode and add decorator example #2729).
• Per-config pass_configs supported in autotuning ([Autotune] Support per-config pass_configs in autotuning #2496).
• Autotuner cache no longer reused across different outputs and validation settings ([BugFix] Prevent autotuner cache reuse across different outputs and validation settings #2793).
• Segfault fixed when tunable parameters default to None ([BugFix] Fix segfault when tunable params default to None #2657).

Examples

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

New Contributors

• @zyy3077 made their first contribution in [BugFix] Preserve guard identity in LoopUnswitching #2585
• @hhy3 made their first contribution in [BugFix] Fix grouped reduce_sum over-counts on straddle layout #2424
• @gvr13n made their first contribution in fix: cast ptxas register usage level to int before building the nvcc command #2641
• @Lfan-ke made their first contribution in [BugFix] Reject non-positive thread extents in T.Kernel #2653
• @UnsettingGalaxy made their first contribution in [BugFix] Fix 1D bulk TMA transfer alignment check #2646
• @Da1L8-X made their first contribution in [Autotune] Support per-config pass_configs in autotuning #2496
• @Lyscoria made their first contribution in [BugFix] Support runtime-dependent vector negative indices #2654
• @xiaoyouPREG made their first contribution in [BugFix] Preserve bf16 NaN and Inf during RNE packing #2690
• @morluto made their first contribution in [BugFix] Map CUDA atomic add consume ordering to acquire PTX #2713
• @arxgy made their first contribution in [Bugfix] Emit st.bulk destination as a shared write to fix missing barrier and potential compilation-introduced races #2700
• @li-ruinan made their first contribution in [BugFix] Fix wrong offset when scanning a non-zero-offset buffer sub-region #2680
• @cla7aye15I4nd made their first contribution in [BugFix] Fix HIP predicated dword copy zero fill #2721
• @mygitljf made their first contribution in [BugFix] Fix FP4 dequant symbolic exponent clamp #2656
• @jayzlee147 made their first contribution in [BugFix] Fix MFMA DataType args causing compilation failure on ROCm #2726
• @zkyue made their first contribution in [BugFix][Transform] Never classify side-effecting binds as replayable (atomics re-executed at every use site since v0.1.11) #2651
• @Butterfingrz made their first contribution in [Example][DeepSeek-V3.2] Adaptive threads for sparse MLA backward #2592
• @perkyfever made their first contribution in [FEATURE] Add block-causal attention for dLLM example #2499
• @GY-Bai made their first contribution in [BugFix] Add threadgroup address space qualifier in Metal codegen for shared memory pointer arithmetic #2770
• @bhaochen made their first contribution in [BugFix] Honor nan_propagate in reduce max/min/absmax clear=False write-back #2788
• @KellyFrog made their first contribution in [BugFix] Resolve partial scalar reduce barrier participation #2777

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

---

This discussion was created from the release v0.1.13.