Research window: 2026-10-21 07:00 ~ 2026-10-22 07:00 (Beijing time, 24 hours).

Sources: GitHub tile-ai organization per-repo verification (Atom commit feed, covering tilelang / TileOPs / TileFoundry / tilelang-ascend [ascendc_pto] / tilelang-hygon [main and feature branches] / tilelang-metax [dev] / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack) + merged and in-review ground-truth searches (is:merged / is:open, grouped by repository) + release page tag verification + docs site commit verification + adopter repository verification (TileKernels / FlashQLA) + new community repository search + Google News multi-query searches in Chinese and English (via proxy) + Hacker News + arXiv + third-party industry analysis reports.

Issue Index

  • Today’s Highlights: TileFoundry 0.0.3 officially released, handing instruction-level scheduling to the agent; TileOPs merged 20 PRs in one day, with MLA backward and the SM89 small-scale path completed on the same day; the eight merges in the main repo are concentrated on compiler defects that “don’t error but produce untrustworthy results”; Ascend DeepGEMM alignment moved from under review to merged (10-21)
  • I. Core Project Progress
    • 1.1 Main repo: batched finalize_reducer previously reduced only the first output block (#3070) (10-22)
    • 1.2 Main repo: BF16→FP32 vectorized conversion switched to inline PTX packed lane conversion (#3595) (10-21)
    • 1.3 Main repo: unroll pass preserves empty stepped loops, ROCm Select lowered per lane (#2887) (10-21)
    • 1.4 Main repo: software pipelining preserves transitive copy dependencies, thread domains participate in region conflict determination (#2822) (10-21)
    • 1.5 Main repo: upfront validation of cast source/target/rounding combinations, rejects silent annotation drops (#3350) (10-21)
    • 1.6 Main repo: CUDA targets without arch are rejected instead of silently defaulting to sm_50 (#2860) (10-21)
    • 1.7 Under review: shift signedness and TIRx loop scope fixes submitted as paired commits in the tvm fork and main repo (#3597 / tvm #87 / #86) (10-21)
  • II. Multi-Backend Adaptation (Ascend / MetaX / Hygon / Moore Threads, etc.)
    • 2.1 Ascend evaluation: SemiAnalysis deep-dive report devotes a dedicated section to TileLang Ascend, stating that among six repos only TileLang was usable at the time (10-21)
    • 2.2 Main repo Ascend: DeepGEMM heuristic and epilogue alignment moved from under review to merged (#3577) (10-21)
    • 2.3 Main repo Ascend: integer signedness fix for warp_reduce_max/min on PTO (#3582) (10-21)
    • 2.4 Ascend repo: sub-32B compact row-stride sources cannot be expressed, keepdim reduce output corrupted (#1795) (10-21)
    • 2.5 Ascend repo: JIT cache switched to atomic publication, recovers when cache is deleted (#1903) (10-21)
    • 2.6 Ascend repo: cann-bench adds depthwise_conv_2d example (#1915) (10-21)
    • 2.7 Ascend under review: indexer_topk_reducesum, GLM-5.3 kpool compress, and WY fast operator migration (10-21)
    • 2.8 Hygon: feature branch advances FP8 decode GEMM, address chain identified as consuming 13% of kernel time (10-21)
    • 2.9 MetaX / Moore Threads / MLIR-Ascend: no pushes in this window, all three lines remain at existing versions (10-21)
  • III. Ecosystem and Adopters (TileOPs / TileFoundry / Operator Libraries)
    • 3.1 TileFoundry: 0.0.3 tagged and released, instruction-level scheduling + refactor analysis + H200 and Apple AMX examples (10-21)
    • 3.2 TileOPs: MLA prefill backward completed at op layer and SM90 warp-specialized kernel (#2678 / #2680 / #2688 / #2686) (10-22)
    • 3.3 TileOPs: dense GQA supports FP32 attention sink, -inf disables, +inf zeroes that head (#2515) (10-21)
    • 3.4 TileOPs: SM89 small-scale paths filled in as a batch, FP8 BMM default tile reduced to 64×64 (#2681 / #2672 / #2669 / #2654) (10-21)
    • 3.5 TileOPs: NSA top-k extended to 2048 compressed blocks on SM89 with output slots completed (#2674 / #2675 / #2673) (10-21)
    • 3.6 TileOPs: GDN forward completes in a single kernel when walk fills the SM, partitioned normalization splits wide rows by cluster (#2682 / #2679) (10-21)
    • 3.7 TileOPs: GQA backward directly serves grouped KV heads in the warp-specialized kernel (#2685) (10-21)
    • 3.8 TileOPs under review: folded edge axis single launch, NSA top-k narrowed column tile (#2689 / #2690) (10-22)
    • 3.9 Adopters: TileKernels and FlashQLA had no pushes in this window (10-21)
  • IV. Community, Tutorials and Events
    • 4.1 Community tutorial: Chinese “TVM / TileLang Cookbook” updated, covering everything from jit to cubin (10-21)
    • 4.2 New community repo: tilelang-wrapper-for-sail created, no description yet (10-21)
    • 4.3 Community tool: remote auto-tuning agent repo continues to receive pushes (10-21)
    • 4.4 Docs and versions: docs site auto-updated seven times, main repo VERSION still 0.1.15 (10-22)
  • V. Trend Observations

Today’s Highlights

This window merged 32 PRs (TileOPs 20, tilelang 8, tilelang-ascend 3, TileOPs docs site 1), plus one official release: TileFoundry tagged 0.0.3 on 10-21 at 10:26. The three lines have different centers of gravity. The eight merges in the main repo are almost entirely cleanup of defects that “don’t error but produce untrustworthy results”—batched reduction handling only the first output block, cast rounding annotations being silently dropped, CUDA targets without arch silently falling back to sm_50, pipelining losing transitive dependencies; TileOPs took MLA prefill backward from the op layer to an SM90 warp-specialized kernel in one day, while advancing four small-scale paths on SM89—FP8 BMM, W4A16, paged DSA, and NSA top-k—bringing the review queue down from 16 to 9; on the Ascend side, the DeepGEMM heuristic and epilogue alignment (#3577) that was still under review yesterday merged today, and the Ascend repo fixed a memory access modeling defect that caused corrupted keepdim reduce output (#1795). On the industry side, SemiAnalysis’s Huawei-DeepSeek deep-dive report devotes a dedicated section to TileLang Ascend, concluding that among the six repos with Ascend support, only TileLang was usable at the time.

I. Core Project Progress

1.1 Main repo: batched finalize_reducer previously reduced only the first output block (#3070)

Merged 10-22 01:16, 2 files +29/-5, fixes #2623. The lowering of the batched finalize_reducer called run_batch only once at the start of the buffer; the remaining consecutive output blocks were never reduced at all: when the output dimension exceeds the batch size, later segments received not the reduction result but uninitialized or stale values from the previous round. The fix emits a batched all-reduce with the correct offset block by block according to ceil(OutputDim / batch)—128 outputs with batch=4 emits 32 calls—while also switching the remapped reducer buffer to a flat view, addressing each batch with a constant linear offset, removing the previous shape and stride reconstruction, and retaining all chunk calls. Verification was done on an L40S: new sum/max/min numerical regressions cover fp16 and fp32 with batch 4/8/16, 23 finalize_reducer-related cases pass, and the external reproduction matrix shows zero mismatches across 16 cases. This is a classic “correct at small sizes, silently wrong at large sizes” class of defect.

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

1.2 Main repo: BF16→FP32 vectorized conversion switched to inline PTX packed lane conversion (#3595)

Merged 10-21 19:28, 4 files +19/-7. Adds tl::__tl_cvt_bf16x2_to_f32x2, which uses inline PTX to convert paired BF16 lanes into f32x2 in one go, replacing the per-lane scalar promotion path; all three lane widths of 2, 4, and 8 remain supported. The expected values for vectorized cast and layout inference were updated accordingly, and the BF16 cast tests were narrowed to the existing coverage without expanding the assertion surface. The change does not affect semantics, only the generated instruction sequence; the impact surface is every BF16 compute path that goes through vectorized cast.

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

1.3 Main repo: unroll pass preserves empty stepped loops, ROCm Select lowered per lane (#2887)

Merged 10-21 19:20. The loop unroll pass previously eliminated stepped loops with a trip count of 0 as empty shells, changing the program structure seen by subsequent passes; the fix preserves such empty loops during unrolling. The same commit also changes the vectorized Select on the ROCm side to lower per lane, avoiding whole-group lowering failures on unsupported targets. One change is in IR structure and the other in backend lowering; being combined in one commit indicates they were prompted by the same reproduction scenario.

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

1.4 Main repo: software pipelining preserves transitive copy dependencies, thread domains participate in region conflict determination (#2822)

Merged 10-21 18:13. Pipeline planning previously followed only direct dependencies, so indirect (transitive) copy dependencies were lost at the planning stage and the waits on dependent copies were deleted along with them; meanwhile, the physical region conflict check had two separate sets of criteria on the planning and lowering sides. The fix uses a backward worklist to follow both scalar and buffer copy dependencies, makes planning and lowering share a single physical region conflict check, keeps replay bindings within scope, and fills in coverage for transitive producers, strided aliases, and the sequential TMA fallback path. A second commit introduces thread domains into region conflict determination: accesses are first expanded to the materialized thread domain before regions are compared, so disjoint shared-memory slices can still be pipelined while cross-thread dependencies and physical overlap of strided buffers are preserved; the same domains are used for pipeline planning, dependency validation, and buffer versioning.

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

1.5 Main repo: upfront validation of cast source/target/rounding combinations (#3350)

Merged 10-21 13:49. round is only lowerable for explicitly listed combinations: PTX cvt.rs requires an f32 source, a 32-bit rbits operand, and a specific packed target. Other combinations were previously let through at branch dispatch, and the rounding annotation silently disappeared; scalar loops already rejected these combinations, but vectorized loops did not, so the same cast behaved differently in the two loop forms. A typical example is f32→E8M0: its helper rounds toward positive infinity and does not accept a rounding parameter at all, yet the vectorized branch never read round. Illegal combinations are now rejected uniformly before dispatch, so no branch can drop the annotation anymore; the accompanying tests pass the parameter only for the rs mode that genuinely accepts rbits.

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

1.6 Main repo: CUDA target without arch rejected instead of silently defaulting to sm_50 (#2860)

Merged 10-21 13:32, 1 file, fixes #1049. On machines with a CUDA toolkit but no usable GPU, a CUDA target without a specified arch degraded to a bare cuda, and TVM then silently filled in arch=sm_50, with compilation proceeding toward a wrong SM and no diagnostics throughout—exactly the trap of tilelang.compile(func) with target="auto". The fix lands in _cuda_target_from_arch, the single entry point shared by auto-detection and an explicit target="cuda", which now raises an actionable error stating the two ways to specify an explicit arch ({"kind":"cuda","arch":"sm_90"} or TILELANG_DEFAULT_TARGET). Headless compilation still works as long as a concrete arch is given, since codegen and nvcc depend only on the toolkit, with device access deferred until kernel launch. Three GPU-less unit tests were added at one point during the PR and removed in the final commit, leaving only the entry-point validation.

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

1.7 Under review: shift signedness and TIRx loop scope fixes submitted as a pair

Main repo #3597 fixes the loss of signedness in mixed-signedness shifts, and the corresponding change at the TIRx layer, “TIRx preserves shift left-operand type”, was submitted in the tile-ai/tvm fork as #87 in the same period; the two are the same logic at the language layer and the TIRx layer. Separately, tvm #86 backports the loop scope fix from upstream apache/tvm#19927. The fork has 25 PRs under review and the main repo 163; the TIRx side going first means the fork carries part of the upfront validation of changes before the main repo follows, the second time this collaboration path has appeared in this window.

Source: https://github.com/tile-ai/tilelang/pull/3597, https://github.com/tile-ai/tvm/pull/87, https://github.com/tile-ai/tvm/pull/86

II. Multi-Backend Adaptation (Ascend / MetaX / Hygon / Moore Threads, etc.)

2.1 Ascend Evaluation: SemiAnalysis Deep-Dive Report Devotes a Dedicated Chapter to TileLang Ascend

SemiAnalysis published “The Huawei-DeepSeek Symbiosis” on 10-21, with the subtitle directly listing “Huawei Ascend 950PR deep dive and microbenchmarks, DSv4.1 Flash analysis, DeepSeek’s Ascend support, AICPU, TileLang Ascend.” Three points are directly relevant to TileLang: First, among the six repositories that gained Ascend support (TileLang, DeepGEMM, DeepEP, TileKernels, FlashMLA, DeepSelect), only TileLang was usable at the time of writing, because CANN 9.2.0 had not yet been released, and DeepSeek requires CANN 9.2.0 plus a C++20 compiler, which 9.2.0-beta.2 does not provide; Second, the Ascend 950 backend has been merged into the TileLang main repository rather than existing only in a separate Ascend adaptation repository, providing native compilation, scheduling, synchronization, and runtime integration. AutoSchedule decides overlap and on-chip buffer slot counts based on data dependencies and estimated execution time, cutting up to 40% for small single operators on the 950PR; the compiler also automatically inserts synchronization, removes redundant waits, and checks conditional branches; Third, when compared with CANN and Triton-ascend, “certain sizes can match or exceed CANN,” with the optimal implementation switching by size (the author notes this is a naive benchmark under hot cache). The article also points out ecosystem gaps: most optimizations target super nodes, and vLLM-Ascend 0.26.0rc1 on x86 hosts still has multiple issues.

Source: https://newsletter.semianalysis.com/p/huawei-deepseek

2.2 Main Repo Ascend: DeepGEMM Heuristics and Epilogue Alignment Moves from Under Review to Merged (#3577)

Merged on 10-21 10:20, 29 files +1392/-1036. Recorded as under review in the previous issue, now merged: completes DeepGEMM-Ascend’s configuration selection and epilogue, covering fused bias, Alpha, BF16 stochastic rounding, and FP8 quantization for dense, batched, and grouped GEMM, and in the author’s own commits removes redundant setup and restores test gating. This corroborates the external judgment in 2.1 — what the Ascend side keeps polishing is precisely the GEMM path’s heuristics and numerical finalization, which is exactly the gap between “usable” and “runs well.”

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

2.3 Main Repo Ascend: Integer Signedness Fix for warp_reduce_max/min on PTO (#3582)

Merged on 10-21 10:21. PTO’s pto.redux_max/min rejects typed integers: PTOAS requires an unsignedness carrier plus an explicit signedness, so int32/uint32 T.warp_reduce_max/min fails to compile outright, reporting “requires signedness=’signed’ or ‘unsigned’ for integer values.” The fix routes warp_reduce_sum/max/min uniformly through contrib.ptodsl’s tl.redux_* wrappers (now exported), with the wrappers providing the carrier and signedness, and gates on value dtype in the manner of AscendAllReduce; the issue references main repo #3541 and PTOAS #1579. Verification is two-layered: lower-only codegen tests over an op × dtype grid, and sparse dual-target (ascend/pto) execution tests — using uint32 lanes ≥2^31 to directly expose signedness mistakes.

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

2.4 Ascend Repo: Sub-32B Compact Row-Pitch Source Cannot Be Expressed, keepdim reduce Output Corrupted (#1795)

Merged on 10-21 17:15, 2 files +302/-3, fixes #1682, and is the most technically substantial item in this window. The root cause is at the template layer: copy_ub_to_gm unconditionally generates a multi-burst DataCopyPad, while the MTE3 engine implements the system-level convention that “UB narrow rows are stored in 32B slots” — the source address advances by ceil(blockLen/32) + srcStride 32B blocks per burst. For producers that write rows on 32B boundaries this is an exact match; but T.reduce_max’s (M,1) fp32 output is written compactly by the V-pipe with a row pitch of only 4B, which the engine cannot express. In practice the second burst reads element 8 instead of element 1 (in [127, 1151, 71, 199], 1151 is exactly the value of element 8); the write side is fine — the error is in the DMA source address sequence. The fix has two halves: when both sides are packed and total ≤ 65535, collapse into a single flat burst (M×L bytes, exploiting the “single burst length is arbitrary” contract); non-expressible row pitches that are not multiples of 32B are changed to an explicit trap() abort, guarded by compile-time if constexpr, generating no extra code for 32B multiples and sub-32B slot layouts. Why not degrade: bisheng’s cross-pipe automatic synchronization only covers straight-line intrinsics; wrapping DataCopyPad in a conditional branch, or placing a GM scalar write in a sibling branch, causes bisheng to generate bad synchronization code and deadlock the device (in practice, per-row DMA wrapped in a runtime branch triggers an aicore exception). 14 test cases cover the flat/multi paths, the keepdim reduce reproduction, UB loop reuse, the legal single-burst form of non-expressible pitches, and the explicit abort; the test environment is Ascend910 A3 (dav-c220) / CANN 9.0.0 local build. The PR itself notes two remaining boundaries: sub-32B compact row pitch with a pitched destination (e.g., writing (M,1) fp32 into one column of a wide tensor) still uses slot stepping; blockCount > 4095 (12-bit field) still silently truncates (a 4200-row test copies only the first 104 rows), to be filed as a separate issue.

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

2.5 Ascend Repo: JIT Cache Switches to Atomic Publication and Recovers When Cache Is Deleted (#1903)

Merged on 10-21 16:20, 7 files +1026/-86, related to #1836. The original implementation overwrote build artifacts in place into kernel_lib.so; if another process had already loaded that shared library, it would read a file being rewritten. The fix first copies artifacts to a separate temporary file in the same directory, then publishes atomically with os.replace(); if the cached library has been deleted before loading, after confirming via the error message and file state it returns a cache miss and recompiles; if the directory disappears during the save phase, it skips saving and retains the compilation or tuning result already obtained; permission errors, missing dependency libraries, invalid ELF, and serialization internal exceptions continue to propagate upward rather than being swallowed as a miss. Coverage includes KernelCache, AutoTunerCache, and AutotuneResult, without changing cache keys, operator computation logic, or CI parallelism. 41 host-only test cases pass (9 for atomic publication, 32 for cache-deletion recovery); the PR itself notes no NPU-side integration verification was done. The atomic publication part follows the implementation of #1853 and retains the original author’s information.

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

2.6 Ascend Repo: cann-bench Adds depthwise_conv_2d Example (#1915)

Merged on 10-21 16:19. examples/cann-bench/ adds a depthwise 2D convolution benchmark. The accompanying smoke-test detail is worth noting: bench_test.sh judges PASSED by the output containing KERNEL OUTPUT MATCH or TEST PASSED!; the script previously printed “smoke test passed” with exit code 0 yet was judged a failure; the fix aligns with the LSTM/GRU smoke tests, preventing the new case from being persistently falsely red in CI.

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

2.7 Ascend Under Review: indexer_topk_reducesum, GLM-5.3 kpool compress, and WY fast Series Migration

The Ascend repo’s under-review queue has 13 updates in this window, concentrated in three groups. #1917 implements the indexer_topk_reducesum operator, corresponding to summing after top-k selection in the sparse attention indexer; #1910 adds GLM-5.3’s kpool compression operator, pointing to the KV compression path of the next-generation model; #1827 / #1828 / #1802 / #1803 / #1799 are a group of WY fast and chunk series operator migrations (WY fast itself and its backward, chunk_o backward, Chunk Scaled Dot KKT, Chunk Delta H), clearly aimed at the training chain of linear attention / gated delta networks. Also #1785 fixes indexing of multi-dimensional scalar BufferLoad in T.tile binary operations, #1854 validates automatic cross-core synchronization configuration, and #1731 fixes incorrect W4A4 int4×int4 results in gemm_v0 (#438).

Source: https://github.com/tile-ai/tilelang-ascend/pull/1917, https://github.com/tile-ai/tilelang-ascend/pull/1910, https://github.com/tile-ai/tilelang-ascend/pull/1827

2.8 Hygon: Feature Branch Advances FP8 Decode GEMM, Identifies Address Chain as 13% of Kernel

tilelang-hygon’s main branch has no new commits in this window (latest is still #25 merged on 10-20), but the feat/hcu-smls-b-vgpr branch has three pushes on 10-21 at 11:04 and 15:27, containing reusable performance-diagnosis and tuning conclusions. The 11:04 push unrolls the decode GEMM’s k-loop and returns M>96 tiles to bn96; the author’s key judgment is that “the fixed per-step overhead comes from the address chain, not the barrier” — the symbolic shared-level index (k+d) % 3 lowers to v_mul_lo_u16 / v_sub_u16 / v_readfirstlane_b32 and lands on matrix_load’s address path, consuming 13% of the kernel, whereas T.Pipelined would emit the same two barriers anyway. The 15:27 push prioritizes splitting the wide-N middle segment of M∈(64,128] by block count, with gate_up taking bm96/bn128/bk128/ns2 for M∈(64,96], and moves the builder into the in-package decode_cfg; it also reorganizes the decode GEMM record document from 3825 lines into a 588-line two-part “results + optimization log,” plus an index of closed levers and a bandwidth model.

Source: https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-smls-b-vgpr

2.9 MetaX / Moore Threads / MLIR-Ascend: No Pushes in Window

tilelang-metax (default branch dev, latest push 10-16, release still at v0.1.14 from 09-17), tilelang-musa (latest 09-30, v0.1.15+musa.1), and tilelang-mlir-ascend (latest 10-19, tag still at v0.1.2.020.sglang.poc from 09-09) — all three repos have no commits in this window and no new tags.

Source: https://github.com/tile-ai/tilelang-metax, https://github.com/tile-ai/tilelang-musa, https://github.com/tile-ai/tilelang-mlir-ascend

III. Ecosystem and Adopters (TileOPs / TileFoundry / Operator Libraries)

3.1 TileFoundry: 0.0.3 tagged and released — instruction-level scheduling + refactored analysis + H200 and Apple AMX examples

The previous issue noted that “the 0.0.3 announcement is already in the README, tag still pending”; this issue, on 10-21 at 10:26, it was officially released, with PyPI as tilefoundry==0.0.3, codename “Agent-driven scheduling”. The core change is where scheduling authority resides: the compiler lists which instructions can implement each HIR operation and what each requires, the agent selects among the candidates, and the compiler then lowers the selected schedule to TIR. The accompanying CLI has tilefoundry schedule candidates and schedule facts to display candidates and constraints, tf.schedule records the selection, and tilefoundry schedule finalize completes the lowering; instructions are declared as AtomSched operations with explicit TIR operand contracts, starting from CUDA MMA atoms, SM90 WGMMA, and async copies. The analysis portion was refactored into memory (segmented allocation, footprint and reuse windows, loop-aware buffer lifetimes), cost (work and traffic counted by execution scope, each access relation declaring what it reads), and structure (definition facts, valuation contexts, and scope snapshots unified through visit context, with projection reuse so each corpus case is analyzed only once). The new tutorial walks a single 8192×5120 by 5120×17408 GEMM all the way to the optimal AtomSched kernel on H200.

The performance side gives three sets of real numbers: Qwen3.8-27B prefill on a single H200 at 1024 / 4096 / 8192 tokens is 80.29 / 309.61 / 631.19 ms respectively, versus SGLang’s 89.06 / 317.11 / 638.25 ms; short sequences at 128 / 256 / 512 / 1024 tokens are 1.88× / 1.41× / 1.10× / 1.03× SGLang respectively; the Qwen prefill example on Apple M2 Pro implements the large projection as an FP16 AMX outer product. Both Qwen3.8-27B runs produce the same first token as transformers at every length. Two known issues: the wheel does not yet package TileFoundry’s CUDA runtime headers, so tf.compile cannot build CUDA kernels from an installed package (a source checkout is required; fixed in the next version), and the performance timeline misprices insert_slice / tf.zeros across mesh locations and does not show mesh lane underfill (#239).

Source: https://github.com/tile-ai/TileFoundry/releases/tag/v0.0.3

3.2 TileOPs: MLA prefill backward completes the op layer and SM90 warp-specialized kernels (#2678 / #2680 / #2688 / #2686)

Four PRs form one line. #2678 (10-21 11:34) adds the two op layers MLADenseBwdOp and MLAVarlenBwdOp, aligned with FA3’s _flash_attn_backward interface; #2680 (10-21 20:58) adds the SM90 kernels, providing two warp-specialized implementations, MLADenseBwdWSKernel and MLAVarlenBwdWSKernel, for D_qk=192, D_v=128; #2688 (10-21 23:28) changes the resident-CTA fetch order of MLAVarlenPrefillWSFwdKernel — previously each head was first given a query tile, and after the first tile finished it could only queue head by head, now only query tiles of the same head are processed at a time; #2686 (10-22 03:08) does final cleanup, moving single-reader constants and helpers to their use sites, with no behavior change. Docs site #78 also adds the two new ops to the Attention page, prompted by check_api_pages.py reporting two missing pages among the 207 ops exported by the main repo.

Source: https://github.com/tile-ai/TileOPs/pull/2678, https://github.com/tile-ai/TileOPs/pull/2680, https://github.com/tile-ai/TileOPs/pull/2688

3.3 TileOPs: dense GQA supports FP32 attention sink (#2515)

Merged 10-21 22:42. Adds an optional per-head FP32 attention sink for dense GQA prefill and decode, covering FP16/BF16/FP8 on SM90, RoPE, softcap, and sliding window; other architectures reject sink calls outright. Semantically each sink is an unscaled, non-softcapped logit in the denominator: -inf disables that sink, +inf directly zeroes that head’s output, and split decode must preserve sink weights under large logits. Also fixes a benchmark distortion — the chunked prefill benchmark was previously being measured as single-query decode; the fused-RoPE tail cases adjust KV length to 271/1025 to cover split execution.

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

3.4 TileOPs: concentrated fill-in of SM89 small-scale paths (#2681 / #2672 / #2669 / #2654)

All four merged PRs address L40S/4090-class gaps. #2681 changes the default tile for FP8 BMM on SM89 to 64×64 — conditioned on the result grid having at most one block per SM; large grids keep 128×128; the author gives single-GPU measurements on PyTorch 2.13.0 + CUDA 13.0 + TileLang 0.1.12: 1×16×64×128 drops from 5.161 to 3.799 us, 16×128×128×1024 from 12.943 to 7.414 us, while 4096³ goes from 512.573 to 583.680 us instead (that row’s tile is unchanged; this is noise). #2672 raises the default M upper bound for 64-column W4A16 from 16 to 64, with replay times at M=17 / 32 / 64 dropping 43.35% / 47.20% / 40.00% respectively, while retaining explicit configuration overrides and numerical tolerance. #2669 implements SM89 paged DSA with MMA, preserving the 656-byte FP8 latent cache layout, BF16 rounding, natural-log LSE, and dead position behavior: batch 4 / 128 head / 256 positions is 0.269 ms versus 0.727 ms for the PyTorch FP32 reference, and 48×64 head×2048 positions is 2.543 versus 2.535 ms; the author explicitly writes that “the large case is essentially at parity and does not constitute a general advantage”. #2654, for small token counts, computes and moves only the covered 8-token accumulator prefix, opening 40-token decode up to M33–40 on a 132-SM H200 with N divisible by 128, 133–225 output tiles, and K∈{7168, 7680, 8192}.

Source: https://github.com/tile-ai/TileOPs/pull/2681, https://github.com/tile-ai/TileOPs/pull/2672, https://github.com/tile-ai/TileOPs/pull/2669, https://github.com/tile-ai/TileOPs/pull/2654

3.5 TileOPs: NSA top-k extended to 2048 compressed chunks on SM89 and output slots completed (#2674 / #2675 / #2673)

Three PRs revolve around the same SM89 path. #2674 lets group48 reuse the existing tiled top-k program: one tile per token (48 query head rows), column-wise GEMM warp layout; #2675 raises top-k selection on SM89 from the previous cap to 1024–2048 compressed chunks, covering head dim 64 / 128 / 256; #2673 fixes output filling when candidates are insufficient — every requested top-k slot must be filled with -1, whereas previously the 32-slot sorting band would leave stale values in slots 32–63. Separately, #2684 adds an upper bound to the NSA sorting capacity for packed requests, using min(total_chunks, ceil(total_tokens / block_size)) and then aligning to the scoring tile.

Source: https://github.com/tile-ai/TileOPs/pull/2674, https://github.com/tile-ai/TileOPs/pull/2675, https://github.com/tile-ai/TileOPs/pull/2673

3.6 TileOPs: GDN single-kernel prefill and quantized normalization splitting (#2682 / #2679)

#2682 targets GDN forward for linear attention: its prefill was slower than FlashInfer on 79 rows in the nightly batch (0.57×–0.98×); the fix is to fold prefill into a single kernel when “the walk can fill the SMs”. #2679 handles FusedAddRMSNormPerTokenQuant: about 25 SASS instructions per element are spent on elementwise math; the change packs the elementwise math and splits wide rows across clusters. Both PRs replace “serial overhead within a single operator” with “wider parallelism and fewer instructions”, the same idea as the SM89 small-scale path tuning group the same day.

Source: https://github.com/tile-ai/TileOPs/pull/2682, https://github.com/tile-ai/TileOPs/pull/2679

3.7 TileOPs: GQA backward serves grouped KV heads directly in the warp-specialized kernel (#2685)

Merged 10-21 22:58. On SM90 with head dim 128, GQABwdOp previously handed grouped KV heads to the pipelined kernel, going CTA by CTA per query head; now the warp-specialized GQA backward kernel consumes grouped KV heads directly, reducing redundant movement and scheduling overhead. The same day, #2687 fixes device backend auto-loading in the PrivateUse1 rename regression test, a test-infrastructure cleanup.

Source: https://github.com/tile-ai/TileOPs/pull/2685, https://github.com/tile-ai/TileOPs/pull/2687

3.8 TileOPs under review: fold edge axis into a single launch, narrow column tiles for NSA top-k (#2689 / #2690)

Two PRs are still under review at the end of the window. #2689 targets reduction: fold the edge axis into a single launch and issue the fold’s tail loads first; #2690 uses narrower column tiles for NSA top-k on SM89. After 20 merges, TileOPs’ review queue dropped from 16 in the previous issue to 9, the only repository in the organization with a clear decline in review count this window.

Source: https://github.com/tile-ai/TileOPs/pull/2689, https://github.com/tile-ai/TileOPs/pull/2690

3.9 Adopters: no pushes from TileKernels or FlashQLA this window

deepseek-ai/TileKernels (1991★) and QwenLM/FlashQLA (732★) had no commits this window, with their latest pushes both stopping on 10-14; FlashQLA’s version remains v0.1.3 from 09-30. The quiet on the adopter side has lasted a week, with momentum still inside the tile-ai organization.

Source: https://github.com/deepseek-ai/TileKernels, https://github.com/QwenLM/FlashQLA

IV. Community, Tutorials, and Events

4.1 Community Tutorial: Chinese “TVM / TileLang Cookbook” Updated

RuneFang/tvm_tilelang_cookbook (34★, repo created 7-21) had a push on 10-21. Its self-description is clear: aimed at readers who “can write some Python and have heard of CUDA but have never read compiler code,” it walks step by step from @tilelang.jit all the way down to .cubin. Chinese-language material that explains the full chain from DSL to compilation artifact remains scarce, and is especially useful for newcomers to new backends such as Ascend and Hygon.

Source: https://github.com/RuneFang/tvm_tilelang_cookbook

4.2 New Community Repo: tilelang-wrapper-for-sail Created, No Description Yet

flytiger-eco/tilelang-wrapper-for-sail was created on 10-21 at 09:04, with no description and no README; the repo name points to providing a TileLang wrapper for SAIL. Small repos of this kind—”vendor-internal runtime + TileLang shell”—are usually outposts for new hardware onboarding, but the contents are not yet public, so this issue only records it.

Source: https://github.com/flytiger-eco/tilelang-wrapper-for-sail

4.3 Community Tool: Remote Autotuning Agent Repo Continues to Receive Pushes

cc-c122/TileLang-Remote-Autotuning-Agent (1★, repo created 7-01) had a push on 10-21 at 22:57. Consistent with TileFoundry’s direction of handing scheduling to an agent, the community side is also trying to make autotuning a remotely drivable agent rather than a local script sweep.

Source: https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent

4.4 Docs and Versions: Docs Site Auto-Updated Seven Times, Main Repo VERSION Still 0.1.15

tilelang.github.io auto-updated seven times within the window (10-21 at 10:18, 11:26, 11:33, 13:54, 18:26, 19:23, and 10-22 at 01:23), all triggered by merges into the main repo. The main repo’s VERSION is still 0.1.15 (released 09-30), and the README’s latest milestone is still the 09-30 Ascend 950 backend; TileOPs has no new tag; TileFoundry added the v0.0.3 tag today, the only new version in the window. Although TileOPs-nightly’s pushed_at changed, its default branch has had no new commits since 08-30—the change came from a non-default branch and is not counted as development activity.

Source: https://github.com/tile-ai/tilelang.github.io, https://github.com/tile-ai/tilelang

V. Trend Observations

The nature of compiler defects is expanding from “computing the wrong number” to “choosing the wrong target.” Of the eight merges into the main repo this issue, four change “what the compiler thinks it is doing”: a CUDA target without an arch silently falls back to sm_50 (#2860), cast rounding annotations are silently dropped by branch (#3350), batched reductions process only the first output block (#3070), and pipeline planning loses transitive dependencies (#2822). What they share is that they do not raise errors and only produce untrustworthy results under certain sizes or certain loop shapes—harder to locate than crashes, and also showing that the project places silent correctness ahead of feature expansion.

The bottleneck on the Ascend line has shifted to the memory model and runtime boundaries. Sub-32B compact row spacing cannot be expressed via DMA and can only be explicitly aborted (#1795); cache concurrency is changed to atomic publication with deletion recovery added (#1903); PTO’s integer signedness has to be patched up by a wrapper layer (#3582). None of these three is “adding an operator”; they are “keeping existing operators from failing under boundary conditions.” Together with the merged DeepGEMM heuristic alignment, the Ascend side is shifting from capability coverage to engineering reliability.

Analyst firms are beginning to treat TileLang as the criterion for Ascend usability. SemiAnalysis uses “only TileLang among the six repos was usable at the time” to define the completeness of DeepSeek’s Ascend support—more than any single technical blog post, this shows TileLang’s position in the domestic-compute narrative: from an experimental DSL to the default carrier layer. The same report also gives a reverse signal—optimization is concentrated on supernodes, x86 host support is still incomplete, and generality lags behind performance.

The center of gravity of the attention stack is moving from inference forward passes to training backward passes. Within a single day, TileOPs filled in the op layer for MLA prefill backward, the SM90 warp-specialized kernel, and data-fetch order optimization, plus grouped KV head handling for GQA backward—fully aligned with the backward-operator backfill on the Ascend side since last month (indexer_bwd, SiTU-GLU backward, WY fast backward). TileLang’s coverage on the training side is being systematically filled in.

SM89 is moving from “it runs” to “it’s tuned.” Four lines—FP8 BMM tiling, the M upper bound for W4A16, paged DSA, and NSA top-k—were all worked on the same day, and all give single-card measurements along with qualifiers such as “does not constitute a general advantage.” Consumer-grade and single-card inference scenarios have become one of the performance targets, and at the same time the authors’ self-discipline on benchmark methodology has clearly tightened.

The division of labor between scheduling and correctness is becoming clear between TileFoundry and the main repo. The main repo continues to tighten semantics and correctness, while TileFoundry 0.0.3 hands “which instruction to choose” to the agent and leaves “guaranteeing semantics” to the compiler, giving reproducible numbers with two examples: H200 and Apple AMX. If this division stabilizes, the TileLang ecosystem will develop a two-layer structure: “the language layer guards correctness, the agent layer fights for performance.”

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
tile-ai/tilelang (main) 8 commits and 8 merges within the window; VERSION still 0.1.15; 163 open PRs, 21 of which were updated within the window
tile-ai/TileOPs (main) 20 commits and 20 merges within the window; 9 open PRs, 3 updated within the window
tile-ai/tilelang-ascend (ascendc_pto) 3 commits and 3 merges within the window; 235 open PRs, 13 updated within the window
tile-ai/tilelang-hygon (main / feat-hcu-smls-b-vgpr) No commits on main; 3 pushes to the feature branch (10-21 11:04, 15:27); 3 open PRs
tile-ai/TileFoundry (main) No new commits; v0.0.3 tagged and released at 10-21 10:26; 1 open PR
tile-ai/tilelang-metax (dev) / musa / mlir-ascend No commits in any of the three repos within the window (latest 10-16, 09-30, 10-19 respectively); no new tags
tile-ai/TileRT / tilescale / DeepStack No commits in any of the three repos within the window (latest 09-24, 08-25, 10-05 respectively)
deepseek-ai/TileKernels / QwenLM/FlashQLA No commits; latest pushes in both remain at 10-14
tilelang.github.io / TileOPs.github.io 7 / 1 automated updates respectively, all triggered by merges to the main repos
TileOPs-nightly No new commits on the default branch since 08-30; pushed_at changes come from non-default branches
Google News (multiple Chinese and English queries) / Hacker News / arXiv Three industry reports within the window center on Huawei-DeepSeek collaboration; only the SemiAnalysis piece contains a dedicated TileLang section; the latest relevant arXiv paper remains 07-24

Additional notes: The Hygon pushes are not on the default branch. They were identified by first pulling the commit stream of the feat/ branch, then cross-checking the PushEvent ref and before→head via the repository events API; pushed_at changes on TileOPs-nightly were excluded the same way. Open-PR counts are the total is:open search results per repo; within-window update counts are the total updated:>=window start search results. Commits landed in the half-day before this window’s start (10-21 07:00) are covered by the previous report and are not counted again here.

Full Source List