Reporting period: Monday, November 2, 2026 to Friday, November 6, 2026, 5 working days in total Sources: 5 issues of this publication’s daily TileLang community updates (11-02, 11-03, 11-04, 11-05, 11-06); the collection windows of the five issues are contiguous, together covering 168 hours from 10-30 07:00 to 11-06 07:00 (the Monday issue used a 72-hour window covering the previous weekend; Tuesday through Friday used 24-hour windows), with no overlap with the previous weekly report (10-26 to 10-30); additionally verified via a supplementary 168-hour window check of the repositories in the tile-ai organization, each repository’s releases page and tags, and a four-source inspection review of TileLang×RISC-V (see appendix)

I. Weekly Highlights

  • Ascend 950’s persistent FlashMLA kernel lands directly in the mainline, written in a compiler-schedulable style rather than ported from a separate repo (merged 11-02, #3678): A single change adding 1,997 lines, deleting 17 lines, across 20 files and 42 commits, covering both the BF16 prefill and compact FP8 decode paths, with support for optional FP8 and FP4 extra KV cache, attention sink, invalid indices, and per-query independent lengths. The attention loop, online softmax, dequantization, and scheduling are all expressed in TileLang, relying on T.Stage to let the compiler manage buffer versions and synchronization; QK and PV share two different layout views of a single L0B storage ring, L1 placement is specified directly via T.annotate_buffer_addresses; prefill splits four disjoint output slices within a single accumulator allocation, draining each slice as soon as its normalization completes. To make this style viable, the compiler side gained three capabilities in the same period: L0C-to-UB copies preserve NZ output, each AIV stream is selected before its own pipeline, and explicitly scheduled stages preserve source order when their physical issue phases coincide. The target shapes are written very narrowly: 64 heads, Dqk and Dv both 512, topk a multiple of 64. Source: https://github.com/tile-ai/tilelang/pull/3678

  • ThreadSync changed three times in one week—first making way for the migration, then restoring performance and proofs (11-01, 11-04, 11-05): #3733 on 11-01 fully reverted the precise thread counting (#3291) merged three weeks earlier, with a blunt rationale—that change moved expensive count queries ahead of full-block barrier checks, so even when all threads reach the barrier the proof still runs; once TVM’s solver budget is exhausted it returns partial counts, and legitimate remap and sequence helper kernels fail to compile outright with “Cannot prove a thread rectangle”; three correctness improvements (quotient guards, exact participating-thread bounds, non-rectangular thread sets) were withdrawn together, with #3732 left as the tracking issue for re-optimization after the migration. #3761 on 11-04 addressed a side effect of the revert: to fix the duplicate named-barrier IDs in #3643, the same group of threads cannot reuse an ID across scopes, so TileKernels’ mhc_pre_big_fuse kernel with 96 threads per block gained an extra ID on its two-stage cp.async pipeline, and the compiled artifact’s declared barrier resources went from 4 to 5, dropping the barrier occupancy ceiling on B300 at hidden dimension 1280 from 16 blocks/SM to 12 blocks/SM; the fix lets ThreadSync accept an ordered scope list and, before sharing an ID across scopes, uses a new analyzer to prove the participation predicate always holds—acceptance criteria being resource declarations back to 4 and occupancy back to 16 blocks/SM, without reverting the #3643 fix. #3787, under review on 11-05, continues to fill prover blind spots: at 192 threads, forms like tx < 160 and tx // 32 % 2 == 0—a boundary term and a warp-index modulo term joined by and—are misreported, the root cause being that the FloorDiv and FloorMod encodings cannot establish the relationship between the two divisions under the “upper bound plus modulo” combination; the fix rewrites threadIdx.x as warp_index * warp_size + lane. Sources: https://github.com/tile-ai/tilelang/pull/3733, https://github.com/tile-ai/tilelang/pull/3761, https://github.com/tile-ai/tilelang/pull/3787

  • The TIRx migration went from “verbal rationale” to three pieces of hard evidence in one week, with the foundation swap entering public preparation (11-03 to 11-06): #3736 on 11-03 deleted tilelang/language/ast/, tilelang/language/parser/, and tilelang/language/tir/entry.py together—leftovers from an early local TVMScript parser fork—adding 1 line, deleting 3,384 lines across 10 files, with the language layer now using the upstream tvm.tirx.script.parser and AST and statement override registration in tilelang.language.overrides. #3752, merged 11-06, together with its prerequisite tile-ai/tvm#91, turns Z3’s rlimit from a hardcoded 100,000 into something configurable per kernel; the author’s words: “this indirectly fixes a lot of Z3-prover-related issues.” #3803 on 11-06 submitted a complete migration draft: replacing the vendored TVM with a standalone Apache TIRx, migrating the frontend, IR transformations, code generation, JIT adaptation, and packaging all to the native object model, pinning apache/tvm-tirx at 7e49a148 and apache/tvm-ffi at 3e906bc6, keeping the dependency tree cleanly upstream with no injected patches. All three lines point to the same thing: the top priority for the mainline at this stage is swapping the foundation, and all other optimizations must make way for it. Sources: https://github.com/tile-ai/tilelang/pull/3736, https://github.com/tile-ai/tilelang/pull/3752, https://github.com/tile-ai/tilelang/pull/3803

  • “Silently generating wrong code” became the week’s dominant defect type, spanning four release versions (11-03 to 11-05): The most serious is #3808—after condition-related scalarization, the shared-memory epilogue is missing a __syncthreads(), so versions 0.1.12, 0.1.13, 0.1.14, and 0.1.15 all silently generate incorrect CUDA with no compile error whatsoever; the reporter provided a complete runnable reproduction script. The same family includes #3800 (with warp specialization disabled, T.Pipelined reloads the next stage before the tcgen05 barrier wait, so the MMA reads a rewritten operand), #3721 (on sm_100a, T.gemm on TMEM with threads/32 not dividing N produces not a single tcgen05.mma in the PTX, and the product is silently dropped), #3775 (in production TP2 sglang with the nvrtc backend, CUDA graph capture is corrupted, with three byte-identical configurations reproducing three different first failure points), and #1991 and #1992 (on Ascend 910_9362 with threads=2, AscendVidReduction unconditionally halves the injected tmp workspace too, and sin/cos on an indirectly addressed tile computes 3978/4096 elements wrong, i.e., 97.1%). This family reports no errors and passes routine tests; it can only be caught by targeted reproduction. Sources: https://github.com/tile-ai/tilelang/issues/3808, https://github.com/tile-ai/tilelang/issues/3800, https://github.com/tile-ai/tilelang-ascend/issues/1991

  • TileFoundry pushes synchronization from “the analysis can prove it” to “emission must match the proof” (11-02 to 11-06): #283 builds “how shared-memory accesses should be ordered” into a provable analysis layer—walking each function in lowering’s emission order, planning commits, waits, and scope-local fences before conflicting accesses; of 83 scheduling fixtures, 52 are planned and 31 rejected (24 use TMA copies, 6 use wgmma, whose completion semantics are not yet declared); because by default no caller requests it, the 75 fixtures that finalize are byte-identical to main. #285, merged 11-04, is its execution side: CUDA lowering emits according to the synchronization plan and verifies each item in reverse, throwing LoweringError on mismatch; on a single sm_89, 6 of 7 buildable programs match the reference exactly with synccheck and memcheck fully clean, while those same 7 lowered the main way produce 4 wrong results, the direct symptom being a maximum deviation of 659 in sm_89 E4M3 GEMM. #286, merged 11-06, adds WGMMA’s asynchronous completion semantics, and #284, merged 11-05, builds FSDP and HSDP training semantics into the HIR evaluator and aligns with real 2×2 PyTorch HSDP on four-card sm90. The author kept an explicit caveat on #286: the matrix instruction emitter is still not wired up, and this change neither proves execution correctness nor delivers performance. Sources: https://github.com/tile-ai/TileFoundry/pull/283, https://github.com/tile-ai/TileFoundry/pull/285, https://github.com/tile-ai/TileFoundry/pull/286

  • TileOPs merged 15, 7, 1, 7, and 9 changes across five windows, with the competitive focus shifting entirely from “operator count” to “selectors and preparation cost” (all week): Of the first window’s 15 changes, 13 landed on a single day, 11-01, nearly every one carrying an H200 measured table and a “manifest falls back when no measured row exists” declaration; the following 11-04 and 11-05 windows had 7 and 9 changes respectively, pushing the pace back to its densest since the holiday. Three hard numbers summarize the week’s direction: GQA dense forward uses a single measured rule to select the decode or prefill program, cutting selection regret relative to the fastest candidate from 1.818 to 1.017 (#2808); warp-specialized decode overlaps softmax with QK and PV, giving a geometric-mean latency of 0.952× main across the 478 affected sweep calls (#2817); in-place reads of broadcast operands cut broadcast-row time for five operator classes—lerp, where, clamp, masked_fill—to 0.13–0.44× the mainline (#2832). Also in the period: a roofline pricing rewrite (three physically impossible bandwidth rows corrected from 7.48, 7.50, and 5.30 TB/s to 0.62, 0.62, and 2.65 TB/s, #2830) and separate fitting of three-stage tile factors (selection regret down from 0.89% to 0.06%, #2836). Sources: https://github.com/tile-ai/TileOPs/pull/2808, https://github.com/tile-ai/TileOPs/pull/2832, https://github.com/tile-ai/TileOPs/pull/2830

  • Domestic backends move from “in-house branches” to “public review,” with Hygon opening its first formal PRs (11-05): The Hygon repo’s main branch was last touched at #32 on 11-02 (treating gfx948 as gfx946-compatible); this week’s activity first accumulated four changes on the feat/hcu-scaled-cast branch (E8M0-scaled FP8/BF8/FP4 conversion, 64-bank LDS read phases, a proof of at most two paths for mixed LDS wrap, and mixed checks by physical read width), then landed three more LDS phase fixes on 11-05 and opened #33 (scaled cast) and #34 (XOR shuffle for reduction using DPP and lane swap) at 09:08 and 19:52 respectively—the first domestic backend to complete the path from in-house branch to public review. mlir-ascend in the same period moved the BishengIR toolchain adapter into contrib (#209) and used three small PRs—#203, #211, #212—to move tl.sub/mul/div/max/min, tl.shl, and tl.shr one by one into the TLIR dialect. Moore Threads landed a README documentation commit on 11-05, ending complete silence since 09-30. Sources: https://github.com/tile-ai/tilelang-hygon/pull/33, https://github.com/tile-ai/tilelang-hygon/pull/34, https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

  • Hardware support expands at both ends: new Blackwell SKUs and domestic accelerators (all week): At one end, new cards. B300’s reported sm_103a trips on two paths—tcgen05.mma.kind::i8 exists only in sm_100a and sm_110a, so after lowering the kernel ptxas reports Feature '.kind::i8' not supported (#3770); CuTeDSL’s TMA helper maintains an architecture whitelist, and B300 is rejected before it even reaches the DSL (#3772; after adding sm_103a and gr100’s sm_107a, B300 shows 8 passed and 1 skipped, gr100 13 passed and 1 skipped). On SM120, FP8 GEMM with fp32 accumulation had been using SM89’s mma.sync atomics; switching to block-scaled MMA with ue8m0 unit scaling raises 8192-cubed on RTX 5090 from 463 to 727 TFLOPS (under review, #3790). At the other end, domestic cards: Hygon does microarchitecture-level tuning on 64-bank LDS phases, while Ascend makes the Cube-to-Vector core ratio configurable instead of fixed at 1:2 (#3783) and tightens AutoSchedule’s cycle lower bound to cut Z3 solve time by 10.04% (#3776). Sources: https://github.com/tile-ai/tilelang/pull/3772, https://github.com/tile-ai/tilelang/pull/3790, https://github.com/tile-ai/tilelang/pull/3783

II. Core Project Progress (main repo tilelang)

Versions and Releases

  • The main repo’s latest tag and release remain v0.1.15 (released 09-30); no new tags and no new release-page entries all week, now 37 days old as of this Friday.
  • Both new versions in the org are stuck at last week: TileOPs at v0.1.0 (10-27) and TileFoundry at 0.0.4rc1 (10-25); neither was promoted to stable this week.
  • Adaptation-line tags were untouched all week: tilelang-ascend v0.1.2.000 (09-09), tilelang-metax v0.1.15 (10-16), tilelang-musa v0.1.15+musa.1 (09-30), and tilelang-mlir-ascend stuck in the v0.1.2.020 series.
  • Merge volume and cadence: Monday’s window covered 72 hours (10-30 07:00 to 11-02 07:00) with 35 merges; the four 24-hour windows from Tuesday to Friday had 18, 8, 12, and 18 merges respectively, totaling 91 across the five windows. Open updates in review were 82, 44, 48, 42, and 57, totaling 273; newly opened issues were 11, 8, 7, 4, and 5, totaling 35. Main-branch commit distribution: tilelang 7, 5, 5, 3, 5; TileOPs 15, 7, 1, 7, 9; TileFoundry 3, 1, 1, 1, 1; tilelang-ascend 5, 4, 0, 1, 2; tilelang-mlir-ascend 1, 0, 1, 2, 0; tilelang-hygon main branch 2, 1, 0, 0, 0 (plus feature-branch commits); tilelang-lsp and the VS Code extension 1 each; tilelang-musa 1.
  • The cadence this week followed a “clear the books early, bottom out midweek, recover by the weekend” pattern: the 11-04 window had only 8 merges, the week’s low point, and the 11-06 window returned to 18; the review queue fell from 82 at the start of the week all the way to 42 before rebounding to 57 on Friday, indicating that the backlog of reviews from early in the week was largely worked off.

Compiler, Language Layer, and Execution Path

  • Ascend 950 FlashMLA kernel and its supporting compiler capabilities (#3678, merged 11-02, 20 files +1997/-17): see the first item in Top Stories. Notably, this change added the three low-level capabilities required for “explicit scheduling” to the mainline at once rather than patching them in the Ascend branch—L0C-to-UB copies preserving NZ output, each AIV stream being selected before the pipeline, and source order being preserved when physical issue phases coincide.
  • Two consecutive fixes for the vector-width regression from cast decoupling (#3734 merged 11-02, #3759 merged 11-05): #3734 fixed Cannot convert type float32x16 to CUDA type in ViT forward on SM100—the vectorization planner chose 16 lanes for 256-bit BF16 storage and deferred cast constraints, and after #3579 made decoupling skip the memory-read guard, final legalization still followed the deferred strategy, stretching the residual FP32 expression into a 512-bit vector; the fix makes defer_cast_constraints a C++ parameter of the planner rather than an IR attribute, so the guard’s FP32-to-BF16 cast drops to 8 lanes while successfully decoupled BF16 copies keep 16. #3759 cleaned up its tail: an int64 cast in a condition previously squeezed the data vector width from 16 lanes to 4 even when the condition was fully uniform across 16 lanes; the fix reuses condition-uniformity analysis for the cast-width exemption, visiting if statements, tirx.if_then_else, and Select conditions with scalar predicates separately, while casts inside clauses remain restricted.
  • The three-stage evolution of ThreadSync (#3733, #3761, #3787): see the second item in Top Stories. #3733 was explicitly a stopgap, stating “making way for the TIRX migration,” and the accompanying TVM fork optimization proposal tile-ai/tvm#89 has been closed. The most noteworthy part of #3761’s 401-line change is the attribution experiment: the author reproduced the same slowdown by changing only the number of barrier resources declared in the compiled artifact, without touching any executable instruction, turning “performance regression” from an impression into a single identifiable variable.
  • Subtraction in the language layer and frontend: #3736 deleted 3384 lines of the legacy TVMScript parser fork (see the third item in Top Stories), with the risk of keeping it stated concretely—if anyone directly imports tilelang.language.parser.operation, it registers conflicting operators on tirx.PrimExpr and tirx.IterVar; #3760 introduced a bounded in-process transformation cache for nested T.macro, storing only context-independent transformation results and rebinding macro instances to the current globals and closure values, solving the problem where “the kernel hits the cache but the macro is still re-parsed and recompiled”; #3802 added T.annotate_host_mapped_buffers, writing the indices of annotated parameters into the PrimFunc’s tl.host_mapped_params attribute, with the runtime validating registration status via cuMemHostGetFlags and resolving the device-visible address via cuMemHostGetDevicePointer, while ordinary pageable memory is still rejected.
  • Debugging and observability: #3703 changed #line directives from “anchoring only the first line of each TIR statement” to tiered emission (0 off, 1 statement-level, 2 and above re-anchoring line by line), implemented in the C-family output stream shared by CUDA, HIP, CPU C, C host, Metal, and Ascend, with TILELANG_EMIT_LINE_DIRECTIVES added as the Python-side default; the author notes that the CUDA and HIP test cases were skipped locally and cuda-gdb’s actual behavior was not directly verified. The LSP repo followed up the same day, aligning CUDA debugging to the new line mapping; the VS Code extension repo followed up again on 11-04, updating docs and renaming the artifact to vsix.
  • Autotuning semantics: #3546 changed the supply closure of set_autotune_inputs from “parameters never passed at the call site” to get_current_device(); previously, on machines without CUDA, the fallback branch moved every captured tensor to CPU, and every tuning round failed due to a launcher device mismatch. On dual-card Ascend950, three Ascend autotuning regression tests all failed before the fix and all passed after. The related #3784 adds a device sync before and after the cross-device copy probe in AscendC CI, bringing asynchronous errors back within the scope of try ... except RuntimeError—previously they surfaced a step late, appearing as an unrelated torch.arange() reporting a 4-byte memcpy size mismatch.
  • Z3 solving resources go from hardcoded to configurable (#3752, merged 11-06, depends on tile-ai/tvm#91): rlimit still defaults to 100,000, and behavior is unchanged when unconfigured. This is the other end of the same direction as the concurrent Ascend-side tightening of the AutoSchedule period lower bound (#3776, Z3 solve time down 10.04%): one changes the search budget, the other changes the search starting point.

Correctness and Bug Fixes

  • ThreadSync and barriers: beyond the full revert in #3733 (see Top Stories) and the cross-domain reuse in #3761 (see Top Stories), #3765 in review explicitly accepts “partial barriers with a single thread-index bound guard” when the thread dimension is a whole number of warps; #3767 replaces the internal StructuralEqual assertion for multi-index fragment access with a ValueError carrying both index expressions; #3786 turns out-of-bounds T.copy destination regions into a compile-time error; #3788 fixes the fast path of ProveFragmentContains—it returned true directly when a large fragment was fully copied, skipping the thread-ownership check, so “small fragment owns threads 0 to 127, large fragment owns only 0 to 95” was misjudged as contained.
  • The silent-wrong-code family: beyond #3808, #3800, #3721, #3775, and #1991/#1992 (see the fourth item in Top Stories), this week also brought #3774 (a 16×16 shared-memory permutation performed only the row shift and dropped the column offset, getting 223 of 256 elements wrong, while the same code at 32×32 was correct), #3773 (three shared-memory copies plus a GEMM inside a uniform while loop: the serial control was 0 errors in 100 runs, while the while version failed 21 of 100 runs with 10 distinct outputs), and #3763 (nested integer floor division (x // -2) // 4 reporting Let cannot bind the same var to two different values at int16, 64 elements, 32 threads). These three, together with #3775, form a concentrated bug-report surface for the “consumer GPU + WSL + pip install” combination.
  • Root cause in dependency analysis and sync insertion: #3749 is the week’s most technically substantial bug report—two consecutive SIMD loads lost their dependency on an earlier GM-to-UB copy, with the root cause being that the logical region reconstructed by MemoryAccessDetector::LogicalPointerRegion missed part of the physical load range: in the minimal repro, work.shape is (2, 128), the generated read descriptor covers physical [64, 192), but the reconstructed region is [0, 128), so dependency analysis considered the input DMA unrelated to the load, AutoSchedule moved the input copy after the load, and InsertSync also emitted no MTE2-to-V set/wait. In the same window the main repo also gained 8 issues pointing at code generation and sync insertion in the Ascend AscendC and PTO backends (#3740 through #3749), 7 of them filed by the same reporter, all with reproducible prim_funcs and line-by-line generated-code comparisons.
  • Internal errors as soon as a boundary condition breaks: #3738 fixes “T.copy from a global region smaller than the padded shared tile (e.g. K=100 copied into BK=112) being judged a full-tile TMA copy and hard-failing”; #3766 fixes TMA byte counting for cta_group::2—the bytes of a CTA pair both land on the even CTA’s barrier, while 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 wrong for (4,1,1) and (2,2,1).
  • Operator and example boundaries: #3709 makes the CPU scalar GEMM clean only the destination region; #3712 makes LegalizeNegativeIndex skip unreachable branches; #3714 adds round-to-nearest-even for fp32-to-tf32 conversion (cvt.rn.tf32.f32; non-saturating requires SM90+, saturating requires SM100+); #3702 makes several attention examples derive tiles from shared_memory_per_block_optin, taking the CUDA 9 items on RTX PRO 6000 (sm_120) from “2 pass, 7 fail” to 9 of 9 passing, while H100 has a 227 KiB optional limit and still uses the original tile.
  • Two semantic deviations on the Ascend side: #3762 reports that T.simd.vcvt(src, "uint8", sat=False) silently ignores sat=False for int32 vectors and always saturates; the documented wraparound semantics (-1 becomes 255, -128 becomes 128, 300 becomes 44) were measured as truncation (-1 becomes 0, -128 becomes 0, 300 becomes 255), affecting paths like the f32-to-int8 tail of int8 quantization; #3789 reports that the PTO backend emits signed pto.si32/si64 for integer T.atomic_add while the PTO atomic interface accepts only unsigned carriers, with fixes in two directions, #3792 and #3794, filed the same day.

Runtime, CI, and Toolchain

  • Both classes of CI stability issues were addressed: #3784 fixes the AscendC CI autotuning supply test case (see above); #2833 bumps TileOPs’ RUNNER_VERSION from 2.336.0 to 2.338.2 (after GitHub deprecated the old runners, all gpu-smoke jobs queued, and entrypoint.sh with --disableupdate could not self-upgrade); #2827 moves TileOPs’ off-CUDA suite to a nightly job (previously these two steps took 13 to 18 minutes on ubuntu-latest in preflight; after the move, preflight keeps only the collection check, the discard-test-report check, and the cuda_only must-fail check, about 3 minutes).
  • Concrete evidence that parallel merges lack a post-merge gate: #2824 merged at 21:06, changing the grouped GEMM roofline row to require declaring sizes and deriving G from it; #2830 merged at 21:08 and still wrote only G without sizes; the two were about 2.5 minutes apart and each had green CI, but right after merging, main’s GPU Smoke reported RowError: row misses ['sizes'], fixed the same day with #2837. This is the week’s most noteworthy process record: a single PR’s CI passing cannot cover “two PRs whose semantics intersect.”
  • Docs and localization: tilelang.github.io had 3, 5, 3, and 3 Update docs automated builds across the week’s four collection windows, corresponding to interface changes such as #line tiered mapping, the ThreadSync scope list, host-mapped annotations, TF32 rounding, and Z3 configuration; TileOPs.github.io had one push, and TileOPs-nightly had only branch pushes with no new commits on the main branch.
  • Toolchain details on the Ascend side: #1956 makes AI Core exception dumps compatible with CANN 9.2.0 (adding required fields to acldumpTensorInfo, skipping null addresses and zero-size tensors, and replacing the deprecated aclrtSetExceptionInfoCallback registration interface with aclrtExceptionInfoCallbackRegister, lowering the version threshold from 9.3 to 9.2); #1932 aligns the A2/A3 performance-tuning skill with the Ascend950 workflow structure and adds two skills: chip-type identification and chip-information retrieval.

Review Queue Observations

  • The largest item remains the full Ascend repo migration (#1975): at its 11-04 00:34 update it reached 607317 lines added, 97595 deleted, and 2887 files, moving TileLang-Ascend into the mainline multi-backend architecture, with the existing Ascend components moved under ascend/a2a3 and split from the mainline A5 path via arch="dav-2201". The author is candid about the TODOs in the description: differences between files in the three directories and the originals still need to be converged, static-check and formatting CI still need to be restored, and docs and Agent-related work still need to be migrated.
  • The TIRx migration draft (#3803) is the week’s most consequential item in review: it determines the cadence at which subsequent TVM-side fixes enter the mainline, and whether the glue layers for Ascend, Hygon, and mlir-ascend need to change along with it.
  • TileOPs’ review queue is concentrated on selectors and layouts: #2836 fits a separate factor (1.36) for coop2_swap’s three-stage tiles; #2838 fixes “treating K as a contiguous dimension and thereby rejecting the legal TN-layout K=1 call”; #2839 makes FP8 grouped GEMM load A’s scales via TMA and switches non-fused tight groups with K ≤ 1536 to 64×256 ping-pong tiles; #2840 adds SM89 prefill sinks and incorporates the presence or absence of sinks into the compiled-instance identity.
  • Trend in open-review counts: 82, 44, 48, 42, 57. The early-week high came from last week’s backlog; after dropping to 42 mid-cycle it rebounded to 57 in the final window, in sync with the rebound in merge volume—a normal clear-the-books/open-new-tickets rhythm.

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

Ascend: From “Adaptation Target” to “First-Release Target”

  • On the first-release side: The FlashMLA kernel (#3678) targets Ascend 950 and is written directly into the mainline, using low-level capabilities such as T.Stage and T.annotate_buffer_addresses that are intended for explicit scheduling, rather than opening a separate repo for porting; the target shape has 64 heads, Dqk and Dv of 512, and topk as a multiple of 64. The accompanying bulk-copy capability was completed on 11-03: #3689 lets a [B, M, N] region be written as a single T.copy, with lowering generating bulk ND2NZ descriptors, preserving the source’s batch and row strides, routing transposed batches through DN2NZ, and skipping source batches that exceed the GM boundary; validation was on Ascend950DT with CANN 9.2, with 50 AscendC example tests and 1733 Ascend unit tests passing, of which 28 BF16/FP32 device cases cover batch and M/N tail blocks, whole-batch out-of-bounds, non-contiguous source strides, destination storage larger than source, and rectangular transpose.
  • Synchronization correctness: #1981 fixes the synchronization annotation of copy_pipe_to_l1 — A2/A3’s PTO cross-core V-to-C communication is relayed through GM, and the three-stage copy belongs to AIV’s MTE3, AIC’s MTE2, and AIC’s MTE1 respectively, yet it was labeled PIPE_MTE3 in operation_config.h, so automatic synchronization waited on the AIC core’s own MTE3, whose queue is empty and whose event is set immediately, allowing the MTE1 read of L1 to be released before the MTE2 write completes; there are two pieces of reproduction evidence: after changing PIPE_MTE3 to PIPE_MTE2 in the generated source, 10/10 runs recovered exactness, while locally on the latest mainline 10/10 runs were wrong, with up to 16384/16384 elements mismatched in a single run. The fix splits the name rather than replacing it outright: copy_pipe_to_l1 is relabeled PIPE_MTE2, and a new copy_pipe_to_l1_bind is labeled PIPE_MTE1 (the A5 direct-connect path), avoiding fixing A2/A3 while breaking A5.
  • Converging on the mainline: The RFC in #1971 lays out the replacement route for the synchronization implementation — port the mainline synchronization core and the standalone Z3 proof service, hook into the existing compilation pipeline through the old TIR adaptation layer, retain TVM, scheduling, memory planning, AscendC/PTO codegen, and the runtime ABI, and after acceptance switch the default implementation and delete the old algorithm; the accompanying #1965 migrates the mainline’s z3-based insert sync to the Ascend repo; #1975 is the migration of the repo-level skeleton. The three are three layers of the same path. In the same period, #3752 makes Z3’s rlimit configurable per kernel, and #3735 proposes proving the disjointness of modular indices within thread sync — precisely addressing the proof-budget problem left by last week’s rollback from both ends: “configuring solver resources” and “reducing the shapes that need proving.”
  • The host-overhead thread: #1948 changes the TVM Analyzer() that is unconditionally constructed on every call in CythonKernelWrapper.forward() to lazy construction with on-demand reuse; on 910B3, vec add single-block drops from 442–467 microseconds to 105.1 microseconds, 32 blocks from 441–493 to 108.0, and kpool compression from 447–450 to 111.3, while a bare ctypes call stays roughly unchanged at about 12.5 microseconds; the follow-up #1967 caches symbol-shape parsing results in the cython wrapper to further cut costs on the symbol path, and the same family includes #3757 under review (keying by Var object rather than printed name, to avoid treating different variables with the same name as the same symbol).
  • Operator supply: #1827 adds WY Fast forward, and #1828 migrates and completes the segmented backward of wy_fast_bwd_split (the two merged within less than a minute of each other around 11-02 02:33), completing this operator chain for linear attention on Ascend to the trainable stage; #1866 migrates Chunk O forward (chunk gated linear attention), using developer mode plus four pass configurations, a single @tilelang.jit kernel plus a grid-stride loop fixed at 20 cores, 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, with the wrapper containing no T.* code at all), with all 20 cann-bench cases passing accuracy, an average speedup of 0.97x against the torch-npu baseline, and a composite score of 73.13 (local 910B2 basis) — 0.97x means there is no net gain yet, but the skeleton is there; #1966, the GLM-5.3 kpool fp8_mqa_logits example under review, is a dual-kernel design with AIV/Cube division of labor, 32 heads, dimension 128, and page size 64.
  • Engineering rigor: #1954 turns the compile-time segfault of T.tile.bitwise_rshift on int64/uint64 into a ValueError carrying the actual dtype and the supported list (only int16/uint16/int32/uint32 allowed), explicitly stating “no new 64-bit right-shift support”; #3783 makes the Cube-to-Vector core ratio explicitly configurable or derived from the device (vector_core_num / cube_core_num, falling back to 2 when unavailable), and carries the selected group count all the way through the sid domain, task assignment, synchronization, and code generation; #3776 adds three necessary conditions to AutoSchedule’s initiation interval — a single-task bound, a resource-group bound, and a self-dependency bound — initializing the binary search with the maximum of the three; on Ascend 950PR with the CANN 9.3.0 weekly build, all examples compile with unchanged scheduling results, and cumulative time in the Z3 solving phase drops 10.04%.
  • Regression status: The Monday-window daily test passed all 2970 items, with one workflow failure on 11-01 (#1976); the Tuesday and Thursday daily test records show all 3030 items passing, a 100% pass rate (#1980, #1988). The test-case count grew from 2970 to 3030 within a week, showing coverage is still expanding, but #1991/#1992 and #3762 — two issues found only through manual reports — show that “pass rate” cannot substitute for correctness.

Hygon: From an In-House Branch into Public Review

  • Main branch: #31 merged on 10-30, improving dynamic MLS boundary inference and MatrixStore stride handling, fixing the sinking of clear_accum in MMAC and blockscaled GEMM, and switching the persistent WASP path to an M-partition consumer with independent A/B fences while retaining slot and phase rotation; #32 merged on 11-02 10:46, treating gfx948 as gfx946-compatible and reusing the existing MLS atoms, MMAC, barrier, and compile options, with the goal of getting the existing builtins running on the new model before integrating gather load. This is a typical two-step approach: first hook the new architecture number into the existing path to get a testable baseline, then add the parts that genuinely need new instructions.
  • Feature branch: feat/hcu-scaled-cast landed four commits within eight hours on 11-04: E8M0-scaled FP8/BF8/FP4 conversion (scalar and packed lowerings, adding a 223-line HCU scaled_cast.h and lower_scaled_cast.cc, with 650 lines changed in the main transform file); deriving plain-copy swizzle from bank-word geometry and providing an explicit inverse transform for narrow b16 tiles (12 files, +1536/-112, 5 new test files); a proof of “at most two paths” for mixed LDS wrap; and switching the mixed-wrap check to physical read width; on 11-05 three more landed (default swizzle adapted to 64 banks, swizzle planning restored on the MN-slice consumer side, and layout-aware reads for K partitions).
  • Review status: #33 (scaled cast, i.e., the branch body) was opened on 11-05 09:08, and #34 on 19:52 (optimizing reduction’s XOR shuffle with DPP and lane swap, plus a fix that “layout inference first propagates known constraints, then assigns free candidates”). Going from “four commits accumulated on a feature branch with no PR opened” to two formal PRs is the first time a domestic backend has completed this path; the topics center on shared-memory bank conflicts and phase alignment, indicating the backend has moved from “it compiles” to “tuning against the microarchitecture.”
  • There is also #19 under review (title only “Distributed,” empty description, updated 11-03 11:06); this report does not speculate on its implementation details.

Moore Threads (MUSA) and MetaX: One Resumes Documentation Maintenance, the Other Still Only Upstream Syncs

  • Moore Threads: On 11-05 19:46 the main branch landed a docs: update README installation and operator development references commit, the repo’s first commit since 09-30; in the same period the timestamps of 6 lingering PRs (dependency upgrades, the TileOPs submodule, the evidence-threshold documentation for MUSA integration) were all refreshed, but their content remains old proposals from April to August. A conservative judgment is warranted: the documentation commit shows someone is maintaining the repo, which is not enough to prove backend development has resumed; the real signal will be whether operator or lowering commits follow.
  • MetaX: The daily report recorded “no pushes” across all five windows, and a review of the 168-hour window confirms that its main branch did have commits this week, but on a commit-by-commit basis they are all syncs from upstream tilelang (#3760, #3752, #3759, #3802, #3714, #3784, #3772, #3534, #3761, #3689, #3546, #3736, #3758, #3748, #3678, #3734, #3702, #3703, etc.), all authored by upstream maintainers, with no MetaX-side commits of its own. The accurate statement is therefore: the main branch is syncing with upstream, while in-house development remains in a quiet period, now for the third consecutive window. The dev branch is still last updated on 10-16.

mlir-ascend and TLIR: Landing Semantics into the Dialect, Layer by Layer

  • Structural cleanup: #206 merged on 10-31, removing references to an internal cache service from CI, a reproducibility cleanup; #209 merged on 11-04 16:35, moving the BishengIR toolchain adaptation code into contrib, separating “the glue layer that must change frequently as the upstream toolchain evolves” from the core lowering, so that future BishengIR upgrades confine the change surface to contrib.
  • Agent workflow: #208 merged into the harness branch (not the mainline; it should not be cited as “merged into mainline”), adding 2715 lines across 66 files and giving the operator harness its first Codex workflow: iterative questioning during project initialization, onboarding existing operator repos, and resuming development, tuning, or analysis from persisted context; the three skill sets share editable Markdown records under .workspace/, and task handoff distinguishes “current candidate” from “verified best candidate,” retaining failed and rolled-back experiments. The author explicitly states there was no real NPU testing this time, with NPU execution adaptation and npm distribution listed as follow-up work.
  • Dialect operators: #203 adds tl.sub/mul/div/max/min to TLIR and the TL-to-HIVM lowering, #211 adds tl.shl, and #212 adds tl.shr, the three advancing on 11-05 at 01:42, 19:44, and 20:58 respectively. This thread is slow but single-directional: landing TileLang’s semantics into the MLIR dialect, layer by layer.

IV. Ecosystem and Adopters

TileOPs: Third Week After 0.1.0, Shifting from Tuning to Selector Engineering

  • Merge cadence: 15 commits in the first window (13 on 11-01), followed by four windows of 7, 1, 7, and 9, totaling 39 for the week. The 11-02 batch introduced two non-performance changes — test infrastructure (#2816’s tileops_test target, 1,281 lines across 59 files, so tests/ops smoke tests no longer require CUDA) and a CI contract (#2818) — indicating the operator library is making structural preparations for multi-backend support and maintainability; the 11-04 and 11-05 windows returned to pure performance work.
  • Attention: #2808 consolidated program selection for GQADenseFwdOp and GQADenseWithLSEFwdOp into a single rule, AttentionCall.dense_decode_program, taking query and key lengths, batch, head grouping, head dim, RoPE, and SM count as inputs; 2 to 64 query positions use warp-specialized decode. Selection regret relative to the fastest available program (geometric mean and maximum) dropped from 1.818 and 25.2 to 1.017 and 1.81, and it fixed two real bugs on mainline: “score scale multiplied by log2(e) vanishes in FP32, causing NaN” and “D256 multi-position calls fail to launch on SM90 due to insufficient shared memory.” #2817 overlaps the softmax of warp-specialized decode with QK and PV (probabilities kept in registers, separate barriers for K and V, shared three-slot ring buffer); on H200, (16, 8, 4096, 64, 8, 128) float16 runs at 0.074610 ms vs. fa3’s 0.077490 (1.039×, 3.654 TB/s), and bfloat16 at 0.074260 vs. 0.077220 (1.040×, 3.671 TB/s). #2807 extends paged MHA decode to every head dim that “fits in one block’s shared memory,” and #2797 removes the old paged decode kernel from GQAPagedFwdOp.
  • MoE and GEMM: #2815 makes tight-layout MoE inputs collect by token, reusing the same hidden row across its multiple routes; on (4096, 256, 8, 7168, 2048), bfloat16 runs at 7.4417 ms vs. flashinfer-cutlass’s 7.6982 (1.034×, 387.844 TFLOPS), and float16 at 7.5040 vs. 7.9869 (1.064×, 384.624 TFLOPS). #2805 changes the partial single-wave W4A8 prefill grid that fills at least three-quarters of SMs from stream-K to one CTA per tile; llama-8b-up-b33 drops from 23.52 µs to 17.98 (vLLM 20.42), and across 44 N/K sweeps for off-list M values from 33 to 128, the reshaped schedule is 5% to 36% faster, median 18%, with none slower. #2834 builds a native-width tile for small-M NT GEMM whose width fills one wave; decode-down (128×7168×2048) uses only 112 of 132 SMs on mainline, and the change leads the fastest comparison by 3%, with gate-up-b32 1.4% faster than mainline.
  • Elementwise and reductions: #2814 switches bfloat16 Sigmoid and 16-bit SiLUAndMul to a quarter-scaled reciprocal, bypassing division with range checks; #2813 processes 8-byte operands as one vector per thread under long launches (after using TileFoundry’s roofline analysis to confirm these bitwise operators are all memory-bound); #2832 stops five operator classes — lerp, where, clamp, masked_fill — from first expanding stride-0 broadcast tensors to the output shape, cutting broadcast-row time to 0.13–0.44× mainline; #2829 collapses unary elementwise ops with “at most two cheap instructions in the body and total read/write no more than three-quarters of L2” into a single 256-thread wave, and incidentally fixes the old bug where ReLU returns 0 for NaN inputs; #2828 gives row-tiled ArgMax/ArgMin eight 16-byte vectors per thread under a “single-wave residency” constraint, cutting bf16 ArgMin on (2048, 4096) from 8.704 µs to 7.393 (FlagGems 8.543). #2835 computes axis permutation and output shape once per instance in reduction kernels; full hot-call time for six FP32/BF16 VarMeanFwdOp configurations drops 4.33% to 6.17%, with generated CUDA and host source identical — the savings are call-preparation cost, not GPU time.
  • Precision-class defects fixed head-on: #2811 adds an invariant to every row-composed fp32 quantity — “constrained only by the row’s spread, not by its offset” — implemented by taking an element within the row as the baseline and subtracting (warp-leading element for warp-dominant kernels, row-leading element for tiled kernels), eliminating overflow near 1e36; #2821 initializes variance merging from the first real partial rather than an empty accumulator, so constant inputs no longer produce NaN variance (tests/ops/test_reduce.py grows from 155 to 158). These “empty accumulator + large offset” numerical traps only surface in shapes that reduce over retained axes simultaneously; ordinary whole-axis reductions miss them.
  • Measurement methodology becomes a product feature: #2830 rewrites roofline around “how much this call actually reads” — GQA dense, varlen, and paged forward count K/V only for “key rows visible to some query” and count the RoPE table only up to the truncation position keys actually reach; GroupedGEMM prices a row’s group by batch_sizes via the new grouped_gemm_roofline. After the fix, three physically impossible rows go from 7.48, 7.50, and 5.30 TB/s to 0.62, 0.62, and 2.65 TB/s; #2822 changes bench_kernel’s L2 flush from “zero_() on an int8 buffer” to “read a private int32 buffer twice the reported L2 size,” because the old flush left about 120 MB of dirty lines whose writeback landed in the next timed call (2048×11008 f32 Sum drops from 27.81 to 21.38 µs).
  • Selectors and cost models: #2825 changes W4A16’s _select_config to take the argmin of _config_cost over _legal_configs, removing three exception rules — tile-first, widest-K tie window, and joint_switch_margin — and adding a stage_fill term; #2826 documents the dense selector’s calibration contract in the docstring (three steps: sweep, fit, check, including soft annealing and holdout validation); #2836 fits a separate factor (1.36) for coop2_swap’s three-stage tiles, cutting regret from 0.89% to 0.06% geometric mean and from 5.04% to 0.99% maximum.
  • Release form and test scope: #2816’s tileops_test detaches smoke tests from CUDA, and the author states the boundary clearly — comparing against a shared reference only proves backend-agnostic contracts and portability, not numerical correctness on a real backend, which remains tracked in #2611; #2827 moves the off-CUDA suite to nightly; #2833 upgrades the runner.

TileFoundry: Synchronization Analysis, Training Semantics, and Naming Governance

  • Synchronization: Beyond #283’s analysis layer and #285’s lowering execution side (see item five in Top Stories), #286 merged on 11-06 13:40, declaring completion conditions for each operand effect, attaching WGMMA’s fence, commit, and wait to the HIR nodes that need them, with lowering consuming these steps at the instruction’s issue point and completing them in the issuing scope if a later scope cannot hold the collective wait; it also fixes register Reshard aliasing so an element stays in the same register on every thread holding a copy. The author retains two caveats: the matrix instruction emitter is still not wired up, so this commit proves neither execution correctness nor performance; TMA completion semantics are handled separately. A separate #288 opened at 11-06 20:12 addresses pricing of “completed partial reductions” in analysis.
  • Training semantics: #284 merged on 11-05 16:36; the previous evaluator treated reshard as a global tensor and missed contribution differences among participants. This commit evaluates locally per participant, GPU coordinates, calls and loops, gather/scatter, and Partial(sum) reductions along axes, making Reshard’s aliasing proof ownership-sensitive. Acceptance: 775 core suite items, 778 backend suite items (13 skipped on dual-card sm90), ten HIR configurations passing on CPU, and two residency strategies running two update steps against a real 2×2 PyTorch HSDP on four-card sm90 with local gradients, parameters, and momentum all matching.
  • Scheduling and release: #280 and #282 consolidate the decisions of “what to emit per call,” “which mesh emits it,” and “how to split windows” into one place, analysis/emission.py, with byte-level acceptance — 75 of 83 scheduling fixtures produce finalize output identical to mainline, and the other 8 fail on mainline as well; the site repo published the 0.3 scheduling article “Scheduling, Designed for Agents” on 10-30, moving “agent-oriented scheduling” from code and release notes into public narrative.
  • Naming governance: #287 is a documentation-oriented issue noting “five names for one concept” in the codebase — a mesh member is called placement in 133 places, program in 110, instance in 92, and participant in 44, while the single word unit refers to the empty type, a hardware compute unit, and a mesh member simultaneously, and Python and the CUDA runtime disagree on the local_layout signature. Such issues rarely change code immediately, but they are a necessary step for a compiler project going from “can be written” to “can be read long-term.”

TileSight: Performance Model Moves from Paper to Repository

  • tile-ai/TileSight (24 stars) 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. On 11-05 the repo saw its first PR (#1): H20 is Hopper/GH100, and plans with requires_arch=("wgmma",) previously threw UnsupportedError because the architecture profile did not declare WGMMA support; this commit sets support_wgmma=True while explicitly keeping support_utcmma=False, so such plans are accepted while tcgen05 and Tensor Memory requirements are still rejected. Validation is 847 items all passing on CPU, and the author explicitly states this is “not GPU performance validation.”
  • The significance of this line is that it forms a triangle with TileOPs’ roofline pricing and TileFoundry’s analysis layer: one provides an analytical model, one provides line-by-line pricing, and one provides provably synchronized plans. All three saw substantive action this week.

Adopters: TileKernels and FlashQLA

  • Neither adopter repo saw commits all week: deepseek-ai/TileKernels is stalled at 10-14 05:49, and QwenLM/FlashQLA at 10-30 03:02 (latest release still v0.1.3 from 09-30).
  • But adopters not pushing does not mean they are unaffected: the barrier-occupancy regression fixed by #3761 on 11-04 was triggered by TileKernels’ mhc_pre_big_fuse kernel — the same kernel lost 25% of its barrier occupancy budget on B300 at hidden dim 1280. This shows adopter kernels have become regression probes on the compiler side: in 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.

Community Branches and Third-Party Repos

  • A RISC-V research chain takes shape: community researcher zishu0131-png organized three repos on the same day, 10-31 — tilelang-riscv as the main entry point (stating research goals, the two-stage compilation chain, and version combinations), numba2tilelang statically parsing Numba CUDA Python and converting it to TileLang source, and tilelang2rvv converting supported real tvm.tirx.PrimFunc into RVV intrinsic C, then compiling via Clang into an RV64GCV ELF running single-hart on QEMU virt. The author states the boundaries clearly: it is not an officially registered TileLang backend; threadIdx-dependent addressing, cross-hart synchronization, and unsupported IR are all rejected; blocks execute sequentially; the research environment is TileLang 0.1.15, Clang 23.1.3, and QEMU 11.1.0. The four-source inspection over the 168-hour window found no new changes this week.
  • PTO-ISA/tilelangGym spun out as its own repo: this sensitivity and regression test suite was split out from PTO-ISA/tilelang’s examples/ascend/t_parallel_st_suite so the suite can independently pin TileLang and PTOAS versions; the three main paths are Ascend SoftLib SimtVF (target=ascend, cython, the main regression gate), VectorizeParallelToPTO and VL’s vload/vstore, and hand-written SIMD DSL reference implementations under benches using @pto.jit; each test case directory carries a config.toml and VERDICT.md.
  • Small third-party repos: sablin39/tilelang-cuda-skills (149 stars, a skill set for writing TileLang and debugging with the CUDA toolchain), xwhzz/tilelang-hexagon (2 stars, bringing TileLang to Qualcomm Hexagon DSP), Index4NekoMeteor/Tilelang-AttnRes (1 star, implementing Kimi K3’s attention residual structure in TileLang), Tianchen-ai/tilelang-weft-bench (0 stars, a unified engine for CPU-side operator benchmarks), and flytiger-eco/tilelang-wrapper-for-sail (0 stars, no description yet). The first three are interesting in direction — one works on toolchain-side developer experience, two port to new target architectures and frontier model structures; the latter two are very early and should go on the watch list rather than be written up as technical progress.

V. Community, Tutorials, and Events

  • Issue reports cluster around installation configurations: Of the 7 new issues on 11-03, four came from the same external reporter using the pip package, running WSL, and using a consumer-grade sm_89 card (RTX 4060 Laptop), all with reproducible scripts and full output attached; the most technically substantive, #3775, was a production-environment combination of “pip package + dual H200 NVLink + multiple CUDA contexts + nvrtc backend + CUDA graph capture,” and the reporter laid out the control conditions for all three runs in full—three byte-identical configurations reproduced three different first points of failure, and the only difference between /proc/environ and the clean control was a single environment variable, TILELANG_EXECUTION_BACKEND=nvrtc. These are not deep compiler defects, but they determine the experience of a new user’s first attempt.
  • A “reporter-as-reviewer” loop has emerged on the Ascend side: The reporters of #1991/#1992 were the reviewers of #1897, and the fix in #1897 covered only the count-side semantics (HasModifiedBuffer skipping workspace slots), while the workspace side remained unaddressed—another facet of the same layout logic was pointed out only two weeks later. This shows that the layout invariants of the Ascend backend have yet to form a checkable written contract.
  • Documentation and localization pipelines are functioning normally: tilelang.github.io received multiple automatic Update docs syncs throughout the week, and TileOPs.github.io received one push; the LSP repo aligned CUDA debugging with the new line mapping (71e762e0), and the VS Code extension repo updated its documentation and renamed its artifact to vsix, following on from 11-02’s “align CUDA debugging with TileLang line mapping”—the line-by-line anchoring merged into the mainline in the previous issue changed the source of line numbers in generated code, so the debugging configuration on the extension side needed to change accordingly.
  • Tutorials and knowledge assets: No new tutorial repos appeared this week; community effort concentrated on maintaining existing repos and on the two small RISC-V-oriented repos mentioned above; the spin-out of PTO-ISA/tilelangGym into standalone repos marks the first time test assets have been versioned separately.
  • Zero substantive news hits all week: Across the five collection windows, Google News queries in both Chinese and English (covering exact terms, GPU kernel, Chinese operators and compilers, Ascend and MetaX/Hygon/Moore Threads) turned up the latest relevant report as SemiAnalysis’s “The Huawei-DeepSeek Symbiosis” from 10-21, with the rest concentrated in the 09-30 wave of reports on “DeepSeek open-sourcing the Ascend version of TileLang, DeepGEMM, DeepEP,” all predating this period; a Hacker News search for tilelang over 24 hours returned 0 results; the latest arXiv hits were “Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving” from 11-03 (a work citing TileLang, outside the window) and TileSight itself from 07-24. The main body of the report therefore comes entirely from the code side, which has been the stable pattern for several consecutive issues.
  • An industry-side discrepancy in framing is worth noting: In the 11-03 window, Wccftech claimed that DeepSeek would train V5 from scratch on an Atlas 950 SuperPoD built from Ascend 950DT, while a contemporaneous report by The Information said “pre-training on 950DT starting no earlier than the end of 2026,” and earlier reports by Bloomberg and the Financial Times said that the 160,000-unit order for 950DT was mainly for inference, after previous training efforts on Ascend faltered and fell back to Nvidia. The three accounts are inconsistent on “whether and when it will be used for training,” and none has been officially confirmed by Huawei or DeepSeek; this report records them as background only and does not infer the pace of Ascend-side deployment from them.

VI. Trend Observations

First, the TIRx migration has effectively become the top priority on the mainline, and it is starting to take the form of “subtraction.” Three pieces of evidence within one week: on 11-01, the entire set of ThreadSync correctness improvements was reverted for the migration (#3733, while also closing the optimization proposal #89 on the TVM fork side); on 11-03, 3,384 lines of the legacy parser fork were deleted (#3736); on 11-06, a complete migration draft was submitted (#3803, pinning two upstream commits and explicitly stating that no patches will be injected). In between, Z3 resources went from hardcoded to configurable (#3752 and tvm #91). What these actions share is that they pull complexity off the critical path: better to temporarily fall back to a conservative implementation than to let the internal structure of vendored dependencies keep growing inside the language layer. The cost is that during the migration window, the mainline will have a batch of paths that are “known not to be optimal”; the benefit is that once the foundation upgrade is complete, all downstream backends can adapt against a stable IR boundary.

Second, the form of correctness risk has changed: from “fails to compile” to “compiles, runs without errors, but produces wrong results,” and it lurks across versions. This week, at least seven issues belong to this family: #3808 is missing a __syncthreads() across four versions from 0.1.12 to 0.1.15; #3800 lets the MMA read an overwritten operand on Blackwell’s tensor memory path; #3721 silently drops the entire GEMM product on sm_100a; #3775 corrupts CUDA graph capture in production; #3749 makes the dependency analysis blind to the physical load range; #1991/#1992 make sin/cos compute 97.1% of elements incorrectly under Ascend with threads=2; #3774 and #3773 produce wrong results on consumer-grade cards. This family hurts users the most because it turns the premise of “whether the compiler is correct” into a hypothesis that needs verification. This week’s response direction is to turn correctness into a refutable proof—TileFoundry’s #283/#285 prove before emitting and reverse-verify after emitting; #3786 and #3787 change silently generating wrong code into compile-time errors or explicit rejection—rather than continuing to increase the number of unit tests.

Third, Ascend’s role has shifted from “adaptation target” to “first-launch target,” with engineering advancing on two legs simultaneously. One leg is architecture: FlashMLA, a core kernel on DeepSeek’s own decoding path, is written directly into the mainline targeting Ascend 950, and it uses low-level capabilities such as T.Stage and annotate_buffer_addresses that are prepared for explicit scheduling; together with three supporting changes on the compiler side, the Ascend path is moving from “it runs” to “it runs according to scheduling intent.” The other leg is the skeleton: #1975 moves the entire TileLang-Ascend into the mainline multi-backend framework with 2,887 files and 600,000 added lines, using arch="dav-2201" for backend dispatch; mlir-ascend’s #208 writes Codex’s three workflows—initialization, onboarding, and recovery—along with task handoff into .workspace/. Looking at “the migration repo” and “using agents to migrate the repo” in the same week, one can see that the bottleneck on the Ascend side has shifted from “missing operators” to “missing a consistent engineering skeleton and a reusable migration process.”

Fourth, the competitive focus of operator libraries has shifted entirely to selector quality and preparation cost; the remaining room for performance gaps is not in “how to compute.” Among the nine TileOPs merges on 11-05, only two actually change the GPU kernel: broadcast in-place read and native-width tiles; the reduction layout cache saves call preparation (4.33% to 6.17%, with generated source unchanged), and the three-stage tile factor changes a constant (regret drops from 0.89% to 0.06%). Combined with the roofline pricing rewrite and the change of L2 flush to read-only, a clear route emerges: first enumerate the legal configurations for each shape, then rank them with a unified cost model, and write the calibration process as a reproducible contract (#2826 has already written the three steps of “scan, fit, check” into the docstring). If this route is done right, performance no longer depends on manual tuning of individual kernels; if done wrong, it degenerates into “some shapes are especially fast, and other shapes rely on luck.”

Fifth, domestic backends are shifting from “merging one big backend at once” to “pushing the smallest reviewable units to the public repository.” Hygon has moved from its in-house branch to two formal PRs, #33 and #34; Moore Threads ends two months of silence with a single README; mlir-ascend uses three small PRs to move tl.sub/mul, tl.shl, and tl.shr one by one into the TLIR dialect; and Ascend fills in engineering details such as core ratios, exception dumps, and performance skills item by item. What the three lines share is that they no longer wait for one big merge: this is good for upstream, because reviewers can comment on the layout and lowering design early instead of seeing it only at the time of a 30,000-line one-shot merge; it also has a cost for downstream—the mainline will be in a state of “known to have several things not working” during such windows, requiring stronger pre-release regression discipline.

One final risk surface: the foundation switch will expose the cost of all glue layers at once. If #3803 moves forward, the vendored structure of 3rdparty/tvm will be replaced by two pinned upstream commits, at which point all backends that depend 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 will be injected,” which is both a commitment and a constraint—this week, 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. To judge whether this line is healthy, look at whether, within two weeks after the migration draft is merged, these backends can complete adaptation without touching the core lowering.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
GitHub tile-ai org, repo by repo (29 repos) The five collection windows had 10, 9, 9, 8, and 10 repos with pushes respectively; a 168-hour window recheck found 14 repos with pushes
Merged ground-truth search (is:pr is:merged) The five windows had 35, 18, 8, 12, and 18 PRs, totaling 91; the first window covers 72 hours (10-30 07:00 to 11-02 07:00)
Open ground-truth search (is:pr is:open updated) The five windows had 82, 44, 48, 42, and 57 PRs, totaling 273
New issue verification The five windows had 11, 8, 7, 4, and 5 issues, totaling 35
Main-branch commit-level verification tilelang 7, 5, 5, 3, 5; TileOPs 15, 7, 1, 7, 9; TileFoundry 3, 1, 1, 1, 1; tilelang-ascend 5, 4, 0, 1, 2; tilelang-mlir-ascend 1, 0, 1, 2, 0; tilelang-hygon main branch 2, 1, 0, 0, 0; tilelang-lsp and the VS Code extension 1 each; tilelang-musa 1
Release pages and tags tilelang latest v0.1.15 (09-30), no new version all week; TileOPs latest v0.1.0 (10-27); TileFoundry latest 0.0.4rc1 (10-25); tilelang-ascend v0.1.2.000 (09-09); tilelang-metax v0.1.15 (10-16)
Adopters TileKernels last active 10-14; FlashQLA last active 10-30, latest release still v0.1.3 from 09-30
TileLang×RISC-V four-source sweep No changes in the 168-hour window
Third-party and ecosystem repos The three research-chain repos, tilelangGym, cuda-skills, hexagon, AttnRes, weft-bench, wrapper-for-sail, and TileSight all had pushes or new creations
Google News (four to six query sets in Chinese and English) No substantive in-window hits across the five windows; the latest related coverage is a SemiAnalysis article from 10-21
Hacker News A 24-hour search for tilelang returned 0 results; what was found were unrelated Show HN posts
arXiv Latest hit is a KV compression paper from 11-03 (cites TileLang, outside the window) and the TileSight paper itself from 07-24

Main References

Notes

  • All facts, version numbers, and dates in this report are taken from the original text of the five daily reports within the period and their primary sources, with no speculative additions; the five collection windows are contiguous (10-30 07:00 to 11-06 07:00, 168 hours total) and do not overlap with the previous weekly report, so the merged, open, and new-issue counts across the five windows are summed directly.
  • Merged counts, open counts, and issue counts come from actual GitHub search output; main-branch commit counts come from repo-by-repo atom verification. The two use different methodologies (one merge may correspond to multiple commits, and commits on non-default branches are not counted toward merged counts), so the appendix lists them separately and does not total them.
  • The “15” for TileOPs appeared within Monday’s 72-hour window, of which 13 fell on 11-01; when citing, it should not be written as “15 on Monday alone.”
  • #1991 and #1992 have identical titles and bodies and are recorded as “two entries with the same title,” not merged into one; the seven Hygon feature-branch items and the Moore Threads documentation commits are described as branch-level activity and do not mean they have been merged into the mainline.
  • MetaX is recorded as “no pushes” in all five windows; the 168-hour recheck found commits on the main branch, but item-by-item verification shows they are all syncs from upstream tilelang (authored by upstream maintainers), so it is described as “main branch syncing with upstream, own development still in a quiet period.” The WGMMA changes in TileSight and TileFoundry are both self-described by their authors as limited to capability metadata or sync plans, and the report follows those qualifications.