Research window: 2026-11-05 07:00 ~ 2026-11-06 07:00 (Beijing time, 24 hours; the previous issue’s cutoff was 11-05 07:00).

Sources: Repository-by-repository check of the GitHub tile-ai organization (29 repositories, 10 pushed within the window) + merged and in-review ground-truth search (18 merged and 57 in-review updates within the window) + 5 new issues + commit-by-commit check of the main branch (tilelang 5, TileOPs 9, TileFoundry 1, tilelang-ascend 2, tilelang-musa 1) + non-default branch check (3 commits on the feat/hcu-scaled-cast branch of tilelang-hygon) + release page check (tilelang / TileOPs / TileFoundry, no new versions within the window) + two adopter repositories + docs site and nightly repository check + four Google News queries in Chinese and English (via proxy) + Hacker News + arXiv. The dedicated four-source RISC-V inspection showed no changes in this window.

Issue Index

  • Today’s Highlights: TileOPs merged nine PRs in one day, advancing three things at once: “in-place reads of broadcast operands, single-wave launch, and tile selection”; beyond the five merges in the main repo, a silent incorrect-CUDA bug that has existed since 0.1.12 was reported; Ascend fixed a segfault, while also being flagged that under threads=2, sin/cos computes 97.1% of elements incorrectly
  • I. Core Project Progress
    • 1.1 Mainline: TF32 conversion supports round-to-nearest-even, with SM90 and SM100 diagnosed separately (#3714) (11-05)
    • 1.2 Mainline: host-mapped buffers can be explicitly annotated; CPU tensors are no longer rejected by device checks (#3802) (11-05)
    • 1.3 Mainline: casts in scalar conditions no longer restrict data vector width (#3759) (11-05)
    • 1.4 Mainline: Z3’s rlimit can be configured per kernel, default remains 100,000 (#3752, depends on tile-ai/tvm #91) (11-05)
    • 1.5 Mainline: nested T.macro reuses conversion templates, reducing frontend recompilation overhead (#3760) (11-05)
    • 1.6 New issue: missing __syncthreads after conditional scalarization; 0.1.12 through 0.1.15 silently generate incorrect CUDA (#3808) (11-05)
    • 1.7 New issue: without warp specialization, T.Pipelined reloads the stage before the tcgen05 wait (#3800) (11-05)
    • 1.8 Under review: full migration draft to native Apache TIRx (#3803); shared transpose switched to ldmatrix/stmatrix (#3809) (11-05)
  • II. Multi-Backend Adaptation
    • 2.1 Ascend: 64-bit dtype for bitwise_rshift now raises an explicit error instead of segfaulting (#1954) (11-05)
    • 2.2 Ascend: A2/A3 performance skills aligned with the Ascend950 workflow (#1932) (11-05)
    • 2.3 Ascend: under threads=2, tmp workspace is unconditionally halved; sin/cos computes 97.1% of elements incorrectly on indirectly addressed tiles (#1991 / #1992) (11-05)
    • 2.4 Hygon: three LDS phase fixes on the branch, and the first two PRs opened (#33 / #34) (11-05)
    • 2.5 Moore Threads: README documentation update, ending the silence since 09-30 (11-05)
    • 2.6 MetaX and mlir-ascend: MetaX has no pushes for a third consecutive window; three TLIR operators under review in mlir-ascend (11-05)
  • III. Ecosystem and Adopters
    • 3.1 TileOPs: nine merges — in-place reads of broadcast operands, single-wave GEMM fill, reduction layout caching (11-05)
    • 3.2 TileOPs: roofline pricing rules rewritten; three impossible bandwidth rows brought back into a reasonable range (#2830) (11-05)
    • 3.3 TileOPs: two merges 2.5 minutes apart caused main’s GPU Smoke failure, fixed the same day (#2837) (11-05)
    • 3.4 TileOPs under review: three-stage tile factors, TN single-element contraction, TMA scaled loads for FP8 grouped GEMM (#2836 / #2838 / #2839) (11-05)
    • 3.5 TileFoundry: WGMMA synchronization plan merged (#286); completion semantics for partial reductions under review (#288) (11-05)
    • 3.6 TileSight: H20 WGMMA capability declaration corrected, the repo’s first PR (#1) (11-05)
    • 3.7 Adopters and third parties: no pushes from TileKernels or FlashQLA; two new third-party repos appeared (11-05)
  • IV. Community, Tutorials, and Events
    • 4.1 Documentation site updated three times; nightly repo has no main-branch commits; Ascend daily tests pass all 3,030 items (11-05)
    • 4.2 News side: four query groups; no in-window hits on Hacker News or arXiv (11-05)
  • V. Trend Observations
  • Appendix: Sources and Verification Notes

Today’s Highlights

The densest output of the day came from the operator library rather than the compiler. TileOPs merged six commits in a row between 21:06 and 22:46 — just over an hour — plus two more at 14:01 and 21:31, nine in a single day, all tagged perf, fix, or ci. Three things advanced at once. First, in-place reads of broadcast operands: five operator classes — lerp, where, clamp, masked_fill — no longer expand stride-0 broadcast tensors into the output shape first; measured broadcast-row latency dropped to 0.13–0.44× that of the mainline (#2832). Second, single-wave launch: unary elementwise ops collapse into one 256-thread wave when “the body has at most two cheap instructions and total read/write stays under three-quarters of L2,” and the long-standing bug where ReLU returns 0 for NaN was fixed along the way (#2829). Third, the tile selector: a native-width tile whose width fills one wave was built for small-M NT GEMM, turning decode-down from a laggard into a 3% lead over the fastest comparison (#2834). What these three share is not “faster” but that each ships a reproducible A/B table and flags the row where it still trails.

The second item is a trustworthiness problem, and it spans two repos. The main repo opened issue #3808 reporting that from 0.1.12 to 0.1.15, after condition-related scalarization, the shared-memory epilogue lost a __syncthreads(); the compiler raises no error and the generated CUDA is wrong. The reporter supplied a full reproduction script and four version numbers — this is silently wrong code, not a performance regression. In the same window, #3800 reports a different silent error: without warp specialization enabled, T.Pipelined reloads the next stage before the barrier wait in T.tcgen05_gemm, so the operand the MMA reads has already been overwritten. These two form a neat contrast with #3759, merged into the main repo the same day — #3759 fixes “an int64 cast inside a scalar condition squeezing the data vector width from 16 lanes to 4.” At the same scalarization/vectorization boundary, one side is fixing width while the other is missing a sync.

The third item is on the Ascend side. tilelang-ascend merged #1954, converting a compile-time segfault in T.tile.bitwise_rshift on int64/uint64 into an explicit ValueError, permitting only int16/uint16/int32/uint32; this is not new 64-bit right-shift support but turning “the compiler crashes” into “the compiler says no.” The same repo opened two issues on the same topic that day (#1991, #1992): under 910_9362 / CANN 9.0.0 with threads=2 (C:V=1:2), AscendVidReduction unconditionally halves every UB allocation in the kernel — including the injected tmp workspace — with the result that sin/cos applied to an indirectly addressed skip-set buffer miscomputes 3978/4096 elements (97.1%), while sigmoid/round violate AscendC’s calCount <= tmp contract and write out of bounds into scratch behind the layout. Set against this is the daily test report from #1988 the same day: 3030 items all passing, 100% pass rate. Together these two sets of numbers say one thing — existing tests do not cover the combination of threads=2 plus indirect addressing, and pass rate cannot substitute for correctness.

The fourth item is ecosystem movement. The Hygon repo went from “four commits accumulated on a feature branch, no PR opened” to formally entering review: the feat/hcu-scaled-cast branch landed three more LDS phase fixes that day (default swizzle adapted to 64 banks, swizzle planning restored on the MN-slice consumer side, layout-aware reads for the K partition), and two PRs were opened — #33 (scaled cast) and #34 (lane swap and constraint propagation for reduction). Moore Threads ended its complete silence since 09-30, landing a README update on installation and operator development on the main branch. Add the first PR in the TileSight repo (correcting the H20 WGMMA capability declaration) and the TIRx migration draft in main-repo #3803, and three organizational threads can be seen unfolding at once: domestic backends moving from in-house branches into public review, performance modeling moving from papers into a repo, and the compiler foundation preparing to replace vendored TVM.

I. Core Project Progress

1.1 Mainline: TF32 conversion supports round-to-nearest-even, with separate diagnostics for SM90 and SM100 (#3714)

Merged at 12:09. Previously, T.cast to tfloat32 only had the default TF32 round-to-nearest path; anyone wanting ties-to-even had to import the CUDA helper themselves. This change adds native lowering for round="rn": the non-saturating form emits cvt.rn.tf32.f32, the saturating form emits cvt.rn.satfinite.tf32.f32, and it bypasses the default TF32 rounding constructor, preserving the converted bit pattern directly. Both scalar and vector paths are covered, with TF32 lane access going through packed carriers. Architecture diagnostics are split by form: the non-saturating form requires SM90 or above, the saturating form requires SM100 or above, and a compile-time error is raised otherwise; casts without an explicit rounding specification keep their existing behavior. The reporter is the requester behind the closed #3713, making this an interface-convergence change of the “users no longer need to bring their own helper” kind.

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

1.2 Mainline: host-mapped buffers can be explicitly annotated, CPU tensors no longer rejected by device checks (#3802)

Created and merged the same day (submitted 12:08, merged 16:02). CUDA kernels can directly access host memory mapped into the GPU address space, but TileLang’s existing device checks assumed all tensor arguments reside on the target device, so such CPU tensors were rejected before launch. The new T.annotate_host_mapped_buffers(*buffers) writes the indices of annotated parameters into the PrimFunc’s tl.host_mapped_params attribute, supports multiple buffers, merges duplicate declarations, and survives IR serialization and kernel cache restoration. Both the Cython and TVM-FFI launchers accept CPU tensors for annotated parameters, while the remaining dtype, shape, rank, and stride checks stay as before; device-side parameter checks still align with the current execution device, and host tensors do not provide a CUDA device number. A new native runtime helper validates registration status via cuMemHostGetFlags and resolves the device-visible address via cuMemHostGetDevicePointer; ordinary pageable memory is still rejected.

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

1.3 Mainline: casts in scalar conditions no longer constrain data vector width (#3759)

Merged at 19:27, tagged ascend. The cause: #3734 correctly tightened the width restriction on data casts after DecoupleTypeCast, but it also counted casts inside conditions: a BF16 copy guarded by T.assume(n % 16 == 0) was squeezed from 16 lanes to 4 lanes on b5cb02633 because the condition contained T.Cast("int64", ...), even though the condition was uniform across all 16 lanes. The fix reuses the planner’s existing condition-uniformity analysis to extend the cast-width exemption from buffer indices to conditions: conditions of if statements, tirx.if_then_else, and Select on scalar predicates are accessed separately, while branches remain restricted; the condition-cast constraint is still retained on targets with lane-level vector predicates. The rbits grouping for stochastic rounding is also preserved when skipping the width constraint. This change and the two new issues in this window fall in the same area, indicating that the scalarization/vectorization boundary is currently a hotspot for defects.

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

1.4 Mainline: Z3 rlimit configurable per kernel, default remains 100,000 (#3752, depends on tile-ai/tvm #91)

Merged at 23:05, preceded by tvm #91 in the fork (merged 23:03). The change turns the Z3 solver’s rlimit from a hardcoded 100,000 into something configurable per kernel; when unconfigured, the old default is kept. The author puts it plainly: “this indirectly fixes a lot of Z3 prover-related issues.” The significance of such changes is turning “proof exceeds budget” from an untunable environmental behavior into a tunable compilation parameter—the same direction as yesterday’s Ascend-side tightening of the AutoSchedule cycle lower bound (cutting Z3 solve time by 10.04%), just two approaches: one changes the search starting point, the other changes the search budget.

Source: https://github.com/tile-ai/tilelang/pull/3752, https://github.com/tile-ai/tvm/pull/91

1.5 Mainline: nested T.macro reuses transformation templates, frontend recompilation overhead converges (#3760)

Merged at 23:38. The problem: each time the outer function runs, nested T.macro definitions redo source parsing, AST rewriting, and Python compilation, even when the kernel itself already hits the cache. The fix introduces a bounded in-process cache that stores only context-independent transformation results (unbound code and rewritten source), with each macro instance then binding the current globals and closure values; context-dependent transformations take the original path, and wrapped functions bypass the cache. Macro call sites and the kernel cache are unchanged. Regression coverage includes new closures, globals, captured buffers, annotations, context-dependent fallback, and closure lifetime; CPU-side frontend checks pass, GPU execution was not verified locally—the author states this boundary explicitly.

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

1.6 New issue: missing __syncthreads after conditional scalarization, versions 0.1.12 through 0.1.15 silently generate incorrect CUDA (#3808)

Opened at 21:14; the heaviest defect report in this window. The repro is a kernel with a shared-memory epilogue: versions 0.1.12, 0.1.13, 0.1.14, and 0.1.15 all drop a __syncthreads() after condition-related scalarization, the compiler reports nothing, and the generated CUDA simply computes wrong results. The reporter provided a complete runnable Python repro script and environment (Python 3.14.9, CUDA 2.15.0+cu134). What makes this damaging is that it spans four releases with no warning whatsoever—any kernel relying on shared memory to hand off between stages can hit it, and all the user sees is “the result is wrong.”

Source: https://github.com/tile-ai/tilelang/issues/3808

1.7 New issue: without warp specialization, T.Pipelined reloads the stage before the tcgen05 wait (#3800)

Opened at 08:44. The trigger is running a T.Pipelined loop with T.tcgen05_gemm after disabling warp specialization (tl.disable_warp_specialized=True): the stage is overwritten by the next round’s copy before mbarrier_wait_parity, so the MMA reads a rewritten operand. The reporter pasted the lowered kernel source and points to the missing ordering constraint between the write via As and the read by tcgen05mma_ss. Like #3808, this belongs to the “generated code looks fine but is actually invalid” category, except this one occurs on Blackwell’s tensor memory path.

Source: https://github.com/tile-ai/tilelang/issues/3800

1.8 Under review: full migration draft to native Apache TIRx (#3803); shared transpose moved to ldmatrix/stmatrix (#3809)

#3803 is the most substantial item under review this period: it replaces the vendored TVM compiler with standalone Apache TIRx, migrating the frontend, IR transformations, code generation, JIT adaptation, and packaging entirely to the native object model, and pins two upstream commits, 7e49a148 of apache/tvm-tirx and 3e906bc6 of apache/tvm-ffi, keeping the dependency tree on clean upstream with no injected patches. The author marks it as a draft—itself a public preparation for a “foundation upgrade.”

#3809 (submitted 01:27) addresses another concrete matter: T.transpose between two shared buffers has always gone through generic SIMT loops, so 8-bit and 16-bit tiles cannot use shared-memory bandwidth, and the FP8 sparse MLA in examples/deepseek_v32 had to hand-write 16x16 register staging to work around it; meanwhile, after T.Pipelined adds a stage axis to the source region, the rank mismatch makes MakeIndices fail outright. The fix first tries warp-matrix lowering on targets supporting stmatrix (SM90 and above, excluding CuTeDSL): 16-bit tiles use one ldmatrix plus one stmatrix.trans per 8x8 block, 8-bit tiles use two ldmatrix.trans plus __byte_perm selection per 16x16 byte block to assemble four rows, and the remaining regions fall back to the original SIMT path.

Source: https://github.com/tile-ai/tilelang/pull/3803, https://github.com/tile-ai/tilelang/pull/3809

II. Multi-Backend Adaptation

2.1 Ascend: 64-bit dtype for bitwise_rshift now raises an explicit error instead of segfaulting (#1954)

Merged at 16:36 (ascendc_pto branch). On A2/A3, passing int64/uint64 operands through the public API T.tile.bitwise_rshift causes the compiler to segfault directly during the OptimizeForTarget stage. The fix checks the dtype of dst and src0 before constructing the intrinsic, allowing only int16/uint16/int32/uint32 and raising a ValueError with the actual dtype and the supported list for anything else; the shift amount scalarValue is not restricted by expression dtype; Buffer and BufferRegion both go through _underlying_buffer(). Adds a direct API diagnostic test and a subprocess crash regression test. The author explicitly states the boundaries: this change does not add 64-bit right-shift support, nor does it modify the path that directly constructs the underlying intrinsic. Related to #1720.

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

2.2 Ascend: A2/A3 performance skills aligned with the Ascend950 workflow (#1932)

Merged at 15:36, a pure skill and process refactor: aligns and unifies the A2/A3 performance tuning skills with the A5 (Ascend950) workflow structure, adding two new skills for chip type identification and chip information retrieval. In isolation it has no functional value, but it means the performance tuning “playbook” is starting to fork by chip generation—the same skill set must be able to identify which generation it is running on, otherwise tuning recommendations will be misapplied across generations.

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

2.3 Ascend: tmp workspace unconditionally halved under threads=2, sin/cos computes 97.1% of elements wrong on indirectly-addressed tiles (#1991 / #1992)

Two issues on the same topic were opened at 17:03 and 17:25, with identical titles and bodies (effectively the same issue submitted twice). The reported environment is 910_9362 / CANN 9.0.0, threads=2 (C:V=1:2). The root cause is in AscendVidReduction: it unconditionally halves all UB allocations within the kernel, including the tmpspace injected by InjectTmpBuffer, whereas the skip-set buffer (indirect addressing where the GM index contains a UB load, e.g. T.copy(KV[b_i, idx_ub[bi_i], g_i, :D], kv_ub)) keeps its full shape and was never meant to participate in vid halving. The consequences fall into two tiers: sin/cos computes wrong results directly (3978/4096 elements wrong on upstream HEAD, 97.1%); sigmoid/round violates AscendC’s calCount <= tmp contract, with numerically correct values but out-of-bounds writes landing on scratch memory behind the layout, which becomes silent data corruption once the layout changes. The reporter notes this was found during the #1897 review, and that the #1897 fix only covered the count-side semantics (HasModifiedBuffer skipping workspace slots), leaving the workspace side unaddressed.

Source: https://github.com/tile-ai/tilelang-ascend/issues/1991, https://github.com/tile-ai/tilelang-ascend/issues/1992

2.4 Hygon: three LDS phase fixes on the branch, and the first two PRs opened (#33 / #34)

The main branch is still at 11-02; all three commits landed on the feat/hcu-scaled-cast branch, authored on the Hygon side (commit domain hygon.cn):

  • 10:06 fix(hcu): adapt default LDS swizzle to 64-bank phases
  • 14:04 fix(hcu): restore LDS swizzle planning for MN slice consumers
  • 19:49 fix(hcu): support K-partitioned layout-aware LDS reads

More important is the change in review status: PR #33 was opened at 09:08 (scaled cast, i.e. the branch’s main body), and PR #34 at 19:52 (optimizing reduction’s XOR shuffle with DPP and lane swap, plus a fix that “propagates known constraints before allocating free candidates” in layout inference). The previous report noted that “all four commits are on the feature branch, no PR yet”; these two PRs in this window move Hygon from a proprietary branch to public review—the first domestic backend to complete this path.

Source: https://github.com/tile-ai/tilelang-hygon/pull/33, https://github.com/tile-ai/tilelang-hygon/pull/34, https://github.com/tile-ai/tilelang-hygon/commit/07451160e5aa758419201adb598725cbef70e3a4

2.5 Moore Threads: README documentation update ends the silence since 09-30 (11-05)

At 19:46 a commit docs: update README installation and operator development references landed on the main branch, the repo’s first commit since 09-30. During the same period, the timestamps of 6 lingering PRs (dependency upgrades among #1 to #7, the TileOPs submodule, and the evidence-threshold documentation for MUSA integration) were all refreshed, but their content remains old proposals from April to August, with no code changes in this window. A conservative judgment is warranted: the documentation commit shows someone is maintaining the repo, but it is not sufficient to prove backend development has resumed; the real signal will be whether operator or lowering commits follow.

Source: https://github.com/tile-ai/tilelang-musa/commit/4cac45b859108a86ac9a0b1cbfdeb33da7711697

2.6 MetaX and mlir-ascend: MetaX silent for a third consecutive window; three TLIR operator PRs under review in mlir-ascend (11-05)

tile-ai/tilelang-metax’s last push is still 10-26, silent for a third consecutive window. tile-ai/tilelang-mlir-ascend’s main branch has no new commits, but three operator PRs based on the TLIR dialect are progressing: #203 adds tl.sub/mul/div/max/min to TLIR and supplements the TL-to-HIVM lowering, updated at 01:42; #211 adds tl.shl and #212 adds tl.shr, submitted at 19:44 and 20:58 respectively. What this line of work is doing is lowering TileLang’s semantics layer by layer into the MLIR dialect—slow in pace but single in direction.

Source: https://github.com/tile-ai/tilelang-metax, https://github.com/tile-ai/tilelang-mlir-ascend/pull/203, https://github.com/tile-ai/tilelang-mlir-ascend/pull/211, https://github.com/tile-ai/tilelang-mlir-ascend/pull/212

III. Ecosystem and Adopters

3.1 TileOPs: Nine PRs merged — in-place reads for broadcast operands, single-wave fill for GEMM, reduction layout caching (11-05)

The nine PRs cluster into two windows, 14:01 and 21:06–22:46. The mainline items, one by one:

  • #2832 In-place reads for broadcast operands: Five operator classes — LerpTensorFwdOp, WhereFwdOp, ClampTensorFwdOp, MaskedFillScalarFwdOp, MaskedFillTensorFwdOp — no longer expand stride-0 broadcast tensors to the output shape first; instead they read directly at stride 0 within a single kernel. One shared builder serves all five, operand shapes collapse into a single loop nest, vectors of 16 bytes or less stay in one inner row, and each operator only writes its own elementwise body. Broadcast rows run at 0.13–0.44x the mainline time, matching or beating torch-compile; the remaining elementwise rows are byte-for-byte identical to mainline.
  • #2829 Single-wave unary elementwise: When the body has at most two operations that each lower to a single instruction, total read/write volume is no more than three-quarters of L2, and the result has the same width as the operand, the kernel collapses to a single 256-thread wave. Also fixes ReLU returning 0 for NaN (now returns NaN: 16-bit uses the torch bit pattern, float32 uses canonical NaN). All 327 GPU regression tests pass on H200, and 13 new cases pass Compute Sanitizer memcheck.
  • #2834 Native-width tiles for small-M NT GEMM: ds-v3-decode-down (128x7168x2048, NT) uses only 112 of 132 SMs on mainline, and the 112-wide tile was split into seven n16 WGMAs on 0.1.12. After switching to issuing a WGMMA of the tile’s own width every 16 K steps, decode-down leads the fastest comparison by 3%, and gate-up-b32 is 1.4% faster than mainline (still 0.5–1% behind cuBLAS). The new fill tile is enabled only for shapes with 96 < M <= 128 and N divisible by 8, with a separately fitted structural factor of 0.58.
  • #2835 Reduction layout caching: The reduction kernel rebuilt the axis permutation and output shape on every call, even though ReduceCall is fixed. Now each kernel instance computes them once and shares pure shape helpers with the existing tensor adapters, with no global cache. Full hot-call time for six FP32/BF16 VarMeanFwdOp configurations drops 4.33%–6.17%, and the generated CUDA and host source are identical — the savings are in call preparation, not GPU time.
  • #2833 CI fix: After GitHub deprecated runner 2.336.0, all gpu-smoke jobs queued, and entrypoint.sh carries --disableupdate so it cannot self-upgrade; RUNNER_VERSION is bumped from 2.336.0 to 2.338.2.

Sources: https://github.com/tile-ai/TileOPs/pull/2832, https://github.com/tile-ai/TileOPs/pull/2829, https://github.com/tile-ai/TileOPs/pull/2834, https://github.com/tile-ai/TileOPs/pull/2835, https://github.com/tile-ai/TileOPs/pull/2833

3.2 TileOPs: Roofline accounting rules rewritten; three impossible bandwidth rows return to plausible ranges (#2830)

Nightly job 41873301494 surfaced three physically impossible roofline rows: two gpt-oss-serve-window configurations of GQAVarlenFwdOp read 7.48 and 7.50 TB/s, and qwen3-30b-ep8-empty of GroupedGEMMFwdOp read 5.30 TB/s. This PR rewrites the accounting rules around “how much this call actually reads”: GQA dense, varlen, and paged forward now count K/V only for key rows some query can see, q only for query rows that can see keys, the RoPE table only up to the truncation position keys actually reach, and FP8 scaling only for requests that see keys; GroupedGEMM uses a new grouped_gemm_roofline to price a row’s group by batch_sizes; NSA compression, NSA sparse, DSA decode, MHA paged, and the FP8 lightning indexer all charge only for positions actually read, and DSA also honors is_causal=False. After the fix the three rows read 0.62, 0.62, and 2.65 TB/s. A side benefit is fixing division by zero for empty K/V storage and empty paged pools. Test nodes: +16.

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

3.3 TileOPs: Two PRs merged 2.5 minutes apart break GPU Smoke on main, fixed same day (#2837)

The PR merged at 22:46 is itself a process lesson. #2824, merged at 21:06, changed the grouped GEMM roofline row to require a declared sizes and derive G from it; #2830, merged at 21:08, still wrote only G with no sizes. The two landed about 2.5 minutes apart, each with green CI, and immediately after merging, GPU Smoke on main reported RowError: row misses ['sizes']. The fix replaces G: 2 in _grouped with sizes: [1, 0]; the asserted byte counts (24, 32) are unchanged. The cause is not that anyone wrote something wrong, but that parallel merges lack a post-merge gate: per-PR CI cannot cover cases where two PRs’ semantics intersect, so only an immediate full run after merge catches them.

Sources: https://github.com/tile-ai/TileOPs/pull/2837, https://github.com/tile-ai/TileOPs/pull/2824, https://github.com/tile-ai/TileOPs/pull/2830

3.4 TileOPs under review: Three-stage tile factor, TN single-element contraction, TMA scale loads for FP8 grouped GEMM (#2836 / #2838 / #2839)

#2836 fits a separate factor (1.36) for the three-stage tile of coop2_swap, because by FLOPs it is not faster than four stages, but by bytes and K steps it was overestimated by about 10%, which squeezed out the 160-wide tile and dropped nightly results from 1.005–1.011x to 0.985–0.995x QuACK. The factor is fitted only on those 42 off-manifest shapes (plateau 1.344–1.392), and on a 44-shape sweep the selector’s regret falls from 0.89% to 0.06% geometric mean and from 5.04% to 0.99% max. #2838 fixes the dispatcher treating K as the contiguous dimension and thereby rejecting valid TN-layout K=1 calls (A=[1,64], B=[1,64] computing AᵀB could use the cp.async kernel); it now falls back to checking that both operands’ physical contiguous dimension has a row width of at least 4 bytes, adding 24 test nodes. #2839 (submitted 00:56) makes FP8 grouped GEMM load A’s scales via TMA (four K steps per box, stage_scales 2, calibrated on three rows) and switches non-fused tight-grouped BF16/FP16 GEMM with K <= 1536 to a 64x256 ping-pong tile; FP8 qwen3-30b-ep8-skew-up rises from 0.973x to 1.015x DeepGEMM, and BF16 skew-down from 0.962x to 1.099x QuACK. #2840 (submitted 06:51) adds prefill sinks for SM89 and folds sink presence into the compiled instance identity.

Sources: https://github.com/tile-ai/TileOPs/pull/2836, https://github.com/tile-ai/TileOPs/pull/2838, https://github.com/tile-ai/TileOPs/pull/2839, https://github.com/tile-ai/TileOPs/pull/2840

3.5 TileFoundry: WGMMA synchronization plan merged (#286); completion semantics for partial reductions under review (#288)

#286, merged at 13:40, was the item under review last issue and addresses WGMMA asynchrony: after an instruction is issued it keeps reading shared memory and registers, so subsequent accesses must be ordered after its effects. The approach declares completion conditions per operand effect, attaches fence, commit, and wait to the HIR nodes that need them, and has lowering consume these steps at the instruction’s issue point; if the subsequent scope cannot hold a collective wait, it is done in the issuing scope. It also fixes aliasing in register Reshard so each element stays in the same register on every thread holding a copy, and tracks register dependencies by this storage identity; on SM90, CUDA units that issue WGMMA group steps are compiled separately with sm_90a. The author keeps two caveats: the matrix instruction emitter is still not wired up, so this PR proves neither execution correctness nor performance; and TMA completion semantics are handled separately. #288 was opened at 20:12 to address how “completed partial reductions” are priced in analysis.

Sources: https://github.com/tile-ai/TileFoundry/pull/286, https://github.com/tile-ai/TileFoundry/pull/288

3.6 TileSight: WGMMA capability declaration fixed for H20 — the repo’s first PR (#1)

Submitted 14:37, updated 17:26, this is the first PR in tile-ai/TileSight (24 stars). The repo is the open-source implementation of the 07-24 arXiv paper “TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters”: starting from tile shapes, data movement, software pipelining, and hardware profiles, it estimates operator cost, cache traffic, and execution time on CPU, with CUDA/HIP probes measuring the target GPU separately. The change is small but pointed: H20 is Hopper/GH100, and plans with requires_arch=("wgmma",) previously threw UnsupportedError because the architecture profile did not declare WGMMA support; it now sets support_wgmma=True while explicitly keeping support_utcmma=False, so such plans are accepted while tcgen05 and Tensor Memory requirements are still rejected. The author also notes that capacity must be handled per SKU (NVIDIA lists both 96 GB and 141 GB H20), and this change touches only capability metadata — capacity, bandwidth, frequency, throughput, and utilization values are all unchanged. Validation is 847 tests passing on CPU, explicitly stated as “not GPU performance validation.”

Sources: https://github.com/tile-ai/TileSight/pull/1, https://arxiv.org/abs/2607.22432

3.7 Adopters and third parties: No pushes to TileKernels or FlashQLA; two new third-party repos appear (11-05)

deepseek-ai/TileKernels is still at 10-14 and QwenLM/FlashQLA at 10-30, with no commits in this window; FlashQLA’s latest release is still v0.1.3 from 09-30. Two new small third-party repos appeared: Tianchen-ai/tilelang-weft-bench (“TileLang-Weft full CPU operator benchmarks using the shared Weft Bench engine”, 0 stars, building a unified engine for CPU-side operator benchmarks) and flytiger-eco/tilelang-wrapper-for-sail (0 stars, no description). Both are very early; they go on the watch list rather than being written up as technical progress.

Sources: https://github.com/deepseek-ai/TileKernels, https://github.com/QwenLM/FlashQLA, https://github.com/Tianchen-ai/tilelang-weft-bench, https://github.com/flytiger-eco/tilelang-wrapper-for-sail

IV. Community, Tutorials, and Events

4.1 Docs site updated three times; no main-branch commits in nightly repo; Ascend daily tests pass 3030/3030 (11-05)

tile-ai/tilelang.github.io pushed Update docs at 16:07, 19:33, and 23:43, corresponding to that day’s interface changes in the main repo (host-mapped annotations, TF32 rounding, Z3 configuration). TileOPs-nightly has only branch pushes, with no new commits on main; TileOPs.github.io had one push at 09:02. Issue #1988 in the Ascend repo records that the 11-05 08:16 daily test run passed all 3030 items with a 100% pass rate, in contrast with #1991/#1992 opened the same day (see 2.3).

Sources: https://github.com/tile-ai/tilelang.github.io, https://github.com/tile-ai/tilelang-ascend/issues/1988

4.2 News side: No in-window hits across four query sets, Hacker News, or arXiv (11-05)

Across the four Google News queries (English "TileLang" OR "tile-lang" and TileLang GPU kernel; Chinese TileLang 算子 and TileLang 编译器), the most recent hit is SemiAnalysis’s “The Huawei-DeepSeek Symbiosis” from 376 hours ago; the rest cluster around the 09-30 wave of coverage on “DeepSeek open-sources Ascend versions of TileLang, DeepGEMM, DeepEP” (about 880 hours ago), all predating this window. A Hacker News search for tilelang returns 0 items in the past 24 hours; the three Show HN hits captured are a terminal data viewer, a workflow engine, and a macOS file manager — unrelated to the topic. The latest arXiv all:tilelang hit is 11-03’s “Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving” (outside the window; a work that cites TileLang), preceded by the TileSight paper itself on 07-24. The report body therefore still comes entirely from the code side.

Sources: https://news.google.com/, https://news.ycombinator.com/, https://arxiv.org/

V. Trend Observations

First, the competitive focus of operator libraries has shifted entirely from “number of operators” to “selector and preparation cost.” Of today’s nine TileOPs items, only #2832 (broadcast in-place read) and #2834 (native-width tiles) actually change the GPU kernel; #2835 saves call preparation (4.33% to 6.17%, with generated source unchanged), and #2836 changes a single fitting factor, cutting regret from 0.89% to 0.06%. This shows that the remaining headroom in performance gaps lies not in “how to compute” but in “which tile to pick and which preparation work to skip.” Taken to its conclusion, this path yields a cost model described by a calibration contract rather than a pile of hand-tuned parameters—today’s #2836, which “fits separately and moves only one constant,” is a prototype of exactly that form.

Second, the dominant defect type on the compiler side is currently “silently generating wrong code,” not compilation failure or insufficient performance. #3808 was missing a __syncthreads() across four versions from 0.1.12 to 0.1.15; #3800 lets MMA read an overwritten operand on the Blackwell path; #1991/#1992 make sin/cos compute 97.1% of elements incorrectly under Ascend with threads=2—none of the three reports an error, all pass routine tests, and all require a person to write a targeted repro to catch. This class of defect hurts users the most, because it turns the premise “is the compiler correct?” into an assumption that must be verified. The corresponding engineering action should be to expand differential testing across shape combinations (comparing element-by-element against a reference implementation as the control group), rather than continuing to add to the number of unit tests.

Third, the rhythm on the Ascend side is shifting from “clearing regressions after architectural changes” to “clearing correctness boundaries.” The four categories of regression brought by structural changes like #3512 were just cleared in the previous window; this window’s two merged items (segfault converted to error reporting, performance skills split by generation) and two new issues (tmp workspace halved, calCount <= tmp violation) are all boundary convergence. Notably, the reporter of #1991 was a reviewer of #1897, yet the fix for #1897 covered only the count side—another facet of the same layout logic was pointed out only two weeks later, indicating that the layout invariants of the Ascend backend have yet to form a checkable written contract.

Fourth, the “communalization” of domestic backends is accelerating. Hygon has moved from a self-maintained branch to two formal PRs, Moore Threads has resumed documentation maintenance, and mlir-ascend is moving arithmetic operators one by one into the TLIR dialect via three small PRs. What the three lines share is that none of them waits any longer for “a one-shot merge of a large backend”; instead they push the smallest reviewable units into the public repository. This is good for upstream: reviewers can comment on layout and lowering design early, rather than only when a 30,000-line one-shot merge arrives.

One final risk surface: if the TIRx migration in main repo #3803 proceeds, the vendored structure of 3rdparty/tvm will be replaced by two pinned upstream commits, at which point every backend that depends on TVM internal APIs (the glue layers of Ascend, Hygon, and mlir-ascend) will have to change accordingly. The draft explicitly states “no dependency patches injected,” which is both a commitment and a constraint—today’s moves, with mlir-ascend moving the BishengIR adapter into contrib and Ascend forking performance skills by chip generation, can both be seen as replaceable surfaces reserved for this kind of foundation switch.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
tile-ai organization (29 repos) 10 repos pushed within the window
Main repo tilelang 5 merges to the main branch; 2 new issues (#3808, #3800); in-review updates clustered around scalarization, CuTeDSL, Profiler, and Ascend
TileOPs 9 merges to the main branch; 4 new or updated in-review PRs; no new issues
TileFoundry 1 merge to the main branch (#286); 1 in review (#288)
TileSight First PR (#1); no commits to the main branch
tilelang-ascend 2 merges to the default branch ascendc_pto (#1932, #1954); 3 new issues (#1991 and #1992 on the same topic; #1988 is a daily test report)
tilelang-mlir-ascend No new commits to the main branch; three TLIR operator PRs (#203, #211, #212) advancing
tilelang-hygon No new commits to the main branch (stalled since 11-02); 3 commits on the feat/hcu-scaled-cast branch; new PRs #33 and #34
tilelang-musa 1 documentation commit to the main branch (first since 09-30); timestamps refreshed on 6 legacy PRs with no content changes
tilelang-metax No pushes within the window (third consecutive window)
TileOPs-nightly / TileOPs.github.io / tilelang.github.io No commits to the nightly repo’s main branch; 1 push to the TileOPs docs site; 3 Update docs pushes to the tilelang docs site
Release pages (tilelang / TileOPs / TileFoundry) No new releases within the window; tilelang’s latest remains v0.1.15 from 09-30, TileOPs v0.1.0 from 10-27, and TileFoundry v0.0.4rc1 from 10-25
TileKernels / FlashQLA No pushes within the window
TileLang × RISC-V four-source sweep No changes (references and PR status across the four upstreams match the previous snapshot)
Google News (four queries, Chinese and English) / Hacker News / arXiv No substantive hits within the window; the most recent hit was 376 hours ago

Full Source List

Verification Notes

All merge, in-review, issue, and commit counts in this issue come from GitHub search and measured output from atomic sources: merges and in-review items via the search API (18 merges, 57 in-review updates, 5 new issues), commit titles via each branch’s commits atom (12 branches, 21 commits falling within the window), and release versions via the releases atom. The organization-level anonymous API hit its hourly limit late in collection; the organization inventory outside the main branches was cross-checked against web pages and atom feeds, with no estimated values used. Issues #1991 and #1992 have identical titles and bodies and are recorded as “two issues on the same topic” rather than merged into one. The three Hygon commits and one Moore Threads commit are unmerged commits, described at the branch level and not as merged into the mainline. The WGMMA-related changes in TileSight and TileFoundry are both qualified by their authors as capability metadata or synchronization plans only, with no claims of execution correctness or performance; this report retains those qualifications. All three news channels — Google News, Hacker News, and arXiv — had no hits in this window.