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

Sources: GitHub tile-ai organization per-repo verification (29 repos, 9 pushed within the window) + merged and in-review ground-truth search (18 merged, 44 in-review updates) + full review of new issues (8) + per-commit review of main branches (tilelang 5, TileOPs 7, TileFoundry 1, tilelang-ascend 4, tilelang-hygon 1, tilelang-lsp 1, docs site 5) + release page and tag verification (5 repos) + two adopter repos + Google News Chinese and English queries (via proxy) + Hacker News + arXiv.

This Issue’s Index

  • Today’s Focus: Ascend 950’s persistent FlashMLA kernel lands (#3678, 1997 lines added, 42 commits); a real regression fix in vectorized planning (#3734); #line directives switched to per-line re-anchoring (#3703); TileOPs merges 7 performance and test-infrastructure PRs in a single day; TileFoundry turns shared-memory ordering into a provable analysis layer (#283), and finds that main’s lowering gets 4 of 7 programs wrong on sm_89; 18 merges, 44 in-review updates, and 8 new issues in 24 hours
  • I. Core Project Progress
    • 1.1 Mainline: Ascend 950’s persistent FlashMLA kernel, dual paths for BF16 prefill and compact FP8 decode (#3678) (11-02)
    • 1.2 Mainline: residual 512-bit vectors after cast decoupling is skipped, ViT forward compilation fails on SM100 (#3734) (11-02)
    • 1.3 Mainline: #line directives switched to per-line re-anchoring, debugger mapping gains optional dense mode (#3703) (11-02)
    • 1.4 Mainline: attention example derives tiles from device shared-memory limit, RTX PRO 6000 goes from 2 pass / 7 fail to 9 pass (#3702) (11-02)
    • 1.5 Mainline: CUDA float32x8 bit conversion switched to builtin bit cast, SASS instruction count drops from 19 to 18 (#3748) (11-02)
  • II. Multi-Backend Adaptation
    • 2.1 Ascend: WY Fast forward operator added, wy_fast_bwd_split backward segmented migration (#1827 / #1828) (11-02)
    • 2.2 Ascend: T.tile.fill and T.tile.clear documentation fact-check merged (#1634 / #1628) (11-02)
    • 2.3 Ascend in review: cython wrapper caches symbol shape resolution, continuing the host-overhead line (#1967) (11-02)
    • 2.4 Ascend in review: z3 version sync migrated from mainline, whole-repo mainline sync advancing (#1965 / #1975) (11-02)
    • 2.5 Ascend in review: two reduction operator defects, first Parallel axis broadcast changed to scalar read (#1978 / #1968 / #1979) (11-02)
    • 2.6 Hygon: gfx948 handled as gfx946-compatible, first making existing instructions testable (#32) (11-02)
    • 2.7 Moore Threads, MetaX, mlir-ascend: no pushes in window (11-03)
  • III. Ecosystem and Adopters
    • 3.1 TileOPs: 7 merges in a single day, 6 performance, 1 test infrastructure (11-02)
    • 3.2 TileOPs: warp-specialized decode softmax overlaps QK/PV, 478 sweep calls geometric mean 0.952x (#2817) (11-02)
    • 3.3 TileOPs: MoE per-token gather pre-permutation, 387.8 TFLOPS on H200 (#2815) (11-02)
    • 3.4 TileOPs: 16-bit elementwise drops range-checked division, 8-byte operands one vector per thread (#2814 / #2813) (11-02)
    • 3.5 TileOPs: Norm statistics changed to subtract in-row elements, eliminating overflow near 1e36 (#2811) (11-02)
    • 3.6 TileOPs: tileops_test target lets tests/ops run without CUDA (#2816) (11-02)
    • 3.7 TileFoundry: shared-memory ordering made into a provable analysis layer, 83 fixture plans: 52 rejected, 31 accepted (#283) (11-02)
    • 3.8 TileFoundry in review: CUDA lowering emits sync plans and reverse-validates, 4 of 7 programs wrong on main (#285) (11-02)
    • 3.9 Adopters: FlashQLA, TileKernels no pushes in window (11-03)
  • IV. Community, Tutorials and Events
    • 4.1 Ascend sync correctness becomes a concentrated bug-report surface: 8 new issues in a single day, two reporters each filing systematically (11-02)
    • 4.2 SIMD loads of cross-row and transposed views drop DMA dependencies, root cause in logical region reconstruction (#3749) (11-02)
    • 4.3 News side: no in-window hits across Google News Chinese/English, Hacker News, and arXiv channels (11-03)
  • V. Trend Observations
  • Appendix: Materials and Verification Notes

Today’s Highlights

There is only one main thread today: TileLang has ported DeepSeek’s MLA decode kernel to Ascend 950, and it has done so in a way that the compiler can schedule. #3678 adds 1,997 lines, 20 files, and 42 commits in one go, covering both the BF16 prefill and compact FP8 decode paths, and also including optional FP8 and FP4 extra KV caches, attention sink, illegal indices, and per-query independent lengths. The implementation approach is noteworthy: the attention loop, online softmax, dequantization, and scheduling are all written in TileLang, relying on T.Stage to let the compiler manage buffer versions and synchronization; QK and PV use two different layout views of the same L0B storage ring; and L1 placement is specified directly with T.annotate_buffer_addresses; prefill splits four disjoint output slices out of a single accumulator allocation, and drains each slice as soon as its normalization is computed. To make this style of writing work, the compiler side simultaneously adds three capabilities — L0C-to-UB copies preserve NZ output, each AIV stream is selected before its own pipeline, and explicitly scheduled stages keep source order when their physical issue phases are the same. The target shapes are written very narrowly: 64 heads, Dqk and Dv both 512, and topk a multiple of 64.

The second thread is fixing a real compiler regression. #3734 deals with ViT forward on SM100: Cannot convert type float32x16 to CUDA type. The path is long but every step is clear — the vectorization planner chose 16 lanes for a 256-bit BF16 store and deferred the cast constraint, expecting DecoupleTypeCast to split computation and memory access; but #3579 had just excluded “guards for reading memory” from decoupling (duplicating these guards would change semantics), so the final legalization step was still choosing widths according to the deferred strategy, and the residual FP32 expression was stretched into a 512-bit vector, which the CUDA backend cannot express. The fix is to make the deferred constraint a planner parameter rather than an IR attribute; only the final legalization step passes false, the guard’s FP32-to-BF16 conversion drops to 8 lanes, and successfully decoupled BF16 copies still keep 16 lanes; address casts are also removed from the data vector type constraint and their legality is judged by the existing access analysis.

The third thread is the debugging experience. #3703 changes #line directives from “anchoring only the first line of each TIR statement” to optional graded emission: level 0 disables it, 1 keeps the original behavior, and 2 and above re-anchor every non-empty line generated for each line, including statements expanded into multiple lines and closing braces. The implementation is placed in the C-family output stream shared by CUDA, HIP, CPU C, C host, Metal, and Ascend, without touching vendored TVM; statements and functions without source spans are explicitly switched to <tilelang-generated>, and user source text is preserved verbatim to avoid breaking macro continuations and string literals. A new environment variable TILELANG_EMIT_LINE_DIRECTIVES serves as the Python-side default and is folded into the pass configuration before cache lookup. Locally, 42 tests pass and 6 are skipped; the CUDA and HIP cases were skipped on the author’s machine, and the actual behavior of cuda-gdb has not yet been directly verified — the author states this explicitly. On the same day, the LSP repo followed up with one commit aligning CUDA debugging to the new line mapping.

The pace on the operator side has not stopped, but the direction has changed. Of the 7 merges in TileOPs in a single day, 6 are performance and 1 is test infrastructure, all from the same author; each still comes with an H200 measured table and a checklist regression gate. The most valuable is #2816: it adds a tileops_test target that exists only during testing, computes a shared workload reference for each operator on PrivateUse1, and means the smoke tests in tests/ops no longer have to stand on CUDA. The change is 1,281 lines across 59 files, and the author also clearly states the boundary — comparing against a shared reference can only prove backend-independent contracts and portability, not the numerical correctness of a real backend; that line is still tracked in #2611.

Finally, there is a hidden thread in TileFoundry. #283 turns “how shared memory accesses should be ordered” into an analysis layer: a synchronization family walks each function in lowering emission order, and for each shared memory access takes the instruction, issue mesh, operands, and byte range, and plans commit/wait and scope fences before conflicting accesses. Of 83 scheduling fixtures, it plans 52 and rejects 31 (24 use TMA copies and 6 use wgmma, whose completion semantics are not yet declared; 1 requires inserting a fence in a loop with unequal cross-thread boundaries). The companion #285 is under review, turning the plan into CUDA lowering and validating it in reverse: of the 7 programs that can be built on a single sm_89 card, 6 match the reference exactly, and synccheck and memcheck are completely clean, whereas lowering the way main does produces wrong results in 4 of the same 7; the direct symptom given in the #283 description is that the E4M3 GEMM on sm_89 is off by up to 659.

I. Core Project Progress

1.1 Mainline: Persistent FlashMLA kernel for Ascend 950, dual paths of BF16 prefill and compact FP8 decode (#3678)

Merged on 11-02 09:26, adding 1,997 lines and deleting 17 lines, across 20 files and 42 commits — the heaviest item in this window. The kernel targets Ascend 950: prefill runs BF16, decode runs compact FP8, with optional extra KV cache in FP8 and FP4, and support for attention sink, invalid indices, and per-query independent lengths. The attention loop, online softmax, dequantization, and scheduling are all expressed in TileLang, driven by T.Stage together with compiler-managed buffer versions and synchronization; QK and PV share two different layout views of a single L0B storage ring, and L1 placement is specified via T.annotate_buffer_addresses; prefill splits four disjoint output slices within the existing accumulator allocation, each with its own synchronization event, draining each slice as soon as its normalization completes; decode uses a full-row output copy. Three compiler-side supports were added: L0C-to-UB copies preserve NZ output, each AIV stream is selected before its pipeline, and explicit scheduling stages keep source order when flattened at the physical issue phase. The target is 64 heads, Dqk and Dv both 512, and topk a multiple of 64; pair gather reuses the documented zero-burst codegen workaround; the PR ships with correctness and benchmark runners, buffer-reuse and query-boundary regressions, and a README carrying the FlashMLA MIT notice.

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

1.2 Mainline: Residual 512-bit vectors after cast decoupling is skipped, ViT forward fails to compile on SM100 (#3734)

Merged on 11-02 09:05. On e8225d7f5, the SM100-targeted ViT forward reports Cannot convert type float32x16 to CUDA type at the guarded FP32-to-BF16 output store: the planner selects 16 lanes for the 256-bit BF16 store and defers the cast constraint, expecting DecoupleTypeCast to separate compute from memory access; after #3579 made decoupling skip guards that read memory, the subsequent LegalizeVectorizedLoop kept deferring under the same policy, so the residual FP32 expression became a 512-bit vector that the CUDA side cannot represent. The fix makes defer_cast_constraints a C++ parameter of the planner (not an IR attribute), chosen by the caller per phase: LayoutInference and LowerTileOp keep true, final legalization always passes false, and the computed width is forwarded as an explicit vectorize_hint so the rewriter does not re-plan with the default policy. The guarded FP32-to-BF16 drops to 8 lanes, successfully decoupled BF16 copies keep 16 lanes, and the guard safety check is retained; address casts and access-pointer parameters are excluded from data vector type constraints, contiguous addresses are lowered to a scalar base plus stride with legality decided by the existing access analysis, and aligned vector accesses with int64 indices and the four-lane atomic case are unaffected.

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

1.3 Mainline: #line directives re-anchored per line, debugger mapping gains optional dense mode (#3703)

Merged on 11-02 04:06, 682 lines added across 11 files. Previously #line emission anchored only the first line generated by a TIR statement, and subsequent lines were advanced by the compiler’s logical line numbers, so cuda-gdb attributed generated loads, arithmetic, and stores to unrelated Python lines; synthesized statements and closing braces could inherit an advanced mapping, and functions without a source span inherited the previous function’s source file. This PR extends tl.emit_line_directives into levels: 0 or False disables it and keeps the default, 1 or True keeps statement-level, and 2 or above re-anchors every non-empty generated line (including multi-line expanded statements and closing braces). The implementation lives in the C-family output streams, covering CUDA, HIP, CPU C, C host, Metal, and Ascend, without modifying vendored TVM; it also records source spans for statements and functions, explicitly switches source-less regions to <tilelang-generated>, and preserves user text verbatim to avoid breaking macro continuations and string literals. A new TILELANG_EMIT_LINE_DIRECTIVES serves as the Python-side default, folded into the pass config before cache lookup, with explicit config taking precedence (including when the default value is invalid). Verified locally with a Metal plus LLVM build: the line-directive and compile-only groups pass 42 and skip 6; with directives removed, the vectorized-add and conditional-select fixtures produce identical output at levels 0, 1, and 2. The author notes that CUDA and HIP cases were skipped on the local machine and cuda-gdb behavior was not directly tested.

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

1.4 Mainline: Attention examples derive tiles from device shared-memory limits, going from 2 pass / 7 fail to 9 pass on RTX PRO 6000 (#3702)

Merged on 11-02 06:59, 345 lines added across 13 files. The non-tuned paths of several attention examples hard-coded Hopper-sized tiles, which fail immediately on sm_86, sm_89, and sm_120 where the per-block dynamic shared memory limit is about 99 KiB: the error asks for 163840 bytes of allowed dynamic shared memory while device 0 allows at most 101376 bytes per block. Each example now derives its tile from shared_memory_per_block_optin, following the existing GQA decode example (example_gqa_decode.py); datacenter configurations are unchanged because each branch only triggers below the previous tile requirement. Landing points on ~99 KiB devices: the four attention sink examples use 64×64 two-stage with 128 threads (forward) and 64×32 one-stage with 128 threads (backward), example_gqa_bwd.py uses 64×64 forward and 64×32 backward, and example_mla_decode.py takes BLOCK_N 32 and BLOCK_H 16. The latter two gain a threads parameter: with 64-row or 16-row tiles, T.gemm needs four warps so each warp owns a full 16 rows, and 256 threads fail to lower outright (reporting a layout conflict for acc_s and acc_s_cast respectively, and no legal warp partition for M=16 N=32). Separately, flash-attn 2.x has no sm_12x kernels, so the four varlen examples previously threw during the reference step; they now probe flash_attn_varlen_supported once and fall back to a pure-PyTorch varlen_attention_ref (covering GQA, bottom-right-aligned causal, and window_size). Verified on RTX PRO 6000 Blackwell (sm_120): CUDA passes all 9 (previously 2 pass / 7 fail), and CuTeDSL also passes all 9; H100 has a 227 KiB opt-in limit, still uses the original tiles, and likewise passes all 9.

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

1.5 Mainline: CUDA float32x8 bit conversion switched to builtin bit cast, SASS instruction count drops from 19 to 18 (#3748)

Merged on 11-02 09:58, 20 lines added and 3 deleted, fixing #3679 and superseding #3681. The change is CUDA-only: the float32x8 bit reinterpretation moves from a bit-manipulation sequence to __builtin_bit_cast(uint64_t, make_float2(v, v)), reducing the generated SASS instruction count from 19 to 18. A small change, but a typical compiler-backend cleanup of the “semantically equivalent, shorter output” kind.

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

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

2.1 Ascend: WY Fast forward operator added, wy_fast_bwd_split backward segmented migration (#1827 / #1828)

Both merged on 11-02 at 02:33, less than a minute apart. #1827 “migrates and adds” the WY Fast operator to the Ascend NPU side, while #1828 migrates and completes the segmented backward pass of wy_fast_bwd_split. WY notation is a common technique in linear-attention-style models for reducing recursive matrix multiplication to a low-rank form; pairing the forward pass with the segmented backward pass completes this operator chain on Ascend to the point of being trainable. The PR descriptions for both contain only the template greeting, so the scope of changes is based on titles and commits; this report does not infer internal implementation details.

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

2.2 Ascend: T.tile.fill and T.tile.clear documentation fact-check merged (#1634 / #1628)

Also merged in this batch were two documentation-side fact-checks, targeting the API behavior of T.tile.fill and T.tile.clear respectively. The discussion in #1628 left a maintenance convention: these two interfaces are fairly similar, so when one is modified, please check whether the other needs a corresponding adjustment. This is part of an ongoing batch of API fact-checks; still under review in the same period are T.tile.bitwise_not (#1656), T.tile.bitwise_and (#1651), T.tile.bitwise_lshift (#1659), and T.tile.bitwise_or (#1655).

Source: https://github.com/tile-ai/tilelang-ascend/pull/1634, https://github.com/tile-ai/tilelang-ascend/pull/1628

2.3 Ascend under review: cython wrapper caches symbolic shape resolution, continuing the host-overhead line (#1967)

Updated 11-02. #1967 caches the resolution results of symbolic shapes in the cython wrapper, aiming to further reduce per-call dispatch overhead, corresponding to defect #1898. This line and #1948, merged in the previous issue, are two landing points for the same goal: #1948 changed the fixed overhead of unconditionally constructing a TVM Analyzer on every call to lazy construction (from about 450 microseconds to about 106 microseconds on 910B3), while #1967 addresses the cost of repeated resolution itself on the symbolic path. In the same family, #3757 is under review on the mainline side, aligning the dynamic ABI and shared-symbol validation: keying by Var object rather than printed name to avoid treating different variables with the same name as the same symbol, and treating strides on singleton dimensions only as a fallback, not letting them preemptively determine the value.

Source: https://github.com/tile-ai/tilelang-ascend/pull/1967, https://github.com/tile-ai/tilelang/pull/3757

2.4 Ascend under review: z3 version migrated in from mainline, whole-repo mainline migration advancing (#1965 / #1975)

#1965 was updated on 11-02 at 12:18, migrating the mainline z3 version of insert sync to the Ascend repo. #1975 was updated the same day at 03:28 and is still about migrating TileLang-Ascend as a whole onto the mainline multi-backend architecture. These two are two layers of the same convergence path: one swaps the implementation of the synchronization algorithm, the other swaps the repo’s skeleton. The accompanying mainline-side actions are #3752 (z3 rlimit configuration) under review, and #3735 proposing to prove the disjointness of modular indices in thread sync—the proof-budget problem left by the ThreadSync rollback in the previous report is being addressed from both ends at once: “configuring solver resources” and “reducing the shapes that need proving.”

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

2.5 Ascend under review: two reduction operator defects, broadcast on the first Parallel axis switched to scalar reads (#1978 / #1968 / #1979)

All three were updated on 11-02. #1978 fixes AscendVidReduction incorrectly halving the tile operator size for the whole [1,N] UB block in the row-splitting loop (corresponding to #1800); #1968 halves the counts for sigmoid and round when threads=2 (corresponding to #1897); #1979 makes broadcast operands on the first Parallel axis use scalar reads, bypassing VEC’s UB alignment trap (corresponding to #1885). All three belong to the same class of problem: the Ascend-side vectorized path makes harder assumptions about shapes and alignment than the CUDA side, and once a boundary condition breaks, it is a hard error at compile time or runtime.

Source: https://github.com/tile-ai/tilelang-ascend/pull/1978, https://github.com/tile-ai/tilelang-ascend/pull/1968

2.6 Hygon: gfx948 handled as gfx946-compatible, first making existing instructions testable (#32)

Merged on 11-02 at 10:46. The change treats gfx948 as gfx946-compatible, 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 require new instructions. The repo’s previous routine-copy LDS swizzle fix was merged in the previous window.

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

2.7 Moore Threads, MetaX, mlir-ascend: no pushes within the window (11-03)

The last push to tilelang-musa was on 09-30, tilelang-metax on 10-26, and tilelang-mlir-ascend on 10-31 08:21 (earlier than the start of this window). This issue’s backend activity is therefore concentrated on the Ascend and Hygon sides, with Ascend accounting for 4 of the 18 merges and the majority of the 44 under-review updates.

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 TileOPs: 7 merges in a single day, 6 performance and 1 test infrastructure (11-02)

All 7 merges landed between 11-02 05:10 and 05:22, all by the same author, lcy-seso, covering six directions: attention, MoE, elementwise, normalization, test targets, and CI contracts. Compared with the previous issue’s “13 merges in one day,” the count has halved, but the structure has changed: the previous issue was per-shape tuning cleanup, whereas this issue includes two non-performance changes—test infrastructure (#2816) and CI contracts (#2818)—indicating that the operator library is beginning to make structural preparations for multiple backends and maintainability. Each merge still includes an H200 measured-results table and states that no already-tested rows in the manifest regressed.

Source: https://github.com/tile-ai/TileOPs/commits/main

3.2 TileOPs: warp-specialized decode softmax overlapped with QK/PV, geometric mean 0.952× across 478 sweep calls (#2817)

Merged at 11-02 05:21. The change overlaps the softmax of warp-specialized GQA decode with the QK and PV matrix multiplications: probabilities stay in registers, K and V use independent barriers, and they share a three-slot ring buffer. The selection conditions are D=128, a tile of 64 keys, one row tile per KV head, no softcap, and a “one CTA per SM” grid that satisfies occupancy and stream-length bounds; calls that do not meet these conditions continue to use the serial program. All nine planner boundaries are validated on the host side using fixed device facts, without requiring a real attention tensor allocation or a 132-SM card. The measured environment is H200, CUDA 13.2, PyTorch 2.13.0+cu132, TileLang 0.1.12, using CUPTI device busy time: for shape (16, 8, 4096, 64, 8, 128), float16 is 0.074610 ms vs. fa3’s 0.077490 (1.039×, 3.654 TB/s), and bfloat16 is 0.074260 vs. 0.077220 (1.040×, 3.671 TB/s); the 478 affected sweep calls recorded a geometric-mean latency of 0.952× that of main.

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

3.3 TileOPs: MoE gathers and permutes rows by token, 387.8 TFLOPS on H200 (#2815)

Merged at 11-02 05:19. The tight-layout MoE input is now gathered by token, so the same hidden row is reused across its multiple routes; when tokens do not fill the GPU, routes are split across different blocks. At the same time, the unused expert-order source-index allocation and writes in the tight sweep are removed. A new scheduling rule is added: for an unfused, K-major tight group, if the average row count is greater than 128 and no more than 192, select the ping-pong 64×256 tile on H200; calls that explicitly specify the group size still use the cost model. Measured on (4096, 256, 8, 7168, 2048), bfloat16 is 7.4417 ms vs. flashinfer-cutlass’s 7.6982 (1.034×, 387.844 TFLOPS), and float16 is 7.5040 vs. 7.9869 (1.064×, 384.624 TFLOPS); the recorded routing and even-group sweeps cover both dtypes and bias modes, with no above-noise regressions in the manifest.

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

3.4 TileOPs: 16-bit elementwise drops range-checked division, 8-byte operands use one vector per thread (#2814 / #2813)

#2814 merged at 11-02 05:18. For bfloat16 Sigmoid and 16-bit SiLUAndMul, it switches to a quarter-scaled reciprocal, bypassing range-checked division, and shares that helper with SiLU; it also removes the unnecessary subnormal guard in the exponential, stores the fused gated output across threads as adjacent 16-byte vectors, and preserves the unrounded SiLU activation that multiplies by infinity when the gate is near negative 88. Measured bfloat16 Sigmoid on 16777216 elements is 0.017504 ms vs. torch-compile’s 0.018656 (1.066×, 3.834 TB/s), and on 268435456 elements is 0.248545 vs. 0.262529 (1.056×, 4.320 TB/s); SiLUAndMul on (2048, 14336) is 0.041856 vs. 0.045440 (1.086×). #2813 merged at 05:17 and makes 8-byte operands process one vector per thread under long launches; before the change, TileFoundry’s roofline analysis confirmed that all such bitwise operators are memory-bound: at 4.8 TB/s HBM bandwidth, the lower bounds for each workload are 41.9 µs for int64 and 21.0 µs for int32 on hidden-state-prefill, 56.4 and 28.2 for llama-7b-ffn, 42.8 for cnn-feat-broadcast int64, and 14.0 and 7.0 for linear-bias.

Source: https://github.com/tile-ai/TileOPs/pull/2814, https://github.com/tile-ai/TileOPs/pull/2813

3.5 TileOPs: Norm statistics subtract an in-row element, eliminating overflow near 1e36 (#2811)

Merged at 11-02 05:10. The root cause is stated plainly: all mean and variance kernels sum the raw row in fp32, and the sum grows with the row’s offset (roughly N times x ); rows near 1e36 overflow outright and all outputs become NaN. Even without overflow, the mean is rounded at the scale of the offset, and x − mean carries that rounding into every element, so the variance computed for a constant row equals the square of the rounding, normalizing to plus or minus weight instead of bias. The fix adds an invariant to every fp32 quantity formed from a row: it is bounded only by the row’s spread, not by its offset—implemented by taking “some element in the row” as the baseline and subtracting it, using the warp’s first element for warp-dominant kernels and the row’s first element for tiled kernels. This is a direct fix for the same class of problem as the four Norm operator defects listed under #2812 in the previous issue’s report.

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

3.6 TileOPs: tileops_test target lets tests/ops run without CUDA (#2816)

Merged at 11-02 05:20, the largest single change in this window: 1281 lines added, 524 removed, 59 files, closing #2803. It adds a tileops_test target that exists only for testing, computes shared workload references for each operator on PrivateUse1, and runs the tests/ops smoke tests through it; unsupported reference calls are explicitly listed as unvalidated and skipped, and exclusions that hold only on CUDA are also validated under the same target. The operator-to-workload registry is shared via workloads/calls, while benchmark-specific bindings, copy accounting, and input overrides remain in benchmarks/_cases; it also adds a test device generator, convolution support, corrected test capability markers, and makes the reference implementation honor the scale and empty-row semantics of GQA backward. The author states the limitations prominently: this target tests backend-agnostic contracts and test portability, and comparing against a shared reference does not prove numerical correctness of a real backend; #2611 remains open; calls without an available reference remain visible as unvalidated. The accompanying #2818 separates the dispatch contract from workload data generation, instantiating only shapes for contract cases without generating or validating unused tensor contents, and runs the five host structural test suites only once in compile-contract-gate.

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

3.7 TileFoundry: shared-memory ordering as a provable analysis layer, 52 of 83 fixture plans accepted and 31 rejected (#283)

Merged at 11-02 02:09. The problem is real: finalize does not emit commit, wait, or barrier between shared-memory accesses, so schedules that stage through shared memory race, and on the card sm_89 E4M3 GEMM deviates from torch by up to 659. The direction set by #267 is “an analysis that orders accesses on the HIR” plus “lowering that emits in target order,” emitting only the minimal correct form and rejecting anything it cannot prove; this change delivers the analysis part. The new synchronization family walks each function in lowering’s emission order: for each shared-memory access, it takes the instruction, emission mesh, and operands from analysis/emission.py, the scope from IterationScope, and the byte count from the placement recorded by the memory family; before a conflicting access it plans one commit plus one wait (for outstanding cp.async writes) and one fence covering “the innermost scope holding both,” numbered by the active position of the child node that runs first. Operators declare how their effects complete via the new registry visitor_registry/completion.py, and anything undeclared is treated as UNKNOWN; the family rejects accesses after UNKNOWN, UNKNOWN accesses still in flight on return, fences inserted in loops that do not wait across thread boundaries, and reserved calls whose callee holds shared memory. By default no caller requests this family, so the 75 scheduling fixtures that finalize are byte-identical to main; of the 83 fixtures, 52 are planned and 31 rejected (24 use TMA copies and 6 use wgmma, with completion semantics not yet declared; 1 requires a fence in a loop at a thread-dependent boundary), and corpus cases assert these rejections individually.

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

3.8 TileFoundry under review: CUDA lowering emits the synchronization plan and validates it in reverse, 4 of 7 programs wrong on main (#285)

Updated 11-02 12:40. This is the execution side of #283: CUDA lowering emits steps according to the synchronization plan and checks each actual emission against the plan, throwing LoweringError on mismatch. Results on the corpus: the default flow is 14 programs plus synchronization, with 31 rejected; for these 14, the TIR differs from main only by the added statements, and the other 30 are identical to main; after the core visitor, the 75 programs that finalize on main are byte-identical. On a single sm_89, 7 programs that build and carry synchronization were run: 6 match the reference exactly and are clean under both synccheck and memcheck; lowering these same 7 the way main does yields 4 incorrect results. The author also reports two pre-existing issues on main: on GEMM and loop_outputs, racecheck reports multiple lanes of the same cp.async writing the same address (copy_async in the HIR does not accept split destinations, so each thread copies the entire tile), and loop_outputs also writes one output to the wrong element.

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

3.9 Adopters: no pushes to FlashQLA or TileKernels in this window (11-03)

QwenLM/FlashQLA last pushed at 10-30 03:02, and deepseek-ai/TileKernels at 10-14 05:49; neither repository had commits in this window. TileKernels has now gone several consecutive issues without a push since 10-14, and all visible adopter activity this issue falls within the two in-house repositories, TileOPs and TileFoundry.

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

IV. Community, Tutorials, and Events

4.1 Ascend synchronization correctness becomes a concentrated bug-report surface: 8 new issues in a single day, systematically filed by two reporters (11-02)

During the window, the main repository added 8 new issues, numbered mainly from #3740 to #3749, all pointing to code generation and synchronization insertion in the two backends AscendC and PTO on Ascend. Seven of them were filed by the same reporter, with an environment of TileLang 0.1.15, PyTorch 2.10.0+cpu, and torch_npu 2.10.0. The theme is behavioral divergence of the same program across the two backends: the generated code disagrees on which cores to launch (AscendC emits mix(1,2), PTO emits vector); with enable_auto_schedule=True, AscendC does not insert synchronization; with auto-scheduling disabled generation succeeds, but with it enabled generation fails with an unclear error message; and the same code that AscendC can generate fails on PTO. What these reports share is that each includes a reproducible prim_func and a line-by-line comparison of the generated code, making them directly convertible into regression cases. Together with the in-review #3735 (proving modular index disjointness in thread sync) and #3580 (predicate and hoisted-index synchronization), synchronization correctness on the Ascend side has become the most concentrated line of work at present.

Sources: https://github.com/tile-ai/tilelang/issues/3747, https://github.com/tile-ai/tilelang/issues/3745

4.2 SIMD loads of cross-row and transposed views lose their DMA dependency, with the root cause in logical region reconstruction (#3749)

The most technically substantive issue in this batch, reported with an environment of main branch e8225d7f5, Ascend950PR, and CANN 9.2.0. The problem is that two consecutive SIMD loads lose their dependency on an earlier GM-to-UB copy: first, a row-major load that starts in the middle of a row and crosses into the next row; second, an access to a transposed view where the physically contiguous dimension is not the last logical dimension. The root-cause path is described in detail: the generated tl.access_ptr covers the physical load range, but the logical region reconstructed by MemoryAccessDetector::LogicalPointerRegion misses part of it, so dependency analysis concludes that the input DMA is unrelated to the load, AutoSchedule moves the input copy after the load, and InsertSync also emits no MTE2-to-V set/wait. In the minimal reproduction, work.shape is (2, 128); first T.copy(A, work[1, 0:64]), then inside T.SimdVF() T.simd.vld(work[0, 64]) reads 128 contiguous half-precision elements: the generated read descriptor is tl.access_ptr(work[0, 64], 128, 1), physically covering [64, 192), of which [128, 192) is exactly the part just written by the DMA, while the reconstructed region is work[0:1, 0:128], i.e. physical [0, 128), entirely missing the written interval; the row count is computed as ceil(128/128) = 1 row, and the starting column 64 is not counted. This sits on the same foundation that this window’s main work on FlashMLA and PTO vectorization is shoring up, and when it hits, it silently produces wrong results.

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

4.3 News side: no in-window hits across Google News (Chinese and English), Hacker News, or arXiv (11-03)

Six groups of Google News queries (English: TileLang, TileLang GPU kernel, tile-ai TileLang; Chinese: TileLang 算子, TileLang 昇腾, 沐曦 or 海光 or 摩尔线程) returned the latest relevant coverage as the 10-21 SemiAnalysis article on Huawei and DeepSeek, with the rest concentrated in the 09-29 to 10-01 batch on DeepSeek’s open-source Ascend components, all earlier than this window. On arXiv, the latest hit for all:"TileLang" is TileSight from 2026-07-24, and the combined query for TileLang and Ascend returns nothing; among Hacker News candidates, the only TileLang-related items are repository push entries. Of the 14 candidates, 12 came from GitHub repository discovery and 2 were an unrelated Show HN and an LSP repository, so this issue’s information surface consists entirely of the repository side.

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

V. Trend Observations

The most noteworthy judgment of the day is this: Ascend is no longer an “adaptation target” but a “first-launch target.” 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 prepared for explicit scheduling such as T.Stage and annotate_buffer_addresses, rather than opening a separate repository for porting. Together with the three compiler-side supports (preserving NZ from L0C to UB, selecting the AIV stream before the pipeline, and emitting explicit stages in source order), it is clear that the Ascend path is moving from “it runs” to “it runs according to scheduling intent.”

The second thread is that the center of gravity of correctness is shifting downward. The previous issue reverted a whole set of ThreadSync correctness improvements for the TIRX migration, and this issue immediately produced two follow-ups: mainline #3752 tunes z3 solver resource limits, and #3735 proposes using modular index disjointness to reduce the shapes that need to be proven; meanwhile #3749 digs a defect where “dependency analysis cannot see the physical load range” down to the level of MemoryAccessDetector::LogicalPointerRegion. Putting these together with the hidden thread from TileFoundry—where the analysis layer can prove shared-memory ordering for 52 of 83 fixtures, while doing lowering the main way would make 4 of 7 programs compute incorrectly—the same signal appears three times: the main risk for TileLang now is not insufficient functionality but “the generated code looks right but is actually wrong,” and the team’s chosen approach is to turn correctness into refutable proofs rather than plugging holes one patch at a time.

The third is that testing and contracts are beginning to form an independent layer. TileOPs’ tileops_test decouples smoke tests from CUDA, and compile-contract-gate separates the dispatch contract from workload data generation; both were merged within the same hour, and the author explicitly noted that “a shared reference cannot prove a real backend correct.” This habit of stating separately “what can be tested” and “what has been proven” says more about the project emerging from its prototype phase than merging a few more performance patches.

Finally, note a contrast: on the operator side, 7 performance merges landed in a single day, 6 of them with H200 measured tables, while on the news side all three channels had zero hits in 24 hours. TileLang’s external visibility remains concentrated on events such as “DeepSeek open-sources Ascend components,” and day-to-day progress happens almost entirely inside the repository. For anyone tracking this project, that means the body of the daily report can only be commits and PRs, not coverage.

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; 8 new issues opened; 5 automated doc-site builds
TileOPs 7 merges to main branch; 1 site build
TileFoundry 1 merge to main branch; 1 new PR under review
tilelang-ascend 4 merges to main branch; 6 PR updates under review
tilelang-hygon 1 merge
tilelang-lsp 1 merge (aligns CUDA debug line mapping)
TileOPs-nightly 1 automated push
tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack No pushes within the window
Release pages (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) No new versions within the window
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