Research window: 2026-10-23 07:00 ~ 2026-10-26 07:00 (Beijing time, 72 hours, covering Friday through Monday).

Sources: Per-repo verification of the GitHub tile-ai organization (tilelang / TileOPs / TileFoundry / tilelang-ascend [ascendc_pto] / tilelang-hygon / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / three documentation sites) + ground-truth searches for merged and in-review items (two sets: is:merged and is:open) + full review of newly opened issues + commit-level diffs and release-page tag verification + PyPI release status verification + adopter-repo verification (TileKernels / FlashQLA) + search for new community repos + multiple Google News queries in Chinese and English (via proxy) + Hacker News + arXiv + industry media coverage.

Index for This Issue

  • Today’s Highlights: 78 merges and 76 open updates in 72 hours; only one item left before TileOPs’ first PyPI release, CHANGELOG already records 0.1.0 (10-26); the TIRx replacement under review in the main repo touches 1043 files; situ_glu merged in the Ascend repo, daily tests pass all 2903 items in a single day
  • I. Core Project Progress
    • 1.1 Main repo: partial-sync named barrier ids collide across two storage domains (#3643) (10-25)
    • 1.2 Main repo: NaN-propagation lowering for clamp reverted, T.clamp back to min/max (#3628) (10-23)
    • 1.3 Main repo: int64 atomic max/min and AtomicStore crashes on int64, bf16, fp16 (#3307) (10-23)
    • 1.4 Main repo: VectorizeLoop no longer silently ignores loop step (#3345) (10-23)
    • 1.5 Main repo: autotuning results not written back after NVRTC disk-cache hit (#3193) (10-23)
    • 1.6 Under review: replace bundled TVM with separately packaged TIRx, 1043 files changed (#3666) (10-25)
    • 1.7 Under review: TMEM release on sm_100 missing CTA barrier, warp-specialization prologue drift (#3645 / #3663 / #3664) (10-25)
    • 1.8 New defects: reduction lowering fails on sm_120, three attention examples exceed shared memory (#3647 / #3646) (10-25)
  • II. Multi-Backend Adaptation (Ascend / Hygon / New CPU and ROCm)
    • 2.1 Ascend main repo: DeepGEMM examples support FP4×FP4 packed input (#3589) (10-23)
    • 2.2 Ascend main repo: DMA preserves buffer base offset, region bounds hoisted first, full lifecycle for cyclic buffers (#3627 / #3625 / #3606) (10-23)
    • 2.3 Ascend main repo: CI image upgraded to cann9.3.0 weekly build (#3630) (10-23)
    • 2.4 Ascend under review: E4M3 UB→L1 copy, early rejection of on-chip memory overflow (#3639 / #3655) (10-24~25)
    • 2.5 Ascend under review: PTO target defaults to tvm_ffi execution backend (#3516) (10-24)
    • 2.6 Ascend repo: situ_glu operator merged, aligned with Kimi K3 gated activation (#1891) (10-23)
    • 2.7 Ascend repo under review: sparse MLA, GLM-5.3 kpool, and three normalization examples (10-23~25)
    • 2.8 Ascend repo: daily tests pass all 2903 items, next day’s test job reports failure (#1941 / #1944) (10-25~26)
    • 2.9 Hygon: three merges to mainline, register pipeline back under review (#26 / #18 / #27) (10-23~25)
    • 2.10 New backends under review: Intel AMX BF16/INT8 GEMM, ROCm FP4 and E8M0 (#3637 / #3665) (10-24~25)
  • III. Ecosystem and Adopters (TileOPs / TileFoundry / Distributions)
    • 3.1 TileOPs: seven of eight first-release cleanup items done, CHANGELOG records 0.1.0 (#2738) (10-25)
    • 3.2 TileOPs: public names and paths finalized in one pass, packaging switched to metadata and supports 3.13 (#2741 / #2742 / #2746) (10-25)
    • 3.3 TileOPs: operator-authoring entry point consolidated as authoring, eager calls converge on a single executor (#2743 / #2744) (10-25)
    • 3.4 TileOPs: shared operator semantics documented in one pass, input validation generated from manifest (#2751 / #2758 / #2760) (10-25)
    • 3.5 TileOPs: release docs, third-party notices, and pip install entry point (#2753) (10-25)
    • 3.6 TileOPs performance: GDN prefill partition width chosen from each call’s own length, KDA native gating moved into kernel (#2745 / #2683) (10-23~25)
    • 3.7 TileOPs performance: merged attention states, paged GQA, and indexer tail wave (#2739 / #2752 / #2735 / #2730) (10-24~25)
    • 3.8 TileOPs performance: GEMM tail-wave narrow tiles, ping-pong BMM tiered staging (#2733 / #2734 / #2720) (10-23~25)
    • 3.9 TileOPs performance: MoE A-box loading and native-width WGMMA, elementwise reciprocal and row broadcast (#2711 / #2715 / #2710 / #2725) (10-23~24)
    • 3.10 TileOPs performance: reductions and normalizations switched to 16-byte evict-last reads (#2757 / #2759 / #2748 / #2737) (10-25)
    • 3.11 TileOPs correctness and CI: SSD scan tail, GLA segment unroll, H200 release validation (#2716 / #2750 / #2749 / #2718 / #2713) (10-23~25)
    • 3.12 TileFoundry: E4M3 warp MMA on sm_89 and per-instruction capability checks (#261 / #262 / #264) (10-23~24)
    • 3.13 TileFoundry: cp.async target-layout addressing and MLA backward layout (#263 / #260 / #257) (10-23)
    • 3.14 TileFoundry: split into three distribution packages core / nvidia / apple, test inputs loaded only once (#268 / #269) (10-25)
    • 3.15 Distribution: Sunrise 0.1.15+sunrise.1.2.0 release PR merged (#9) (10-23)
  • IV. Community, Tutorials, and Events
    • 4.1 Community: tilelang-puzzles-ascend builds PTO quantization kernel tutorial for Ascend (pushed 10-23)
    • 4.2 New repo: Chinese TileLang operator interview practice site, 9 modules and 18 problems (10-23)
    • 4.3 Community skills and tutorial repos continue to be pushed, docs site follows first release with four revision PRs (10-23~25)
    • 4.4 Discussion: TileScale user asks whether the TileLang version in use can be upgraded (#72) (10-25)
  • V. Trend Observations

Today’s Highlights

In the past 72 hours, the tile-ai organization merged 78 PRs (TileOPs 51, tilelang 10, TileFoundry 10, docs site 4, tilelang-ascend 1, tilelang-hygon 1, tilelang-sunrise 1), updated 76 open PRs, and opened 19 new issues. The heaviest item is in the release queue: the TileOPs “first-release cleanup” tracking issue (#2738) has seven of eight items done, with only the TestPyPI dry run and the first tag remaining. Meanwhile, the CHANGELOG already recorded its first official version entry, “0.1.0 - 2026-10-26”, at 23:17 on 10-25, stating that operators are imported as tileops.<family>.<Op>, that tileops.backend and tileops.authoring serve backend vendors and third-party operator authors respectively, and that the validation configuration is locked to SM90 (H200), Python 3.12/3.13, PyTorch 2.13, TileLang 0.1.12, and CUDA 13.2. Yet the latest version of the package on PyPI is still 0.0.1.dev2 from 2025-10-09, and the repository has no tags or releases, so the first release sits in its final stage: “docs ready, artifacts not uploaded.” The 10 merges to the main repo center on silent errors: named barriers colliding across storage domains (#3643), the NaN-propagation lowering of T.clamp being reverted (#3628), int64 atomic max/min and AtomicStore crashes (#3307), and VectorizeLoop dropping the loop step (#3345). The most notable item under review is #3666: it deletes 3rdparty/tvm and migrates the frontend and lowering wholesale to a separately packaged TIRx, with 1043 files and +71428/-104176. On the same day, the Ascend line merged the Kimi K3 gated activation operator situ_glu (#1891); that repo’s daily tests reported all 2903 items passing on 10-25, then reported a test job failure in the early hours of 10-26 (#1944).

I. Core Project Progress

1.1 Main repo: partial-sync naming barrier IDs collide across two storage domains (#3643)

Merged 10-25 14:30, 2 files +57/-1, fixes #3642. ThreadSync runs once each for shared and shared.dyn, and each run creates a fresh rewriter that renumbers from kFirstUsedBarrier; the second run skips syncs already rewritten by the first but has no knowledge that those IDs are taken, so it assigns the same bar.sync id, count to two disjoint groups of threads. The fix makes numbering continuous across both runs and adds an assertion for cross-domain reuse. This class of defect does not error at compile time and only shows up intermittently on large kernels.

Source: https://github.com/tile-ai/tilelang/pull/3643

1.2 Main repo: clamp’s NaN-propagating lowering reverted, T.clamp back to min/max (#3628)

Merged 10-23 16:51. Reverts the entire #3205 chain: removes the tl.clamp builtin, per-backend lowering registrations, codegen branches, and device-side helper functions (including the Ascend registration); T.clamp(x, lo, hi) is restored to T.min(T.max(x, lo), hi), on the grounds that the extra NaN handling blocked min/max simplification and vectorization. The NaN requirement did not go away: #3605 was opened the same day to implement it separately as T.clamp_symmetric on CUDA SM86+, with the division of labor being “T.clamp stays general-purpose, the new interface preserves semantics.”

Source: https://github.com/tile-ai/tilelang/pull/3628

1.3 Main repo: int64 atomic max/min and AtomicStore crashes on int64, bf16, fp16 (#3307)

Merged 10-23 18:44. Since #1716, normalize_atomic_type has normalized int64_t to unsigned long long to support atomic add, which made Atomic max/min and their reference forms crash outright on int64; AtomicStore was hit the same way and additionally affected bf16 and fp16. The fix selects types separately for int64 per operation, with AtomicStore going through its own type dispatch. A textbook regression of “changing a shared normalization for one operation and breaking the rest along the way.”

Source: https://github.com/tile-ai/tilelang/pull/3307

1.4 Main repo: VectorizeLoop no longer silently ignores loop step (#3345)

Merged 10-23 19:02, part of F319 tracked by #3301. The vectorizer only received (loop_var, extent, body) — the step never made it in: T.vectorized(0, 8, step=2) (iteration set 0/2/4/6) was lowered into contiguous accesses covering the whole [0, 8). The fix rejects non-unit steps before entering vectorization and emits a diagnostic. The common trait of this batch of sub-items: the frontend accepts a form that cannot be lowered, and the backend emits machine code for a different program.

Source: https://github.com/tile-ai/tilelang/pull/3345

1.5 Main repo: autotuning results not written back after NVRTC disk cache hit (#3193)

Merged 10-23 15:04. When NVRTC loads a kernel from disk it does not set adapter.libpath, but after autotuning selects a configuration it needs that field to save the cached kernel, so every cache hit threw AttributeError and the conclusion was discarded — what users saw was “it retunes on every startup.” The fix restores libpath on the cache-loading path to match the fresh-compile path, and adds cache-reload and negative-control test cases.

Source: https://github.com/tile-ai/tilelang/pull/3193

1.6 Under review: replace built-in TVM with separately packaged TIRx, 1043 files changed (#3666)

Submitted 10-25 19:16, 1043 files +71428/-104176. Removes 3rdparty/tvm and its submodule entry (keeping CUTLASS and HIP headers), and migrates the frontend, native tensor and scoped IR, lowering passes, backend pipelines, and execution adapters wholesale to the independently built and released tvm_tirx; the compatibility layer is confined to TileLang internals, with no global tvm alias and no second copy of TIRx retained. This takes the earlier “tvm fork first, main repo follows” pattern to its endpoint: the foundation moves from vendored source to an external dependency, which can henceforth be versioned and rolled back independently, with the merge impact spanning the entire repo.

Source: https://github.com/tile-ai/tilelang/pull/3666

1.7 Under review: sm_100 TMEM release lacks CTA barrier, warp-specialization prologue statements drift (#3645 / #3663 / #3664)

Three related items, all in automatic warp specialization and Blackwell’s new storage. #3645: on sm_100, tensor memory release has no CTA-level barrier, so a producer warp can free the accumulator while consumer tcgen05.mma/tcgen05.ld are still reading it. #3663: specialization moves prologue statements outside the T.serial outer loop into the loop body, so initialization that should run once is replayed every iteration, producing wrong output. #3664: when an if wraps a specialized pipeline and statements follow the if, the setmaxnreg-related lowering makes ptxas fail. Taken together, there is still no complete semantic model for “which statements belong to which warp, and where their lifetimes end.”

Source: https://github.com/tile-ai/tilelang/pull/3645, https://github.com/tile-ai/tilelang/issues/3663, https://github.com/tile-ai/tilelang/issues/3664

1.8 New defects: reduction lowering fails on sm_120, three attention examples exceed shared memory (#3647 / #3646)

Two reports filed the same day, 10-25 09:08. The former is a ReduceOp lowering failure
for gdn wy_fast_bwd_split on sm_120 — the source index depends on a thread-held reduction segment, which existing rules do not cover; the latter is three examples — attention_sink, gqa_bwd, and
block_causal_attention_varlen — exceeding the 101376-byte per-block shared memory limit. The corresponding fix #3662 is already open: pick smaller tiles, fewer pipeline stages, and 128 threads according to the device limit, while large devices keep the original defaults. Consumer Blackwell’s capacity constraints are becoming a hard boundary for examples and tuning tables.

Source: https://github.com/tile-ai/tilelang/issues/3647, https://github.com/tile-ai/tilelang/issues/3646

II. Multi-Backend Adaptation (Ascend / Hygon / New CPU and ROCm)

2.1 Ascend Main Repo: DeepGEMM Examples Support FP4×FP4 Packed Inputs (#3589)

Merged 10-23 10:26. Enables packed E2M1 FP4×FP4 inputs for the three DeepGEMM interfaces—dense, M-grouped, and K-grouped—reusing the existing MX GEMM kernel and heuristics, deriving logical dimensions from packed-byte storage and aligning to FP4’s 64-element alignment; on the backend side only the MX K-tail padding is changed, and the batched constructors also accept FP4. The public batch/einsum interfaces still keep FP8 to align with DeepGEMM-Ascend-Nightly, with FP8×FP4 dequantization left for later.

Source: https://github.com/tile-ai/tilelang/pull/3589

2.2 Ascend Main Repo: DMA Preserves Buffer Base Offset, Region Bounds Promoted First, Full Lifecycle of Cyclic Buffers (#3627 / #3625 / #3606)

All three merged 10-23; they are layout and lifecycle fixes. #3627: DMA planning dropped Buffer.elem_offset, so when a UB alias was constructed with elem_offset=256, T.copy(alias, out) actually read from element 0; the fix computes the base offset into the address and preserves the parent offset in strided slices (including nested ones). #3625: region minimum and length are promoted with Int64Promoter before constructing the half-open interval, avoiding int32 casts masking dynamic tail bounds, which exhausted the Z3 budget in the overlap proof and introduced spurious cross-iteration dependencies for in-place persistent RoPE. #3606: a sub-buffer lifecycle summary must fully cover accesses before it can be promoted; values preloaded externally live across the entire parent cycle, preventing reuse under nesting and manual pipelining from overwriting pending readers.

Source: https://github.com/tile-ai/tilelang/pull/3627, https://github.com/tile-ai/tilelang/pull/3625, https://github.com/tile-ai/tilelang/pull/3606

2.3 Ascend Main Repo: CI Image Upgraded to cann9.3.0 Weekly Build (#3630)

Merged 10-23 19:20, CI-only change: the ascend job’s test container is bumped from cann9.3.0-weekly.20260930 to cann9.3.0-weekly.20261021-ubuntu22.04-py3.12. It determines which CANN runtime version all Ascend-side regression cases actually validate against, and is the precondition for judging the validity period of conclusions on this line.

Source: https://github.com/tile-ai/tilelang/pull/3630

2.4 Ascend Under Review: E4M3 UB→L1 Copy, Early Rejection of On-Chip Memory Overflow (#3639 / #3655)

#3639 enables float8_e4m3fn UB→L1 copy for Ascend 950: ND inputs are reordered through the existing scatter/post-copy pipeline, pre-packed NZ is moved directly, and footprint and latency analysis add the FP8 element width. #3655 makes MergeUBAllocations validate planned occupancy against the respective capacities of UB, L1, L0A, L0B, and L0C during the lowering phase and emit diagnostics rather than generating out-of-bounds allocations—the trigger was two independent loops each legitimately selecting two versions of a 32KB L0A tile, exceeding the limit after merging. On-chip memory management is shifting from runtime errors to compile-time rejection.

Source: https://github.com/tile-ai/tilelang/pull/3639, https://github.com/tile-ai/tilelang/pull/3655

2.5 Ascend Under Review: PTO Target Defaults to tvm_ffi Execution Backend (#3516)

Under-review change updated 10-24: adds tvm_ffi support for pto and makes it the default, so execution_backend="auto" now resolves to tvm_ffi. The shared host-side codegen handles MakePackedAPI and the C host build; the Ascend runtime module launches PTODSL-compiled aibin via aclrtLaunchKernelWithHostArgs, consistent with the AscendC path. This aligns with main-repo CPU-side #3634 (unifying on TVM-FFI, removing the CPU-specific Cython wrapper): the execution backend is being simplified.

Source: https://github.com/tile-ai/tilelang/pull/3516

2.6 Ascend Repo: situ_glu Operator Merged, Aligned with Kimi K3 Gated Activation (#1891)

Merged 10-23 16:04. Semantics and dtype/shape/attr constraints are aligned item by item with CANN ops-nn’s activation/situ_glu, with accuracy benchmarked against the CPU fp64 PyTorch golden reference; a pure Vector single kernel, where the host merges axes into a (rows, cut*inner) zero-copy view by dim, and non-last-dim splits are non-strided. The operator corresponds to the forward gated activation of Kimi K3 Stable LatentMoE; together with the previous issue’s situ_glu gradient and sparse MLA, the Ascend repo’s operator list is being rolled out one by one according to specific conversational model architectures.

Source: https://github.com/tile-ai/tilelang-ascend/pull/1891

2.7 Ascend Repo Under Review: Sparse MLA, GLM-5.3 kpool, and Three Normalization Examples (10-23~25)

Among the 24 under-review PRs, the operator-side representatives are sparse_mla_fwd (#1929, DSA sparse MLA forward, with varlen packed random-row gather + online softmax + dual QK/PV GEMM, fast path restricted to conditions such as H=128, G=1, dim/tail=(512,64)) and GLM-5.3 kpool compression (#1910, per-dimension softmax pooling within the pool, crossing the BF16 boundary, Hadamard-128 rotation, E4M3 quantization, and scattering rows and scales into the paged cache). On the example side, rms_norm (#1558), NMS (#1748, pairwise IoU compressed to 1 bit into a uint16 bitmap, 2MB replacing a 64MB matrix at N=4096), and GroupNorm (#1607) are added.

Source: https://github.com/tile-ai/tilelang-ascend/pull/1929, https://github.com/tile-ai/tilelang-ascend/pull/1910, https://github.com/tile-ai/tilelang-ascend/pull/1748

2.8 Ascend Repo: All 2903 Daily Tests Pass, Next Day’s Test Job Reports Failure (#1941 / #1944)

The bot account reported at 10-25 05:15 that all 2903 daily tests passed with a 100% pass rate; at 10-26 04:30 it reported the workflow’s test job failed, with another failure record at 10-24 05:08. #1944 does not include failing case details, but from the time series, the daily full suite has stabilized at around 2900 and serves as a regression gate, so a single failure means some newly merged change or environment change broke an existing case.

Source: https://github.com/tile-ai/tilelang-ascend/issues/1941, https://github.com/tile-ai/tilelang-ascend/issues/1944

2.9 Hygon: Three Merges to Main, Register Pipeline Re-Entered Under Review (#26 / #18 / #27)

Main landed three PRs at once on 10-23 17:57 (author dates 10-14, 10-16, 10-20; a batch merge): fixing MLS slot and linear DS reads, CI support for DTK 26.04/26.10, and switching to wave64 shuffle with vectorized BF16 stores; the merged #26 also fixes the FP4 MLS warp-partition handle. Under-review #18 reintroduces the register pipeline (coexisting with shared-memory pipelining, inserting commit and barrier per pipeline stage) and enables warp diverge when the warp count exceeds 4, aiming to bring example_gemm on par with ROCm; updated 10-25 06:00. #27 handles builtin lowering, precision, and integer vector codegen. Contrasting with the previous issue’s “merge branch removes enable_register_pipeline,” the same capability is being redone as a PR.

Source: https://github.com/tile-ai/tilelang-hygon/pull/26, https://github.com/tile-ai/tilelang-hygon/pull/18, https://github.com/tile-ai/tilelang-hygon/pull/27

2.10 New Backends Under Review: Intel AMX BF16/INT8 GEMM, ROCm FP4 and E8M0 (#3637 / #3665)

#3637 adds an Intel AMX backend to the CPU path: target detection, compile options, GEMM dispatch, with BF16 and INT8 kernels using AVX-512/VNNI packing and aligned reuse of scratch, and the execution side adapting OpenMP across 1, 40, and 120 cores. TileFoundry had previously run the prefill example on Apple AMX; this PR brings AMX directly into the main repo’s CPU path. #3665 fixes two gaps in the HIP template: T.cast(x, T.float4_e2m1fn) was printed on HIP as (fp4_e2_t)x, but that type only has an explicit constructor, so floats were truncated into integer storage (1.0 became 0.5, 2.0 became 1.0); it also adds the E8M0 type and enables software FP4 on gfx94x.

Source: https://github.com/tile-ai/tilelang/pull/3637, https://github.com/tile-ai/tilelang/pull/3665

III. Ecosystem and Adopters (TileOPs / TileFoundry / Distributions)

3.1 TileOPs: Seven of eight first-release clean-up items complete; CHANGELOG pins 0.1.0 (#2738)

On 10-25 at 10:45 the “Release clean-up” tracking issue was opened, with seven of eight items checked off: naming, packaging and repository files, Python 3.13, public API and extension points, operator behavior, input validation, and the docs site; the only incomplete item is “third-party notices, PyPI install docs, dev dependency group, CHANGELOG and release steps, TestPyPI dry run and first tag, to be executed by the maintainer.” The CHANGELOG landed its first version entry in the same window, “0.1.0 - 2026-10-26,” noting the first release and that docs/stability.md defines what counts as a public interface. Checking PyPI: the latest tileops is still 0.0.1.dev2 from 2025-10-09; checking GitHub: no tags, no releases. Separately, #2756 opened a “Post-release clean-up” ahead of time.

Sources: https://github.com/tile-ai/TileOPs/issues/2738, https://github.com/tile-ai/TileOPs/blob/main/CHANGELOG.md

3.2 TileOPs: Public names and paths finalized in one pass; packaging moved to metadata and Python 3.13 supported (#2741 / #2742 / #2746)

#2741 is a one-shot rename across 739 files, +6452/-6740: tileops, the per-operator-family modules, tileops.backend, tileops.manifest, tileops.trace and the new tileops.authoring become the public surface, everything else moves down into _ops, _kernels, _utils, _perf, the tileops.ops lazy re-export layer is removed, and each operator keeps exactly one import path. #2742 fixes packaging metadata (bench moved from an extra to a dependency group, dev deduplicated with pre-commit and pytest-timeout added, Linux/CUDA 13/Python 3 declared, py.typed shipped with the package), and adds a CHANGELOG and a bug report template. #2746 sets requires-python to >=3.12 with no upper bound and runs all 43 packaging test cases on an H200 with Python 3.13.15 + torch 2.13.0 + PyPI tilelang 0.1.12.

Sources: https://github.com/tile-ai/TileOPs/pull/2741, https://github.com/tile-ai/TileOPs/pull/2742, https://github.com/tile-ai/TileOPs/pull/2746

3.3 TileOPs: Operator authoring entry point consolidated into authoring; eager calls converge on a single executor (#2743 / #2744)

#2743 defines the extension point: out-of-repo code writes operators by subclassing tileops.authoring.Op and declaring the manifest entry as the class attribute spec, which the validator checks at class creation, granting signature checks, shape inference, roofline and compilation boundaries; only classes internal to tileops take an entry of the same name and are served directly by the backend. #2744 collapses the signature resolution, call spec, kernel program and call record that were re-derived on every eager call into “one plan per specialization,” resolving only once under concurrency and rolling back published tuning configs on build failure. Both the operator author’s and the caller’s interfaces go from implicit convention to explicit contract.

Sources: https://github.com/tile-ai/TileOPs/pull/2743, https://github.com/tile-ai/TileOPs/pull/2744

3.4 TileOPs: Shared operator semantics written down once; input validation generated from the manifest (#2751 / #2758 / #2760)

#2751 and #2758 turn “one convention per concept” into documentation enforced per operator: LSE for all attention is uniformly float32 natural log, dense [B,H,S] and packed [H,T] are stored in that layout (GQADenseBwdOp previously read base-2), rows with no visible keys output 0 and -inf; out= is always keyword-only and returns that object, and backward operators are paired with their corresponding forward. #2760 changes the packed-offset check in validate_inputs from hand-written per operator to manifest-generated, incidentally fixing a latent issue where four linear attention operators took cu_seqlens by argument position.

Sources: https://github.com/tile-ai/TileOPs/pull/2758, https://github.com/tile-ai/TileOPs/pull/2760

3.5 TileOPs: Release docs, third-party notices and the pip install entry point (#2753)

11 files, +103/-32: THIRD_PARTY_NOTICES includes the MIT notice for code from the Qwen team (two GDN prefill kernels are adapted from there) and ships as a license file with the wheel and sdist; the README first screen switches to pip install tileops with verified configuration (Linux, Python 3.12/3.13, PyTorch 2.13, TileLang 0.1.12, CUDA Toolkit 13.2 and nvcc, SM90). The commit note requires merging just before tagging, sequencing it with the last item on the tracking issue.

Source: https://github.com/tile-ai/TileOPs/pull/2753

3.6 TileOPs performance: GDN prefill partition width chosen from the call’s own length; KDA native gating moved into the kernel (#2745 / #2683)

#2745 fixes GDNFwdOp prefill: the partition width previously came from the key head count, and q35-397b-tp8-32768 cut each walk into 64 segments across 4 waves, leaving the nightly benchmark at only 0.988 of FlashInfer; after solving the partition from the call’s own sequence length and value heads, bf16 drops from 289.3 to 200.6 microseconds (1.434x FlashInfer) and fp16 from 285.0 to 195.2 microseconds (1.468x). #2683 lets KDA forward consume raw gates directly (with optional dt_bias and bounded gating), 16 files +3095/-458, adding 28 workloads and exceeding the FLA baseline on every KDAFwdOp row, with strongly decaying chunks recomputed as paired Gram terms to avoid NaN.

Sources: https://github.com/tile-ai/TileOPs/pull/2745, https://github.com/tile-ai/TileOPs/pull/2683

3.7 TileOPs performance: Merged attention states, paged GQA and indexer tail waves (#2739 / #2752 / #2735 / #2730)

#2739 adds a row-free guarded program and one-shot LSE storage after merging to MergeAttentionStatesFwdOp, which previously reached only 0.985 of FlashInfer’s native merge_state on llama-70b-cascade, with the other two rows at 0.988 to 0.999. #2752 makes it support both natural log and base-2 (exp2/log2), changing the single-wave 16-row grid to 512-thread CTAs handling 64 rows. #2735 fixes the RoPE scratch overflow in qwen35, changing the head dim 256 FP8 cache from a single-warpgroup 64-row tile to a dual-warpgroup 128-row tile (previously only 0.16 to 0.57x the fastest comparison implementation), and the flat dim-64 rows to a three-consumer warpgroup. #2730 splits the final sub-wave block of the FP8 indexer into up to 16 CTAs by powers of two, but in practice long rows run 2 to 7 percentage points slower due to ptxas FMNMX ordering.

Sources: https://github.com/tile-ai/TileOPs/pull/2739, https://github.com/tile-ai/TileOPs/pull/2735, https://github.com/tile-ai/TileOPs/pull/2730

3.8 TileOPs performance: GEMM tail-wave narrow tiles, ping-pong BMM staged storage (#2733 / #2734 / #2720)

#2733 targets the last wave of the persistent grid: the tile column count of the busiest CTA determines overall time, and several ds-v3 rows previously trailed cuBLASLt or quack by 0.6 to 2.3 percentage points; the fix runs wide tiles first, then handles the last row with narrower tiles, choosing the layout with the fewest columns on the busiest CTA, with 176-wide tiles paired with a 5-stage ring and a double-buffered epilogue. #2734 stages BmmPingPongKernel stores alternately through two buffers (16 columns when K spans 16 or more segments, 32 otherwise), avoiding the last tile’s epilogue serially waiting on cast and store. #2720 switches short-K batched GEMM within two waves to 128x128 ping-pong tiles, hiding tile switching and the epilogue under the other group’s main loop.

Sources: https://github.com/tile-ai/TileOPs/pull/2733, https://github.com/tile-ai/TileOPs/pull/2734, https://github.com/tile-ai/TileOPs/pull/2720

3.9 TileOPs performance: MoE A-box loading and native-width WGMMA, elementwise reciprocal and row broadcast (#2711 / #2715 / #2710 / #2725)

#2711 fixes decode-time MoE: an expert holds only 1 to 7 rows, but a full 64-row TMA load also pulls the next expert’s rows into L2, making gate/up in gpt-oss-120b-decode L2-bound (3.56GB of DRAM reads correspond to 6.11GB of L2 traffic); the fix splits A into 16-row boxes at 128-byte K atoms and issues only boxes reachable by active rows. #2715 makes the 192- and 224-wide template tiles each issue a single 16K WGMMA — previously T.wgmma_gemm with a width not divisible by 256 was split into multiple gcd(width,256)-wide pieces, each re-reading A, running 27 to 35 percentage points slower. #2710 changes the reciprocal to one approximate instruction plus one vector per thread (previously 1 to 1.4 percentage points behind torch.compile on 16 million fp16/bf16 rows). #2725 has one thread handle four rows sharing a single bias vector when the grid occupies only one wave.

Sources: https://github.com/tile-ai/TileOPs/pull/2711, https://github.com/tile-ai/TileOPs/pull/2715, https://github.com/tile-ai/TileOPs/pull/2710

3.10 TileOPs performance: Reductions and normalization switch to 16-byte evict-last reads (#2757 / #2759 / #2748 / #2737)

#2759 routes read-once 16-byte loads uniformly through tileops::load16_evict_last (argreduce shared-memory segments, per-thread full-vector argreduce, InstanceNorm row normalization, row folding within 8MiB), falling back to evict-first only above the threshold. #2757 changes the down pass of leading-axis reductions so each of 1024 threads holds one 16-byte vector, bringing ProdFwdOp on dim 0 of (4,128,4096) fp16 from 3.33 microseconds to parity with torch.compile. #2748 selects LayerNorm thread count by row count (256 when above 8192 and at most one block per SM, otherwise 128), compiling once each for decode and prefill serving. #2737 stops tuning AdaptiveMaxPool2d — configurations selected by CUDA event timing could be 1.7x slower between rounds — running the derived default configuration instead and adding a register-based global pooling kernel.

Sources: https://github.com/tile-ai/TileOPs/pull/2759, https://github.com/tile-ai/TileOPs/pull/2757, https://github.com/tile-ai/TileOPs/pull/2748

3.11 TileOPs correctness and CI: SSD scan tail, GLA segment unrolling, H200 release validation (#2716 / #2750 / #2749 / #2718 / #2713)

#2716 unrolls the fixed tail of the SSD chunk scan pipeline before CUDA codegen, fixing an integer underflow compile error caused by a residual arithmetic expression in the async copy wait argument (the Mamba case previously failed to compile). #2750 explicitly unrolls the two-stage pipeline of packed GLA, avoiding rejection of symbolic wait counts by newer CUDA code generators. #2749 aligns the optional bias storage before the vectorized causal convolution prefill load — contiguous views may retain an unaligned base address. #2718 tightens H200 release validation: missing cases, empty reports, unsuccessful tests or unverified rows all fail, each install-round verification uses an independent cache, and CUDA Graph covers 23 tests and 9 operators. #2713 makes nightly benchmarks collect serially on an exclusive GPU — previously early warmup initialized a CUDA context while another benchmark held the exclusive device, causing 1061 case errors in one night.

Sources: https://github.com/tile-ai/TileOPs/pull/2716, https://github.com/tile-ai/TileOPs/pull/2718, https://github.com/tile-ai/TileOPs/pull/2713

3.12 TileFoundry: E4M3 warp MMA on sm_89 and per-instruction capability validation (#261 / #262 / #264)

#261 declares FP8 warp MMA on sm_89 (the m16n8k32 E4M3 form of mma.sync): the sm_89 target established in #259 could previously select only BF16, leaving E4M3 products in registers with no instruction to land on. #262 extends capability validation from “MMA only” to every instruction: on sm_89, bulk copy, TMA copy or mbarrier transaction barrier forms previously passed validation and failed only at codegen or ptxas with errors that did not mention the missing hardware; now admit_fields checks each against target capabilities at admission time. #264 runs that atom from scheduling to the GPU with a 16x8x128 E4M3 product, staging K slices through shared memory — precisely the case that exposed the cp.async addressing defect later fixed in #263.

Sources: https://github.com/tile-ai/TileFoundry/pull/261, https://github.com/tile-ai/TileFoundry/pull/262, https://github.com/tile-ai/TileFoundry/pull/264

3.13 TileFoundry: cp.async destination layout addressing and MLA backward layout (#263 / #260 / #257)

#263 fixes the inconsistency in ops::copy_async between addressing the destination by element index and iterating the source by vector: when a thread slice contains multiple vectors and the contiguous segment is not on the first pattern, later vectors land at the wrong address — in the sm_89 E4M3 GEMM, with per-thread slice (2,1,1,8) and strides (256,-,-,1), the second 8-byte vector was written at byte 4 instead of byte 256. #260 closes the layout and analysis gaps in MLA backward (preserving the author’s affine mapping, using CuTe primitives for slicing and instruction projection, adding StMatrix and accumulating CopyReduceAsyncTensor semantics), closing #251. #257 fixes three defects in the post-lowering TIR memory access path: ops::reduce misidentifies the thread axis as cta-level, warps_per_group is taken as the whole grid instance count (256 instead of 8 in the RMSNorm tutorial), and cross-warp workspace indexing reads out of bounds.

Sources: https://github.com/tile-ai/TileFoundry/pull/263, https://github.com/tile-ai/TileFoundry/pull/260, https://github.com/tile-ai/TileFoundry/pull/257

3.14 TileFoundry: Split into core / nvidia / apple distribution packages; test inputs loaded once (#268 / #269)

#268 is a structural change across 739 files, +5489/-4518: core and vendor distributions are split apart under a shared version contract, with CUDA and AMX implementations, hardware docs, specs, templates and runtime headers packaged separately, emission and linking implementations resolved via target binding, and installed targets loaded after core initialization; model resources go to targets and shared causal language model orchestration to core, netting 274 lines removed after review. #269 targets test runtime (314 to 509 seconds on main, 651 seconds when runtime headers change) by running the same input through the CLI loader only once — previously a single scheduling HIR fixture was loaded 12 times, validated 9 times and analyzed once per family, and the test-private loader accepted programs the CLI would reject.

Sources: https://github.com/tile-ai/TileFoundry/pull/268, https://github.com/tile-ai/TileFoundry/pull/269

3.15 Distributions: Sunrise 0.1.15+sunrise.1.2.0 release PR merged (#9)

Merged 10-23 at 10:43. The PR updates the Sunrise distribution to 0.1.15+sunrise.1.2.0, incorporating TileLang 0.1.15 and the accompanying reviewed TileOPs, TileKernels, TVM and tvm-ffi changes, with source import and merge adaptation committed separately, finally adapting the TANG backend wheel and fixing a sliced GEMM regression (both accumulator flags validated against the public template, retaining stride, origin and offset assertions). The previous issue noted that after installing this PR on S2, the import smoke test failed at tangGetDeviceCount; there have been no further commits in that repo this window, and no fix for that failure has been recorded.

Source: https://github.com/tile-ai/tilelang-sunrise/pull/9

IV. Community, Tutorials, and Events

4.1 Community: tilelang-puzzles-ascend offers a PTO quantization kernel tutorial for Ascend (pushed 10-23)

Pushed 10-23, self-described as “a PTO quantization kernel tutorial and exercises for Ascend,” following the format of tile-ai/tilelang-puzzles. The same organization’s tilelangGym is used to measure the performance sensitivity of different benchmarks and microbenchmarks under PTO lowering, and was covered in the previous issue. Both repos belong to the PTO-ISA organization and, together with the PTO vectorization under review in the main repo (#3512, vectorizing T.Parallel into PTO VMI instructions), form a “compiler—exercises—sensitivity” trio, indicating that the PTO Ascend instruction line is filling in its developer entry points.

Source: https://github.com/PTO-ISA/tilelang-puzzles-ascend

4.2 New repo: Chinese TileLang operator interview practice site, 9 modules, 18 problems (10-23)

Created 10-23, self-described as a Chinese operator interview practice site: 9 modules, 18 problems, 21 teaching examples, including interactive exercises, solutions, and performance analysis. Previously, community content was mainly cookbooks and skill packs; this is the first artifact organized “by interview/course.”

Source: https://github.com/zhachar14h-cell/tilelang-operator-interview

4.3 Community skill and tutorial repos continue to push; four doc-site commits follow the first-release revamp (10-23~25)

sablin39/tilelang-cuda-skills (146 stars), RuneFang/tvm_tilelang_cookbook (36 stars), and cc-c122/TileLang-Remote-Autotuning-Agent (pushed 10-24) continue to update; kevin3113/models-tilelang (created 10-22, pushed 10-23), whose author also submitted schedmode support for Ascend cross-core synchronization to the main repo (#3614) in the same period. On the docs side, four commits to the TileOPs documentation site follow the main repo’s revamp: #79 removes the retired GQAPrefillPagedWithKVCacheFwdOp and posts a release announcement on the homepage, #81 follows the renaming of 26 operator classes and the underscore prefixing of internal packages, #82 follows changes to the public API contract and bench dependency groups, and #83 tidies the homepage and installation page for the first PyPI release.

Source: https://github.com/tile-ai/TileOPs.github.io/pull/83, https://github.com/sablin39/tilelang-cuda-skills, https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent

4.4 Discussion: TileScale user asks whether the TileLang version in use can be upgraded (#72, 10-25)

On 10-25 at 05:17, a user asked in tile-ai/tilescale whether the TileLang version that TileScale depends on can be upgraded. The repo’s last push was on 08-25, and its docs and examples still point to an earlier version. Such questions typically appear after upstream interfaces enter a period of rapid renaming—the TIRx replacement under review in the main repo (#3666) and the renaming of 26 operator classes in TileOPs in this window will both amplify downstream version-tracking costs.

Source: https://github.com/tile-ai/tilescale/issues/72

V. Trend Observations

Release engineering has become the main battleground. Of the 78 merges, TileOPs accounts for 51, of which more than ten are renames, packaging, semantic unification, docs, and CI gates—not performance or new operators, but freezing the interface. The tracking issue is 70% complete, the CHANGELOG is finalized for 0.1.0 with a 10-26 date, and from then on the 0.x compatibility commitment takes effect; in the same period, TileFoundry is split into three packages (core/nvidia/apple) and Sunrise merges the 0.1.15 distribution. All three lines do the same thing in the same week: turn “it runs” into “it ships.”

The focus of correctness is shifting from shape validation to new hardware and new scheduling. Last week’s main-repo merges were mostly shape and geometry validation; this window shifts to lifetimes and barriers: duplicate numbering of cross-storage-domain barriers, missing CTA barrier on sm_100 tensor memory deallocation, warp specialization moving the wrong prologue statement, and setmaxnreg triggering ptxas failures. The common thread is that compilation succeeds but results are intermittently wrong, and all of them land on automatic warp specialization and new-architecture storage—the bill that must be paid once scheduling is handed to the compiler.

The Ascend line moves from getting it running to laying out operators per model plus daily regression. situ_glu aligns with the Kimi K3 Stable LatentMoE gated activation, and the sparse_mla_fwd and GLM-5.3 kpool under review are likewise components of specific model structures. The daily full run of 2903 items all passing, followed by failures reported the next day, shows that the test suite is already large enough to serve as a gate; combined with the CI image tracking the weekly cann9.3.0 release, the conclusions on this line have a shelf life measured in weeks.

Execution backends subtract, while the backend list adds. CPU unifies on TVM-FFI, PTO defaults to tvm_ffi, plus CuTeDSL fixes and NVRTC compilation—all in the same direction: fewer execution-path forks. The backend list expands in the opposite direction—Intel AMX enters the CPU path, ROCm adds FP4 and E8M0, and Hygon re-proposes the register pipeline. The distribution form is shifting from “one codebase with multi-backend branches” to “one IR, multiple packages,” on the same path as #3666 making TVM an external dependency.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
tile-ai organization repositories 29 repos, 11 pushed within the window; 78 merged, 76 under-review updates, 19 new issues
tilelang main repo 10 merged (9 on 10-23, 1 on 10-25); VERSION still 0.1.15, no new tag; 50 under-review updates, including the TIRx replacement (#3666)
TileOPs 51 merged, centered on first-release cleanup and performance; no tag, no release; PyPI latest still 0.0.1.dev2 (2025-10-09)
TileFoundry 10 merged (#257 to #269); the v0.0.3 release notes were already covered in the 10-22 issue, no new tag in this window
tilelang-ascend 1 merged (situ_glu, default branch ascendc_pto), 24 under-review updates; daily tests passed all 2903 items, but the next day’s job failed
tilelang-hygon 3 merged on main (author dates 10-14 to 10-20, part of a batch merge); 2 under review (#18 / #27)
tilelang-sunrise 1 merged (the 0.1.15+sunrise.1.2.0 release PR), no subsequent commits after merge
tilelang-mlir-ascend 0 commits on the default branch main within the window; pushes occurred on three branches update0.1.15 / ci-0.1.15 / probe-ci, for version adaptation and CI probing
tilelang-metax / musa No pushes within the window, stalled at 10-16 / 09-30 respectively
TileRT / tilescale / DeepStack No pushes; tilescale received one version-upgrade inquiry (#72)
Adopters TileKernels / FlashQLA Both stalled at 10-14, no pushes
CANN community ops-tilelang operator repo Last commit on the default branch was 10-22, no additions within the window
Google News (multiple Chinese and English queries, via proxy) 1 exact hit (SGLang blog); full-text verification found no mention of TileLang, excluded as irrelevant
Hacker News 2 hits, both AI investment discussions, unrelated to the topic
arXiv Searched with TileLang as title or abstract keyword, latest was 2026-07, 0 within the window
PyPI tilelang 0.1.15 (09-30), tilefoundry 0.0.3 (10-21); tileops still 0.0.1.dev2
New community repos 1 new repo (tilelang-operator-interview), 1 first-push repo (tilelang-puzzles-ascend), and 4 active repos; creation dates and push times all verified

Full Source List