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

Sources: Repository-by-repository check of the GitHub tile-ai organization (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon branch / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / three documentation sites) + ground-truth search of merged and under-review items (two sets, is:merged and is:open, tallied by repository) + commit-level diff and release-page tag checks + adopter repository checks (TileKernels / FlashQLA) + community new-repository search + multiple Google News queries in Chinese and English (via proxy) + Hacker News + arXiv + industry media coverage.

This Issue’s Index

  • Today’s Highlights: Among 35 merged PRs, the main repo shifts focus to analyzer lifecycle; DeepSeek V3.2’s top-k selector out-of-bounds write traced to CUDA_ERROR_ILLEGAL_ADDRESS; Ascend NZ layout enables partial UB→L1 copies (18 files); CANN community opens new ops-tilelang operator repo
  • I. Core Project Progress
    • 1.1 Main repo: Clone analyzer before vectorization rewrite (#3612) (10-22)
    • 1.2 Main repo: Old analyzer binding invalidated after parallel loop splitting (#3226) (10-22)
    • 1.3 Main repo: deepseek_v32 top-k selector candidate buffer out-of-bounds (#3280) (10-22)
    • 1.4 Main repo: CUDA 16-bit inverse trig and binary math fallback completed (#3172) (10-22)
    • 1.5 Main repo: sm8x sparse metadata non-bijective shapes now rejected (#3125) (10-22)
    • 1.6 Main repo: Eager frontend preserves Python source locations (#3558) (10-22)
    • 1.7 Under review: Pass capability checks, unsigned mulhi, and assume predicates (10-22)
  • II. Multi-Backend Adaptation (Ascend / Hygon / MetaX / Moore Threads)
    • 2.1 Ascend: NZ buffer supports partial UB→L1 copies (#3530) (10-22)
    • 2.2 Ascend PTO: MAD direction predicates grow exponentially, stalling compilation (#3581) (10-22)
    • 2.3 Ascend PTO: SIMT vector local variables allocated per lane (#3523) (10-22)
    • 2.4 Ascend repo: cann-bench adds Dilation2D, reverts depthwise_conv_2d same day (#1916 / #1919) (10-22)
    • 2.5 Ascend repo under review: sparse_mla forward and backward (#1929 / #1928) (10-22)
    • 2.6 Hygon: Merge branch removes enable_register_pipeline (10-22)
    • 2.7 MetaX / Moore Threads / MLIR-Ascend: No pushes in window (10-22)
    • 2.8 Industry: SCMP republishes SemiAnalysis report; TileLang the only usable among six repos (10-22)
  • III. Ecosystem & Adopters (TileOPs / TileFoundry / Operator Libraries)
    • 3.1 TileOPs: FP8 paged warp-specialized kernel shared memory aliasing causes NaN (#2708) (10-22)
    • 3.2 TileOPs: Dense prefill switched to persistent traversal, D64 three consumers (#2706) (10-22)
    • 3.3 TileOPs: SM90 dense GQA decode uses warp specialization (#2694) (10-22)
    • 3.4 TileOPs: Causal sliding window serves entire query head group with two key tiles from one KV head (#2696) (10-22)
    • 3.5 TileOPs: Paged decode autotuning previously shared a single block table (#2703) (10-22)
    • 3.6 TileOPs: Masked softmax FMA overflow and varlen RoPE odd half-width (#2699 / #2700) (10-22)
    • 3.7 TileOPs: FP8 grouped GEMM staged scale buffering and gate/up ping-pong warp group (#2704 / #2693) (10-22)
    • 3.8 TileOPs: Predicate broadcast and boolean operands switched to 16-byte vector reads (#2702 / #2697) (10-22)
    • 3.9 TileOPs: Three SM89 paths advanced same day (#2705 / #2691 / #2690) (10-22)
    • 3.10 TileOPs: Reduction folding single launch and empty DSA selection validation (#2689 / #2698) (10-22)
    • 3.11 TileOPs: Benchmark reference computed in float32 (#2692) (10-22)
    • 3.12 TileFoundry: Install sm_89 architecture and RTX 6000 Ada device (#259) (10-22)
    • 3.13 TileFoundry: gmem pricing changed to footprint and runtime mesh, bound to callee dimensions (#258) (10-22)
    • 3.14 TileFoundry: Must raise error when device refuses to launch kernel (#256) (10-22)
    • 3.15 Distribution: Sunrise opens release PR for 0.1.15+sunrise.1.2.0 (#9, under review) (10-22)
    • 3.16 Adopters: TileKernels and FlashQLA no pushes in window (10-22)
  • IV. Community, Tutorials & Events
    • 4.1 CANN community opens ops-tilelang operator repo (10-22)
    • 4.2 tilelang/agent: Operator optimization written as agent SKILL workflow (10-22)
    • 4.3 New community repos: Sail wrapper, remote autotuning agent, and tutorial repo (10-22)
    • 4.4 Docs & versions: Docs site updated three times, VERSION still 0.1.15 (10-22)
  • V. Trend Observations

Today’s Highlights

In this window, the tile-ai organization merged 35 PRs (TileOPs 20, tilelang 9, TileFoundry 3, tilelang-ascend 2, TileFoundry docs site 1), with another 71 open PRs updated in the window. The main repo’s nine PRs shift focus from last week’s “shape and geometry validation” to analyzer lifecycle: cloning the analyzer before vectorization rewrite (#3612, 2 lines in a single file), old bindings invalidated after parallel loop splitting (#3226, 6 files), and eager frontend preserving Python source locations (#3558)—none of these cause kernels to compute incorrectly, but rather cause the compiler to read stale state after rewrites. On the numerical side, the most notable is #3280: deepseek_v32’s top-k selector stages radix candidates in a fixed [2, 4096] shared buffer without constraining the size of the threshold bucket selected in the first round; FP32 scores with identical low mantissa bits or fully identical inputs cause out-of-bounds writes, crashing with CUDA_ERROR_ILLEGAL_ADDRESS. On the Ascend side, #3530 is the day’s largest PR (18 files, +993/-256), preserving compressed NZ padding through storage normalization and explicitly providing C0 group row stride for partial UB→L1 copies; #3581 fixes a predicate expression exponential growth on PTO that could stall compilation for over ten minutes. Two new ecosystem threads: the CANN community created the ops-tilelang operator repo on 10-20, with its first fix pushed on 10-22; the Sunrise distribution opened a release PR for 0.1.15+sunrise.1.2.0, incorporating the TANG backend wheel into distribution.

I. Core Project Progress

1.1 Main repo: Clone analyzer before vectorization rewrite (#3612)

Merged on 10-22 23:43, changing only the vectorization source location, +2/-1. Loop vectorization reuses the passed-in arithmetic analyzer when rewriting IR, but the analyzer caches constant and range constraints from before the rewrite; once vectorization changes loop bounds or index expressions, subsequent passes still receive stale conclusions. The fix is to copy the analyzer before entering the vectorization rewrite, so the rewrite only affects the copy. This commit ships no regression test; its value lies in plugging the same class of problem as #3226 from another entry point—that commit, landed the same day, addressed stale bindings after parallel loop partitioning. Taken together, the two commits show the lifetime of analysis contexts being cleaned up as an independent defect class.

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

1.2 Main repo: Stale analyzer bindings after parallel loop partitioning (#3226)

6 files, +55/-19, touching the loop_partition, loop_vectorize, and lower_tile_op passes and their headers. After a parallel loop is partitioned, the loop body’s iteration variables are rebound, but the pass still queries the analyzer using the pre-partition variables, retrieving range and divisibility information for the old variables, which distorts simplification and out-of-bounds checks. The fix isolates analysis contexts per pass and simplifies the initialization path of loop analysis snapshots, with a set of minimal parallel-loop analysis regression cases added. The commit includes a Codex co-signoff.

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

1.3 Main repo: deepseek_v32 top-k selector candidate buffer out-of-bounds (#3280)

2 files, +91/-10, fixing #1351. tl_topk_impl packs radix candidates into a fixed [2, 4096] shared buffer, but nothing constrains the threshold bucket selected by the first radix round to at most 4096 elements—when FP32 scores share low mantissa bits or inputs are all equal, atomic appends write past the buffer end and crash with CUDA_ERROR_ILLEGAL_ADDRESS. The fix allocates a (batch, 2, seq_len) global workspace in the tl_topk wrapper so both radix buffers are bounded by seq_len; the public tl_topk() signature is unchanged. It ships a 79-line reproduction case built from the #1351 inputs.

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

1.4 Main repo: CUDA 16-bit inverse trig and binary math lowering completed (#3172)

3 files, +463/-8. asin, acos, asinh, acosh, atanh, log1p, atan2, hypot, copysign, ldexp, and nextafter lacked CUDA-specific lowering rules; the generic rules only handle float32/float64, leaving float16 and bfloat16 calls unhandled at codegen. The fix registers them in CUDAMath and adds half_t / bfloat16_t overloads: the ten operators evaluate in float32 and convert back, while nextafter steps via integer comparison on the 16-bit encoding to preserve bfloat16 subnormals under fast math. float32/float64 function names are unchanged, and 182 lines of template implementations are added on the common.h side.

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

1.5 Main repo: Reject non-bijective shapes in sm8x sparse metadata (#3125)

Single file tilelang/layout/gemm_sp.py, +6/-0, fixing #2606. The interleave in the SM8x cutlass metadata layout builder is bijective only when M is a multiple of the group size and K is even; otherwise it previously generated incorrect metadata silently, making sparse GEMM results wrong without warning. Both conditions are now explicit rejections. The sparse metadata shape test originally added with the PR was requested to be removed during review and did not enter the main branch.

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

1.6 Main repo: Eager frontend preserves Python source locations (#3558)

4 files. After the eager frontend lowers Python for / with statements into generated calls, errors and debug mappings point back to generated-code line numbers. The fix maps generated loop calls back to for statements and context wrappers back to with statements, and adds locations for unpacked iterable arguments; as with #3125, review required removing the accompanying source-location regression test.

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

1.7 Under review: pass capability checks, unsigned mulhi, and assume predicates (10-22)

The main repo’s review queue saw 50 updates in the window. Notable ones to follow: [Transform] Locate VerifyBufferInit reads outside tile ops (#3620), [Codegen] Add unsigned mulhi intrinsic (#3617), [Transform] Keep assumption checks with local dependencies in scope (#3611) and [Transform] Preserve tl.assume predicates that simplify to true (#3610), [CPU] Support clz in C and C host code generation (#3609); on the Ascend side, [BugFix][Ascend] Reject odd FP4 view geometry in GM->L1 copies (#3608) and [Ascend] Support exact placement of versioned on-chip buffers (#3607) were updated the same day. Also #3613, a native argmin/argmax reduction draft explicitly marked “do not review.”

Source: https://github.com/tile-ai/tilelang/pulls

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

2.1 Ascend: NZ Buffer Support for Partial UB→L1 Copies (#3530)

The largest single change of the day: 18 files, +993/-256. Padding of the compact NZ layout was previously flattened at the logical layout stage, so partial UB→L1 copies could not express the respective C0 group row strides of source and destination. The fix preserves compact NZ padding through the storage normalization stage and then materializes it exactly once by capacity; when the copy region is lowered, it explicitly provides source/destination C0 group row strides in 32-byte blocks, and threads the row stride through the scatter template, Ascend/PTO DMA codegen, and access footprints. It also adds geometry, alignment, and out-of-bounds checks, preserves the physical view owner, and migrates callers to logical allocation shapes plus explicit row strides. src/ascend/op/nd2nz.h gains 83 lines.

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

2.2 Ascend PTO: Exponential Growth of MAD-Direction Predicates Stalls Compilation (#3581)

2 files, +48/-5. A PTO kernel with repeated mutable guards can spend over ten minutes in analysis even without any matrix multiplication — a compile-time stall before NPU execution. The root cause is that PTOMmadDirectionAnalyzer conservatively preserves the incoming state when a guard reads mutable memory; if both branches retain the MAD direction, Merge receives the same predicate S, and the implementation always constructs and simplifies S || S. The arithmetic simplifier cannot eliminate this duplication in the reproducer, so consecutive branches cause the predicate expression to double repeatedly. The fix skips rebuilding when both sides’ predicates are identical, and adds 20 lines of targeted tests.

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

2.3 Ascend PTO: SIMT Vector Local Variables Allocated Per Lane (#3523)

4 files, +273/-29. SIMT vector local variables were allocated as scalar PTODSL buffers without accounting for lane count, yet were accessed with Python subscripts; AllocatedBufferValue does not support per-item assignment, so loop-index writes failed even earlier during the PTODSL AST rewrite stage. The fix allocates enough scalar elements per vector slot and translates slot subscripts into scalar offsets, reusing the existing contiguous load/store path, preserving AscendC’s vector-local semantics, without changing loop unrolling or adding a new lowering pass. It adds two codegen test groups, float32x2 and float32x4, totaling 209 lines.

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

2.4 Ascend Repo: cann-bench Adds Dilation2D, Reverts depthwise_conv_2d the Same Day (#1916 / #1919)

#1916 merges Dilation2D forward (a cann-bench level3 task): NHWC, fp16/fp32, arbitrary kernelSize/stride/rate, and SAME/VALID; the contribution of out-of-bounds taps takes -inf + filter and matches the golden bit-for-bit, NaN inputs propagate, and a +inf filter tap turns out-of-bounds contributions into NaN. It is implemented as a pure Vector kernel plus a Python wrapper, dispatched according to the Ascend copy engine’s 32B row-stride alignment constraint, with the grid adapting to geometry rather than hardcoding core counts. The same day, #1919 fully reverts depthwise_conv_2d (#1915), which had just been merged on 10-21; the commit message gives no reason, leaving the examples repo in a “one in, one out” state.

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

2.5 Ascend Repo Under Review: sparse_mla Forward and Backward (#1929 / #1928)

#1929 adds sparse_mla_fwd (DSA sparse MLA forward: varlen packed random-row gather, online softmax, fused QK/PV dual GEMM), placed under examples/dsa_sparse_finetune/, with dual kernels (a fast specialization plus a portable general one) and host dispatch; the fast path’s applicability domain is hardcoded to H=128 with G=1, dim/tail of 512/64, topk a multiple of 256, and NI between 1 and 6, with anything outside the domain taking the general path. #1928 adds sparse_mla_bwd (dQ GEMM chain, P/dP vector chain, dKV random-row scatter), placed under examples_experiment/dsa/ and registered in the operator test list; the forward result k1 has already been folded into this operator. The Ascend repo has 11 PRs under review in total.

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

2.6 Hygon: Merge Branch Removes enable_register_pipeline (10-22)

The merge branch gcy_merge_main pushed a 6-file, +10/-24 cleanup on 10-22 at 11:11, deleting the enable_register_pipeline switch entirely, touching the HCU register soft-pipelining injection and pipeline planning passes, the pipeline utility header, and the language-level loop extension, with examples and performance scripts updated to match. Removing a switch usually means the register pipelining path is now on by default and dual-track maintenance is no longer worthwhile; however, main is still at the 10-20 commit titled “preserve constant K origins in hoisted MLS windows” (#25), and this batch is still waiting to land on the merge branch.

Source: https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main

2.7 MetaX / Moore Threads / MLIR-Ascend: No Pushes in the Window (10-22)

tilelang-metax’s last push was on 10-16, tilelang-musa’s on 09-30, and tilelang-mlir-ascend’s on 10-19; none of the three adaptation lines had any commits in the window. The Ascend and Hygon lines accounted for all backend changes in this window.

Source: https://github.com/tile-ai/tilelang-metax

2.8 Industry: SCMP Reposts SemiAnalysis Report, TileLang the Only Usable of Six Repos (10-22)

South China Morning Post published “Jensen’s nightmare” on 10-22, relaying SemiAnalysis’s Thursday Huawei-DeepSeek report: DeepSeek’s latest software release plus more experiments on Ascend 950PR show the gap with the US is narrowing, and it judges that by 2027 or 2028 DeepSeek will “primarily” rely on Huawei’s “already quite mature” technology stack. The report’s original text is the same one covered in the previous issue, in which the judgment that “only TileLang is usable among the six Ascend-supported repositories” is unchanged; it also notes that the Ascend 950 backend was merged into the main TileLang repo rather than remaining in a separate adaptation layer. This is secondary propagation from the same source and does not constitute new technical progress, but it shows the judgment is entering mainstream media framing.

Source:
https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report

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

3.1 TileOPs: FP8 paged warp-specialized kernel shared-memory aliasing causes NaN (#2708)

2 files +80/-30. GQAPagedFP8WSKernel returns NaN under the combination of “seven single-token requests placed before a 300-token request, 4099 rows of FP8 cache, 48-row pages.” The root cause is that the shared-memory planner read the warp-role buffers as active in different time windows, so it placed the split handover onto key or value ring addresses still being filled by TMA, and the handover’s running maxima were also squeezed onto the ring flag bits. The fix explicitly separates these buffers in shared memory and adds a 24-line regression case (under the qwen3-fp8-ragged configuration).

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

3.2 TileOPs: dense prefill switched to persistent traversal, D64 three consumers (#2706)

6 files +590/-5, adding GQADensePersistentWSKernel (external name gqa_dense_persistent, taking priority over gqa_dense): when the number of query tiles exceeds the number of SMs, one CTA occupies one SM and walks its tiles in alternating rounds, loading the next tile’s query during the current tile’s epilogue. On D64 each tile is 192 rows of query spread across three 64-row consumer warpgroups, cutting K/V re-reads by one third; D128 keeps 128 rows.

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

3.3 TileOPs: SM90 dense GQA decode goes warp-specialized (#2694)

Adds GQADecodeWSKernel: a TMA producer warp fills a two-slot K/V ring, and one WGMMA warpgroup maintains the online softmax for the head group; the number of splits is chosen to fill the resident CTA slots in one pass, with each split guaranteed at least four KV tiles, and with a single split the normalization happens within the CTA and writes output directly without going through the combine kernel; KV heads vary fastest across CTAs.

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

3.4 TileOPs: causal sliding window serves an entire query head group with two key tiles of one KV head (#2696)

Adds GQADenseSlidingWindowBandKernel: GQADenseFwdOp previously reloaded the same two key/value tiles for every query head, with one CTA doing only one (query tile, head); the new kernel has one CTA per (query tile, KV head) load the band’s two key/value tiles only once, reduce over that group of query heads in turn, and keep the next head’s query tile in flight. It targets scenarios where the sliding-window row count does not exceed 128 keys.

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

3.5 TileOPs: paged decode autotuning previously shared a single block table (#2703)

4 files +24/-13. The paged decode autotune supply gave every request the same block table, meaning it only measured one request’s pages, and those pages stayed resident in L2, so the candidate ranking reflected the traffic actually paid by the served batch. Measured on a 1500 MHz H200, the MHA decode tuner on opt-1.3b-b16 chose block_N 128 (cold 80.0 us) over block_N 64 (72.5 us), dropping to 0.994 of flashinfer. The fix makes tuning read different cache pages per request.

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

3.6 TileOPs: masked softmax FMA overflow and varlen RoPE odd half-width (#2699 / #2700)

#2699 (4 files +54/-6) fixes NaN rows in the dense GQA sliding-window forward: the 128x64 tuning tile under the causal window makes each query tile start from the first-row window, and the last few rows encounter two fully masked key tiles in a row; after the first block both running maxima sit at the -1e38 clamp, and fast math fuses prev * scale - curr * scale into an FMA, whose rounding error (about 2 to the 99th power) overflows to NaN. The fix subtracts the running maxima before the scaling rescale. #2700 fixes the loop bound of varlen query RoPE under odd half rotary widths (rotary_dim 10, 14, etc.): when the parallelism rows times chunks is not divisible by the thread count of the rotary branch, TileLang drops the loop upper bound, subsequent query rows in the warp-specialized path are computed incorrectly, and the generic path (dim 256) hangs outright.

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

3.7 TileOPs: FP8 grouped GEMM staged scale staging and gate/up ping-pong warp groups (#2704 / #2693)

#2704 adds the config key stage_scales: it lets the producer copy each stage’s A/B scales into shared memory via cp.async and arrive at that stage’s barrier on landing; stage_scales=0 still generates the old kernel. The background is that masked FP8 grouped GEMM reads each step’s A scale for K from global memory in the math warpgroup, and deepening the ring compresses L1 and stalls these loads, leaving ds-v3-e8-up-masked at only 0.979 of DeepGEMM. #2693 adds a ping-pong form to GemmTemplate: two math warpgroups each run a full 64-row tile and take turns taking the CTA’s tiles, with one group’s epilogue falling under the other group’s mainloop, ordered by rotating barriers.

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

3.8 TileOPs: predicate broadcast and boolean operands switched to 16-byte vector reads (#2702 / #2697)

#2702 makes row-broadcast predicates (Eq/Ne/Gt/Lt/Ge/Le and LogicalAnd/Or) read each full 1024-column block’s stride-1 operand as 16-byte vectors via load16_evict_last, while partial blocks and other staged bodies keep fragment copies; linear-bias-float32 (1024x4096 against 1x4096) improves the six comparisons from 0.99 to 1.06, and the logical operators to 1.04. #2697 changes the full-size operands in the comparison and logical binary kernels to 16-byte evict-first vector reads; float32 rows were previously only at 0.99 to 1.00 of torch.compile.

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

3.9 TileOPs: three SM89 paths advanced the same day (#2705 / #2691 / #2690)

#2705 lowers the default pipeline stage count of the SM89 SSD chunk scan to two, reducing the scan kernel’s buffer requirements; the FP16/BF16 measurements on six shapes take the median of three rounds of 200 CUPTI samples each. #2691 splits the SM89 paged DSA key range when the existing planner selects multiple partitions, reusing the existing FP32 partial-output and log-sum-exp merge; with a single partition it keeps the BF16 output and single kernel, and folds the compile target into the merge cache identity, building a separate JIT builder per target. #2690 switches NSA top-k selection on SM89 for query head groups of 32/64 and head dims of 64/128 to 32-column tiles, keeping the 64-row tile and the same scoring/sorting implementation, reducing median GPU busy by 2.1% to 15.0% across eight FP16/BF16 cases.

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

3.10 TileOPs: reduction folding in a single launch and empty DSA selection validation (#2689 / #2698)

#2689 adds ReduceFoldKernel, folding the contiguous suffix of each retained index’s lead segment in a single launch (previously reducing over the axis prefix and suffix took two launches), and has ReduceEdgeKernel and VectorNormEdgeKernel hand off to it all calls whose rows can be held in flight by a 512-thread block; the background is that Sum lagged torch-compile on 3d-multidim-float16 (0.977) and mlp-intermediate-float32 (0.990). #2698 fixes paged DSA validation: when a query row has no valid selected keys, the reference implementation returns zero output and negative-infinity LSE, and subtracting the reference LSE yields a NaN error bound that marks correct results as wrong; the fix gives empty rows an exact zero LSE bound and adds a control case requiring the validator to reject nonzero output and finite LSE when all rows are empty.

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

3.11 TileOPs: benchmark references computed in float32 (#2692)

ref_program previously served both as the correctness reference and as the object timed for the torch baseline, and many references called torch directly with the 16-bit input dtype, so results depended on how torch rounded each step on the device side. CUDA masked this difference; CPU torch and torch_musa did not — GroupNorm and InstanceNorm failed on the MTT S5000, while the same kernel passed against a float64 reference. The fix makes ref_program compute in float32 or wider, while the torch baseline is timed at the input dtype. This is a rare issue surfaced by a domestic GPU backend but that had to be fixed first at the reference-implementation layer.

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

3.12 TileFoundry: install sm_89 architecture and RTX 6000 Ada device (#259)

11 files +392/-22. Adds two target descriptions, nvidia_sm89.toml and nvidia_rtx_6000_ada.toml (114 and 124 lines respectively), with docs/spec/target.md updated by 42 lines in step. This change extends TileFoundry’s device modeling from Hopper to Ada, and is two sides of the same coin as #256 from the same day, which exposed that tests compiled for the sm_90 target only produce wrong values on a single sm_89 card; it is accompanied by an sm_89 e4m3 warp mma declaration (#261, under review).

Source: https://github.com/tile-ai/TileFoundry/pull/259

3.13 TileFoundry: gmem accounting switched to footprint and running mesh, with callee dimensions bound (#258)

The widest-reaching change of the day: 74 files +1128/-553, fixing #239, #252, #253. Previously, when 132 CTAs probed only 40, accounting assumed “one move per position per window,” so under-fill was invisible; dimensions that could only be declared by the callee could not be bound; and the meaning of analyze’s output targets and per-function memory rows was unclear. The fix centralizes dimension binding for check and analyze into instantiate, so dimensions declared by the callee or submodules can now be bound. The commit was initiated by a community contributor and taken over and rewritten by a maintainer.

Source: https://github.com/tile-ai/TileFoundry/pull/258

3.14 TileFoundry: must throw when the device refuses to launch a kernel (#256)

7 files +69/-11. Previously the launcher issued kernel<<<grid, block, smem, nullptr>>>(...) and never read the launch status; when the device refused the launch, the call returned as usual, the output kept its old contents, and the error never reached the caller. On a single sm_89 card, tests compiled for the bundled sm_90 target only showed wrong values, and only compute-sanitizer revealed that cudaLaunchKernel returned cudaErrorNoKernelImageForDevice. The fix makes the CUDA launcher throw after reading cudaGetLastError(), with the CPU-side module changed in step and the spec docs and tests updated alongside.

Source: https://github.com/tile-ai/TileFoundry/pull/256

3.15 Distribution: Sunrise opens the 0.1.15+sunrise.1.2.0 release PR (#9, under review)

At 16:48 on 10-22 it created the candidate/sunrise-0.1.15-1.2.0-20261019 branch and opened the PR, updating the Sunrise distribution to 0.1.15+sunrise.1.2.0 and incorporating TileLang 0.1.15 plus the accompanying reviewed TileOPs, TileKernels, TVM, and tvm-ffi changes; the source import and merge and the distribution adaptation are committed separately, and the final adaptation configures the TANG-backend Sunrise wheel and fixes the sliced GEMM regression (both accumulator flags validated against the public template, with stride, origin, and offset assertions retained). Self-reported verification: 55 local contract tests pass, 54 TANG pipeline codegen tests pass, both lint jobs and the trust-boundary check pass; however, after the S2 wheel build produced the package, the post-install import smoke check failed at tangGetDeviceCount. The three commits on the branch are the 0.1.15 distribution adaptation, narrowing CI to the TANG backend, and the sliced GEMM accumulator flag test.

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

3.16 Adopters: no pushes from TileKernels or FlashQLA in the window (10-22)

deepseek-ai/TileKernels last pushed on 10-14, and QwenLM/FlashQLA likewise stopped on 10-14; neither adopter had commits in this window.

Source: https://github.com/deepseek-ai/TileKernels

IV. Community, Tutorials, and Events

4.1 CANN Community Launches ops-tilelang Operator Repository (10-22)

Atlas-Ascend-Fork/ops-tilelang was created on 10-20, positioned as the CANN community’s TileLang high-performance operator repository for Ascend NPU: it centrally manages Kernel implementations, automated test cases, standard workloads, and performance baselines, with the goal of forming reusable, verifiable, and sustainably evolving operator assets. The repository already has four layers—src, tests, scripts, docs—along with open atom community specification files such as OAT.xml and SECURITY.md. On 10-22 a fix was pushed for “overflow after combining a valid upper-bound seed with a test-case hash, causing CPU test failures”; earlier commits include “replace Ascend A5 references with Ascend 950” and “sessions where benchmark fails or is interrupted still overwrite the performance baseline.” A community-level operator repository using a single DSL as the sole implementation language has not been seen before; this line is worth tracking weekly.

Source: https://github.com/Atlas-Ascend-Fork/ops-tilelang

4.2 tilelang/agent: Writing Operator Optimization as an Agent’s SKILL Workflow (10-22)

tilelang/agent (created 08-28, 3 stars) pushed on 10-22; the README itself is an executable prompt template: it injects domain knowledge via @tilelang-operator-dev/SKILL.md, with inputs limited to the operator name, GPU/CPU reference implementation, shape, and the constraints “tilelang npu delivery, Chinese delivery”; comments state that the agent must execute in the A5 environment and, after loading, follow a 10-step closed loop. The Skill contains a selection decision tree, 8-CSV rules, L0 to L4 strategies, and 11 pitfalls. Looking at the operator repository in 4.1, the instruction-level scheduling in 3.15, and this repository together, a division of labor is forming on the Ascend side: “humans write rules, agents tune, and baselines accumulate in the repository.”

Source: https://github.com/tilelang/agent

4.3 New Community Repositories: Sail Wrapper, Remote Autotuning Agent, and Tutorial Repositories (10-22)

flytiger-eco/tilelang-wrapper-for-sail was created on 10-21 and pushed on 10-22 (Python, no description); cc-c122/TileLang-Remote-Autotuning-Agent (Python, pushed 10-22) and easy-tilelang/tilelang-devkit, easy-tilelang/tilelang-tutorials were pushed on the same day. There are many of them with low stars per repository; this is natural growth in the tutorial and tooling layer, and no technical direction can yet be judged.

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

4.4 Docs and Versions: Documentation Site Updated Three Times, VERSION Still 0.1.15 (10-22)

tilelang.github.io had three Update docs commits, TileFoundry.github.io had one render TileFoundry v0.0.3 commit, and TileOPs.github.io had one commit. The main repository’s VERSION remains at 0.1.15, and the latest release is still v0.1.15 from 09-30; the repository has 8742 stars, 915 forks, and 389 open issues.

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

V. Trend Observations

1. The defect types in the main repo have shifted to a different layer. Last week the main repo was clearing out “shape and geometry validation” issues—cast rounding annotations being silently dropped, CUDA targets without an arch silently falling back to sm_50, sparse metadata shapes not being validated. In this window, four of nine commits landed on compile-time state such as analyzers and source locations: analyzers being reused after rewriting, variable rebinding after splitting, and the eager frontend losing Python locations. The common thread across these defects is that “the generated kernel may still be correct, but what the compiler tells you and what it actually does are inconsistent”—and as TileLang moves toward agent-driven automatic kernel writing (4.2), compile-time state consistency will directly determine whether the agent iteration loop can converge.

2. Progress on the Ascend side has shifted from operator count to the underlying data path. #3530 touches 18 files in a single commit, implementing row-stride expression for the NZ compressed layout in some UB→L1 copies; #3581 and #3523 respectively fix compile-time predicate explosion and the allocation semantics of SIMT vector local variables. None of the three adds new operators—all of them make existing paths compilable and correctly executable across more shapes. Combined with the CANN community spinning up a separate ops-tilelang operator repo (4.1), the division of labor becomes clear: operator assets are carried by the community repo, while the main repo focuses on backend semantics and codegen.

3. TileOPs saw 20 consecutive commits over two days, with the focus on SM90 specialization and SM89 gap-filling. In this window, three new SM90 warp-specialized kernels were added (dense decode, causal sliding-window band, FP8 varlen overlap), while three small-scale paths on SM89—SSD chunk scan, paged DSA, and NSA top-k—were advanced together. Notably, two NaN fixes (#2708, #2699) both appeared on warp-specialized paths: liveness determination in the shared-memory planner and FMA fusion under fast math are the two places where this kind of hand-orchestrated kernel is most error-prone—and that is precisely the part TileLang promises to manage on the user’s behalf.

4. A counter-indicator is the benchmark layer. #2692 changed the reference implementation to compute in float32, because GroupNorm and InstanceNorm were failing on MTT S5000—the problem was that the reference implementation followed the input dtype rather than the kernel itself. On the same day, TileFoundry #256 turned “device refuses to start” from silent into a thrown error. Both point in the same direction: turning “looks like it passed” into “actually passed.” For teams relying on these tools for chip adaptation, the real value of such fixes exceeds a few performance numbers.

5. Distribution and industry are both accelerating. The Sunrise distribution packaged 0.1.15 and the TANG backend wheel into a release PR, indicating that TileLang now has more than one downstream distribution form; on the industry side, SCMP republished SemiAnalysis’s assessment, pushing “only TileLang is usable among the six repos” into mainstream media. On the technical cadence, this remains a quiet day (no releases, no breaking changes in the main repo), but the ecosystem layer shows noticeably more organizational activity than the code layer.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
tile-ai organization repos 29 repos, 10 pushed within the window; 35 merged, 71 updates under review
tilelang main repo 9 merged, 1 single-line fix, 1 Ascend change spanning 18 files; VERSION 0.1.15, no new release
TileOPs 20 merged, all kernels and tests for Attention / MoE / Elementwise / Reduction
TileFoundry 3 merged, 4 more under review; the 0.0.3 release on 10-21 was already covered in the previous issue
tilelang-ascend 2 merged (1 new example, 1 reverted example), 11 under review
tilelang-hygon No commits on main; one switch cleanup on merged branch gcy_merge_main
tilelang-sunrise New candidate branch created and a 0.1.15+sunrise.1.2.0 release PR opened, not merged
tilelang-metax / musa / mlir-ascend No pushes within the window; last activity on 10-16 / 09-30 / 10-19 respectively
TileRT / tilescale / DeepStack No pushes within the window
Adopters TileKernels / FlashQLA Both last active on 10-14, no pushes
Google News (multiple Chinese and English queries, via proxy) 3 hits; 2 were media reprints of the same SemiAnalysis report, 1 was DeepSeek funding and was removed as off-topic
Hacker News No TileLang-related discussion
arXiv Searched with TileLang as title or abstract keyword, 0 hits within the window
New community repos 5 new or active repos, all with creation dates and push times verified
Member or vendor official sites Not applicable to this topic

Full Source List