Research window: 2026-10-28 07:00 ~ 2026-10-29 07:00 (Beijing time, 24 hours).

Sources: GitHub tile-ai organization per-repo verification (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / docs site) + merged and in-review ground-truth search (24 merged, 48 in-review updates) + full review of new issues (6) + commit-level diff and file-list verification + release page and tag verification + third-party backend repo verification + six Google News queries in Chinese and English (via proxy) + Hacker News + arXiv.

This Issue’s Index

  • Today’s Highlights: The NPUIR backend for Ascend 910 A2/A3 has been merged into the mlir-ascend mainline, with all 504 self-tests passing; the TileOPs selector now calibrates against busy-hour conditions, cutting manifest regret from 5.79% to 0.48%; 24 merges, 48 in-review updates, and 6 new issues in 24 hours
  • I. Core Project Progress
    • 1.1 Main repo: Index coefficients for scalar reads no longer mix into lane counts; compile crash for row widths above 32768 closed (#3453) (10-28)
    • 1.2 Main repo: Batched all-reduce resolves thread ranges via the scalar branch; fixed the issue where all 2048 outputs were wrong when batch>1 (#3341) (10-28)
    • 1.3 Main repo: Shape validation added for SIMT atomic regions; out-of-bounds regions rejected at compile time (#3465) (10-28)
    • 1.4 Main repo: Sparse MMA K tail blocks are no longer silently let through (#3367) (10-28)
    • 1.5 Main repo: atomic_addx4 target dtype narrowed to float* (#3369) (10-28)
    • 1.6 Main repo: Handle suffixes for non-restrict parameters are no longer trimmed (#3366) (10-28)
    • 1.7 Main repo: TMA descriptors record only validated names (#3484) (10-28)
    • 1.8 Main repo: Unified validation for profiler output indices; scalar and list forms share one set of boundary rules (#3486) (10-28)
    • 1.9 Main repo: T.Parallel’s coalesced_width uniformly converted to IntImm (#3374) (10-28)
    • 1.10 Main repo: Test infrastructure routes cache-sensitive cases through without_kernel_cache (#3695 / #3696) (10-28)
    • 1.11 Three new defects: 13 examples fail on SM120, IfStmtBinding compiles exponentially, without_kernel_cache does not force recompilation (#3701 / #3700 / #3699) (10-28)
  • II. Multi-Backend Adaptation (Ascend / Hygon / ROCm)
    • 2.1 Ascend: The NPUIR backend for 910 A2/A3 has been merged into the mlir-ascend mainline; all 504 self-tests pass, and only 1 of 72 examples fails for environmental reasons (#207) (10-28)
    • 2.2 Ascend main repo: Bisheng build identity added to the binary cache key (#3691) (10-28)
    • 2.3 Ascend main repo: Handwritten pipelines preserve configuration register phases (#3624) (10-28)
    • 2.4 Ascend repo in review: Root cause and fix for incorrect W4A4 int4×int4 GEMM results (#1731) (10-28)
    • 2.5 Ascend repo in review: a2a3 scheduler migration, 100 files, 18,453 lines (#1951) (10-28)
    • 2.6 Ascend repo in review: Carver reads A3 resource budgets according to the selected SDK (#1955) (10-28)
    • 2.7 Ascend repo in review: AI Core Exception Dump dependency lowered to CANN 9.2.0 (#1956) (10-28)
    • 2.8 Ascend repo in review: Docs switch to T.Tensor, deprecating the T.Buffer DSL interface (#1957) (10-28)
    • 2.9 Hygon: LDS swizzle feature branch for HCU normal copies, 18 files including 5 new tests (#30) (10-28)
    • 2.10 ROCm: ROCm 10 support for the pip SDK merged; MI355X is no longer misidentified as gfx900 (#3549) (10-28)
  • III. Ecosystem and Adopters (TileOPs / TileFoundry)
    • 3.1 TileOPs: Selector now calibrates against busy-hour conditions, cutting manifest regret from 5.79% to 0.48% (#2784) (10-28)
    • 3.2 TileOPs: Non-NT layouts now use a transposed persistent kernel; 2112×7168×4096 TN drops from 195.3 to 175.3 microseconds (#2789) (10-28)
    • 3.3 TileOPs: cumsum long rows switch to two-stage seed scan; (8,65536) drops from 12.75 to 5.18 microseconds (#2790) (10-28)
    • 3.4 TileOPs: extrema reduction switches to PTX max.NaN; NaN propagation unchanged while regression disappears (#2792) (10-28)
    • 3.5 TileOPs: Two commits on effective row tiles for parallel cumsum and GDN narrow prefill tiling (#2788 / #2786) (10-28)
    • 3.6 New TileOPs issue: Large GEMM rows lag the fastest baseline under back-to-back calls (#2791) (10-28)
    • 3.7 TileFoundry: Ring stage binding and compilation flow unified, 70 files (#277) (10-28)
    • 3.8 TileOPs in review: Release image bundles TileOPs 0.1.0 with TileFoundry 0.0.3 (#2793) (10-28)
  • IV. Community, Tutorials, and Events
    • 4.1 Two automated updates to the documentation site (10-28)
    • 4.2 Batch of Ascend repo doc fact-checks and installation fixes in review, seven items advancing together (10-28)
    • 4.3 New operator examples in the Ascend repo in review: GroupNorm and fla_chunk (#1607 / #1953 / #1952) (10-28)
    • 4.4 Third party: Community auto-tuning agent refreshes Chinese documentation (#74); the Hexagon backend repo’s main branch has no new commits in the window (10-28)
    • 4.5 News side: Six groups of Chinese and English Google News queries, Hacker News, and arXiv all had no substantive hits in the window (10-28)
  • V. Trend Observations
  • Appendix: Sources and Verification Notes

Today’s Highlights

Today’s window centers on the Ascend side. tilelang-mlir-ascend merges the entire NPUIR backend for Ascend 910 (A2/A3) into the mainline in a single commit touching 55 files with 21,490 added lines: on the Python side, the tilelang/ascend910 package (backend / adapter / codegen / compiler / execution_backend / jit / pipeline / runtime / target, plus the kernel, copy, gemm, and npuir dialect surfaces under language); on the C++ side, the developer and expert emitters under src/ascend910/target, the op layer, and three passes — LowerNpuirBlock, NpuLoopVectorize, and RewriteFDivToMul; the runtime gains npu_launcher.h and npu_utils; installation goes through install_npuir.sh; and the target alias accepts ascend910 / ascend-910 / a2 / a3 / npuir. The author’s acceptance criteria are the repo’s two bundled test suites: all 504 pytest testing/npuir cases pass, and 71 of 72 examples pass with 0 timeouts and 0 errors — the sole failure, deepseek_v32/fp8_lighting_indexer.py, fails identically on the old backend on the same machine and is an environment and tolerance issue. A second commit aligns behavior differences with the old backend item by item, focusing on bf16 legalization (HIVM does not accept bf16 operands, so computation is done in fp32 with explicit conversion) and copy sinking for dynamic extents.

TileOPs merged seven commits today, with both main threads on the performance side. The selector no longer fits against the manifest; instead it is recalibrated on busy time (device_busy_ms) across 160 dense shapes, 60 shapes with m in 9–32, and 264 shapes with m≤8, using coordinate descent on softmin expected log regret. End-to-end busy-time regret on the held-out manifest drops from 5.79% to 0.48%, and the m≤8 holdout set drops from 14.18% to 1.79%. On the GEMM side, coop2_swap is added: non-NT layouts are converted to a transposed persistent kernel, with two consumer warpgroups each holding 128 rows, tail writes landed back via stmatrix.trans and TMA transpose, and mixed wide/narrow tiles supported (e.g., 2112 = 13×160 + 32 split into 160-wide tiles plus a 128 narrow tile); 2112×7168×4096 TN drops from 195.3 µs to 175.3 µs, and another H200 on the same machine reads 174.3 down to 158.5 µs — both beating QuACK’s 187.9 and 167.0 µs.

The reduction side broke and was fixed the same day. In the morning, #2787 implemented NaN propagation for extrema reductions using int32 keys, and #2788 changed the row tiles of the parallel cumsum; both landed around 07:00, after which #2790 and #2792 respectively closed out the two regressions: cumsum was changed to seed-based segmentation with two launches ((8,65536) float32 from 12.75 µs to 5.18 µs, also fixing mainline launch failures above 65535 tiles), and extrema switched to PTX max.NaN / min.NaN with a separate kernel for rows that are not whole 16-byte vectors (amax (2048,4096) bf16 back from 8.352 µs to 6.976 µs, faster than FlagGems). This “introduce and fix within the same day” cadence shows that TileOPs’ nightly benchmarks are now catching regressions the same day.

I. Core Project Progress

1.1 Main repo: Index coefficients of scalar reads no longer leak into lane counts; compile crash for row widths above 32768 closed (#3453)

Merged 10-28 11:06, single file +8/-5. A kernel that reads a full row in vectorized fashion while also reading individual elements of that row by index failed to compile with InternalError: Can’t fetch the lanes of a scalable vector when the row width modulo 65536 was at least 32768. Root cause was in VectorTypeAccessChecker::OnArrayAccess: it recorded scalar reads as access_dtype.with_lanes(me->coeff), where me->coeff is the modulo coefficient of the read index — for x[row, 0] that is the row width. DataType’s lane count is only 16 bits and is read back as int16, so 32768 through 65535 was treated as a scalable vector and lanes() threw directly in get_preferred_dtype, while wider multiples silently wrapped. The fix stores the coefficient separately in an int64_t, uses it for the modulo in compatibility checks, and keeps only the element type in the dtype of scalar reads. The original trigger was TileOPs’ GroupNorm: both the 40960 and 163840 group widths in Stable Diffusion hit it. Verified with two wheel builds compared: mainline fails to compile at both 32768 and 40960 widths, while this change builds at all tested widths.

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

1.2 Main repo: Batched all-reduce resolves thread ranges via the scalar branch; fixes all 2048 outputs wrong when batch>1 (#3341)

Merged 10-28 05:46, 2 files +170/-8. The scalar all-reduce branch resolves the participating thread range before emitting the call (ResolveAllReduceThreadRange when reducing_threads exceeds 32 and the target supports named barriers), but the batched branch took block boundaries directly as offset and barrier width. When the reduction threads are a partial range not starting at 0, such as [32, 96), the batched call indexed from thread 0 and barred on whole blocks, folding uninitialized registers of non-participating threads into the reduction: batch=1 was correct, batch=2 on the same kernel produced all 2048 output elements wrong with a maximum absolute error of 1083 and no diagnostics. The fix makes the batched branch use the same resolution as the scalar branch, passing the resolved offset and range to MakeBatchAllReduce, with workspace stride still the block size. A new test runs batch=1 and batch=2 on the same kernel and compares against the torch reference.

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

1.3 Main repo: Shape validation added for SIMT atomic regions; out-of-bounds regions rejected at compile time (#3465)

Merged 10-28 08:14. SIMT atomic operations previously did not validate the shape of the passed region, so mismatched shapes only surfaced as wrong results or crashes in codegen. This change adds region shape validation on both the CUDA and ROCm paths, rejecting invalid regions at compile time.

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

1.4 Main repo: Sparse MMA K tail blocks no longer silently accepted (#3367)

Merged 10-28 16:01. When sparse MMA had unhandled tail blocks along K, it previously did not error but reused whole-block copy and compute parameters, causing silent numerical deviation. This change explicitly rejects such K tail blocks, moving the problem to compile time.

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

1.5 Main repo: atomic_addx4 target dtype narrowed to float* (#3369)

Merged 10-28 16:08, 3 files +27/-9. The generic CUDA implementations of atomic_addx4 and atomic_addx4_ret previously accepted any target pointer, so int32 or int64 targets fell through to float semantics and computed wrong results. This change restricts both generic overloads to float* targets, leaving the existing fp16 and bf16 specializations untouched; the Python-side API stays type-neutral, and integer targets on the CPU backend are unaffected. A regression case for compile failure with integer targets is added on the CUDA side, and accumulation correctness cases for two integer widths are added on the CPU C backend.

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

1.6 Main repo: handle suffix of non-restrict parameters no longer stripped (#3366)

Merged 10-28 13:19. Parameter deduplication compared names with the _handle suffix removed, so two distinct parameters A and A_handle in the same kernel were treated as one and codegen emitted only one. The fix deduplicates by exact name plus variable identity, and adds a CUDA-side regression case with A and A_handle coexisting.

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

1.7 Main repo: TMA descriptors record only validated names (#3484)

Merged 10-28 17:12. TMA descriptor registration previously happened before validation, so names of descriptors that failed validation stayed in the table and later valid descriptors with the same name were misjudged as duplicates. This change moves registration after validation.

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

1.8 Main repo: Unified validation of profiler output indices; scalar and list forms share one set of bounds rules (#3486)

Merged 10-28 19:24, 2 files +15/-9, fixes #2884. The profiler and adapter previously handled result_idx inconsistently: the scalar form had bounds checks, the list form did not, and negative indices were not normalized. This change makes BaseKernelAdapter and Profiler share one set of rules, accepting None, an integer, or a list of integers, judged over -len(params) to len(params)-1, with valid negative indices normalized and out-of-bounds or non-integer elements raising ValueError. 28 new CPU-side regressions cover both the scalar and list paths.

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

1.9 Main repo: T.Parallel’s coalesced_width uniformly converted to IntImm (#3374)

Merged 10-28 16:27. When coalesced_width was passed a Python integer, some paths fed the native integer straight into passes, and type assertions failed in later transformations. This change converts it uniformly to IntImm at the language layer, eliminating failures caused by differences in how it is written.

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

1.10 Main repo: Test infrastructure moves cache-sensitive cases to without_kernel_cache (#3695 / #3696)

Merged 10-28 01:59 and 09:19. Following up on yesterday’s issue #3694 (kernel cache hits breaking tests that assume cold compilation), these two changes move cases that inspect compilation artifacts, source, and diagnostics to without_kernel_cache() or explicit .compile(…), so assertions see the results of the current build rather than an old kernel in the in-process cache.

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

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

1.11 Three new defects filed: 13 examples failing on SM120, exponential IfStmtBinding compile time, without_kernel_cache not forcing recompilation (#3701 / #3700 / #3699)

Three newly opened issues point at testing and compile time. First, #3701 runs the full examples suite on an RTX PRO 6000 Blackwell (sm_120, 99 KiB dynamic shared memory limit): both the CUDA and CuTeDSL backends show 13 failed, 199 passed, 27 skipped, with an identical failure set — 7 are examples hardcoding Hopper-sized tiles on the non—tune path (e.g. attention_sink needs 163840 bytes, deepseek_mla decode needs 231424 bytes), 4 are because the reference implementation flash-attn 2.5.8 has no sm_12x kernels, and the other two are ReduceOp sinking and ThreadSync failing to prove thread ranges. Second, #3700 reports that repeatedly writing x = x + x in a loop with an i < N guard makes tl.LayoutInference go from 0.087 seconds to 118.5 seconds, dropping back to 0.100 seconds when replayable bind inlining is disabled. Third, #3699 notes that without_kernel_cache() only flips CacheState (read only by the disk cache), while the @tilelang.jit wrapper also checks the in-process _call_form_cache and _kernel_cache before JITImpl.compile(), so an ordinary call inside that block does not guarantee recompilation.

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

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

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

II. Multi-Backend Adaptation (Ascend / Hygon / ROCm)

Merged 10-28 08:04, 55 files +21490/-1, author CeleNewYear, based on a35f8ddf of v0.1.15; the two commits handle backend landing and behavior alignment respectively. The backend surface includes backend, adapter, codegen, compiler, execution_backend, jit, pipeline, runtime, and target under tilelang/ascend910, plus kernel, copy, gemm, and the npuir dialect surface under language; on the C++ side are the developer and expert emitters under src/ascend910/target, the operator layer under src/ascend910/op, and LowerNpuirBlock, NpuLoopVectorize, and RewriteFDivToMul under src/ascend910/transform; the runtime requires npu_launcher.h, npu_utils.cpp, and npu_utils.py to generate a torch-compatible launch wrapper and estimate per-block scratch via *_infer_workspace_shape_function; the build side wires into CMakeLists and adds the –bishengir-path option to install_npuir.sh; target aliases accept ascend910, ascend-910, a2, a3, npuir, and the mode switch TILELANG_ASCEND_MODE takes expert or developer, defaulting to expert. Acceptance data: pytest testing/npuir 504 passed, 0 failed; examples/run_all.py –sequential 71 of 72 passed with 1 failure, 0 timeouts, 0 errors, and the sole failure deepseek_v32/fp8_lighting_indexer.py also fails on the same machine with the old backend. Differences aligned in the second commit include bf16 legalization (HIVM rejects bf16 operands, so computation is done in fp32 followed by explicit conversion) and copy sinking for dynamic-extent regions.

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

2.2 Ascend main repo: Bisheng build identity added to binary cache key (#3691)

Merged 10-28 10:26. The Ascend-side binary cache previously keyed only on kernel and target, so switching Bisheng toolchain versions would still hit stale binaries. This change folds build identity into the cache key, so a toolchain change invalidates and triggers recompilation.

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

2.3 Ascend main repo: manual pipeline preserves config register phase (#3624)

Merged 10-28 09:57. The manual pipeline lost config register phase information during the rewrite stage, causing the same configuration to fail to take effect on the correct cycle. This change preserves the phase during rewriting, restoring timing semantics consistent with the automatic pipeline.

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

2.4 Ascend repo under review: root cause and fix for incorrect W4A4 int4×int4 GEMM results (#1731)

Under review, 7 files +634/-24, fixes #438. The root cause is that the entire C++ template stack derives fractal geometry from sizeof(AscendC::int4b_t) == 1, whereas int4b_t is a nibble-packed type: the correct element count for a 32-byte C0 block is 64, not 32, so all copy strides, repeat counts, and byte counts are off by a factor of 2. Three layers compound: the LoadData parameters from L1 to L0A/L0B are derived from that constant, so data piles up in the wrong place; the zN-to-nZ transpose of the B operand is silently disabled by the s4-specific instruction load_cbuf_to_cb_s4, requiring a switch to LoadDataWithTranspose to load_cbuf_to_cb_transpose_s4, whose parameter semantics are undocumented in the CANN headers, forcing the author to reverse-engineer them on-device with probe kernels (correct parameters repeatTimes = N/64, srcStride = roundUp16(K)/64, dstGap = 3, and the s4 version works in 2048-byte blocks rather than the 512-byte fractal stated in the docs); the UB-to-GM byte length is computed at 1 byte per element, moving twice as much per row and mixing in adjacent rows. The mad_s4 instruction itself is correct; the symptom is 99.6% mismatch, with output -501/627 versus the correct answer 641/1785.

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

2.5 Ascend repo under review: a2a3 scheduler migration, 100 files, 18453 lines (#1951)

Under review, 100 files +18453/-666, author erhsh. Migrates the A2/A3 scheduler to the new architecture, the largest under-review change in the window. The PR description table is unfilled; specific boundaries need to be checked against the file list.

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

2.6 Ascend repo under review: Carver reads A3 resource budget per selected SDK (#1955)

Under review, 3 files +203/-27. On Ascend910_9392, Carver falls back to Ascend910A, placing a 256×256 FP32 accumulator against a 256 KiB L0C budget, whereas the selected SDK provides only 128 KiB. The fix identifies the two specific SKUs 9392 and 9382, uses the existing ASCEND_HOME selection to read each platform’s INI (with an additional explicit file option for cross-compilation), and trims SDK capacity to the current backend arena while preserving original capacity, config identity, BF16 support, core count, and L2 attributes. Six targeted cases cover device, target, and explicit chip probing, oversized and legal L0 tiles, and a separate 9382 configuration with BF16 disabled; the author reports regressions of 1352 passed, 3 skipped on the native Ascend CI main suite and 153 passed on the examples suite.

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

2.7 Ascend repo under review: AI Core Exception Dump dependency lowered to CANN 9.2.0 (#1956)

Under review, 5 files +49/-32. The feature previously declared a dependency on CANN 9.3.0; the author verified it works on 9.2.0 and fixed two compatibility issues: acldumpTensorInfo now fills in the four required fields addrType=ACL_DUMP_ADDR_RAW, placement, shapeNum, and originShapeNum (9.2.0’s acldumpSaveExceptionInfo strictly validates addrType, and the original implementation was zeroed by calloc to ACL_DUMP_ADDR_PTR and rejected), and tensors with NULL addresses or zero size are skipped; the registration interface moves from the deprecated aclrtSetExceptionInfoCallback to aclrtExceptionInfoCallbackRegister. The parsing side supports both msaicerr.py output forms (9.2.0 produces .npy, older versions produce .bin), and the documented version requirement is updated to CANN no lower than 9.2.0.

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

2.8 Ascend repo under review: docs switch to T.Tensor, deprecating the T.Buffer DSL interface (#1957)

Under review, single file +124/-122. The kernel parameter definitions in section 3.1 of the programming guide are aligned with the upstream API: text and examples uniformly use T.Tensor, with an added note that the deprecation of T.Buffer(…) and T.Buffer[…] applies only to the TileLang DSL interface, while the underlying TVM Buffer type and concept are not deprecated. This is based on the @deprecated decorator in tilelang/language/proxy.py on the ascendc_pto branch.

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

2.9 Hygon: LDS swizzle feature branch for HCU plain copy, 18 files including 5 new tests (#30)

Under review, 18 files +2000/-84, branch Feat/hcu plain copy lds swizzle pushed 10-28 13:59. Changes concentrate in the new gemm_lds_strategy_utils (+148/-0 and +29/-0), annotate_mls_gemm_dep.cc (+332/-45), and ds_read_format.cc (+61/-22), with small accompanying changes to copy and merge_shared_memory_allocations; on the test side, five new tests are added at once: test_hcu_plain_copy_lds_swizzle (396 lines), _formula (237), _formula_b8 (220), fp8 (200), and _b8_write_phases (80). The direction is LDS write phasing and swizzle strategy for the plain copy path, covering both b8 and fp8 data widths.

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

2.10 ROCm: ROCm 10 support for the pip-based SDK merged, MI355X no longer misdetected as gfx900 (#3549)

Merged 10-28 07:38, 2 files +106/-19. ROCm 10.x’s rocm/pytorch images ship ROCm as Python wheels (rocm-sdk and _rocm_sdk_devel), with neither an /opt/rocm directory tree nor rocm-* apt packages. get_rocm_arch() previously hardcoded the /opt/rocm prefix and fell back to gfx900 when the path did not exist, so rocwmma’s static assertion failed outright on MI355X (gfx950); setting ROCM_PATH did not help either, because TVM’s tvm_callback_rocm_get_arch calls the function with the default prefix. The fix resolves in the order of caller prefix, find_rocm_path(), torch’s gcnArchName, and gfx900 fallback, with a 10-second timeout on each rocminfo probe; it also adds docker/Dockerfile.rocm10 based on rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12, removing the rocm-* apt installs, pointing ROCM_PATH at rocm-sdk path –root, retaining the base image’s torch build, and adding an import self-check. The author measured architecture detection changing from gfx900 to gfx950 on MI355X, with the testing/python/amd suite building and passing against the correct architecture.

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

III. Ecosystem and Adopters (TileOPs / TileFoundry)

3.1 TileOPs: Selector recalibrated on busy-hour conditions, manifest regret cut from 5.79% to 0.48% (#2784)

Merged 10-28 10:51, 4 files +259/-203. The dense selector’s scoring now uses a tensor_core_tflops multiplied by a structural factor; previously structural differences only showed up in a single max(), so memory-bound shapes scored basic, simple, and coop2 equally. The launch_us term was removed (it fits to 0 under busy-hour conditions). The m≤8 NT band replaces hand-tuned thresholds with a streaming model: per-CTA rate is min(ns·rows·bk·2/latency, cta_gbps), waves are computed from SMEM-limited residency, plus a DRAM floor of n·k·2/dram_tbps and a split-K reduction term. The TN/TT persistent kernel path is now wired up (A/B SMEM tiles follow the storage layout, transpose_A/B); 176/192 widths, narrow tiles, ping-pong, and split-K remain NT-only. Calibration is done under busy-hour conditions across three grids of 160/60/264 shapes, 7 interleaved CUPTI rounds taking the median, with softmin expected log regret via coordinate descent and tau annealing 0.1→0.05→0.02. End-to-end results: held-out manifest 5.79% to 0.48%, calibration grid 9.28% to 2.65%, m 9–32 3.13% to 2.09%, m≤8 11.20% to 1.17%, m≤8 held-out set 14.18% to 1.79%.

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

3.2 TileOPs: Non-NT layouts moved to transposed persistent kernel, 2112×7168×4096 TN from 195.3 to 175.3 µs (#2789)

Merged 10-28 21:06, 4 files +620/-15, resolves #2785. Adds coop2_swap: for TN/TT/NN it computes Cᵀ = op(B)ᵀ·op(A)ᵀ, runs a persistent kernel over 256×block_n tiles, with two consumer warpgroups each holding 128 rows, and writes the tail back transposed via stmatrix.trans and TMA. Mixed tile widths (narrow_n and row_mix) keep M from being wasted on padding—e.g. 2112 = 13×160 + 32 splits into width 160 plus narrow 128, wide runs first and narrow finishes, with the split point taken from coop2’s narrow-tile _busiest_cta metric. The selector enumerates coop2_swap block_n 128 and 160 outside NT, structural factor 1.2, refit under busy-hour conditions. H200 busy hour: 2112×7168×4096 TN from 195.3 to 175.3 µs (QuACK 187.9), 24576×1536×4096 TN from 421.9 to 416.4 µs; on another H200, 174.3 to 158.5 (QuACK 167.0) and the 384–402 µs range respectively. The row-group rule changes from fixed 4 to 16, within 0.989–1.022× of group 4 across 24 TN/TT/NN calibration shapes; on another 9 TN shapes, mixed widths are 3.6% to 11% faster than uniform coop2_swap.

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

3.3 TileOPs: cumsum long rows switched to two-phase seeded scan, (8,65536) from 12.75 to 5.18 µs (#2790)

Merged 10-28 20:45, 4 files +255/-246. cumsum with few long rows now uses a seeded segmented scan, cutting launches from three to two, eliminating the regression introduced by #2788 and fixing its layout failure. The rule: use the segmented scan when rows are at least 64 KiB and row count is small; cumprod unchanged. Also fixes a launch failure on main when row count exceeds 65535 tiles (e.g. (1, 2**24) float32). H200 @1500 MHz busy hour: (8,65536) float32 12.75 to 5.18 µs, (16,65536) 19.14 to 7.01 µs, (1,65536) bf16 9.03 to 3.45 µs, (3,8193) fp16 7.52 to 3.29 µs. Across a 384-shape sweep, nothing is slower than before #2788, and the CumsumFwdOp and CumprodFwdOp manifest rows are both faster than the fastest comparison implementation; the new few-long-seq workload (8×65536 float32) runs 5.20 µs vs FlagGems 6.30 µs.

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

3.4 TileOPs: extrema reduction switched to PTX max.NaN, NaN propagation unchanged and regression gone (#2792)

Merged 10-28 23:31, 4 files +346/-119. amax/amin now use PTX max.NaN and min.NaN on fp32 instead of #2787’s int32 key scheme, keeping NaN propagation semantics unchanged while removing the regression from earlier that day; rows that aren’t full 16-byte vectors go through a separate extrema row kernel, since T.reduce_max drops NaN on fp32. H200 @1500 MHz: amax (2048,4096) bf16 was 7.008 µs before #2787, 8.352 on main, 6.976 with this change; fp16 7.008 to 7.775 to 6.944; amax (64,32768) bf16 3.552 to 4.544 to 3.520; amin same shape 3.552 to 4.224 to 3.520.

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

3.5 TileOPs: Two changes—effective row tiles for parallel cumsum, narrow prefill tiling for GDN (#2788 / #2786)

#2788 merged 10-28 07:13, making parallel cumsum process only effective row tiles; it was the earliest change in the window and was later superseded by #2790’s two-phase scan on the problematic path. #2786 merged 07:09, 2 files +30/-2, switching the 32-column GDN prefill plan on SM89 to the existing state scan with fused output, handling the full-block segment and tail block separately and keeping the three interfaces of prepare, output, and final state; selected when input is dense, batch 1, 16 value heads, state width 128, key major, and sequence length greater than 16384 and at most 65536. At H=4, median device execution time drops 13.65% to 18.31% versus the TileOPs baseline, and device active span drops 13.48% to 17.07%; across four interleaved rounds against FLA, the span difference at H=16, length 65536, BF16 is between -2.57% and +1.77%, which the author notes falls within measurement noise.

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

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

3.6 New TileOPs issue: Large GEMM rows lag the fastest baseline under back-to-back calls (#2791)

Newly opened, following up on #2133 and #2789. bench_gemm reads a single call with an L2 flush and sync in between, so calls start below the power cap and the energy cost is masked; switching to back-to-back consecutive calls as in training, six of the ten manifest large rows lag the fastest baseline. A typical row: prefill-o-proj 4096×7168×16384 NT runs 1412.7 µs here vs QuACK 1371.1, cuBLAS 1360.1, DeepGEMM 1359.2, a ratio of 0.962. Method: single H200, unlocked clocks, 1 s warmup then 1.5 s timed, four interleaved rounds taking the median.

Source: https://github.com/tile-ai/TileOPs/issues/2791

3.7 TileFoundry: Ring stage binding and compile flow unified, 70 files (#277)

Merged 10-28 16:55, 70 files +1567/-1298, supersedes #276. Two motivations: unbound ring views and stage selection could carry views without SSA names into CUDA operands; and the three compile entry points each assembled different pass lists. The fix stores ring stages as a set of bound TensorViews, binding once per consuming scope, and shares default_pipeline across finalize, compile, and scaffold; the core host entry pass is resolved via Target.resolve with CUDA providing the @bind implementation; Module.entry becomes the host starting point, and the analysis’s shared-memory peak is carried into host launch; all CPU hosts in the call chain participate in linking, with the root host first. Contract changes: link_program drops entries, the loader binds the first linked host, and build requires a CPU host root. Remaining risks are noted in the PR: #278 tracks Broadcast layout normalization, #255 tracks SymbolRef and function references.

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

3.8 TileOPs under review: Release image bundles TileOPs 0.1.0 with TileFoundry 0.0.3 (#2793)

Under review, 9 files +120/-44. .github/runner/Dockerfile adds a release target that installs the PyPI releases of tileops and tilefoundry on top of the dev image, names the image by TILEOPS_VERSION and TILEFOUNDRY_VERSION, and pushes it as ghcr.io/tile-ai/tileops-release:tileops0.1.0-tilefoundry0.0.3-cuda13.2-torch2.13-tilelang0.1.12; a companion scripts/verify_release_image.sh pulls the image and runs stack and baseline checks on GPU, packaged tests, and the specified –bench files, with cases taken from the git tag matching the in-image tileops version (the wheel ships without tests or benchmarks). Also removes the sgl-kernel baseline: no image installs it anymore, and bench_moe_permute_align has always skipped it.

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

IV. Community, Tutorials, and Events

4.1 Two automated documentation site updates (10-28)

tilelang.github.io received one “Update docs” commit each at 08:20 and 13:31 on 10-28, both submitted by github-actions. The content is documentation build artifacts synced with the main branch, with no independent content changes.

Source: https://github.com/tile-ai/tilelang.github.io

4.2 Batch of Ascend repo documentation fact-checks and installation fixes under review, seven advancing in the same batch (10-28)

The tilelang-ascend documentation fact-check pipeline updated seven items within the window: fact-checks for five APIs — T.tile.bitwise_lshift, bitwise_or, bitwise_and, fill, and clear — plus clamp taking effect on PTO with documentation verified (#1620), and strengthened docstrings and tests for axpy (#1600). Separately, #1943 changes the SHMEM installation to stop on build failure, and together with the snapshot analyzer validation in #1942, both are bot submissions.

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

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

4.3 New operator examples under review in the Ascend repo: GroupNorm and fla_chunk (#1607 / #1953 / #1952)

Three example-type PRs were updated within the window: #1607 adds 20 cann-bench cases for GroupNorm, #1953 adds the fla_chunk_o operator and tests, and #1952 adds the fla_chunk_intra_token_parallel operator and tests. This is also the original scenario behind the scalar-read crash defect in 1.1 (GroupNorm reading the first element of each row).

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

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

4.4 Third party: community auto-tuning agent refreshes Chinese documentation (#74); Hexagon backend repo has no new commits on the main branch within the window (10-28)

The community repo cc-c122/TileLang-Remote-Autotuning-Agent merged #74 at 09:21 on 10-28, refreshing the Chinese README per the current V2 workflow; the repo previously had a merge on 10-24 covering baseline cancellation status and mcprofiler attribution. Another third-party backend repo, xwhzz/tilelang-hexagon, shows a push timestamp of 06:56 on 10-29 in the organization view, but the latest commit on the main branch is dated 08-24 (CI dependency upgrade); there were no code commits on the main branch within the window, and yesterday’s 15 commits were feature-branch activity, so this item is not counted as new content.

Source: https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent/pull/74

Source: https://github.com/xwhzz/tilelang-hexagon

4.5 News side: six Google News queries in Chinese and English, Hacker News, and arXiv all yield no substantive hits within the window (10-28)

The six queries covered the exact term TileLang over two days, Chinese operator terms over seven days, English GPU kernel DSL over three days, organization names over seven days, Chinese vendor terms (Ascend / MetaX / Hygon) over seven days, and operator library terms over two days. Among the returned results, the only TileLang-related item falling within the window was a reposted article (an Ascend 950PR matrix compute comparison, which is secondhand market commentary with no original data and is not included); the rest were general DeepSeek funding and technology coverage whose titles carry no TileLang-related content. All English queries returned empty. The Hacker News item of the day was the launch of a new language called Vu, unrelated to TileLang. arXiv had no new tilelang or tiled DSL papers within the window. The gap in signal strength between the news side and the GitHub side continued to widen today, and the main body of the daily report remains repo-based.

Source: https://news.google.com/

V. Trend Observations

The Ascend side’s approach over the past two days is now clear: the main repo merges small fixes by the hour (cache keys, register phases), the dedicated repo mlir-ascend moved the 910 NPUIR path from a feature branch into the main branch with a single 20,000-line merge, and tilelang-ascend is filling in documentation and operators on its established dual-backend cadence of PTO and AscendC. The three operate at different tempos but point to the same thing: the 910 compiler path is shifting from “it runs” to “it’s acceptance-ready,” and the 504 self-tests and 71/72 examples are the acceptance criteria given this time. Notably, the a2a3 scheduler migration in the review queue spans 100 files, indicating that the next phase of change is at the scheduling layer rather than the operator layer.

On the performance side, the focus is on the measurement methodology itself. These two TileOPs changes move both the selector and GEMM conclusions to busy time (device_busy_ms), and in the same PR present the difference between single-call and back-to-back invocation: under the single-call reading, this implementation leads QuACK, while in back-to-back, power-capped training scenarios, six of the ten large rows lag behind. This is a more important signal than “how many microseconds a given row is faster by,” because training is in practice back-to-back. #2791 makes this a standalone item, showing the team now treats the methodology issue as a product issue rather than glossing it over as benchmark bias.

On engineering hygiene, three things are worth noting today: first, self-inflicted breakage and self-repair within the same batch (#2787/#2788 merged in the morning, #2790/#2792 reverted in the afternoon), which shows the nightly benchmark’s fallback is working but also that the semantic boundaries of this reduction code remain fragile; second, the semantic hole in without_kernel_cache() (#3699) was caught during the migration — the two in-process cache layers do not honor CacheState, so tests believe they are cold-compiling while actually getting an old kernel, and such defects render “all our tests pass” meaningless; third, 13 example failures on SM120 (#3701) expose examples hardcoding Hopper dimensions, and as Blackwell installations grow, this kind of “examples are not the product” debt will be repeatedly charged.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
GitHub tile-ai org, repo by repo (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / docs site) Pushes within window: tilelang, TileOPs, TileFoundry, tilelang-mlir-ascend, tilelang-hygon, tilelang.github.io; tilelang-metax, tilelang-musa, TileRT, tilescale, DeepStack, tilelang-puzzles had no main-branch commits within window
Merged ground-truth search (is:pr is:merged) 24 PRs: tilelang 15, TileOPs 7, TileFoundry 1, tilelang-mlir-ascend 1
Open ground-truth search (is:pr is:open updated) 48 PRs, including the Ascend repo docs-verification batch and a2a3 migration
New issue check 6 issues, 1 of which is the Ascend repo’s daily-test bot
Main-branch commit-level check tilelang 15, TileOPs 7, TileFoundry 1, docs site 2; merged counts match main-branch commit counts
Releases page and tags tilelang latest remains v0.1.15 (released 09-30), no new version within window; TileOPs latest is v0.1.0
Third-party backend repos Hexagon backend main branch’s latest commit dates to 08-24, no commits within window; autotuning agent merged Chinese docs refresh; six branches of the MetaX repo had no commits within window
Google News (six queries, Chinese and English) Only one secondhand reprint touched on it, no primary data, not included
Hacker News No TileLang-related items that day
arXiv No new tilelang or tiled DSL papers within window

Full Source List