Research window: 2026-11-03 07:00 ~ 2026-11-04 07:00 (Beijing time, 24 hours; previous cutoff at 11-03 07:00).

Sources: GitHub tile-ai organization per-repo check (29 repos, 9 pushed within the window) + merged and in-review ground-truth search (8 merged, 48 in-review updates) + full check of newly added issues (7) + per-commit check of main branches (tilelang 5, TileOPs 1, TileFoundry 1, VS Code extension 1, docs site 3) + release page and tag check (5 repos, no new versions within the window) + two adopter repos + Google News Chinese and English queries (via proxy) + Hacker News + arXiv.

This Issue’s Index

  • Today’s Highlights: ThreadSync reuses partial barriers across storage domains, restoring the performance regression introduced when #3643 fixed correctness (#3761); cleanup of 3,384 lines of legacy TVM parser branches (#3736); autotuning input feeding switched to the current device (#3546); TileFoundry’s sync plan formally lands in CUDA lowering (#285); 8 merges, 48 in-review updates, and 7 new issues in 24 hours
  • I. Core Project Progress
    • 1.1 Mainline: ThreadSync allows partial barriers to be reused across storage domains; barrier occupancy on B300 returns from 12 to 16 blocks/SM (#3761) (11-04)
    • 1.2 Mainline: Remove legacy TVMScript AST and parser branches, deleting 3,384 lines in one go (#3736) (11-03)
    • 1.3 Mainline: Autotuning input feeding now uses the current device; tuning on Ascend and MPS is usable again (#3546) (11-03)
    • 1.4 In review: Scalar casts in conditions no longer compress data vector width, following up on yesterday’s #3734 (#3759) (11-03)
    • 1.5 In review: Two user-visible errors changed to explicit error reporting (#3765 / #3767) (11-04)
    • 1.6 In review: Fallback when padded shared tiles cannot be copied wholesale; 2-CTA TMA bytes accounted per peer leader (#3738 / #3766) (11-03)
    • 1.7 In review: B300 sm_103a adaptation across two lines, tcgen05 int8 selection and CuTeDSL TMA architecture gate (#3770 / #3772) (11-03)
  • II. Multi-Backend Adaptation
    • 2.1 Ascend: Batched ND2NZ copy from 3D GM to L1 merged, one descriptor covers the whole batch (#3689) (11-03)
    • 2.2 Ascend in review: TileLang-Ascend entire repo migrated into mainline multi-backend architecture, 2,887 files (#1975) (11-04)
    • 2.3 Ascend in review: copy_pipe_to_l1 sync annotations split by execution engine, fixing wrong results reproducing 10/10 on A2/A3 (#1981) (11-03)
    • 2.4 Ascend in review: First Parallel axis broadcast uses scalar reads, Chunk O forward migration, MHA example (#1979 / #1866 / #1879) (11-03)
    • 2.5 mlir-ascend: Codex workflow and task handoff merged into harness branch (#208) (11-03)
    • 2.6 Hygon: Distributed in-review PR updated, no new commits on main branch (#19) (11-03)
    • 2.7 MetaX and Moore Threads: no pushes during the window (11-03)
  • III. Ecosystem and Adopters
    • 3.1 TileFoundry: Sync plan implemented as CUDA lowering with reverse validation, merged into mainline (#285) (11-04)
    • 3.2 TileFoundry in review: wgmma sync plan, FSDP and HSDP training semantics into HIR evaluator (#286 / #284) (11-03)
    • 3.3 TileOPs: Variance merging starts from the first partial, variance of constant inputs no longer yields NaN (#2821) (11-03)
    • 3.4 TileOPs in review: GEMM with optional bias, FP8 grouped GEMM supports contiguous layout (#2823 / #2824) (11-03)
    • 3.5 TileOPs in review: W4A16 and dense GEMM selectors refitted as a single cost model (#2825 / #2826) (11-03)
    • 3.6 TileOPs in review: Short varlen packed by KV group, L2 flush changed to read-only (#2819 / #2822) (11-03)
    • 3.7 Adopters: No pushes from FlashQLA and TileKernels, but #3761 fixes TileKernels’ fused kernel (11-04)
  • IV. Community, Tutorials and Events
    • 4.1 7 new issues in a single day: one production-grade CUDA graph issue, four from the same external reporter (11-03)
    • 4.2 Ascend 950 vcvt saturation semantics and mlir-ascend daily tests (#3762 / #1980) (11-03)
    • 4.3 Documentation site updated three times, VS Code extension renamed to vsix (11-03)
    • 4.4 News side: no hits across three channels within the window; foreign media reports on Ascend training vary in framing (11-03)
  • V. Trend Observations
  • Appendix: Materials and Verification Notes

Today’s Highlights

The main thread today is a fix that “takes back the performance.” #3643, merged in the previous window, addressed a real correctness issue: separate calls to ThreadSync("shared") and ThreadSync("shared.dyn") could assign the same named barrier ID to disjoint thread groups. The fix was to have the later call preserve IDs already consumed by the earlier one—but this preservation rule also blocked something else: the same group of threads could originally reuse the same ID across two scopes. The cost landed directly on TileKernels’ mhc_pre_big_fuse kernel: this kernel has 96 threads per block, with threads 32 to 95 running a two-stage cp.async pipeline. The sync after the async wait gets ID 3 on the first pass, and the buffer-dependency sync for the same 64 threads gets ID 4 on the second pass. That extra ID raised the barrier resources declared in the compiled artifact from 4 to 5; when running hidden dimension 1280 on B300, the barrier occupancy ceiling thus dropped from 16 blocks/SM to 12 blocks/SM, while register, shared memory, and sync instruction counts were unchanged. #3761’s approach is to have ThreadSync accept an ordered list of scopes, with the CUDA pipeline changed to ThreadSync(["shared", "shared.dyn"]). Sync is still planned and inserted in the original scope order, but some barrier IDs are allocated uniformly only after both scopes have been processed. Before sharing an ID across scopes, a brand-new analyzer proves that the participation predicate holds over the entire candidate thread range; branch and Bind contexts are provided by ConstrVisitor, thread-related predicates are treated as proof obligations rather than assumed true, proven-dense groups use the bounding-box volume directly and skip enumeration of partial thread counts, unproven groups keep the scope in the key and include the participant count in the key, and explicitly numbered barriers (including IDs allocated by an earlier ThreadSync) continue to be preserved. The change is 401 lines across 4 files.

The author did a very clean attribution experiment in the description: changing only the number of barrier resources declared in the compiled artifact, without touching any executable instruction, reproduces the same slowdown. This kind of experiment is worth noting—it turns “performance regression” from an impression into an identifiable single variable, and it makes the acceptance criterion for the fix explicit: not a vague “benchmarks got faster,” but resource declarations back to 4 and occupancy back to 16 blocks/SM.

The second item is supply-side subtraction and consistency of conventions. #3736 deletes tilelang/language/ast/, tilelang/language/parser/, and tilelang/language/tir/entry.py together. These are modules left over from an early local fork of the TVMScript parser, about 3380 lines of Python, with no remaining callers in the repo: the language layer now uses the upstream tvm.tirx.script.parser interface through tilelang.language.tir.common, AST and statement override registration lives in tilelang.language.overrides, and T.prim_func and T.macro are exported by tilelang.language.eager. The actual change is 1 line added, 3384 lines deleted, 10 files, while also removing the ignore rule for *language/ast* in docs/conf.py; public language-layer exports are covered by dialect tests and frontend port tests. #3546 fixes the device convention for autotuning: the supply closure installed by set_autotune_inputs resolved the benchmark device from a parameter that no call site ever passes—what is actually passed is the kernel argument list—so the CUDA or CPU fallback branch moved every captured tensor to CPU on machines without CUDA, and every tuning round failed due to a launcher device mismatch. The closure now uses get_current_device(), the same policy as default tensor supply: CUDA behavior is unchanged, Ascend and MPS now supply on the current device, and other backends still fall back to CPU. Acceptance was on dual-card Ascend950: three Ascend autotuning regressions all failed before the fix and all passed after; in the main repo, 12 tests in testing/python/autotune/ passed and 28 were skipped, with another 3 failing because that machine has no CUDA, which the author explicitly left for CUDA CI to verify.

The third item is on the Ascend side: GM-to-L1 copies expand from “one descriptor per matrix” to whole batches. #3689 allows a [B, M, N] region to be written as a single T.copy; lowering generates batched ND2NZ descriptors, preserving the source-side batch and row strides, deriving the destination-side batch pitch from the physical layout, and routing transposed batches through DN2NZ; source batches beyond the GM boundary are skipped directly while destination storage is left untouched, and M/N tail zero-padding applies only to valid batches. Existing 2D and single-matrix behavior is unchanged; the implementation targets only the AscendC backend and requires the L1 destination to be in canonical matrix layout, with PTO out of scope. Validation was on Ascend950DT with CANN 9.2: 50 AscendC example tests and 1733 Ascend unit tests passed, of which 28 device cases cover batch and M/N tail blocks, whole-batch out-of-bounds, non-contiguous source strides, larger destination storage, and rectangular transpose; CUDA and PTO cases were excluded, and the CI image’s CANN 9.3 was not tested locally, which the author noted in the description. The change is 579 lines across 10 files.

The fourth item is the landing of a hidden thread in TileFoundry. In the previous report it was still listed as “under review”; it merged in this window: #285 makes the commits, waits, and fences planned by sync analysis actually emitted by CUDA lowering, and reverse-validates each one against the plan, throwing LoweringError on mismatch. On the corpus, the default flow yields 14 programs plus sync and rejects 31; for those 14, the TIR differs from mainline only by the newly added statements, and another 30 are identical to mainline; on a single sm_89, 7 programs that build and carry sync were run, 6 matched the reference exactly with clean synccheck and memcheck, while lowering those same 7 the mainline way produced 4 wrong results. The direct symptom the author cites in the motivation is a maximum deviation of 659 for E4M3 GEMM on sm_89. This change adds 788 lines, deletes 214, across 25 files, advancing from “the analysis can prove it” to “emission must match the proof.”

I. Core Project Progress

1.1 Mainline: ThreadSync allows partial barriers to be reused across storage scopes; barrier occupancy on B300 returns from 12 to 16 blocks/SM (#3761)

Merged 11-04 00:50, +401 / -72 lines, 4 files. Background is #3643 from the previous window: when ThreadSync("shared") and ThreadSync("shared.dyn") are called separately, they may assign the same named barrier ID to disjoint thread groups — a real correctness issue — and the fix was to reserve IDs already consumed by earlier calls. The side effect was that the same thread group could no longer reuse an ID across scopes either; on TileKernels’ mhc_pre_big_fuse kernel this meant one redundant allocation per 96-thread block: the two-stage cp.async pipeline for threads 32 to 95 got ID 3 on the first pass and ID 4 on the second, the compiled artifact declared 5 barriers instead of 4, and the barrier occupancy ceiling for B300’s hidden dimension 1280 dropped from 16 blocks/SM to 12 blocks/SM, while registers, shared memory and sync instruction counts were unchanged. This PR makes ThreadSync accept an ordered scope list, and the CUDA pipeline now emits ThreadSync(["shared", "shared.dyn"]): syncs are still scheduled and inserted in the original order, and partial barrier IDs are allocated uniformly after both scopes are processed; before sharing an ID across scopes, a new analyzer proves the participation predicate holds over the candidate thread range, with branch and Bind contexts supplied by ConstrVisitor and thread-dependent predicates treated as proof obligations rather than assumptions — dense groups already proven are sized by bounding-box volume, skipping partial thread-count enumeration, unproven groups keep the scope in the key with participant counts folded into the key, explicitly numbered barriers (including IDs allocated by earlier calls) are still preserved, and disjoint thread groups remain isolated — the #3643 fix is not reverted. Legacy string-form calls remain compatible, and CuTeDSL goes through the same CUDA pass pipeline.

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

1.2 Mainline: Remove the legacy TVMScript AST and parser fork, deleting 3384 lines in one go (#3736)

Merged 11-03 13:05, +1 / -3384 lines, 10 files. Removed were tilelang/language/ast/, tilelang/language/parser/ and tilelang/language/tir/entry.py — leftovers from an early local fork of the TVMScript parser, with no callers anywhere in the repo: the language layer uses the upstream tvm.tirx.script.parser interface via tilelang.language.tir.common, TileLang’s own AST and statement overrides are registered in tilelang.language.overrides, T.prim_func and T.macro are exported from tilelang.language.eager, and only four modules under tilelang/language/tir/ are consumed: common, exports, ir, op. The risk of keeping this orphaned parser fork is stated concretely: anyone importing tilelang.language.parser.operation directly would register conflicting operators on tirx.PrimExpr and tirx.IterVar. The PR also drops the unused re-export of prim_func in tilelang/language/tir/__init__.py and the ignore rule for *language/ast* in docs/conf.py (*parser* is kept because tilelang/language/overrides/parser.py still uses it). Acceptance criterion is that public exports are unchanged, covered by test_tilelang_language_dialect.py and the language frontend tests.

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

1.3 Mainline: Autotuning input supply uses the current device; tuning works again on Ascend and MPS (#3546)

Merged 11-03 13:19, +153 / -12 lines, 3 files, fixes #3545. The frozen supply closure installed by set_autotune_inputs resolved the benchmark device from a parameter that is never passed at the call site — per the supply program’s contract, what is actually passed in is the kernel argument list — so on machines without CUDA, the CUDA-or-CPU fallback branch moved every captured tensor to CPU and every tuning round failed with a launcher device mismatch. The closure now uses get_current_device(), the same policy as the default tensor supply: behavior on CUDA is unchanged and multi-GPU benchmark workers still copy per card; on Ascend and MPS, captured tensors are supplied on their respective current devices, so tuning is usable there; other backends still fall back to CPU. New Ascend cases live in testing/ascend/autotune/, covering the decorator path, the benchmark supply device of AutoTuner, and cross-device captured inputs; on a dual-card Ascend950 all three failed before the fix and pass after; in the main repo testing/python/autotune/, 12 pass and 28 are skipped, with 3 more (eager mode, scalar inputs, with inputs) failing because that machine has no CUDA — the author notes these are left for the CI CUDA runner.

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

1.4 Under review: Scalar casts in conditions no longer narrow data vector width, following up on yesterday’s #3734 (#3759)

Updated 11-03 14:09, +148 / -29 lines, 2 files. This closes out yesterday’s merged #3734: #3734 correctly tightened the width limit for data casts after DecoupleTypeCast, but it also made casts inside scalar conditions start limiting data vector width — on b5cb02633, a guarded BF16 copy could be squeezed from 16 lanes to 4 because the condition contained an int64 cast, even when that condition is uniform across all 16 lanes. The fix reuses the planner’s existing condition-uniformity analysis for the cast width exemption, visiting if statements, tirx.if_then_else and Select conditions of scalar predicates separately from their clauses: casts in conditions no longer impose a data-type width limit, casts in clauses still do, and the condition-cast constraint is retained on targets with per-lane vector predicates; the memory-guard behavior of #3579 and the guarded FP32-to-BF16 limit fixed by #3734 are both unchanged. Regression coverage includes uniform guards at 8 and 16 lanes, all three condition forms, FP32 casts inside clauses, vector predicates and rounding-bit grouping; the nine int64 guard cases fail before the fix, and the final legalization and cast-rounding suites pass 74 tests in total.

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

1.5 Under review: Two user-visible errors become explicit diagnostics (#3765 / #3767)

Both updated around 11-04 01:10, filed by the same reporter. #3765 corresponds to #3653: for a partial barrier guarded by a single thread-index bound, when that bound is a multiple of the warp size, all threads in a warp agree on the condition; but the general arithmetic proof missed thread_binding < 448 when the enclosing branch uses bitwise &, and ThreadSync aborted outright. This PR explicitly accepts this case when the thread dimension is a whole number of warps (also covering bounds equal to or beyond that dimension), while other conditions still go through the existing proof path. #3767 corresponds to #3005: when a fragment is distributed across threads and the same fragment is read and written with two structurally different indices in one T.Parallel body (e.g. a cross-lane pattern like buf[i] and buf[(i + 1) % N]), it cannot be lowered without cross-thread exchange; previously ParallelOpNode::RecordBufferAccess hit an ICHECK in StructuralEqual and surfaced as a compiler internal assertion; it is now a ValueError that includes both index expressions and explains that the fragment cannot be accessed this way. Both are “replace an internal assertion with a diagnostic the user can understand” changes whose perceived impact outweighs their size.

Source: https://github.com/tile-ai/tilelang/pull/3765, https://github.com/tile-ai/tilelang/pull/3767

1.6 Under review: Fall back when a padded shared tile cannot be copied wholesale; account 2-CTA TMA bytes to the peer leader (#3738 / #3766)

Both updated between 11-03 21:31 and 21:47. #3738 fixes #3704: a T.copy from a global region smaller than the padded shared tile (e.g. K=100 into BK=112) was classified as a full-tile TMA copy by ClassifyWarpSpecializedCopy and SelectCopyInstForLowering, then failed hard in lowering with s_extent == g_extent (112 vs. 100). The cause is that the selection logic only checked last-dimension alignment, dtype pairing and global stride, not whether the padded extent can form a single whole-tile copy; this PR adds that criterion and falls back to the normal copy path when it does not hold. #3766 fixes TMA byte counting in cluster scenarios: a cta_group::2 TMA load clears the peer bit of the barrier address, so the bytes of a CTA pair both land on the even CTA’s barrier, whereas tma_copy with a cluster barrier previously expected cluster_size × tile_bytes only on rank 0 — which happened to be right for (2,1,1) and (1,2,1), but for (4,1,1) and (2,2,1) rank 0 received only 2 shares while expecting 4, with the other bytes landing on rank 2’s barrier. Neither is a new feature; both are the “internal error the moment a boundary condition breaks” class of defect.

Source: https://github.com/tile-ai/tilelang/pull/3738, https://github.com/tile-ai/tilelang/pull/3766

1.7 Under review: Two tracks for B300’s sm_103a — tcgen05 int8 selection and the CuTeDSL TMA arch gate (#3770 / #3772)

B300 reports its target architecture as sm_103a, and two PRs under review each address one of the blockers it creates. #3770 corresponds to #3769: tcgen05.mma.kind::i8 exists only on the architecture-specific sm_100a and sm_110a — neither on sm_103a (B300) nor in the family-portable sm_10xf — yet AllowTcgen5Mma accepts int8/uint8 operands for any TargetIsSm100 (arch 100 to 110), so the kernel passes lowering and then ptxas reports Feature '.kind::i8' not supported on .target 'sm_103a'; this PR stops TileLang from selecting kind::i8 on these targets. #3772 handles the CuTeDSL TMA arch gate: B300’s sm_103a is rejected before entering the DSL, with an error listing sm_90/90a/100a/110/120/120a/121/121a but not 103a; on mainline fa69f7fa with nvidia-cutlass-dsl 4.7.0 and CUDA 13.3, six CuTeDSL tests fail on B300, including several JIT cases for tma_load. The same batch includes #3768, fixing the SM100 flash attention backward example: T.tcgen05_gemm now makes mbar a required keyword-only argument and no longer waits implicitly, while the example still calls it like a synchronous GEMM, so both scripts fail before compilation with missing 1 required keyword-only argument; the author also notes there is no test_*.py under examples/flash_attention_sm100/, which is why the problem went unnoticed.

Source: https://github.com/tile-ai/tilelang/pull/3770, https://github.com/tile-ai/tilelang/pull/3772, https://github.com/tile-ai/tilelang/pull/3768

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

2.1 Ascend: Batched ND2NZ Copy from 3D GM to L1 Merged, One Descriptor Covers the Whole Batch (#3689)

Merged on 11-03 17:07, +579 / -39 lines, 10 files. Previously, Ascend’s GM-to-L1 copy handled only one matrix per DMA descriptor, so copying a batch required for b in T.serial(B) with per-batch calls; this PR allows a [B, M, N] region to be written as a single T.copy, with lowering generating a batched ND2NZ descriptor that preserves the source’s batch and row strides, derives the destination-side batch pitch from the physical layout, routes transposed batches through DN2NZ, and reuses the existing matrix-axis boundary handling: source batches beyond the GM boundary are skipped and the destination storage is left untouched, while M/N tail zero-padding applies only to valid batches. Existing 2D and single-matrix behavior is unchanged; the implementation targets only the AscendC backend and requires the L1 destination to be in canonical matrix layout, with PTO support out of scope for this change. Validation was on Ascend950DT with CANN 9.2: 50 AscendC example tests and 1733 Ascend unit tests passed, including 28 BF16/FP32 device cases covering batch and M/N tail blocks, whole-batch out-of-bounds, non-contiguous source strides, destination storage larger than source, and rectangular transposed copies. The author notes that CUDA and PTO cases were excluded and that the CANN 9.3 used by the CI image was not verified locally. This pairs with the Ascend 950 FlashMLA kernel merged in the previous issue: both attention and GEMM batched transfers need one descriptor to cover the whole batch, since per-batch calls incur substantial launch overhead at production shapes.

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

2.2 Ascend Under Review: TileLang-Ascend Whole-Repo Migration into Mainline Multi-Backend Architecture, 2887 Files (#1975)

Updated 11-04 00:34, with a change size of +607317 / -97595 lines across 2887 files — the largest item this issue and the most worth tracking among all org PRs under review. The goal, per proposal #1911, is to move TileLang-Ascend onto the mainline multi-backend architecture: existing Ascend components move under the corresponding mainline directories’ ascend/a2a3 (tilelang/ to tilelang/ascend/a2a3, src/ to src/ascend/a2a3), and the original Ascend C and PTO A2A3 backends register independently in the mainline, orthogonal to the mainline Ascend and PTO A5, so passes and codegen do not interfere and only a small amount of infrastructure is shared; both sides share target="ascend" and "pto", but the A2A3 backend carries an arch="dav-2201" tag (auto-detected from the NPU name) for dispatch; A2A3 frontend primitives switch to import tilelang.ascend.a2a3.language as T, coexisting with A5’s import tilelang.ascend.language as T, with operator code unchanged relative to the original repo; passes and codegen use versions migrated directly from TileLang-Ascend to preserve the original compilation flow, the execution backend reuses the mainline cython with Ascend-specific adaptations such as workspace indexing added, examples and testing land in examples/ascend_a2a3 and testing/ascend_a2a3 respectively, and CI temporarily uses the original repo’s ci_cd.yml, confirmed to still run after adapting to the new paths. The author is candid about the TODOs in the description: differences between files in the three directories and the originals remain to be reconciled, static-check and formatting CI remain to be restored, docs and Agent-related items remain to be migrated, and notes that “given the number of files involved, batch Agent operations will inevitably have omissions,” requesting review assistance. Laying migration uncertainties out directly in the PR description is more review-friendly than waiting for a “done” conclusion.

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

2.3 Ascend Under Review: copy_pipe_to_l1 Sync Annotation Split by Execution Engine, Fixing Wrong Results Reproducing 10/10 on A2/A3 (#1981)

Updated 11-03 20:36, corresponding to #1780. The root cause is clearly stated: A2/A3’s PTO cross-core V-to-C communication relays through GM, and the three transfer stages belong to three engines — AIV’s MTE3 does the UB-to-GM send, AIC’s MTE2 does the GM-to-L1 receive (TPOP’s internal TLOAD), and AIC’s MTE1 does L1-to-L0A; but copy_pipe_to_l1 is tagged PIPE_MTE3 in operation_config.h, so the auto-sync pass generates set_flag(PIPE_MTE3, PIPE_MTE1, …) plus the corresponding wait_flag between it and copy_l1_to_l0a — yet it waits on the AIC core’s own MTE3, whose queue is empty and whose event is set immediately, so MTE1’s read of L1 may be released before MTE2 finishes writing, and L0A gets stale data, producing wrong MMA results. The mis-annotation was introduced with #1238; in the same table copy_pipe_to_ub and copy_l0c_to_pipe are both annotated by execution-side engine, and only this one retained the sender’s perspective. Two pieces of reproduction evidence: after changing PIPE_MTE3 to PIPE_MTE2 in the generated source, the issue author saw 10/10 exact recovery, while locally on the latest main, 10/10 runs were wrong, with up to 16384/16384 mismatched elements in a single run. The fix follows the issue’s suggestion to split the name rather than replace it outright: copy_pipe_to_l1 is annotated PIPE_MTE2 (A2/A3’s GM-to-L1 load), and a new copy_pipe_to_l1_bind is annotated PIPE_MTE1 (waiting for L1 readiness and binding addresses on the A5 direct path), following the existing precedent of copy_pipe_to_ub and copy_pipe_to_ub_V to avoid fixing A2/A3 while breaking A5. The same root issue also covers the WAR direction: MTE2’s TLOAD may overwrite the same L1 buffer while MTE1 is still reading.

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

2.4 Ascend Under Review: First Parallel-Axis Broadcast via Scalar Read, Chunk O Forward Migration, MHA Example (#1979 / #1866 / #1879)

All three were updated in a burst on the afternoon of 11-03, forming a milestone batch for Ascend operators. #1979 corresponds to #1885: the failing family reduces to one pattern — a 1D UB operand indexed by the first T.Parallel axis (or an outer serial variable) appears on the left side of a binary expression, exactly the shape of the issue’s reproducer dk_i = sum_j v[j] * A[j, i]; all 5 FAILs in the 21-point bisection matrix belong to this family, independent of reduction direction or broadcast form. The root cause is that when AscendLowerParallelToVector lowers a 2D parallel nest into a row loop, the entry HandleSimpleCase only performs scalar-access detection on the right operand, so the left side takes the vector path and hits the UB alignment trap; this PR makes the first-axis broadcast operand take the scalar read as well. #1866 migrates Chunk O forward (chunk-based gated linear attention O forward, GDN) to the Ascend NPU, using developer mode with four pass configurations (automatic C/V merge, automatic C/V sync, auto-sync, memory planning), a single @tilelang.jit kernel with a fixed 20-core grid-stride loop, achieving on-chip fusion, zero GM workspace, and zero host ATen operators. #1879 adds an MHA forward example: BSHD layout, Expert-mode three-dispatch MMA pipeline, split into 5 Python files across wrapper, shared layer, and three-form modules, with the wrapper containing no T.* code and doing only metadata validation and JIT dispatch; is_causal uses a bottom-right-aligned mask, and fully masked rows are protected by sumexp + 1e-30 to output 0; all 20 cann-bench cases pass accuracy, with an average speedup of 0.97x versus the torch-npu baseline and a composite score of 73.13 (measured locally on 910B2). 0.97x means no net gain yet, but as an example it provides the skeleton for “how to write a complete attention operator on Ascend.”

Source: https://github.com/tile-ai/tilelang-ascend/pull/1979, https://github.com/tile-ai/tilelang-ascend/pull/1866, https://github.com/tile-ai/tilelang-ascend/pull/1879

2.5 mlir-ascend: Codex Workflow and Task Handoff Merged into harness Branch (#208)

Merged 11-03 17:50, +2715 / -39 lines, 66 files, targeting the harness branch rather than main (the mainline’s latest commit is still 10-31’s ci: remove internal cache-service references) — do not describe this as “merged into mainline” when citing it. The content adds the first Codex workflow to the operator harness: project initialization via iterative questioning, onboarding existing operator repos, and resuming development, tuning, or analysis from persisted context; the three skills share editable Markdown records under .workspace/, and task handoff distinguishes “current candidate” from “verified best candidate,” retaining failed and rolled-back experiments so sub-agents get focused context while the main agent maintains shared state; a local installer registers the skills and appends project entries to AGENTS.md (preserving existing content, ignoring .workspace/, resolving Git projects to the worktree root). The README, dependencies, reuse notes, and roadmap all describe a TileLang-first, Ascend-scoped setup; the workflow references GSD and retains source attribution and the MIT notice, with the NPU execution adapter and npm distribution still TODO. This item and the Agent-related directories in the main and Ascend repos are three faces of the same thing: distilling the “agent writes operators” process into reusable in-repo assets.

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

2.6 Hygon: Distributed PR Under Review Updated, No New Commits on Main (#19)

tilelang-hygon had a push in the window (11-03 11:06 Beijing time), but no new commits on the main branch — the latest is still 11-02’s #32 (treating gfx948 as gfx946-compatible so existing instructions can be tested first). The push came from the under-review #19 (title just “Distributed,” empty description), updated 11-03 11:06; judging by timing and title it is advancing distributed-related adaptation, but with no PR description this report makes no inference about implementation details. The repo’s recent cadence is one HCU (Hygon Compute Unit) side fix every two or three days: 10-30’s #31 changed MLS boundaries and persistent WASP sync, #30 added LDS swizzle for normal copies, and 10-27’s #29 did bank partitioning for the scale buffer.

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

2.7 MetaX and Moore Threads: No Pushes in Window (11-03)

tilelang-musa (Moore Threads) last pushed 09-30 09:50, and tilelang-metax (MetaX) last pushed 10-26 07:08; neither repo had commits or under-review PR updates in this window. This issue’s backend activity is therefore concentrated on Ascend and Hygon: of the 8 merges, 1 belongs to Ascend (#3689) and 1 to mlir-ascend (#208, merged into the harness branch); of the 48 under-review updates, the Ascend repo accounts for 15 and the main repo for 23.

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

III. Ecosystem and Adopters (TileOPs / TileFoundry / TileKernels / FlashQLA)

3.1 TileFoundry: Sync plan lands CUDA lowering with reverse validation, merged to mainline (#285)

Merged 11-04 00:19, +788 / -214 lines, 25 files. The sync analysis planned the commits, waits, and fences required for shared-memory accesses, but finalize previously emitted none of them, so schedules staged through shared memory had races—on-device this showed up as an E4M3 GEMM on sm_89 deviating from torch by up to 659. This PR attaches the analysis steps to the HIR nodes they lead: SyncMetadata indexes them by tag—the node’s own statements carry an empty tag, copies added by InsertSlice use base and update, and loop-carried copies use ("entry", slot) and ("backedge", slot). The core lowering calls a new hook before(owner, ctx, tag) before lowering a node’s statements, and CudaLowering emits that node’s steps there; schedule finalize runs the sync analysis. The contract is written into four spec documents, with analysis spec §1.2.5 describing SyncMetadata, tags, rules, and the cases the analysis still rejects. The author also flagged the risk surface: every program in the corpus that can finalize on mainline still finalizes, only with extra statements—the exception is hir/staged_dynamic_inner, whose lanes read and write the same shared tile inside a loop where “trip count equals lanes,” and no fence can converge them; additionally, reserved calls touching shared memory are treated as “one access of unknown bytes,” and two such calls in a region without a thread grid need a fence TIR cannot place, so the analysis rejects them even when the bytes are disjoint; CopyAsyncTensor, CopyReduceAsyncTensor, and wgmma have no emitters yet.

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

3.2 TileFoundry under review: wgmma sync plan, FSDP and HSDP training semantics into the HIR evaluator (#286 / #284)

#286 updated 11-04 02:28, +1907 / -310 lines, 52 files, a direct follow-up to #285: WGMMA may keep reading shared memory and registers after issue, so the CUDA sync plan needs to order subsequent accesses after those effects; this PR declares completion semantics for each operand effect, attaches WGMMA’s fences, commits, and waits to the HIR nodes that need them, and lowering consumes those steps at the instruction’s issue point—if a later scope cannot fit a set wait, the operation completes in the issuing scope. It also fixes register Reshard aliasing so an element stays in the same register on every thread holding a replica, and the plan tracks register dependencies by this storage identity. The author explicitly states the risk: the matrix instruction emitter is still unavailable, this PR does not establish WGMMA execution correctness or performance, and TMA completion semantics remain separate. #284 updated 11-03 14:34, handling training semantics: the previous evaluator treated reshard as a global tensor and missed contributions from different participants; this PR evaluates participant-local values, GPU coordinates, calls and loops, gather/scatter, and per-axis Partial(sum) resolution, and gives Reshard alias proofs ownership semantics, extending the structured lifetimes of #260 and #170. It adds a modded-nanogpt-inspired two-block GQA study with explicit forward/backward, momentum SGD, and an independent PyTorch reference, covering FSDP and 2×2 and 4×4 HSDP residency strategies, reporting collective groups, payloads, and storage. Acceptance numbers: 775 core-suite and 778 backend-suite tests after the refactor (13 skipped on two-card sm90), ten HIR configurations passing on CPU, and both strategies running two update steps on four-card sm90 against real 2×2 PyTorch HSDP with local gradients, parameters, and momentum all matching.

Source: https://github.com/tile-ai/TileFoundry/pull/286, https://github.com/tile-ai/TileFoundry/pull/284

3.3 TileOPs: Variance merge starts from the first partial, no more NaN for constant inputs (#2821)

Merged 11-03 18:04, +43 / -4 lines, 3 files. The defect was that variance merging started from an empty accumulator: merging a large finite mean into an empty accumulator overflowed the squared difference before it was multiplied by zero weight, so the variance of constant inputs became NaN; affected were var, std, and var_mean calls with a reduction on each side of the retained axis. The fix initializes the merge from the first real partial, with subsequent partials following the existing Chan recurrence; row reduction, dispatch, output layout, and the public API are unchanged. Acceptance: three new regressions fail on the original kernel and pass after the fix, the candidate passed 155 regression tests plus 24 extra numerical checks, and tests/ops/test_reduce.py grew from 155 to 158; three autotuning tests were excluded because earlier baseline attempts exceeded the 60-second test limit. This kind of “empty accumulator + large offset” numerical trap is not rare in the reduction family; what’s worth noting is that it only surfaces on shapes that reduce on both sides of the retained axis—ordinary whole-axis reductions won’t catch it.

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

3.4 TileOPs under review: GEMM with optional bias, FP8 grouped GEMM supports contiguous layouts (#2823 / #2824)

#2823 adds an optional bias[N] to GEMMFwdOp.forward: bias is added to the float32 accumulation and the output is rounded only once, matching torch.nn.functional.linear; each dense structure adds bias in its own epilogue—warp-specialized (TMA and row-store), coop2 (including 176/192-wide and narrow tiles), coop2_swap, ping-pong, coop2s, simple, swap_ab, GEMV each with cp.async kernels—and the split-K path adds it only once in the reduction. #2824 lets MoEGroupedGEMMFP8FwdOp accept contiguous expert rows (compact or aligned packing, with segment-end offsets physical_psum or per-row expert ids) in addition to masked blocking, with semantics matching deep_gemm.m_grouped_fp8_gemm_nt_contiguous; the SM90 kernel walks each expert’s run segment by segment-end, compact per-row ids are converted to segment-ends by binary search before the producer/consumer split, aligned per-row ids walk tile by tile reading the tile’s expert from the first row, and a compact expert’s final block with fewer than block_m rows is stored row by row, never writing onto the next expert’s rows; SM89 rejects contiguous layouts. Masked calls still build the same programs as before—35 masked-built TIRs are byte-identical to mainline. The benchmark side adds three Qwen3-30B-A3B EP8 shape rows using group sizes measured from real routing (the model has 128 experts, top-8; layer 44 of 48 was chosen for its mid-range skew, and the 16 experts of the EP8 rank with the most rows were selected): prefill rows route 4096 tokens (gate/up and down), and decode rows take the last token of 128 sequences.

Source: https://github.com/tile-ai/TileOPs/pull/2823, https://github.com/tile-ai/TileOPs/pull/2824

3.5 TileOPs under review: W4A16 and dense GEMM selectors refit to a single cost model (#2825 / #2826)

Both belong to subtask 2 of #2133, updated 11-03 20:54 and 20:54 respectively—one covers W4A16, one covers dense, and the latter only changes documentation. #2825 changes _select_config to take the argmin of _config_cost over _legal_configs, removing the three exception rules: tile-first, widest-K tie window, and joint_switch_margin; _config_cost gains a stage_fill term—deeper rings hide more per-K-tile latency but fill more slowly, and without this term the model ranks 256-wide four-stage ahead of 512-wide two-stage at K=4096, while the latter measures up to 11% faster; mma_big, iter_ws, and launch_ms were removed because the data was insufficient to pin them down. Constants were refit on busy time at the benchmark’s SM clock, fitting on shapes outside the list with the list itself held out for validation, and the _Calibration docstring spells out the scan, fit, and check steps—rerunning from neutral constants reproduces the factory values. On the kernel side there is a real memory saving: split-K CTAs stage only their own slice’s scale and zero groups rather than the whole K, using vectorized loads when the slice is aligned; _smem_bytes and _legal_configs count shared memory per split, whereas mainline has every split’s CTA staging groups for the entire K. Another detail: _ConfigSpace.smem_bytes sits 1 KB below the optional cap—right at the cap, 17 scanned configurations failed to launch because the allocator rounds dynamic shared memory up while the barrier is static. #2826 documents how the dense selector’s constants are refit (scan, fit, check, including soft annealing and holdout validation) without code changes.

Source: https://github.com/tile-ai/TileOPs/pull/2825, https://github.com/tile-ai/TileOPs/pull/2826

3.6 TileOPs under review: Short varlen packed by KV group, L2 flush made read-only (#2819 / #2822)

#2819 updated 11-03 23:37, part of #2496, adds a packed warp-specialized program to GQAVarlenFwdOp on SM90: when a request’s KV-head grouping fills a tile bit-exactly and saves tiles, keys and values are read once per KV head; the packed kernel’s rejection logic uses a cost model weighing the call’s batch, packed query count, and group size, taking over only when the shape-guaranteed savings exceed the measured 9% single-head tile overhead; wide and narrow items are claimed alternately in the same launch, with no host readback and no extra synchronization; when the on-device count of unsplit items is below the SM count, single-tile packed requests split the key scan into segments merged through an fp32 workspace. The window side also has two measured gains: window rows save 1% to 9%, interleaved RoPE rows save 3% to 6%. It also fixes three mainline issues: rows with no visible keys returning NaN when sinks are present and score scale exceeds about 2.4 (both the single-head warp-specialized and FP8 warp-specialized programs), the fallback query rotation reading the RoPE table on rows before the request’s first key, and a rejected negative-control case leaving the caller frame in a reference cycle so one test’s CUDA graph was destroyed during another test’s capture. Test nodes grew from 173 to 197. #2822 updated 11-03 19:01, changes bench_kernel’s L2 flush from “zero_() on an int8 buffer” to “read a private int32 buffer twice the reported L2 size,” with the buffer drawn once from a fixed-seed CPU generator; on CUDA the read uses a Triton kernel (one load per element, one atomic_max per 16 KB block)—torch’s own reduction isn’t used because treating it as a flush slows the next Marlin w4a16 call by 1.3 to 1.5 microseconds. Gains on H200 at 1500 MHz, CUPTI kernel time (median of 30 calls across 3 runs): Sum on 2048×11008 f32 drops from 27.81 to 21.38 microseconds, Ge (torch-compile) on 1024×11008 f32 from 31.31 to 25.67, and 1024×4096 f32 from 6.30 to 5.47; the old flush left about 120 MB of dirty lines whose writeback landed in the next timed call, while the new flush itself writes only 71 to 445 L2 sectors (2 to 14 KB). It also fixes a mainline bug: when L2 size is unknown, the log said 256 MB while 512 MB was allocated.

Source: https://github.com/tile-ai/TileOPs/pull/2819, https://github.com/tile-ai/TileOPs/pull/2822

3.7 Adopters: FlashQLA and TileKernels have no pushes, but #3761 fixes a TileKernels fused kernel (11-04)

QwenLM/FlashQLA last pushed 10-30 03:02 and deepseek-ai/TileKernels 10-14 05:49; neither repo has commits in this window, so there is no new adopter-side code. However, this issue’s mainline headline is directly relevant to adopters: the performance regression fixed by #3761 was triggered by TileKernels’ mhc_pre_big_fuse kernel—the same kernel lost 25% of its barrier occupancy budget on B300 at hidden dimension 1280. This shows adopter kernels have become regression probes for the compiler side: during these weeks when TileKernels pushes no code, its kernels are still passively “being tested” through mainline changes. For those tracking the ecosystem, this is a reverse signal worth noting—during adopter quiet periods, compiler changes can still break on their shapes.

Source: https://github.com/QwenLM/FlashQLA, https://github.com/deepseek-ai/TileKernels, https://github.com/tile-ai/tilelang/pull/3761

IV. Community, Tutorials, and Events

4.1 Seven new issues in a single day: one production-grade CUDA graph issue, four from the same external reporter (11-03)

The main repo added 7 new issues during the window (numbered #3762 to #3775, plus one automated daily test report from the Ascend repo). The most technically substantive is #3775, reported by a team running production workloads: the environment is tilelang 0.1.12 installed from pip, Python 3.12.3, two H200s with NVLink, driver 580.178.04, a CUDA 13 container, torch 2.14.1+cu130, with sglang on the serving side running single-process TP2 (one worker thread per device, each with its own CUDA context), a GLM-family MoE model with DSA attention, and the TileLang kernels on the path being mhc_pre_big_fuse_with_norm, mhc_post, and the DSA prefill kernel. After forcing the execution backend to TILELANG_EXECUTION_BACKEND=nvrtc (the default auto resolves to tvm_ffi), running tilelang’s JIT kernels during CUDA graph capture corrupts the capture, and the corruption surfaces in downstream, unrelated kernels after tilelang has launched. Three byte-identical configurations reproduced three different first failure points: one was cudaError700 (invalid TMA descriptor handle) thrown at TMA descriptor initialization and then echoed at an unrelated launch point; one was cublasGemmEx reporting CUBLAS_STATUS_EXECUTION_FAILED (pure BF16 in and out); and one was an illegal memory access inside an elementwise kernel JIT-compiled by the host framework. All three occurred during decode CUDA graph capture (batch-size bucket 8), all about 100 seconds after service startup, and the only difference between /proc/environ and a clean control was that single environment variable. This falls into the hardest-to-localize category: random symptoms, drifting failure points, and a single environment variable as the trigger—and the reporter documented the control conditions for all three runs in full.

The other four come from the same external reporter, with a uniform environment of the main branch or the 0.1.15 pip package, Linux under WSL, Python 3.13.5, an RTX 4060 Laptop (sm_89), driver 566.26, and the CUDA 12.6 runtime, and all come with reproducible scripts and full output: #3774 is a 16×16 shared-memory permutation that applies only the row shift and drops the column offset (223 of 256 elements wrong, exactly matching the result of “dropping the column shift”), while the same code at 32×32 is correct; #3773 is a uniform while loop doing three shared-memory copies plus a GEMM that produces unstable results (a serial-loop control gives 0 errors in 100 runs, while the while version gives 21 errors in 100 runs with 10 distinct outputs); #3763 is a nested integer floor division (x // -2) // 4 that reports Let cannot bind the same var to two different values at int16, 64 elements, and 32 threads. These three plus #3775 form a concentrated failure surface for the “consumer-grade card + WSL + pip install” combination. The remaining two: #3769 is tcgen05 int8 GEMM on B300 rejected by ptxas (a fix, #3770, is already under review), and #3762 is T.simd.vcvt on Ascend 950 still saturating when sat=False, detailed in the next item.

Sources: https://github.com/tile-ai/tilelang/issues/3775, https://github.com/tile-ai/tilelang/issues/3774, https://github.com/tile-ai/tilelang/issues/3773

4.2 Ascend 950 vcvt saturation semantics and the mlir-ascend daily test (#3762 / #1980)

#3762 was filed by an Ascend-side user, with an environment of Ascend950PR (NpuArch 3510 / dav-3510), CANN 9.2.0-beta.2, driver 25.7.rc1, torch_npu 2.12.0, and TileLang installed from source (e84ce3de from 10-21). The problem is that T.simd.vcvt(src, "uint8", sat=False) silently ignores sat=False for int32 vectors and always saturates: negative values become 0 and values above 255 become 255, with no error and no warning, while the caller believes they are getting wraparound semantics. The docstring is explicit about this case—out-of-range values wrap when sat=False—with expected values of -1 → 255, -128 → 128, and 300 → 44, but the measured values are -1 → 0, -128 → 0, and 300 → 255. The impact lands on f32-to-int8 paths like the tail end of int8 quantization: once a negative intermediate value appears, the semantics shift from wraparound to clamping to zero. #1980 is the Ascend repo’s automated daily test; the 11-03 09:35 round passed all 3030 items for a 100% pass rate, workflow run number 41553942797. The two numbers are interesting side by side: the Ascend repo’s daily test is fully green, yet semantic discrepancies the daily test cannot cover (such as the actual behavior of the sat parameter) still depend on manual reports.

Sources: https://github.com/tile-ai/tilelang/issues/3762, https://github.com/tile-ai/tilelang-ascend/issues/1980

4.3 Three doc-site updates, VS Code extension renamed vsix (11-03)

The docs site tilelang.github.io had three Update docs commits during the window (11-03 13:11, 13:28, and 11-04 00:59 Beijing time), routine updates from the automatic build that follows mainline merges, corresponding to several merges in this window—the #line hierarchical mapping, the ThreadSync scope list, and the corresponding doc locations for Ascend bulk copy. TileOPs.github.io had no push this window; the latest is still 11-01’s “follow TileOPs mainline: overlap and alignment guarantees, build instructions.” The VS Code extension repo had a commit at 11-04 02:19 titled “update docs and rename vsix,” following 11-02’s “align CUDA debugging to TileLang line mapping”: the line-level anchoring merged into mainline last issue changed the source of line numbers in generated code, so the extension’s debug configuration had to follow suit, and this commit adjusts the artifact name and docs together.

Sources: https://github.com/tile-ai/tilelang.github.io, https://github.com/tile-ai/tilelang-vscode-edition

4.4 News side: no in-window hits across three channels, foreign media inconsistent on Ascend training reports (11-03)

The Google News combined Chinese/English query, Hacker News, and arXiv channels all had no substantive TileLang-related hits this window: of 15 candidates, 12 came from GitHub repo discovery, 2 were unrelated Show HN posts (a detective game and a hardware-design prompt tool), and 1 was a Wccftech report on DeepSeek and Ascend. The latest all:tilelang hit on arXiv is still TileSight from 2026-07-24 (a tile-centric GPU performance analysis model), and the joint query with TileLang found no new papers in the window. The only industry-side item worth noting has conflicting accounts: Wccftech says DeepSeek will train V5 from scratch on an Atlas 950 SuperPoD built from Ascend 950DT, while a contemporaneous The Information report says “pre-training on 950DT starting as early as end of 2026,” and earlier Bloomberg and Financial Times reports said the 160,000-unit 950DT order is mainly for inference, after previous training efforts on Ascend faltered and fell back to Nvidia. The three accounts disagree on whether and when it will be used for training, and none has been officially confirmed by Huawei or DeepSeek. It is noted here only as background: whether or not the training materializes, the compiler and operator investment on the Ascend side (this issue’s #3689, #1975, #1981) is laying the groundwork for that path.

Sources: https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods, https://arxiv.org/

V. Trend Observations

First, performance regressions are beginning to be treated as first-class citizens, and the methodology of the fix is more noteworthy than the conclusion. The ID retention rule introduced by #3643 to fix correctness is itself necessary and correct; #3761 did not revert it, but instead allowed retention and reuse to coexist at different levels—before sharing an ID across scopes, first prove that the participating predicate always holds; if it cannot be proven, fall back to scoped counting. Together with the author’s isolation experiment that “only changes the declared barrier resource count,” this fix chain separates two things—”correctness” and “resource declaration”: the latter is a variable that can be reproduced and validated independently. This ability to separate concerns is precisely the mark of a compiler project moving from “it runs” to “it is maintainable.”

Second, Ascend’s engineering is advancing on two legs simultaneously. One is architecture: #1975 moves the entire TileLang-Ascend into the mainline multi-backend framework with 2,887 files and 600,000 added lines, using identifiers such as arch="dav-2201" for backend dispatch, with the goal of letting A2A3 and A5 coexist orthogonally in the same repository; the other is pipeline: mlir-ascend’s #208 writes Codex’s three workflows—initialization, onboarding, and recovery—along with task handoff into .workspace/, and the main repo and the Ascend repo each have their own Agent directory. Looking at “migrating the repo” and “using an agent to migrate the repo” on the same day reveals that the bottleneck on the Ascend side has shifted from “missing operators” to “missing a consistent engineering skeleton and a reusable migration process.”

Third, the culture of self-justification in operator libraries is deepening. Nearly all of the TileOPs PRs under review this issue write the “evidence” into the PR: the EP8 group size of Qwen3-30B-A3B measured on a real route (taking layer 44, the median-skew layer among the 48 layers), the TIR of 35 mask constructions matching the mainline byte for byte, the L2 flush change accompanied by a comparison table of the 46-line implementation and an explanation that the old flush left 120 MB of dirty lines, and the W4A16 cost model explicitly stating “which terms were removed because the data cannot pin them down.” TileFoundry takes another path—turning synchronization into a provable analysis, then having lowering verify it in reverse, and directly listing the cases the analysis rejects (TMA and wgmma completion semantics not yet declared, loops with unequal cross-thread boundaries). Together, the two paths point to the same orientation: better to admit “not proven” than to treat “it runs” as “it is correct.”

Fourth, the risk surface is concentrating on installation form. Of the 7 new issues in one day, four came from the same external reporter using the pip package, running WSL, and using a consumer-grade sm_89 card, all with reproducible scripts attached; #3775 is a production combination of “pip package + multiple CUDA contexts + nvrtc backend + CUDA graph capture.” These are not deep compiler defects, yet they determine the experience of a new user’s first attempt. Contrasting with the Ascend repo’s daily test of 3,030 items all green while the semantic deviation of sat=False still relies on manual reports, it is clear that there is still a gap between the shape of test coverage and the shape of real usage.

One final observation: the news side remains blank this issue, with 12 of the 15 candidates being repository discoveries. TileLang’s external visibility is still tied to events such as “DeepSeek and Ascend,” while day-to-day progress happens only in commits. For those tracking this project, this means the main body of the daily report can only be PRs and issues, not coverage—this has been the stable pattern for several consecutive issues.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
tile-ai organization (29 repos) 9 repos pushed within the window
Main repo tilelang 5 merges to main branch; 6 new issues opened; 23 updates under review; 3 automated doc-site builds
TileOPs 1 merge to main branch; 7 updates under review
TileFoundry 1 merge to main branch; 2 updates under review
tilelang-ascend No new commits on default branch ascendc_pto within the window; 15 updates under review; all 3030 daily tests passed
tilelang-mlir-ascend 1 merge to harness branch (non-mainline)
tilelang-hygon No new commits on main branch; update on #19 under review
tilelang-vscode-edition 1 merge (docs and vsix rename)
tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / tilelang-puzzles No pushes within the window
Release pages (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) No new versions within the window; tilelang latest remains v0.1.15 from 09-30
FlashQLA / TileKernels No pushes within the window
Google News (six queries in Chinese and English) / Hacker News / arXiv No substantive hits within the window

Full Source List