PyPI · #2341 most downloaded on PyPI

tilelang

0.1.153.0M downloads/motile-ai/tilelang ↗New releases in your daily digest, changelog included.

A tile level programming language to generate high performance code.

Last release 5 days ago
30 Sep 2026
Ships fairly regularly
a new release about every 4 weeks
Nearly every release is documented
notes for 22 of 22 stable releases
1 version withdrawn
withdrawn after publishing
2 years old
24 releases · first in 2025

Release timeline

24 releases since 2025

20252026

One column per month.

Releases

0.1.1530 Sept 2026Deprecated
T.symbolic remains available as a deprecated alias; use T.dynamic for new code ( #3216 ).
Highlights

Native Huawei Ascend 950 support (#3308): an end-to-end NPU backend with native code generation, automatic Cube/Vector scheduling and synchronization, and mixed SIMD/SIMT programming.

Automatic CUDA warp specialization (#3059, #3185): an opt-in role-based scheduler that assigns TMA loads, MMA computation, TMA stores, and worker operations to specialized warp groups.

Unified block-scaled GEMM (#3237, #3257, #3284): common T.gemm_blockscaled semantics with dedicated backend dispatch, improved SM100 instruction selection, and expanded SM120 fragment support.

More expressive Python frontend (#3230): compile-time iteration over Python iterables, enumerate, zip, comprehensions, and generator expressions.

Ascend 950

Add tilelang.ascend.language and target="ascend" for Huawei Ascend 950 (dav-3510).

Combine Cube GEMM and Vector computation in one kernel, with T.SimdVF and T.SimtVF regions.

Support explicit UB/L1/L0 storage, tiled copies, cross-core transfers, and MXFP8/MXFP4 block-scaled GEMM.

Add automatic scheduling, pipelining, multi-buffering, layout inference, and synchronization insertion.

Integrate Bisheng compilation, tvm_ffi and Cython execution, PyTorch NPU tensors and streams, and NPU profiling.

Include GEMM, DeepGEMM-style kernels, FlashAttention forward/backward, RMSNorm, and FP8 quantization examples.

See the Ascend 950 guide for installation and usage. Ascend A2/A3 support remains in the community-maintained TileLang-Ascend projects.

CUDA

Enable role-based automatic warp specialization with TL_ENABLE_AUTO_WARP_SPECIALIZATION: "role_based"; add GEMM and FlashAttention examples (#3059, #3185).

Extend SM120 block-scaled GEMM to fragment-resident A operands, row-major scale fragments, and odd per-warp atom grids; diagnose conflicting scale-fragment layouts (#3257, #3284).

Add round-to-nearest T.fma and T.fmul, and complete additional FP16/BF16 math bridges (#3134, #3132, #3163).

Preserve packed FP8 vector copies and fuse exact FP4-to-FP8 conversions through FP32 (#3276, #3204).

Fix WGMMA/UMMA K-panel strides for operand layouts and insert async-proxy fences for sparse MMA (#2965, #3139).

Correct atomic vectorization for invariant or non-contiguous destinations; keep shared FP32 atomics scalar on SM90 (#3129, #3219, #3238).

Fix NVRTC warp reductions, kernel-body assertions, boolean bitwise negation, and non-constant int4/uint4 broadcasts; restrict 256-bit global loads/stores to SM100+ (#3260, #3206, #3228, #3114, #3248).

Add an FP8 sparse MLA forward example for DeepSeek V3.2 on Hopper (#3224).

Language and Compiler

Make kernel launch encoding target-neutral, with launch options and operation hints owned by their backend dialects (#3186, #3203).

Support Python compile-time iteration and comprehensions while preserving device-loop semantics for for ... in range(...) (#3230).

Support unary plus on symbolic expressions, fix loop-variable binding, remove warnings for immutable rebinding, and reject unsupported loop else clauses (#3141, #3232, #3262, #3142).

Treat T.copy / T.async_copy coalesced_width as a hint and clamp it to the achievable vector width (#3246).

Preserve dynamic reduction tail guards, fix parallel-loop lowering with let inlining disabled, and restore async-copy lowering with partitioned layouts (#3294, #3269, #3278).

Improve reducer layout planning and symbolic layout validation; reject unsupported AllReduce thread strides and reduction NaN-propagation dtypes (#3171, #3233, #3266, #3273).

Fix NaN-propagating clamp lowering and eliminate unused bindings during simplification (#3205, #3293).

ROCm, CPU, and Metal

ROCm: support the DeepSeek V3.2 Top-K selector, wave64 Hadamard transforms, and GLM-5.3 k-pool examples, including Top-K index transformation (#3147, #3154, #3254).

Expand ROCm validation for attention kernels and portable examples, and document pip installation (#3146, #3148, #3165, #3168).

CPU: compute scalar GEMM products in the accumulator dtype and legalize BF16 arithmetic (#2917, #3201).

Metal: support 32-bit integer atomic add and respect GEMM buffer-region offsets (#3211, #3209).

Runtime, Build, and Tooling

Store cached kernel parameters as JSON instead of cloudpickle (#3143).

Publish CUDA binaries and metadata atomically in immutable cache directories, preventing readers from observing partially published entries (#3177).

Fix dynamic-output allocation when the sizing input follows the output, and add missing uint64 argument mappings (#3207, #3229).

Add wall-clock benchmarking and MPS-compatible timing helpers (#3234).

Improve CUDA/ROCm compiler discovery and library loading for symlinked installations (#2839, #3166, #3227).

Enable optimization for default single-config native builds and restore effective Windows wheel-build caching (#3191, #3133, #3305).

Validate built wheels on GPU runners and make performance regressions fail CI (#3167, #3182).

Compatibility Notes

Remove the unused sync and group parameters from T.Pipelined, and k_pack from T.gemm_sp (#3202).

ROCm kernels using T.gemm(k_pack=...) should import tilelang.rocm.language (#3203).

Kernels querying thread extents during tracing must specify threads= explicitly in T.Kernel (#3186).

T.symbolic remains available as a deprecated alias; use T.dynamic for new code (#3216).

Remove TILELANG_CACHE_VERIFY_HASH; binary artifact hash verification is now mandatory. Legacy cache formats are rebuilt automatically (#3143, #3177).

Full Changelog: v0.1.14...v0.1.15

What's Changed

[CI] Cap torch<2.14 for CUDA tests until flash-attn supports torch 2.14 by @LeiWang1999 in #3136

Bump transformers from 5.5.0 to 5.10.1 in /examples/bitnet-1.58b by @dependabot[bot] in #3131

[CI] Fix slow Windows wheel builds by making ccache effective by @LeiWang1999 in #3133

[Language][CUDA] Add T.fma and T.fmul round-to-nearest intrinsics by @LeiWang1999 in #3134

[CUDA] Complete 16-bit bridges for CUDA-lowered unary math by @Chennesxu in #3132

[Refactor] Restore LoopUnswitching test altered by #3121 by @Yongqi-Zhuo in #3124

[BugFix][CUDA] Support non-constant int4/uint4 broadcast by @jjppp in #3114

[BugFix][CPU] Compute scalar GEMM products in accum dtype instead of input dtype by @Dino1844 in #2917

[BugFix][Hopper][Blackwell] Read the WGMMA/UMMA K-panel stride from the layout instead of the operand extent by @bigSheep123 in #2965

[BugFix][Carver] Gate _legalize_info fallback on sm_version by @mocusez in #3145

[ROCm] Validate block-causal attention kernels by @andyluo7 in #3146

[ROCm] Validate attention sink kernels by @andyluo7 in #3148

[ROCm] Support DeepSeek-V3.2 Top-K selector by @andyluo7 in #3147

[Bugfix][Language] Reject loop else clauses in the eager frontend by @rishabhsinha17 in #3142

[Example] Fix hadamard reference precision on TF32 and add pytest entry for CI by @Guan-jeans in #3126

[BugFix][Carver] Compare sm_version numerically in plan_rasterization by @mocusez in #3144

[Misc] Expose .agents skills to Claude Code via .claude/skills symlinks by @LeiWang1999 in #3161

[Language] Cache dtype.as_torch and demote storage-dtype fallback logs to debug by @LeiWang1999 in #3160

[Refactor][CUDA] Replace per-lane-count vector constructors with variadic packers by @LeiWang1999 in #3158

[CUDA] Add 16-bit hpow and hfmod bridges by @Chennesxu in #3163

[BugFix][ROCm] Resolve hipcc robustly instead of trusting bare PATH lookup by @LeiWang1999 in #3166

[CI] Validate built wheels on GPU runners in the Dist workflow by @LeiWang1999 in #3167

[Doc] Document pip installation on AMD GPUs (ROCm) by @LeiWang1999 in #3168

[Bugfix][Language] Support unary plus on PrimExpr in the eager frontend by @rishabhsinha17 in #3141

[CUDA] Add async-proxy fences for sparse MMA intrinsics by @ZenAlexa in #3139

[Feature] [CUDA] Role-based automatic warp specialization by @Yongqi-Zhuo in #3059

[CUDA] Resolve symlinked nvcc before deriving CUDA_HOME by @morluto in #2839

[Language] Remove deprecated T.symbolic alias by @SiriusNEO in #3122

[ROCm] Support wave64 Hadamard transforms by @andyluo7 in #3154

[Cache] Store cached kernel params as JSON instead of cloudpickle by @rishabhsinha17 in #3143

[BugFix][Transform] Keep shared FP32 atomics scalar on SM90 by @ZenAlexa in #3129

[Layout] Consider scalar reducer plans in register-count search by @LeiWang1999 in #3171

[Transform] Add snapshot mode to FreshenMutableReads by @LeiWang1999 in #3174

[CUDA][Cache] Publish immutable binary cache directories by @LeiWang1999 in #3177

[CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in #3181

[CI] Fail the perf bot on a regression instead of only reporting it by @cklxx in #3182

[Refactor] [CUDA] Rename AutoSchedule to AutoWarpSpecialization by @Yongqi-Zhuo in #3185

[Refactor][Language] Make T.Kernel target-neutral with dialect-owned launch annotations by @LeiWang1999 in #3186

[Cleanup][Language] Remove dead T.Pipelined(sync/group) and T.gemm_sp(k_pack) parameters by @LeiWang1999 in #3202

[Refactor][Language] Move backend-specific op hints into their owning dialects by @LeiWang1999 in #3203

[BugFix][CPU] Legalize BF16 arithmetic in the CPU pipeline by @anerli in #3201

[Language] Restore deprecated T.symbolic alias by @LeiWang1999 in #3216

[Refactor][JIT] Replace callee-allocated-output target sniffing with a backend capability flag by @LeiWang1999 in #3218

[Transform] Restore bind-before-use order when merging constraint sets by @LeiWang1999 in #3221

[BugFix][Metal] Respect GEMM buffer region offsets by @anerli in #3209

[Runtime] Fix library loading for symlink installs by @sepcnt in #3227

[CUDA] Fix boolean bitwise negation codegen by @sepcnt in #3228

[Example] FP8 sparse MLA forward for DeepSeek V3.2 on Hopper by @xuebozhang525-alt in #3224

[Profiler] Split timing helpers and add wall-clock benchmarking by @SiriusNEO in #3234

[CPU] Import the CPU dialect in CPU tests by @penguin-wwy in #3236

[BugFix][Vectorize] Keep atomic_add scalar for invariant/non-contiguous destinations by @Dino1844 in #3219

[Transform][CUDA] Plan atomic vector widths from destination addresses by @LeiWang1999 in #3238

[Do not review][Op][Language][CUDA] Add common block-scaled GEMM semantics and backend dispatch by @LeiWang1999 in #3237

[JIT] Add missing uint64 argument type mappings by @sepcnt in #3229

[BugFix] Restore symbolic loop-layout injectivity proof; reject layouts on symbolic shared tiles by @sepcnt in #3233

[ROCm] Run portable example validation in CI by @andyluo7 in #3165

[BugFix][CUDA] Only emit 256-bit global load/store on SM100+ targets by @penguin-wwy in #3248

[Metal] Support 32-bit integer atomic add by @anerli in #3211

[Testing] Drop duplicated codegen-smoke tests; fix fastmath self-comparison assertions by @penguin-wwy in #3252

[ROCm] Add GLM-5.3 k-pool Top-K transform by @andyluo7 in #3254

[Frontend] Support Python iterables and comprehensions by @sepcnt in #3230

[Frontend] Remove warnings for immutable variable rebinding by @LJC00118 in #3262

[BugFix] Reject non-power-of-two AllReduce thread strides by @LeiWang1999 in #3266

[NVRTC] Fix warp reduction compilation by @sepcnt in #3260

[BugFix][CUDA] Lower a kernel-body assert to a device-legal check by @xy200303 in #3206

[BugFix][JIT] Allocate a dynamic-shape output that precedes its sizing input by @xy200303 in #3207

[CUDA] Fuse exact FP4 to FP8 conversion through FP32 by @ZenAlexa in #3204

[Fix]Clamp T.copy/T.async_copy coalesced_width to achievable vector size instead of LOG(FATAL) by @edragain2nd in #3246

[BugFix] Bind loop targets independently of mutable scalar variables by @sepcnt in #3232

[BugFix][Language] Lower NaN-propagating clamp through device templates by @xy200303 in #3205

[Fix] Reject unsupported reduction NaN propagation dtypes by @ZenAlexa in
Note truncated.

0.1.142 Sept 2026

Reducer v2 ( #2940 , #3093 , #3043 , #3044 , #3079 , #3100 ): T.alloc_reducer reworked into first-class deferred reduction epochs whose physical lower

Highlights

Reducer v2 (#2940, #3093, #3043, #3044, #3079, #3100): T.alloc_reducer reworked into first-class deferred reduction epochs whose physical lowering is planned by layout inference (via a first-class PartialFragment layout). Adds loop-scoped epochs, conditional reducer finalization, and automatic vectorization of contiguous reducer updates.
Warp specialization schedules (#2892): new scheduling and materialization mechanism for warp-specialized kernels.
Layout inference cost models (#2960, #3055, #3061): new IO-aware cost model for free-mode layout selection; register-count restored as the default, with an environment override to switch models.
Unified backend resolution policy (#2318) plus backend split-up (#2855, #2870, #2850): backend selection is now resolved through a single policy, and builtin ops / Python op proxies are split per backend (CUDA/ROCm/Metal).
Compilation speed: up to ~4x faster cold parallel/AOT compilation (#2809); Z3 solvers materialized lazily (#3105) and analyzer contexts isolated per kernel compilation (#2890).
TMA rework: TMA copy lowering unified on CuTe algebra (#3106); TMA layouts made region-aware to keep slices contiguous (#3089).

Language

Recycle T.unroll(explicit=True) for early explicit unrolling (#2859)
Make the region bridge a builtin intrinsic (#2983)
Expose cluster_mask on T.tma_copy (#2932)
Unify contiguous stride construction under a single implementation (#3016); honor declared strides in pointer helpers (#3073)
Stricter validation: reject symbolic T.gemm tile dimensions with a clear message (#3113), validate T.gemm k_pack arguments (#3094), reject non-positive arrive_count in alloc_barrier/alloc_cluster_barrier (#3112), reject T.Parallel indexing of local buffers (#3041), reject break in fully expanded loops (#3078)

CUDA

tcgen05: pack logical TMEM buffers into shared tcgen05.alloc arenas (#2831); support half-subpartition (M=64) TMEM tiles in tcgen05.ld/st (#2880); fix ld/st segment pointer advancement in b32 columns (#2952)
Select the widest legal WGMMA N instead of gcd (#2931)
FP32x2 accumulation for reductions: per-reduce control (#3057) and a global PassConfig (#3128)
Pre-SM80 fallback for bf16 atomic add (#2938); int4x2/uint4x2 codegen (#3036); 16-bit CUTLASS type overloads for fast-math, __ldg, and htan intrinsics (#3097, #3077, #3028, #2894)
Fixes: warp shuffle for half/bfloat16/FP8 (#3056), FP8 min/max codegen (#3047), UB in packed 8-bit vector stores (#3092), logical not for vectorized bool (#3117, also HIP), vectorized Select codegen (#2843), ldmatrix source offsets wrapped within shared-memory regions (#3110), NVRTC kernel handles isolated per adapter (#2950), flat CUDA include discovery for NVRTC (#2829), masked warpsync in in-warp allreduce (#2865)
Remove T.{reads,writes} for T.tma_{gather4,scatter4} (#3053); separate TMA atomic-add dtype support from layout encoding (#2846)

ROCm and other backends

Remove the Composable Kernel dependency (#3111); ROCm CI re-enabled on a gfx942 runner (#2874, #2910)
Fixes: preserve FP8 bits in warp shuffles (#3104), lower vector Select conditions lane-wise (#2889), emit a compiler barrier for tl.sync_warp on HIP (#2872), reject sub-wavefront block sizes instead of crashing (#2918), resolve versioned device properties in the HIP stub (#2919); emit #line directives for the HIP target (#3058)
CPU backend: support atomic ops (#2941) and reduce ops (#2893)
Metal: preserve pointer address spaces for byte offsets (#2925); resolve auto backend to torch and skip disk cache for torch (#2856)
CuTeDSL: port backend intrinsics to CUTLASS DSL primitives (#2871)

Compiler / Transform

Refactor the loop vectorization plan with ConstraintKind (#2935); always vectorize T.Parallel loops (#3121); scalarize Select in automatic vectorization (#3060)
Add VerifyBufferInit, a general buffer-initialization check (#2956)
Debug info: preserve source spans across lowering passes (#2966); emit #line directives from TIR spans (#3048)
Fixes: don't drop syncs from the other if branch (#3085), fix wait parity for explicit mbarriers in pipelined loops (#3087), avoid int32 overflow in vector analysis (#3066), fix non-divisible nested modulo simplification (#3065), fix ties-away-from-zero round compile on bfloat16/float8 (#2873), fix absmax/abssum for uint dtypes (#2845), carry memory_order through vectorized atomic_add (#2924), fix unsigned zero-point decode underflow (#3118), keep cp.async operands in their address spaces (#2869), handle grid barriers and unbounded pointer ranges (#3050), deduplicate replicated reducer updates (#2881), bind symbolic coordinate ranges in FragmentThreadIndexProbe (#3096), reject non-round-tripping inferred layout inverses (#3090)
Cleanup: remove obsolete compiler and runtime paths (#3086); remove the obsolete disable-fast-math pass config (#3098)

Runtime / JIT / Build

Kernel cache: detect and repair corrupted cache entries (#3074), export libraries after disk-cache hits (#3116), remove the separate cache temporary directory (#3069)
Allocate kernel outputs through the packed API (#2937); fix get_parent_locals frame self-reference leak (#2934)
Support Cython 3.3 with the Python 3.9 limited API (#3068); fix CMake reconfigure aborting in FindPipCUDAToolkit before project() (#3102)

Tooling / Ecosystem

Official compile-only CLI (#3045)
Unified pass instrumentation per compilation (#2923); Pass Visualizer driven by PassInstrument (#2866); show kernel name in the pass timing report (#2905)
Data race check disabled by default, opt-in via env var (#2851)
Open-source TileLang LSP announced (#2862); new agent skills: simplification (#3080), semantic validation (#3054), PR submission (#3082), backend architecture (#2900)
Examples: generalize TCGEN05 GEMMs for Thor (sm110a) and add Stream-K scheduling (#2902); adopt multi-staged buffers in examples (#2836)
Docs: add Sunrise-AI TANG (#3107), HYGON (#3101), and MetaX MACA (#3095) to supported platforms; link the multi-backend architecture design (#3123)

What's Changed

[CUDA] Pack logical TMEM buffers into shared tcgen05.alloc arenas by @Rachmanino in #2831
[Enhancement] Speed up cold parallel/AOT compilation up to ~4x by @cklxx in #2809
[Docs] Refresh README news and onboarding by @LeiWang1999 in #2849
[Enhancement] Disable data race check by default, opt-in via env var by @KellyFrog in #2851
[BugFix] Fix absmax and abssum for uint dtypes by @jjppp in #2845
[CUDA][TMA] Separate atomic-add dtype support from layout encoding by @LeiWang1999 in #2846
[TIR][Python] Trim redundant op proxy wrappers by @SiriusNEO in #2850
[CUDA][ROCm][Metal] Split backend-specific builtin ops by @LeiWang1999 in #2855
[CUDA] Adopt multi-staged buffers in examples by @Yongqi-Zhuo in #2836
[CI] [pre-commit.ci] autoupdate by @pre-commit-ci[bot] in #2861
[Docs][LSP] Announce the open-source TileLang LSP by @LeiWang1999 in #2862
[BugFix][Metal] Resolve auto backend to torch and skip disk cache for torch by @oraluben in #2856
[Typo] Correct source spelling errors by @morluto in #2858
[Refactor] Recycle T.unroll(explicit=True) for early explicit unrolling by @Yongqi-Zhuo in #2859
[BugFix] Use masked warpsync in in-warp allreduce by @jjppp in #2865
[Debug][TIR] Drive Pass Visualizer with PassInstrument by @LeiWang1999 in #2866
[Doc] Fix wrong loop bound in FlashAttention README example by @shanyi0228-web in #2868
[Examples] Gate CUDA-only and flash_attn-dependent example tests by @andyluo7 in #2864
[Testing] Gate CUDA-only tests so non-CUDA backends can run the suite by @andyluo7 in #2863
[TIR][Python] Split backend-specific op proxies by @SiriusNEO in #2870
[BugFix][ROCm] Emit a compiler barrier for tl.sync_warp on HIP by @andyluo7 in #2872
[BugFix] Handle vectorized SelectNode in codegen_cuda by @jjppp in #2843
[CI] Re-enable ROCm CI on a gfx942 runner by @andyluo7 in #2874
[Cleanup] Replace root reproducers with CPU regression coverage by @GY-Bai in #2878
[Compiler][Z3] Isolate analyzer contexts per kernel compilation by @LeiWang1999 in #2890
[Backend] Add unified backend resolution policy by @SiriusNEO in #2318
[Doc] ROCm CI is no longer disabled by @andyluo7 in #2896
[BugFix] Deduplicate replicated reducer updates by @KellyFrog in #2881
[Test] Add regression test for issue #2883 by @SiriusNEO in #2899
[CPU] Support reduce ops on CPU by @penguin-wwy in #2893
[Docs] Define backend architecture and integration skill by @SiriusNEO in #2900
[Testing] Gate the issue #2883 regression test on CUDA by @andyluo7 in #2908
[BugFix][ROCm] Lower vector Select conditions lane-wise by @morluto in #2889
[CI] Use stable torch for the ROCm leg by @andyluo7 in #2910
[Testing] Run portable regression tests on auto targets by @SiriusNEO in #2914
[BugFix][Carver] Parse lettered SM arch strings in check_sm_version by @adityasingh2400 in #2891
[Enhancement] Show kernel name in pass timing report by @penguin-wwy in #2905
[CuTeDSL] Port backend intrinsics to CUTLASS DSL primitives by @cherichy in #2871
[BugFix][ROCm] Reject sub-wavefront block sizes instead of crashing by @andyluo7 in #2918
[Fix] Discover flat CUDA includes for NVRTC by @morluto in #2829
[CI]: Bump pypa/cibuildwheel from 4.1 to 4.2 by @dependabot[bot] in #2930
[ROCm] Resolve versioned device properties in HIP stub by @skyguan92 in #2919
[BugFix] Skip DecoupleTypeCast on Evaluate roots to keep cp.async operands in their address spaces by @li-ruinan in #2869
[Debug][TIR][JIT] Unify pass instrumentation per compilation by @LeiWang1999 in #2923
[Cherry-Pick][BugFix] Fix get_parent_locals frame self-reference leak by @SiriusNEO in #2934
[BugFix] Fix failed ties-away-from-zero round compile on bfloat16/float8 by @edragain2nd in #2873
[JIT][FFI] Allocate kernel outputs through the packed API by @LeiWang1999 in #2937
[Refactor][BugFix] Refactor the loop vectorization plan with ConstraintKind by @SiriusNEO in #2935
[BugFix][CUDA] Provide htan overloads for fp16/bf16 tangent by @Ruihan11 in #2894
[Layout] Support half-subpartition (M=64) TMEM tiles in tcgen05.ld/st by @Rachmanino in #2880
[Feature] Warp specialization schedules and materialization by @Yongqi-Zhuo in #2892
[CUDA] Add pre-SM80 fallback for bf16 atomic add by @Chennesxu in

Note truncated.

Open source →

Additional notes

v0.1.14

Compare

Choose a tag to compare

Open source →

0.1.133 Aug 2026Breaking

Breaking changes : this release removes several legacy APIs and packages. See Backend, API & Refactors before upgrading.

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 st

Note truncated.

Open source →

Additional notes

v0.1.13

Compare

Choose a tag to compare

Open source →

0.1.128 Jul 2026

TileLang v0.1.11 → v0.1.12 Changes

Summary of the main changes between v0.1.11 and v0.1.12 (91 commits).

New Features

LLVM backend support (#2409), with follow-up fixes for auto backend resolution (#2519) and module export (#2467)
Tile scheduler introduced (#2441)
Backend registry architecture: host and device CodeGen are now dispatched through a backend registry (#2442, #2446), target detection/normalization is registration-based, and ExecutionBackend was merged into the backend module (#2323); kernel launch is materialized per backend (#2387)
Developer tooling: pass_visualizer structure-tree pass browser (#2449) and pass-diff display for debugging (#2375)
New CUDA intrinsics exposed (#2473), st.bulk shared-memory zero fill on SM100+ (#2403), stmatrix m16n8 on Blackwell (#2417), SM75 MMA dispatchers for FP16 accumulation and UINT8 (#2392)

CUDA / Codegen Improvements

Optimized fp8↔half/bf16 casts: vectorized and scalar cast codegen (#2511, #2475), plus a fix for vectorized fp16↔bf16 cast compilation (#2407)
TMA lowering for arbitrary/swizzled SMEM layouts (#2380) and GMMA/UMMA lowering for sliced (arbitrary-layout) SMEM (#2452)
Reduced CUDA template include overhead (#2474), swizzled TMA buffer alignment (#2391), RNG state kept in kernel scope (#2540)
Warp-specialization fixes: register over-subscription (#2406), register reallocation for 1P1C (#2440)

JIT / Caching / Build

Cross-host CUDA binary cache (#2459); compile options now included in the cache key (#2532)
PyTorch extensions and perf wheels are cached (#2509); frontend disk cache removed (#2363); lazy kernel lookup caching improved (#2357)
JIT diagnostics and configurable NVCC timeout (#2350), -ccbin support for choosing the C++ compiler (#2348), TILELANG_VERBOSE env var to control compile output (#2453)

Notable Bug Fixes

Pipeline: fixed physical async wait counts (#2505) and simplified async copy lowering (#2444)
Layout/Transform: divide-by-zero in LayoutInference on non-power-of-two broadcast (#2469), avoided thread-indexed replicated fragment readback (#2514), kept all-rep reducers from scalarizing vector plans (#2507), reducer workspace only allocated for cross-warp AllReduce (#2494)
Correctness: T.Persistent dropping tiles when the last dim isn't a multiple of group_size (#2455), sign-extension bugs in packed uint32 decode (#2500) and make_int negative int8 lanes (#2438), PTX v4 atomics for fp16/bf16 atomic_addx4 (#2492), vectorized atomic_add dtype mismatch (#2414), bf16 exp self-recursion (#2402) and rsqrt overload (#2386)
DeepSeek V3.2 topk threshold on exact-boundary inputs (#2513); flash attention bwd varlen NaN fix (#2461); SM100 CLC GEMM schedule-state lifetime (#2423)
Autotuning benchmarking stabilized across devices (#2370); do_bench gained a cache_size option (#2531)

Other

T.view / T.reshape enhancements (#2450), better T.assume/loop-bound handling to eliminate redundant boundary checks (#2502), improved diagnostics for T.serial fragment access (#2462)
C++ style guide added and API naming/namespace normalization across the C++ codebase (#2430, #2434–2436)
Auto target arch now detected from the current device instead of device 0 (#2517)

Overall

This release centers on the new LLVM backend and backend-registry refactor, major TMA/GMMA layout flexibility on CUDA, fp8/fp16/bf16 cast performance, and a large batch of correctness fixes across layout inference, atomics, and pipelining.

What's Changed

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

Note truncated.

Open source →

Additional notes

v0.1.12

Compare

Choose a tag to compare

Open source →

tilelang

Release timeline

Releases

Highlights

Ascend 950

CUDA

Language and Compiler

ROCm, CPU, and Metal

Runtime, Build, and Tooling

Compatibility Notes

What's Changed

Choose a tag to compare

Highlights

Language

CUDA

ROCm and other backends

Compiler / Transform

Runtime / JIT / Build

Tooling / Ecosystem

What's Changed

Choose a tag to compare

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

Choose a tag to compare

TileLang v0.1.11 → v0.1.12 Changes

New Features

CUDA / Codegen Improvements

JIT / Caching / Build

Notable Bug Fixes

Other

Overall

What's Changed

Choose a tag to compare

What's Changed

New Contributors

Choose a tag to compare