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

Sources: GitHub tile-ai organization per-repo check (29 repos, 8 pushed within the window) + merged and in-review ground-truth search (12 merged, 42 in-review updates) + 4 new issues + per-commit check of main branches (tilelang 3, TileOPs 7, TileFoundry 1, tilelang-ascend 1, mlir-ascend 2) + non-default branch check (4 commits on the feat/hcu-scaled-cast branch of tilelang-hygon) + release page and tag check (5 repos, no new versions within the window) + two adopter repos + four Google News queries in Chinese and English (via proxy) + Hacker News + arXiv.

This Issue’s Index

  • Today’s Focus: An Ascend PTO framework change merged via a 10-14 PR (#3512) has entered a concentrated cleanup period, with five mutually independent patches appearing in a single day; TileOPs merged 7 performance changes in one day, the densest day since the holiday; all three merges in the main repo relate to “measurement methodology,” with L2 flush changed from zero-write to read-only
  • I. Core Project Progress
    • 1.1 Mainline: CuTeDSL’s TMA architecture gate accepts sm_103a and sm_107a, six test cases restored on B300 and gr100 (#3772) (11-04)
    • 1.2 Mainline: L2 flush changed from zero-write to read-only reduction, do_bench drops the fast_flush parameter (#3534) (11-04)
    • 1.3 Mainline: AscendC CI’s autotuning supply test case explicitly synchronizes two devices, async copy errors no longer masquerade as arange failures (#3784) (11-04)
    • 1.4 Under review: ThreadSync’s warp consistency proof adds a composite predicate assembled from warp indices (#3787) (11-04)
    • 1.5 Under review: T.copy destination region out-of-bounds now reports a compile-time error, the containment proof for full-copy fragments adds thread-ownership validation (#3786 / #3788) (11-04)
    • 1.6 Under review: SM120’s FP8 GEMM switches to block-scaled MMA, 8192 cubed goes from 463 to 727 TFLOPS on RTX 5090 (#3790) (11-04)
    • 1.7 Under review: Profiler and autotuning methodology validation submitted as a group, six PRs in one day (#3797 / #3798 / #3799 / #3716 / #3717 / #3718) (11-04)
  • II. Multi-Backend Adaptation
    • 2.1 Ascend PTO: One architecture change triggers four categories of regressions, five mutually independent patches submitted the same day (#3792 / #3793 / #3794 / #3795 / #3796) (11-04)
    • 2.2 Ascend: Cube-to-Vector core ratio becomes configurable, no longer fixed at 1:2 (#3783) (11-04)
    • 2.3 Ascend: AutoSchedule’s cycle lower bound tightened, Z3 solve time drops 10.04% (#3776) (11-04)
    • 2.4 Ascend: AI Core exception dump compatible with CANN 9.2.0, version threshold lowered from 9.3 (#1956) (11-04)
    • 2.5 Hygon: Four commits on the feat/hcu-scaled-cast branch, E8M0 scaled conversion and 64-bank LDS read planning (11-04)
    • 2.6 mlir-ascend: BishengIR toolchain adapter migrated into contrib (#209) (11-04)
    • 2.7 MetaX and Moore Threads: no pushes within the window (11-04)
  • III. Ecosystem and Adopters
    • 3.1 TileOPs: 7 merges in one day, GQA varlen packing, ArgMax/ArgMin eight vectors per thread, default cache policy (11-04)
    • 3.2 TileOPs under review: unary elementwise single-wave launch, FP8 grouped GEMM contiguous layout, GEMM optional bias (#2829 / #2824 / #2823) (11-04)
    • 3.3 TileFoundry: FSDP and HSDP training semantics enter the HIR evaluator (#284) (11-04)
    • 3.4 TileFoundry under review: wgmma synchronization plan (#286), plus a separate issue for naming unification (#287) (11-04)
    • 3.5 Adopters and third parties: no pushes from TileKernels and FlashQLA, three third-party repos active within 24 hours (11-04)
  • IV. Community, Tutorials, and Events
    • 4.1 4 new issues: proof for warp size 64, PTO signed atomics, CuTeDSL’s NVFP4 quantizer (11-04)
    • 4.2 News side: four query groups, no in-window hits on HN or arXiv (11-04)
  • V. Trend Observations
  • Appendix: Sources and Verification Notes

Today’s Highlights

The main thread today is not “what got merged again,” but “yesterday’s foundation is starting to be validated.” #3512, merged on 10-14, vectorized T.Parallel into PTO’s VMI pipeline—a structural change to the Ascend backend: integer types moved from unsigned-semantics i* to explicitly signed si*, SIMD lowering was taken over by a new three-stage PTO flow of validation, legalization, and vectorization, and the determination of explicit SIMD parallel loops was rewritten. Three weeks later, this change is being tripped up one by one on real operator shapes, surfacing four classes of regressions: integer atomic operations fail to compile because pto.si32 is not accepted (#3789 report, #3792 fix), T.clear inside T.SimdVF hits a broadcast path that only supports floats (#3795), scalar cast boundary facts are bound too early in VMI memory legalization (#3793), explicit SIMD parallel loops are rejected as scalar lane accesses (#3796), and there is an additional patch adapting integer atomics to an unsigned carrier (#3794). All five patches are based on main and independent of each other; the author repeatedly writes the same line in the descriptions—”This PR is based directly on main and does not depend on the other PTO regression fixes”—and this is deliberate: each of the five can be merged independently, avoiding the past practice of accumulating a giant rollback bundle. Validation is no longer “it runs,” but hardware testing with specific case counts: 47, 18, 16, and 37 items, with four independent operators (SwiGLU, Top-k, Cache, Engram backward) as representative cases.

The second item is a self-correction of measurement methodology, and two repos did the same thing on the same day. Main repo #3534 changed do_bench’s L2 flush from “writing zeros into the buffer” to “reading through the buffer”: writing zeros leaves entire dirty lines in L2, and those lines must be written back during the next timing round, effectively folding the cost of flushing the cache into the measured kernel. #3534 uses cache.view(-1, 4096).sum(dim=1) for chunked reads and also removes the public fast_flush parameter—that parameter only chose whether to build the buffer with int32 or int8. The Ascend-side msprof backend changes more thoroughly: previously a SimtVF loop wrote zeros into UB and then stored via WTS_FV; now it reads from global memory into UB via NORMAL_FV, with the type fixed to int32. TileOPs’ #2822 is a twin change along the same lines, and additionally records the flush method, comparing against historical readings only when the method matches. Taken together, these three changes are a signal: performance numbers are starting to require “cross-version comparability,” not just “faster than last time.”

The third item is catching up on the hardware frontier. B300 reports a target architecture of sm_103a, but CuTeDSL’s TMA helper functions block the architecture with a whitelist, so six CuTeDSL test cases on B300 throw ValueError before even entering the DSL; #3772 adds sm_103a together with the sm_107a reported by gr100 (cc 10.7) to the whitelist, and the author reproduced 8 passed and 1 skipped on a B300 node, and 13 passed and 1 skipped on gr100. In the same window, #3790 takes another path: FP8 plus fp32 accumulation GEMM on SM120 had been using SM89’s mma.sync atomics, reaching only 68% of cuBLASLt on an RTX 5090; after switching to block-scaled MMA with ue8m0 unit scaling, the 8192 cube rose from 463 to 727 TFLOPS, and the output is bitwise identical to the old atomics. The two changes are different in nature—one is “let the new card run,” the other is “let the new card run at full speed”—but both point to the same thing: the boundary of architecture support is expanding from Hopper toward both Blackwell and domestic accelerators at the same time.

The fourth item is left to TileOPs’ commit cadence. Of the 7 merges in a single day, 5 are performance, 1 is CI, and 1 is documentation: GQA varlen packs requests by KV group on SM90, reducing key/value reads for short requests to once per KV head; ArgMax/ArgMin, provided they can reside in one wave, let each thread hold eight 16-byte vectors, and bf16 ArgMin on (2048, 4096) drops from 8.704 microseconds to 7.393 microseconds, faster than FlagGems’ 8.543; unary elementwise and normalization kernels revert operands that are “read only once” to the default cache policy, retaining evict-last/evict-first only where reads are genuinely repeated. What these changes share is not “faster,” but “knowing why it is fast”—each comes with a comparison table and marks the row where it still trails the comparison target (Softmax’s (32768, 4096) regresses from 121.792 to 131.648 microseconds).

I. Core Project Progress

1.1 Mainline: CuTeDSL’s TMA architecture gate accepts sm_103a and sm_107a, six test cases restored on B300 and gr100 (#3772)

Merged on 11-04 at 12:46 (Beijing time, same below), the earliest merge in this issue. The target architecture reported by B300 is sm_103a, while CuTeDSL’s TMA helper functions maintain an architecture whitelist and throw an error before entering the DSL whenever they encounter a value not in the table; the error message lists the allowed values as sm_90, sm_90a, sm_100a, sm_110, sm_120, sm_120a, sm_121, sm_121a. On mainline at fa69f7fa, with nvidia-cutlass-dsl 4.7.0 and CUDA 13.3, six test cases failed on B300 as a result: three on the tma_load path in test_tilelang_jit_cutedsl.py (GEMM JIT, multi-stream kernel, and dynamic shape, one each), plus three tmem_copy_roundtrip variants on the tma_store path in another tcgen05 CuTeDSL test file. The DSL itself accepts sm_103a—its tcgen05 operations accept the entire sm_100f family—so the problem lies purely in this gate. With this addition, the two files report 8 passed and 1 skipped on B300 (cc 10.3, CUDA 13.3, cutlass-dsl 4.8.1). The same PR also adds sm_107a, because gr100’s bring-up board (cc 10.7) reports this value and would hit the same gate; with both added, the same files report 13 passed and 1 skipped on gr100, including the gather/scatter file introduced by #3685.

Worth noting is the verification boundary the author disclosed: no sm_103 hardware was on hand, and the change was verified on a B300 machine provided by another reporter (0z5a, who first reported and bisected in #3685); CI does not cover sm_103a. Such cases of “code merged, verification borrowed hardware” will become more common during the architecture catch-up period, and readers need to know which capabilities have empirical backing and which were merely added to the whitelist.

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

1.2 Mainline: L2 flush changed from zero-write to read-only reduction, fast_flush parameter removed from do_bench (#3534)

Merged on 11-04 at 11:27, changing 7 files, adding 36 lines, and deleting 49 lines, of which three are real implementation changes and four are tests updated accordingly. Previously it was cache.zero_(): writing zeros to a cache_size MB buffer to evict L2. The problem was that after zeroing, the buffer is full of dirty lines, and when the next round of the benchmarked kernel executes, these lines must be written back to memory, so the write-back bandwidth is counted into the benchmarked kernel’s time—the larger the buffer, the more pronounced this bias. Now it is changed to cache.view(-1, 4096).sum(dim=1, dtype=cache.dtype): reading in chunks of 4096 columns, pulling rows into L2 and producing a reduction output, so the dirty-line problem does not exist; chunking avoids the global synchronization of a whole-block reduction; retaining dtype avoids triggering a full-buffer int64 conversion. This change lands simultaneously in the estimate loop of bench.py, the event, cupti, and cudagraph paths of torch_bench.py, and the Ascend-specific flush kernel in msprof.py—the latter previously used SimtVF to write zeros into UB and then store them back to global with WTS_FV, a typical write-style flush; it now reads from global into UB with NORMAL_FV, and the element type is fixed to int32 instead of choosing between int8/int32.

A public interface change to note: the fast_flush parameter of do_bench is removed. It originally only switched between int32 and int8 buffer element types (int32 builds the buffer faster); now the buffer is fixed to int32, so the parameter is meaningless. For scripts that directly call do_bench(..., fast_flush=False), this is a breaking change; tests in the repository have been updated accordingly, and external users must be reminded via release notes.

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

1.3 Mainline: AscendC CI autotuning supply case explicitly synchronizes two devices, async copy errors no longer masquerade as arange failures (#3784)

Merged on 11-04 at 17:01. The symptom was that in the AscendC CI job, the [other_npu] parameterization in the autotuning set-inputs case reported srclen=4, dstlen=4 on the next call to torch.arange()—a 4-byte memcpy size mismatch, exactly the shape of a cross-device copy probe. The real cause was that a cross-device copy failure was reported asynchronously, so the error surfaced one step later than where it occurred, making the error point look like a completely unrelated arange. The fix is straightforward: add a device synchronization before and after the probe (current device and npu:other), so the async error falls back within the scope of try ... except RuntimeError. Verification was running that case file on a single machine with two cards, all four items passing, and when the cross-device copy genuinely fails asynchronously, the [other_npu] parameterization is skipped rather than failing the entire case.

This PR itself is small, but it picks up the tail of #3546 from the previous issue—that PR changed autotuning input supply to resolve by current device, and acceptance mentioned a regression on dual-card Ascend; this one supplements stability of the same scenario on CI. The same feature being fixed once in each of two consecutive windows shows that autotuning on the Ascend side is still in a “runs but boundaries are unstable” state.

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

1.4 Under review: ThreadSync’s warp-consistency proof gains composite predicates built from warp indices (#3787)

This PR takes over the same area left by #3761 and #3765 from the previous issue: ThreadSync must prove that “barrier-participating threads are consistent across the entire warp,” and refuses lowering if it cannot. The triggering shape is a composite predicate tx < 160 and tx // 32 % 2 == 0 under 192 threads—warps 0, 2, and 4 participate as whole groups, warps 1, 3, and 5 do not participate as whole groups, and no warp is split, so it should be provable, yet mainline directly throws the internal error “participating threads are not warp-consistent.” The author built a comparison table that pins down the defect scope precisely: the single predicates tx // 32 % 2 == 0, tx % 64 < 32, and tx < 160 all lower on mainline; only the form “boundary term and warp-index modulo term joined with and” fails; while tx < 168 (which genuinely splits the 5th warp) is correctly rejected on both sides. The root cause is in the prover’s encoding: it renames the thread ID into two hypothetical threads and asks the arithmetic solver to prove that linear1 / warp_size == linear2 / warp_size implies the predicates on both sides are equal, but the encoding of FloorDiv and FloorMod cannot establish the relationship between the two divisions under the combination of “upper bound + modulo.” The fix rewrites threadIdx.x as warp_index * warp_size + lane, with the two hypothetical threads sharing the warp index and each having a free lane—exactly the “same warp” relationship; thus predicates that read the thread ID only through the warp index collapse to the same expression on both sides and consistency is provable, while predicates that genuinely depend on lane still differ, and tx < 168 listed in #3786 is still rejected—the proof is not relaxed, only a missed class is added.

The same author also submitted #3786 on the same day: when the destination region of T.copy exceeds the buffer, it now reports a compile-time error instead of generating an out-of-bounds access. Together the two PRs share the same orientation—turning “the compiler silently generates wrong code” into “the compiler says no.”

Source: https://github.com/tile-ai/tilelang/pull/3787, https://github.com/tile-ai/tilelang/pull/3786

1.5 Under review: Containment proof for fully replicated fragments gains thread-ownership validation (#3788)

When determining whether “a large fragment covers a small fragment,” ProveFragmentContains has a fast path: as long as the large fragment is fully replicated, it returns true directly. This path returns before checking thread ownership, leading to misjudgments such as: the small fragment owns threads 0 to 127, the large fragment is fully replicated but owns only threads 0 to 95, so 96 to 127 are not covered at all, yet the function says containment holds; with check_forward_index enabled, two fully replicated fragments with different physical slots are also accepted. The consequence is that layout inference treats incompatible mappings as compatible—a “silently letting errors through” type, which the reporter found incidentally while investigating #2399. The fix first looks at the physical index, then converts the small fragment’s thread mapping into the large fragment’s coordinate system (accounting for bounded thread-range offsets), and uses the arithmetic analyzer to prove it falls within the replication range and thread range, returning false if it cannot. Verification was on a fresh build of upstream mainline on A100, calling the real C++ implementation directly through a test-time-compiled bridge, covering rejection of 128 vs 96 threads, acceptance of 96 vs 128 and 96 vs 96, physical slots, non-contiguous ownership, bounded ranges, symbolic indices, and partial replication.

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

1.6 Under review: SM120 FP8 GEMM switches to block-scaled MMA, 8192-cube on RTX 5090 goes from 463 to 727 TFLOPS (#3790)

This is the most substantial performance item of this issue. The problem lies in instruction selection: on SM120, FP8 with fp32 accumulation T.gemm reaches only about 465 TFLOPS on RTX 5090 at the 8192-cube shape, while cuBLASLt reaches 681; tuning tiles and pipeline stages does not help, because the bottleneck is the instruction itself—TileLang emits the SM89-era mma.sync ... f32.e4m3.e4m3.f32, and this atom is only half-speed on this card. The author ran a set of register-only probes: e4m3 to f32 via mma.sync is 512 TFLOPS, while via kind::mxf8f6f4.block_scale (unit scaling) it is 1021, and via mma.sync with f16 accumulation it is 1022—the gap is entirely in the instruction. Since block-scaled and ordinary MMA produce identical results when the ue8m0 scale takes 0x7f (i.e., 2 to the power of 0), it is possible, when TL_SM120_BLOCK_SCALED_MMA_ENABLED is on and the target is sm_120a/f or sm_121a/f, to route the four fp32-accumulation FP8 dispatch paths through the block-scaled atom, while other targets, f16 accumulation, and explicit sm_120 without the a suffix remain unchanged. Bitwise consistency was verified at the instruction level: all four combinations of e4m3 and e5m2, with random register contents including NaN and Inf encodings.

Numbers: under e4m3 GEMM, bf16 output, 128×128×128 tile, 3 pipeline stages, 256 threads, 4096-cube goes from 408 to 673 (cuBLASLt 530), 8192-cube from 463 to 727 (681), 4096×14336×4096 from 449 to 696 (498), and 256×8192×8192 from 348 to 606 (457). The only slowdown is the shape with M=16 and a 128-row tile, from 31.5 microseconds to 33.5 microseconds; the author notes this shape is bandwidth-limited and would be faster with a 16- or 64-row tile. New test cases check PTX directly: sm_120a must get the block-scaled atom, and sm_89 must get the ordinary atom.

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

1.7 Under review: Profiler and autotuning consistency checks submitted as a group, six PRs in one day (#3797 / #3798 / #3799 / #3716 / #3717 / #3718)

All six come from the same author (Afloat16), with the theme of aligning “the validator’s semantics with the object under test.” assert_allclose now rejects a kernel returning None where a tensor is expected—the original implementation skipped comparison whenever either side was None, so a kernel returning nothing could “pass” (#3797); assert_consistent previously did zip directly on return values, so scalar tensors threw TypeError, a missing or extra output was silently ignored, and differing shapes could be glossed over by truncating rows or broadcasting; it now normalizes to an output list first and then checks count and shape, including empty tensors (#3798); suppress_stdout_stderr in the CUPTI benchmark previously called fileno() on Python streams, crashing in notebooks and with StringIO where there is no file descriptor; it now redirects native descriptors 1/2 and Python streams separately, using ExitStack to guarantee restoration even on exception paths (#3799). The other three are on the autotuning side: when the target dict is used as a cache key, its string representation meant equivalent dicts with different insertion orders produced different keys; it now serializes as sorted JSON (#3716); Tensor Core detection now prefers explicit compute capability, then explicit or current target architecture, and only last checks the local device, so cross-compilation without a GPU no longer falsely reports unsupported (#3718); Profiler’s dynamic shape substitution now handles composite dimensions like m + 1 and raises ValueError instead of TypeError for missing variables (#3717).

The technical weight of this group is not in the changes themselves but in the shape they expose: validators have long been a weak link of “passing means correct,” and once infrastructure like autotuning cache keys and target resolution drifts, it shows up as “results not reproducible,” which is very hard to attribute.

Source: https://github.com/tile-ai/tilelang/pull/3797, https://github.com/tile-ai/tilelang/pull/3798, https://github.com/tile-ai/tilelang/pull/3799

II. Multi-Backend Adaptation

2.1 Ascend PTO: One Architectural Change Leads to Four Classes of Regressions, Five Mutually Independent Patches Submitted the Same Day (#3792 / #3793 / #3794 / #3795 / #3796)

First, the background. #3512 (merged 10-14, [Ascend][PTO] Vectorize T.Parallel to PTO VMI) migrated vectorization of parallel loops from the old SIMD channel to the PTO VMI pipeline: integer types changed from unsigned i* to explicitly signed si*, Lowering goes through the three stages Verify, Legalize, Vectorize, and the detection of explicit SIMD loops was reimplemented. Each of these five patches corresponds to one place where it tripped over a real operator:

  • #3792 Integer atomic compilation failure: codegen_pto emits pto.si32 / pto.si64 for integers, but PTO’s atomic interface only accepts unsigned i32 / i64, so T.atomic_add, T.atomic_max, and T.atomic_min fail to compile directly on integer buffers (corresponding to issue #3789). The fix wraps tl.atomic_* in a compatibility layer: the pointer element type and operands are rewritten to the same-width unsigned type, integer min/max are passed the correct signedness, and the return value is converted back to the original type. Test cases expanded from 15 to 18, covering signed and unsigned 32/64-bit add/subtract/extremum, negative operands, previous-value return, and counters exceeding 32 bits.
  • #3794 Signed carrier adaptation: Same origin as #3792 but takes a different path—instead of changing the semantics of tl.atomic_*, it rewrites the pointer view and operands at the generated PTODSL call boundary, preserving the address space, with floating-point atomics keeping the original path. The representative case is SwiGLU’s clamp counting (bf16, hidden 2048, clamp 0.5, psum layout); all 18 atomic test cases pass.
  • #3795 SIMD fill lost: int32 T.clear inside T.SimdVF fell into a broadcast path that only supports floating point, because the new PTO pipeline no longer runs the old SIMD pass, and that pass was responsible for both fill and parallel loops. The fix extracts a fill-only entry point that reuses the existing LowerFill, placed after buffer version normalization and before PTO parallel verification, still emitting the original vdup/vsts sequence, and retaining the checks for contiguous regions, alignment, full vectors, and multi-buffer version slicing. 16 test cases pass; the representative case is Top-k (512 tokens, 1 expert, top-1).
  • #3793 Boundary facts of scalar casts bound too early: Scalars unrelated to VMI addresses such as u = uint32(i - 1) with i in [0, 63] throw during memory legalization; eagerly calling Bind also leaves an underestimated interval that throws later during IntSet’s cast analysis. The fix checks the representability of the integer cast before installing facts, giving the proof engine a safe scalar definition and IntSet a conservative numeric interval, avoiding directly evaluating cast endpoints. 47 test cases pass; the representative case is Cache (block_size=64, n=65536).
  • #3796 Explicit SIMD parallel loops treated as scalar lanes: Engram backward’s initialization uses a Parallel nesting in which explicit SIMD stores advance by full-vector blocks, and the new lane layout inference treats them as scalar accesses and rejects them (M=2, N=32, L=128). The fix shares the classification of explicit SIMD between layout inference and parallel unit discovery; such nestings retain the generic layout path from before #3512, while automatic element-wise units still take the new PTO lane path. 37 test cases pass; the representative case is Engram backward (512 tokens, hc 4, hidden 2048). The author also notes a limitation: the explicit-plus-automatic mixed case with FP32/128 lanes only has generated source-level coverage; the current PTOAS/Bisheng device compiler exceeds the 6 KiB VF stack budget on this combination, so there are no device run results.

Looking at these five side by side reveals the engineering rhythm on the Ascend side: structural changes like #3512 land first, then boundaries are exposed one by one against real operators, with each patch carrying hardware measurements and test counts and deliberately kept independent of the others. The “merge the framework first, clean up regressions later” playbook was already used in the previous architecture migration (#1975) and has now become standard procedure—at the cost of the mainline sitting in a “known to have several broken spots” state during such windows, with a non-trivial chance of external users hitting them.

Source: https://github.com/tile-ai/tilelang/pull/3792, https://github.com/tile-ai/tilelang/pull/3793, https://github.com/tile-ai/tilelang/pull/3794, https://github.com/tile-ai/tilelang/pull/3795, https://github.com/tile-ai/tilelang/pull/3796

2.2 Ascend: Cube-to-Vector Core Ratio Configurable, No Longer Fixed at 1:2 (#3783)

Ascend’s AI Core is split into Cube and Vector types; AscendC supports groupings such as mix(1,2), mix(1,1), mix(1,0), and mix(0,1), whereas TileLang previously always launched at 1:2, effectively locking down the upper bound of a hardware degree of freedom and making it hard to adapt to future chips with different core ratios. This change turns the core ratio into an explicit or device-derived parameter: T.MixedKernel(..., sids=1) or sids=2 follows the requested group size with capability validation; when omitted or passed None, it uses the current device’s vector_core_num / cube_core_num, falling back to 2 if device information is unavailable; T.Kernel keeps its original interface and core-type determination, using the prepared vector core count when recognized as a mixed kernel, with the single-core path unchanged and the grouping semantics of explicit T.Vector(1/2) preserved. The selected group count is carried all the way through the sid domain, task assignment, synchronization, and code generation, avoiding repeated device queries and repeated normalization; conflicting declarations and requests exceeding device capability are rejected outright.

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

2.3 Ascend: AutoSchedule’s Cycle Lower Bound Tightened, Z3 Solve Time Down 10.04% (#3776)

AutoSchedule uses binary search to find the pipeline’s initiation interval (II); the original lower bound took the maximum of all task IIs, which was too loose and required extra binary search rounds. This change adds three necessary conditions under the resource model: a single-task bound (retaining P >= max_i II_i), a resource-clique bound (for cliques of pairwise resource conflicts, sorting by residual and summing to obtain P >= sum_i II_i), and a self-dependency bound (S_u - S_u >= wait - d*P yields P >= ceil(wait/d), with version-dependent distances truncated by pipeline stage count). The binary search is initialized with the maximum of the three, with all other logic unchanged. The test environment is Ascend 950PR with the CANN 9.3.0 weekly build, caches disabled, one warm-up and three alternating measurements: all of examples/ascend compiles and scheduling results are unchanged, and cumulative Z3 solve-phase time drops 10.04%. Changes like this—”don’t change results, only change the search starting point”—carry the lowest risk yet deliver global benefit: compile-time solve time is something every user pays.

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

2.4 Ascend: AI Core Exception Dump Compatible with CANN 9.2.0, Version Threshold Lowered from 9.3 (#1956)

This merged into the ascendc_pto branch of the tilelang-ascend repo (11-04 10:18) and is the only merge in the Ascend repo this period. The AI Core exception dump feature previously declared a dependency on CANN 9.3.0 or above; the author verified feasibility on 9.2.0 and fixed three things: filling in the required fields of acldumpTensorInfo (addrType=ACL_DUMP_ADDR_RAW, placement, shapeNum, originShapeNum), because 9.2.0 strictly validates addrType and the original implementation’s ACL_DUMP_ADDR_PTR obtained by zeroing via calloc is rejected; skipping null-address and zero-size tensors, since 9.2.0 does not accept a NULL tensorAddr; and switching the registration interface from the deprecated aclrtSetExceptionInfoCallback to aclrtExceptionInfoCallbackRegister. On the tool side, parse_exception_dump() supports both the .npy format output by 9.2.0 and the older .bin format, and the version requirements in the Chinese and English tutorials are updated to CANN 9.2.0 or above. For users still on 9.2, this is a change from “feature unavailable” to “available.”

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

2.5 Hygon: Four Commits on the feat/hcu-scaled-cast Branch, E8M0 Scaled Conversion and 64-bank LDS Read Planning (11-04)

The Hygon repo’s main branch stopped at 11-02; all activity in this window is on the feature branch, with four commits concentrated within eight hours (16:26 to 20:16 Beijing time), authored by Hygon contributors (huangteng@hygon.cn and zhouweiwang@hygon.cn):

  • feat(hcu): add scaled_cast and unify GEMM LDS read planning: Adds E8M0-scaled FP8/BF8/FP4 conversion with both scalar and packed lowering, adding src/tl_templates/hcu/scaled_cast.h (223 lines) and lower_scaled_cast.cc; also shares GEMM consumer-side metadata between matrix read selection and LDS strategy, supporting single-level LDS slicing, with 650 lines changed in the main transform file.
  • fix(hcu): align GEMM LDS strategies with 64-bank read phases: Derives plain-copy swizzle from bank-word geometry, provides an explicit inverse transform for the narrow b16 tile shared by matrices and the two native pack sizes; asynchronous consumers with mixed or constrained phases prefer compatible common storage; when selecting native wrap, the full b128 read phase is accounted for. 12 files, 1536 lines added, 112 lines removed, 5 new test files.
  • fix(hcu): prefer mixed LDS wraps proven at most two-way: Adds an “at most two-way” proof for mixed LDS wraps, preferring them only when the proof holds.
  • fix(hcu): use physical read width for mixed wrap checks: Mixed wrap checks now use the physical read width, removing the previous inference of 8 bytes per lane from kPack.

The theme is shared-memory bank conflicts and phase alignment—the critical path for GEMM performance on Hygon DCU, and a sign that this backend has moved from “can compile” to “tuning against the microarchitecture.” All four are on the feature branch with no PR opened yet; timing suggests they may be submitted after a sync with the TileLang mainline.

Source: https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-scaled-cast

2.6 mlir-ascend: BishengIR Toolchain Adapters Moved into contrib (#209)

Merged into the main branch (11-04 16:35), moving the BishengIR toolchain adapter code from its original location into the contrib directory. This is pure structural cleanup: 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. The same branch also carries a main commit refactor(contrib): move BishengIR toolchain adapters into contrib, with the merge commit as the second. Such cleanup has no functional value on its own, but it marks the beginning of managing the “adapter layer” as an independent lifecycle—Ascend toolchain versions iterate quickly, and replaceability of glue code matters more than elegance.

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

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

tile-ai/tilelang-metax last pushed at 10-26 15:08, and tile-ai/tilelang-musa at 09-30 17:50; neither repo has commits in this window, nor updates to PRs under review. Among the four domestic backends, Ascend and Hygon are the two active lines this window, while MetaX and Moore Threads are in a quiet period. MetaX also had no pushes in the previous report, now two consecutive windows; Moore Threads has shown no activity since 09-30, worth watching in subsequent windows to see whether it has shifted to internal development.

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

III. Ecosystem and Adopters

3.1 TileOPs: 7 merges in a single day — GQA varlen packing, ArgMax/ArgMin eight vectors per thread, default cache policy (11-04)

All 7 were merged within the window, making this the densest day of the quarter. Item by item:

#2819 adds a packed warp-specialized program to SM90’s GQAVarlenFwdOp: when a request’s KV head group can fill a tile positionally, requests fill tiles by KV group so each KV head reads key/value only once; whether to pack is decided by a cost model, used only when the “savings guaranteed by shape” exceed the measured 9% per-head overhead; wide and narrow requests are claimed alternately within the same launch, with no host readback and no added synchronization; when the device-counted non-splittable items are fewer than the number of SMs, single-tile packed requests split the key scan into multiple segments and merge via an fp32 workspace. Three mainline defects were also fixed: rows with no visible keys returning NaN when sink plus score scaling exceeds about 2.4, the fallback path reading the RoPE table before the request’s first key, and a negative control leaving the caller’s stack frame in a reference cycle that caused cross-test CUDA graphs to be destroyed early. Test nodes grew from 173 to 197.

#2828 gives row-tiled ArgMax/ArgMin eight 16-byte vectors per thread under a “one-wave residency” condition, and adds 64-thread blocks to the tuning candidates; residency is determined jointly by SM thread count, resident CTA limit, and registers. bf16 ArgMin on (2048, 4096) drops from 8.704 µs to 7.393 (vs. FlagGems 8.543, 1.155× faster), and back-to-back from 7.808 to 6.000; f16 ArgMax drops from 6.816 to 6.400 (vs. 9.344, 1.46× faster). Of 12 benchmark rows, 4 got faster and 0 got slower.

#2820 reverts operands of “read each element only once” operations to the default cache policy: elementwise, rope, norm, softmax, and reduction kernels no longer uniformly use evict-last/evict-first, retaining it only where reads genuinely repeat (broadcast rows in binary kernels, rope cos/sin rows shared across batches and heads, the two reads of two-pass RMSNorm, the first two passes of fused add). It also adds paired-kernel and staged-kernel implementations for RMSNorm, shrinks the CTA of warp-row Softmax to exactly the warp count that fills the SM, switches 2D neox RoPE to 128 threads, and folds per-architecture resident CTA limits into RESIDENT_CTAS. The cost: Softmax on (32768, 4096) regresses from 121.792 to 131.648 µs (vs. quack 126.080), which the author honestly flags in the table.

#2822 is a read-only change to L2 flushing, sharing provenance with main-repo #3534; 168 tests pass, and it additionally records the flush method, comparing against historical readings only when methods match. #2825 ranks legal W4A16 configurations with a single refitted cost model, has split-K CTAs stage only their own quantization metadata, and accounts for launch headroom in shared-memory bookkeeping. #2826 is companion documentation spelling out the calibration contract for dense and tiny-M (measurement method, fitting constraints, held-out validation), with no runtime changes. #2827 moves the off-CUDA suite to a nightly job: these two preflight steps previously took 13–18 minutes on ubuntu-latest and reported ten minutes later than other checks, while the tests/ops smoke run uses operator reference implementations and checks device-agnosticism and dispatch rather than kernel results; after the move, preflight keeps only collection checks, drops test reports and the cuda_only must-fail check (about 3 minutes), and 1856 plus 2789 test nodes move to nightly, with the release gate still requiring a green nightly.

Sources: https://github.com/tile-ai/TileOPs/pull/2819, https://github.com/tile-ai/TileOPs/pull/2828, https://github.com/tile-ai/TileOPs/pull/2820

3.2 TileOPs under review: unary elementwise single-wave launch, FP8 grouped GEMM contiguous layout, optional GEMM bias (#2829 / #2824 / #2823)

#2829 makes unary elementwise launches whose body contains only one or two single-instruction operations run as a single wave: read/write bytes no more than three quarters of L2, result the same width as the operand, and at most two body operations each reducible to one instruction — all three must hold to shrink to a one-wave 256-thread block; body conditions are judged by node counts over TIR variables, and bodies with casts, division, math calls, guards, or anything not evaluable on bare variables keep the default launch. Currently qualifying are relu, abs, neg, bitwise_not, and logical_not on bool. A real bug was also fixed along the way: ReLU returned 0 for NaN inputs and now returns NaN (torch’s bit pattern for 16-bit, canonical NaN for float32). In the benchmark, three ReLU rows and one LogicalNotFwdOp row got faster with other configurations unchanged; the author also flags that the decode row with flushing on GPU 6 merely breaks even.

#2824 adds contiguous layouts to SM90’s FP8 grouped GEMM (compact and aligned packing, segment-tail or per-row expert metadata), keeps SM89 on masked layout only, and fills in measured workloads skewed by expert size plus a comparison against DeepGEMM’s contiguous layout. #2823 adds an optional bias[N] to GEMMFwdOp, covering the dense, GEMV, and split-K paths, with a cuBLASLt BIAS comparison. Both correspond to gaps listed in #2130, showing TileOPs is advancing along a “align item by item with the torch ecosystem” checklist rather than adding operators at random.

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

3.3 TileFoundry: FSDP and HSDP training semantics enter the HIR evaluator (#284)

Merged 11-04 16:36. The previous evaluator treated reshard as a global tensor, losing the differing contributions of different participants, so training semantics (especially HSDP, which mixes sharding and replication) could not be evaluated correctly at the HIR level. This implements the training subset of #255: per-participant local values, GPU coordinates, calls and loops, gather/scatter, and axis-wise Partial(sum) reductions all go through shared projection rules; Reshard alias proofs become ownership-sensitive, wired to #170’s structured lifetimes, with tuple arguments validated from tensor leaves. A companion modded-nanogpt-inspired two-block GQA study was added: explicit forward/backward, momentum SGD, an independent PyTorch reference, covering FSDP, 2×2 and 4×4 HSDP, and two residency strategies, reporting collective communication groups, payloads, and storage footprint, with reusable SVGs embedded so every table in the README is derivable from code.

Validation: the refactored core suite at 775 items and the backend suite at 778 items (13 skipped, two sm90 cards), with 10 HIR configurations passing on CPU; both residency strategies match real 2×2 PyTorch HSDP on four sm90s (two update steps, comparing local gradients, parameters, and momentum).

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

3.4 TileFoundry under review: wgmma synchronization plan (#286), plus a naming-consistency issue (#287)

#286 addresses WGMMA’s asynchrony: after the instruction is issued it may still read shared memory and registers, so subsequent accesses must be ordered after its effects. The approach declares completion conditions per operand effect, attaches WGMMA’s fence, commit, and wait to the HIR nodes that need them, and lowering consumes these steps at the instruction’s issue point; if the following scope cannot fit a collective wait, the work is completed within the issuing scope. It also fixes aliasing for register Reshard, keeping each element in the same register on every thread holding a replica, with the plan using this storage identity to track dependencies. The author explicitly notes a risk: the matrix-instruction emitter is not yet wired up, so this neither proves WGMMA execution correctness nor provides performance, and TMA completion semantics are handled separately.

#287 is a documentation-oriented issue noting that the codebase has “five names for one concept” (a mesh member is called placement in 133 places, program in 110, instance in 92, participant in 44, plus wave_units-style spellings with unit), while the single word unit simultaneously means an empty type, a hardware compute unit, and a mesh member; Python and the CUDA runtime also disagree on the local_layout signature. Such issues rarely change code immediately, but they are a necessary step for a compiler project going from “writable” to “readable long-term by people.”

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

3.5 Adopters and third parties: no pushes from TileKernels or FlashQLA, three third-party repos active within 24 hours (11-04)

Neither adopter repo had commits in this window: deepseek-ai/TileKernels last at 10-14 05:49, QwenLM/FlashQLA last at 10-30 03:02 (latest release still v0.1.3 from 09-30). Note that adopters not pushing does not mean they are unaffected — the previous issue’s #3761 fixed exactly a barrier-occupancy regression in TileKernels’ mhc_pre_big_fuse kernel on B300, showing compiler changes reach their shapes as “passive regressions.”

On the third-party side, repo discovery surfaced three active items, all created or pushed within 24 hours: 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), and Index4NekoMeteor/Tilelang-AttnRes (1 star, implementing Kimi K3’s attention residual structure in TileLang). All three are small, but the directions are interesting: one works on toolchain-side developer experience, two port to new target architectures and frontier model structures — the kind of community-driven work that extends TileLang’s boundaries.

Sources: https://github.com/deepseek-ai/TileKernels, https://github.com/QwenLM/FlashQLA, https://github.com/sablin39/tilelang-cuda-skills

IV. Community, Tutorials, and Events

4.1 Four new issues: proof for warp size 64, PTO signed atomics, CuTeDSL’s NVFP4 quantizer (11-04)

Three in the main repo. #3791 is “[BUG][ice-on-valid-code]”: ThreadSync can only prove warp-uniform predicates at warp size 32, and the same IR is rejected when thread_warp_size is 64—this issue is tightly coupled with #3787, currently under review in this cycle, indicating that the prover’s assumption about warp size is hardcoded, whereas AMD and some domestic accelerators use 64-wide warps. #3789 reports that the PTO backend emits signed pto.si32/si64 for integer T.atomic_add, causing validation to fail; two fix directions, #3792 and #3794, appeared the same day. #3781 reports that CuTeDSL’s NVFP4 block-scaled quantizer fails to compile: the reinterpret for packed FP4 and missing ceil/floor. Together with TileFoundry’s #287 (naming unification), this window has 4 new issues in total, all of the “internal implementation details leaking into the user-visible surface” type—not one is a performance complaint.

Sources: https://github.com/tile-ai/tilelang/issues/3791, https://github.com/tile-ai/tilelang/issues/3789, https://github.com/tile-ai/tilelang/issues/3781

4.2 News side: four query groups, no in-window hits on HN or arXiv (11-04)

Across Google News’s four query groups (English TileLang GPU kernel, "TileLang" OR "tile-lang"; Chinese TileLang 算子, TileLang 编译器), the most recent hit is 352 hours old (SemiAnalysis’s Huawei and DeepSeek report on 10-21); the rest cluster around the 09-30 wave of coverage on “DeepSeek open-sources Ascend version of TileLang, DeepGEMM, DeepEP” (856 hours ago)—all earlier than this window. A Hacker News search for tilelang over the last 24 hours returns 0 results; the latest arXiv all:tilelang hit is still TileSight from 07-24, with no new papers in the window. Of the 9 candidates surfaced by GitHub repository discovery, 7 are the repository itself, 1 is an unrelated Show HN like Swarm World, and 1 is an unrelated 2048-type project. It is safe to conclude that the news side is zero for this window, and the report body can only come from commits and PRs—this has now been the stable pattern for several consecutive issues.

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

V. Trend Observations

First, the clearing-the-backlog pattern for Ascend PTO has taken shape, and both its benefits and costs are worth recording. #3512 merged on 10-14, and within three weeks real operators kept pushing it past its boundaries; today’s five patches are just the latest batch. This rhythm of “land the architecture first, then clear regressions operator by operator” also appeared during the #1975 whole-repo migration. The upside is fast architectural iteration and controllable individual changes; the downside is that the mainline sits in a “known to have several things not working” state for a considerable window. The five patches deliberately remain independent of one another, showing that the maintainers are consciously reducing rollback granularity—more mature than bundling everything into one big rollback package, but it also means stronger pre-release regression discipline is needed, or “independently mergeable” will become “independently leakable.”

Second, performance measurement is starting to be treated as a product feature. Main repo #3534 and TileOPs #2822 both changed L2 cache flushing from write-zero to read-only on the same day, for the same reason: “dirty-line writeback pollutes the next timing round.” TileOPs goes further, recording the flush method and comparing historical readings only when the method matches. This is not a minor tweak: if a benchmark number is not comparable across versions, then every “how much faster” conclusion has to be recomputed. Bringing the measurement methodology itself under version control is a prerequisite for performance engineering to become trustworthy.

Third, the competitive focus of operator libraries is shifting from “number of operators” to “quality of the selector.” TileOPs’ W4A16 single-refit cost model and the dense GEMM calibration contract document merged today, together with the FP8 grouped GEMM contiguous layout and GEMM bias under review, form a clear line: 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. If this line is done right, operator library performance no longer depends on hand-tuning individual kernels; if done wrong, it degenerates into “some shapes are especially fast, others rely on luck.”

Fourth, the boundaries of architecture support are expanding at both ends. At one end are new cards: B300’s sm_103a and gr100’s sm_107a enter the whitelist, and SM120’s FP8 switches to block-scaled MMA. At the other end are domestic accelerators: Hygon does microarchitecture-level tuning on the 64-bank LDS phase, while Ascend fills in engineering details like core ratios and exception dumps. The workload structures of these two kinds of expansion are completely different—new cards are mainly about “getting the path to work,” while domestic cards are more about “getting the same path to run at near-hardware efficiency.” The latter requires far higher measurement density, which is exactly why test files make up such a large share of Hygon’s four commits.

One final risk dimension: within a single window, 42 PR updates are under review, of which 12 carry the Ascend/PTO label, 7 carry the CUDA/SM/CuTeDSL label, and 4 cluster in Profiler. This structure shows that the repository’s change pressure is shifting from “adding features” to “fixing boundaries.” For users, the good news is that defects are being identified one by one; the bad news is that most of these defects fall in places that are normally invisible—validators, cache keys, provers—and by the time they surface, they have often already affected the credibility of results.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
tile-ai organization (29 repos) 8 repos pushed within the window
Main repo tilelang 3 merges to main; 3 new issues; 38 updates under review (Ascend/PTO 12, CUDA and CuTeDSL 7, Profiler 4)
TileOPs 7 merges to main; 3 updates under review; no new issues
TileFoundry 1 merge to main; 1 update under review; 1 new issue
tilelang-ascend 1 merge to default branch ascendc_pto; the under-review list could not be retrieved due to anonymous API rate limits, so verification relied on main-branch and commit-level atomic sources instead
tilelang-mlir-ascend 2 commits on main (including merge of #209)
tilelang-hygon No new commits on main (last activity 11-02); 4 commits on the feat/hcu-scaled-cast branch
tilelang-metax / tilelang-musa No pushes within the window
TileOPs-nightly Branch pushes only; no new commits on main
Release pages (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) No new versions within the window; tilelang latest remains v0.1.15 from 09-30
FlashQLA / TileKernels No pushes within the window
Google News (four queries in Chinese and English) / Hacker News / arXiv No substantive hits within the window; most recent hit was 352 hours ago

Full Source List

Verification Notes

All merge, under-review, issue, and commit counts in this issue come from actual GitHub search and atomic-source output, with no estimates. The under-review PR lists for the Ascend and Hygon repos could not be retrieved because the anonymous API quota was exhausted; the related conclusions are based only on main-branch commits and branch-level atomic sources, as noted in the source verification table. The four commits in tilelang-hygon are on a feature branch with no PR opened; they are branch-level activity and do not indicate that the mainline has merged them. The three news channels — Google News, Hacker News, and arXiv — had no hits in this window, so the body of the report is drawn entirely from the code side.