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

Sources: Per-repo verification of the GitHub tile-ai organization (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / docs site / nightly mirror repos) + ground-truth search of merged and in-review items (12 merged, 54 in-review updates) + full check of newly added issues (4) + commit-by-commit verification of the main branches (tilelang 5, tilelang-ascend 5, TileOPs 1, TileFoundry 1, docs site 5) + release page and tag verification + verification of third-party backends and ecosystem repos (10 repos) + six Google News queries in Chinese and English (via proxy) + Hacker News + arXiv.

This Issue’s Index

  • Today’s Highlights: Five Ascend repo merges fill in algorithm-side operators (GLM-5.3 kpool compression, Chunk Scaled Dot KKT, W4A4 GEMM fix, depthwise convolution example, performance profiling skill); the main repo switches full-warp reduction to redux.sync and changes the autotuner’s independent configs to concurrent compilation; 12 merges in 24 hours, 54 in-review updates, 4 new issues
  • I. Core Project Progress
    • 1.1 Main repo: full-warp reduction switched to redux.sync; four 32-lane FP32 max reductions go from 20 shuffles to 4 hardware reductions (#3692) (10-29)
    • 1.2 Main repo: autotuner’s independent JIT configs changed to concurrent compilation; locks retained only for same-key candidates (#3476) (10-29)
    • 1.3 Main repo: Ascend buffer view layout and pointer units no longer propagate across aliases (#3623) (10-29)
    • 1.4 Main repo: GM-to-L1 copy and out-of-bounds fill merged into one scheduled task, removing a redundant MTE2 fence (#3543) (10-29)
    • 1.5 Main repo: CI pins ROCm and Metal jobs back to compatible PyTorch wheels (#3706) (10-29)
    • 1.6 Three new defects: shared buffer with padding via bulk copy reports shape mismatch, negative-scale iterative graph miscompiles reverse indexing, view bit width overflows in int32 (#3704 / #3705 / #3707) (10-29)
    • 1.7 Three fixes in review: bit width changed to int64, CPU scalar GEMM clears only the target region, two in-process cache layers bypassed when caching is disabled (#3708 / #3709 / #3710) (10-29)
  • II. Multi-Backend Adaptation (Ascend / Hygon / Moore Threads / MetaX)
    • 2.1 Ascend: kpool compression operator for GLM-5.3 sparse attention indexer merged, single-kernel pure Vector path (#1910) (10-29)
    • 2.2 Ascend: Chunk Scaled Dot KKT forward operator for chunked gated linear attention merged, static four-block unrolling, zero GM workspace (#1803) (10-29)
    • 2.3 Ascend: fix merged for W4A4 int4-by-int4 GEMM returning all-wrong results; root cause was misreported type sizes and a specialized transpose instruction being silently disabled (#1731) (10-29)
    • 2.4 Ascend: depthwise_conv_2d example resubmitted, switched to manifest-based test channels to avoid concurrent batches (#1920) (10-29)
    • 2.5 Ascend: new operator performance profiling skill added, using msprof op output to determine optimization direction (#1931) (10-29)
    • 2.6 Ascend in review: three view and broadcast fixes at the NPUIR language layer (#1962 / #1960 / #1961) (10-29)
    • 2.7 Ascend in review: fix for PTO backend bitwise XOR unsigned byte workspace view (#1959) (10-29)
    • 2.8 Hygon: two HCU items in review updated, community DCU fused-operator repo appears; no pushes from Moore Threads or MetaX during the window (10-29)
  • III. Ecosystem and Adopters (TileOPs / TileFoundry / TileKernels / FlashQLA / Third Parties)
    • 3.1 TileOPs: release image merged; operator library and TileFoundry packaged into the same image for the first time (#2793) (10-29)
    • 3.2 TileOPs in review: GEMM recomputes grid grouping and resident block count under a power cap; q-b DRAM reads drop from 168 MB to 94 MB (#2794) (10-29)
    • 3.3 TileFoundry: AMX runtime operators changed to 39 static TIR contracts, a breaking change (#279) (10-29)
    • 3.4 Adopters: no pushes to TileKernels or FlashQLA main branches during the window (10-29)
    • 3.5 Independent community operator repo ops-tilelang appears on the Ascend side, with pushes during the window (10-29)
    • 3.6 PTO-ISA companion repos active: problem-set repo migrated to PTOAS 0.70, benchmark repo adds SV2 chunked scaling and per-column analysis (10-29)
    • 3.7 Observation: tilelang branch repo under the T-HEAD name has pushes during the window; branch names point to PPU matrix instruction layouts (10-29)
  • IV. Community, Tutorials, and Events
    • 4.1 Docs site adds applicability boundaries of built-in reductions to the instruction page, synchronized with the redux merge (10-29)
    • 4.2 Dev tools repo completes three follow-ups after extracting sink tracking, aligned with upstream in-review tooling changes (10-29)
    • 4.3 New repo tilelang-knowledge: four commits in one day — 33 operator design cases, 20 sourced knowledge cards, 7 optimization patterns (10-29)
    • 4.4 Illustrated course companion repo and distributed layout experiment repo created on the same day (10-29)
    • 4.5 News side: six groups of Chinese and English Google News queries, Hacker News, and arXiv all had no substantive hits during the window (10-29)
  • V. Trend Observations
  • Appendix: Sources and Verification Notes

Today’s Highlights

Today’s merge volume fell from yesterday’s 24 to 12, but the structure is more concentrated: five in the Ascend-specific repo, five in the main repo, and one each in TileFoundry and TileOPs. The Ascend batch is entirely algorithm-side operators, and all of them are the kind that can be used directly by upstream models. The kpool compression operator for the GLM-5.3-Flash sparse attention indexer packs per-dimension softmax pooling, BF16 boundaries, normalized Hadamard-128 rotation, and E4M3 quantization into a single pure Vector kernel; each block covers 64 pools, with two vector subkernels handling 32 each, zero host-side reordering, and strided transfers directly following the caller’s 3D layout. E4M3 encoding is implemented with pure integer bit operations (round-half-even, including subnormals, saturating to 448), relying on neither the FP8 type nor approximation. Chunk Scaled Dot KKT is the A = gate(K K^T) in the chunked gated linear attention forward pass, with static four-block unrolling, zero GM workspace, and zero host operators; match rates for all three L0 and L1 cases are 1.0.

The main repo’s theme today is the compilation and launch layer. Full-warp reduction now uses redux.sync: on the main branch, four independent 32-lane FP32 max reductions expand to 20 shuffles on CUDA 13.1, sm_103a; after the change they become four redux.sync.max.f32 instructions with no shuffles. Integer reductions use hardware reduction on SM80 and above; FP32 max/min is available only on sm_100a, sm_100f, sm_103a, and sm_103f; narrow types (including CUTLASS’s FP16 and BF16) are promoted to FP32; FP64 still uses the shuffle path. The author explicitly does not claim end-to-end speedup and provides only instruction-level comparisons. The autotuner’s concurrent compilation, meanwhile, splits the compile lock that was previously held for the entire section: each coverage config is copied into an independent JITImpl via dataclasses, only candidates sharing the first-stage key are serialized, and debug output writes to disk remain serialized. The author’s cross-backend validation is 1797 items on Ascend, 3220 on ROCm, 95 on Metal, and 5500 on CUDA.

Another detail worth noting: today’s newly opened reverse-indexing miscompilation defect (#3705) was reproduced at runtime using the tilelang branch repo under the T-HEAD name. This indicates that a full-fledged team is already tracking this upstream code as the foundation of its own hardware stack, rather than treating it merely as an external dependency.

I. Core Project Progress

1.1 Main repo: full-warp reduction switched to redux.sync; four 32-lane FP32 max reductions go from 20 shuffles to 4 hardware reductions (#3692)

Merged 10-29 15:08, 4 files +337/-35. Previously, ordinary cross-thread reductions (such as T.reduce_max) were still expanded into shuffle trees even when the whole group was an aligned full warp: on main branch 837486a, four independent 32-lane FP32 max reductions compiled to 20 shuffles and zero redux instructions with sm_103a and CUDA 13.1; after the change they are four redux.sync.max.f32 and zero shuffles. The approach reuses the existing warp_reduce helper from scalar and batched full reductions: callers only check thread geometry (whether aligned to a warp, unit stride, full warp), while architecture, data type, and operator dispatch are left to the helper. Integer reductions use hardware redux on SM80 and above, with smaller integers promoted to 32 bits; FP32 max/min are only available on sm_100a, sm_100f, sm_103a, and sm_103f; narrow floats are promoted to FP32 before entering the helper, while FP64 keeps its precision and stays on the shuffle path. Sub-warp, strided, and packed reductions, as well as unsupported operators and targets, all fall back to shuffle; cross-warp shared-memory exchange and barriers are unchanged. A compile failure is also fixed along the way: custom unsigned integer reducers previously failed to compile because inline dispatch returned an empty type; they now fall back to shuffle. PTX regression assertions check for the real instructions and the absence of shuffles, and runtime coverage includes signed/unsigned boundary values, batched non-zeroing, NVRTC execution, non-finite values, signed zero, and exact FP64 results. The author explicitly does not claim end-to-end speedup. An existing known issue is also noted in the description: batched reductions fail for FP16 and BF16 at batch=4 with two outputs per thread and batch=8 with four outputs per thread; the original template on main branch reproduces it as well, and the failure occurs in the lowering stage, unrelated to this change.

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

1.2 Main repo: autotuner’s independent JIT configs now compile concurrently; the lock is kept only for candidates with the same key (#3476)

Merged 10-29 16:24, 2 files +54/-46. On main branch, _make_jit_compile_func held a single lock for the entire compilation period, serializing even configs that only used default compile options, so frontend, lowering, and device compilation were all queued. The fix constructs a candidate JITImpl for each override config with dataclasses.replace, keeping backend, target, output selection, log level, compile flags, and debug options at their original values; default options still go through the original lazy call path, preserving its kernel cache and object identity; only candidates sharing the first-stage key (computed by the argument binder used by get_tir, preserving positional arguments, defaults, and factories with tensor annotations) are serialized; shared output for debug files is still serialized inside compilation, covering debug_root_path, IR dump settings, and layout visualization. Acceptance data comes from the cross-backend pipeline: the Ascend job merged 3be11be6 with main branch c58a86f0, Python 3.12.14, torch_npu 2.10.0, CANN 9.3.0, with 1797 checks passed, all 50 AscendC examples passed, and 44 PTO examples passed with 6 skipped; the ROCm 7.2 job on gfx942 passed 3220 and skipped 2765; the Metal job on macOS arm64 passed 95 and skipped 5; the CUDA job on H100 PCIe with cu130 passed 5500 checks and skipped 491, with examples at 76 passed and 14 skipped; the CuTeDSL job had 75 passed, 17 skipped, and 2 unexpected passes.

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

1.3 Main repo: Ascend buffer view layouts and pointer units no longer propagate across aliases (#3623)

Merged 10-29 13:17, 23 files +1139/-438. Multiple buffers created by T.view or T.reshape can share the same storage while requiring different layouts; previously layouts were resolved uniformly by the shared data variable, which could select the wrong Cube mapping or DMA address, and stale tiling geometry could drop version offsets or interpret typed pointer offsets with the wrong unit. This change attaches layout to the individual Buffer view throughout layout inference, fractal storage normalization, multi-buffer materialization, and copy and matmul lowering; view identity is preserved through layout mapping and compact L0 spacing application, and raw pointer offsets are remapped according to their own element units. Layouts no longer propagate across storage aliases; each view uses its own annotation, an operator-inferred layout, or plain addressing. On the same versioned allocation, Cube views can now use different MajorK and MajorMN layouts, provided the logical strides are dense, the base offset is zero, the operand types are compatible, and the padded storage sizes are equal. Multi-buffer aliases share one physical slot size, converted by each view’s element unit without propagating index mappings; a view merely reinterprets existing bytes, so writers must write in the representation the consumer expects, and the access-pointer interface can be used when a view must be selected for built-in instructions.

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

1.4 Main repo: GM-to-L1 copy and out-of-bounds padding merged into one scheduling task, removing a redundant MTE2 barrier (#3543)

Merged 10-29 11:39. The global-memory-to-L1 copy and its out-of-bounds padding jointly prepare the same L1 data; previously the padding could be split into a separate scheduling task, so synchronization inserted a redundant MTE2 pipeline barrier between the copy and the padding. This change compiles each copy and its padding into the same Ascend task, so consumers wait for the complete data including padding; FP8 MX K-tail padding is likewise merged in and retains conditional padding within the task, and existing explicit task scopes remain unchanged with no nesting introduced; the former MX deferred-padding queue is removed, and each operand’s padding immediately follows its own copy. 47 focused tests pass, covering both ordinary out-of-bounds padding and MX tail blocks.

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

1.5 Main repo: CI pins ROCm and Metal jobs back to compatible PyTorch wheels (#3706)

Merged 10-29 21:38. ROCm CI exposed a matching PyTorch ROCm source via an index variable, but the workflow-level “unsafe best match” strategy still let later PyPI candidates win: in the ROCm job of #3534, the resolver chose PyPI’s PyTorch 2.15.0 (CUDA 13.2) instead of the available 2.14.1 plus ROCm 7.2 build, and both resolution results were reproduced directly on the official runner. The same 2.15.0 rollout also broke the Metal job: on a macOS 26.6.2 ARM64 runner, 2.15.0 rejects Apple’s paravirtual GPU because it does not expose Metal 2.0, while the last passing run on the same image used 2.14.1. There are three fixes: set the backend variable according to the configured stable ROCm version so uv keeps PyTorch on the ROCm source while other dependencies still use best match; switch nightly PyTorch installation to the first index so nightly jobs cannot escape the configured source; and add a fast-fail validation step that confirms before the test suite starts that PyTorch is a HIP build, not a CUDA build, and can access the runner GPU.

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

1.6 Three new defects filed: padded shared buffer via bulk copy reports shape mismatch, negative-scale iterative graph miscompiles reverse indexing, view bit width overflows in int32 (#3704 / #3705 / #3707)

All three issues point to boundaries of static checking. First, #3704 reports that when a global tensor whose last dimension is not 16-aligned (for example K=100) is copied into a padded shared buffer (BK = cdiv(K,16) times 16 = 112), once this copy is routed to the bulk copy path, lowering directly reports a shape mismatch between the shared block and the global block (112 vs 100), at the bulk lowering of the CUDA copy; padded buffers work fine on the elementwise path, and padding is exactly the common pattern for satisfying matmul alignment, so any kernel using it together with bulk copy can hit this. Second, #3705 reports that flatten-buffer simplification miscompiles reverse indexing on a negative-scale iterative graph; the reproduction environment is v0.1.15 and a source build of the tilelang branch repo under T-HEAD, and the report also verified that the affected split function is unchanged in upstream main branch and the corresponding TVM revision, making it a shared upstream path. Third, #3707 reports that views and reshapes multiply in int32 when checking buffer bit width, wrapping above 2^31: the Qwen3.8 lm_head weight is 248320 times 2560 bytes of NVFP4, totaling 5,085,593,600 bits, and the bit-width product wraps to 790,626,304, after which converting the static view bit count to an int32 constant throws a literal-out-of-range error.

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

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

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

1.7 Three fixes under review: bit width switched to int64, CPU scalar GEMM clears only the target region, cache disable bypasses both in-process caches (#3708 / #3709 / #3710)

Corresponding to the three defects above and yesterday’s cache semantics topic, the fixes went under review the same day. The bit-width change multiplies static dimension products as Python integers, converts symbolic dimensions to int64, and converts literals exceeding int32 to int64 constants; the regression case directly uses the buffer from that issue (view shape 248320 times 5120), which errors on main branch and passes after the change, and same-shape views with symbolic dimensions are also covered. The CPU change fixes scalar matmul zeroing the entire underlying buffer when clearing the accumulator: for a 4-by-4 output buffer initialized to 7, calling a subregion of rows and columns 1 to 3 would zero the 12 elements outside the region; the fix initializes only with the output region’s own offset and extent, and the regression case accordingly uses an all-ones 2-by-2. The cache change fixes #3699 recorded yesterday: disabling the kernel cache only flipped the cache state, while the jit wrapper still checked two in-process caches before actually compiling, so ordinary calls within a block would get a kernel compiled earlier in the same process; the fix reads the cache switch at call time (also covering environment-variable disabling), skips both in-process caches while disabled, and restores afterward. The test went from 1 failed and 4 passed before the change to all 5 passing.

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

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

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

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

2.1 Ascend: kpool compression operator for GLM-5.3 sparse attention indexer merged, single-kernel pure Vector path (#1910)

Merged at 15:41 on 10-29. For each pool (pool_size = 4), the operator adds the score tensor to the positional encoding, then performs per-dimension softmax pooling; the pooled vector passes through the BF16 boundary, is multiplied by a normalized Sylvester Hadamard-128 rotation (divided by the square root of 128), and is finally E4M3-quantized with per-vector FP32 power-of-two absolute-max scaling: the E4M3 rows and scaling factors are written to a contiguous scratch area, from which the host adapter scatters them into the paged cache (the key cache carries E4M3 as unsigned bytes; the scale cache is written in place as FP32). The implementation is a single-kernel pure Vector path; developer mode enables the three switches for auto-sync, auto vector-and-matrix core merging, and memory planning. Each block covers 64 pools, with two vector subcores each handling 32 independent parallel pools; the host performs zero reordering, and the kernel reads directly with strides according to the caller’s 3D layout, with positional encoding folded into pooling. E4M3 encoding is pure integer bit manipulation: round-half-even, including subnormals, saturating to 448, with no FP8 types and no approximation.

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

2.2 Ascend: Chunk Scaled Dot KKT forward operator for chunked gated linear attention merged, static four-chunk unrolling, zero GM workspace (#1803)

Merged at 15:01 on 10-29. Adds a scaled dot-product forward for chunked inner products: for the forward computation of chunked gated linear attention, A equals the gate-gated K times K transpose. Implemented as a single compiler-decorated kernel, statically unrolling four chunks for on-chip fusion, with zero global memory workspace and zero host operator calls, and all four developer-mode switches enabled. Tests are tiered into L0, L1, L2, and boundary levels, selectable from the command line; L0 covers three cases (with gating, without gating, small-scale with gating), with accuracy checked against a PyTorch CPU FP32 golden reference using dual mixed-tolerance gates. The reported test shape is batch 1, sequence length 32768, 32 heads, head dimension 128, chunk 64, gating enabled, BF16 input/output, FP32 accumulation; all three cases each of L0 and L1 pass, with a match rate of 1.0000 and maximum absolute error between 0 and 0.25.

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

2.3 Ascend: Fix merged for W4A4 int4-by-int4 GEMM producing entirely wrong results, root cause being misreported type size and silently disabled dedicated transpose instruction (#1731)

This one, still under review yesterday, was merged today at 15:08, corresponding to #438. The root cause is three layers of compounding errors. First, fractal geometry was derived assuming type size equals 1, but the nibble-packed type occupies only 4 bits per element with two elements per byte, so the correct element count for a 32-byte C0 block is 64 rather than 32; all copy strides, repeat counts, and byte calculations were off by a factor of two, and data piled up in the wrong locations in L0A and L0B. Second, the transposed load of the B operand was silently dropped: the generic load path sets a transpose flag for the zN-to-nZ conversion, but the 4-bit dedicated instruction hardcodes transpose off and ignores that flag, so B never gets the layout the matrix multiply instruction expects; the correct path is the dedicated load with transpose, which has no documentation whatsoever in the CANN headers, so probe kernels were used to empirically reverse-engineer the parameters on device (move the full block in one repeat, repeat count = N divided by 64, source stride = K rounded up to 16 then divided by 64, destination gap = 3), and that version actually works in 2048-byte blocks, differing from the generic documentation’s 512-byte fractal unit. Third, the byte count between vector and global memory was likewise computed as 1 byte per element rather than 0.5, making each row copy twice as long and mixing in adjacent row data. The compute instruction itself was always correct, faithfully multiplying the misplaced data; the symptom was 99.6% mismatch, with outputs of -501 and 627 versus correct values of 641 and 1785. The report also documents the global memory int4 packing convention: two nibbles per byte, low nibble for even elements and high nibble for odd elements, with row r starting at byte position r times N divided by 2, consistent with CANN’s own format.

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

2.4 Ascend: depthwise_conv_2d example resubmitted, switched to manifest-based test channel to avoid the concurrent batch (#1920)

Merged at 15:43 on 10-29. This is a resubmission of #1915: after the original merged, the nightly full CI failed and it was reverted via #1919, with the root cause being that the L0 largest-shape case triggered an unknown error under the eight-way concurrent legacy script batch; this PR slims that case by a quarter and registers the operator into the continuous-integration operator test manifest via the Pytest channel, the same route as the other twenty-five migrated operators, fully withdrawing it from the concurrent batch, with the kernel body byte-for-byte identical to the original. The operator itself supports fp16, fp32, bf16 and arbitrary kernel size, stride, padding, and dilation; strides 1, 2, and 4 take the quantized-region path, others fall back to scalar inner product; the implementation is padding plus convolution as two kernels plus a 1-by-1 fast path, a pure Vector path, with the grid adaptive to geometry rather than hardcoded core counts. Several measured constraints in the report are worth noting: the padding phase costs roughly the number of statements times 35 nanoseconds, independent of byte count, so the tail band is closed with truncation; on the shared memory side, a copy destination offset not aligned to 32 bytes triggers a 0x26 exception in the vector core; synchronization uses a hand-written champion model (auto-sync disabled in the convolution section), with adjacent depth-1 transfers and a compute event protocol between channels, while a depth-2 protocol or inserting a full barrier midway deadlocks.

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

2.5 Ascend: New operator performance profiling skill added, using msprof op output to set optimization direction (#1931)

Merged at 17:17 on 10-29. Adds a development workflow asset for operator performance profiling: before optimizing performance, run an operator-level pass with msprof, then use the output to decide the next optimization direction. The author has validated the workflow on three operators: elementwise add, reduce-min example, and dynamic quantization. The value of such an asset is that it writes “measure first, then change” into the process, rather than re-deriving the profiling methodology each time.

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

2.6 Ascend under review: Three NPUIR language-layer view and broadcast fixes (#1962 / #1960 / #1961)

Three PRs entered review the same day and depend on each other. First, #1962 fixes the view descriptor for local unified buffers: the bridge layer looks up buffers by shared data variable, which may return an allocation memory reference without the descriptor applied, so same-type offset aliasing reads the allocation prefix, different-type aliasing enters copy sinking with the allocation’s element type, and same-type reshaping may enter reductions with wrong dimensions; the fix introduces a unified view-acquisition helper, using checked reinterpret casts for same-type non-identity views, and single-block byte allocation plus typed view operations for multi-element-type accesses to preserve common storage identity. Second, #1960 fixes views dropping the source offset: when rebuilding a buffer from a data variable, the view lost the element offset, so offset aliasing started from allocation offset zero rather than the source’s first element address; the fix preserves the offset for same types, converts by element width when the byte-level type changes and requires divisibility, and rejects unaligned or unsupported conversions outright, without introducing storage or data copies. Third, #1961 changes unary vector broadcast inference to use the access region: previously broadcast dimensions were inferred from the original buffer shape, so writing an 8-by-8 source into rows 8 to 16 of a 24-by-8 target failed; the fix infers from the generated memory reference type, preserving the access segment’s shape and existing compatibility checks. All three include independent clean-baseline builds and bitwise verification notes.

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

2.7 Ascend under review: Fix for bitwise XOR unsigned-byte workspace view in the PTO backend (#1959)

Entered review at 16:36 on 10-29, closing #1721. In the PTO backend, the temporary workspace for bitwise XOR was passed to the instruction as a raw 1D byte region rather than the required 2D view: at generation time the destination’s 2D shape was copied into the workspace’s shape info but not marked as a slice, so a 1D block descriptor was resolved for unsigned bytes (the workspace does no type reinterpretation), and the instruction read and wrote out of bounds at runtime. The workspace itself is a 1D region allocated by source element count, and its capacity suffices for aligned 2D blocks; the failure was in the descriptor, not the capacity. The fix inherits the destination shape, then forcibly constructs a 2D view and marks it as a slice; tests add multiple coverage groups for unsigned and signed bytes (2D automatic temporary, explicit 1D temporary, 1D, aligned non-zero-offset slice with sentinel, signed bytes), and retain the 16-bit regression.

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

2.8 Hygon: Two HCU PRs under review updated, community DCU fused-operator repo appears; no pushes from Moore Threads or MetaX in the window (10-29)

On the Hygon side, still under review: the shared-memory vertical interleaving feature branch for ordinary copies (#30) and the distributed one (#19) both saw updates in the window, while the main branch’s latest commit remains the buffer-grouping feature from 10-27. On the community side, a DCU-oriented repository appeared the same day, preserving Hygon-platform fused operators together with changes to the TileLang compiler, mostly C++, about 15 MB, with attention operator support and variable-length support visible in the history; it is a personal practice repo rather than official Hygon. The Moore Threads adaptation repo (tilelang-musa) and the MetaX adaptation repo (tilelang-metax) saw no pushes in the window, with their latest on 09-30 and 10-26 respectively.

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

Source: https://github.com/SealSoul117/TileLang_DCU_FuseCmpComm

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

3.1 TileOPs: Release image merged; operator library and TileFoundry packaged into the same image for the first time (#2793)

Merged on 10-29 at 11:18, 9 files, +123/-49. The runner image gains a new release target: on top of the development image, it installs the PyPI release versions of the operator library and the scheduler, and names the image after both version numbers; the image tag is tileops0.1.0 plus tilefoundry0.0.3 plus three environment identifiers — cuda13.2, torch2.13, and tilelang0.1.12 — published under the operator library release path in the container image registry. An accompanying script pulls the image and runs stack checks, baseline checks, packaging tests, and optional benchmark files on the GPU; because the wheels do not include tests or benchmarks, these are all taken from the tag corresponding to the image’s own version. The docs update the quick start, and the runner docs add build, verification, and push steps. The sgl-kernel baseline is also removed: no image installs it anymore, and the related benchmarks had long been skipped.

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

3.2 TileOPs under review: GEMM recomputes grid grouping and resident block count under a power cap; q-b DRAM reads drop from 168 MB to 94 MB (#2794)

Entered review on 10-29 at 21:44, 2 files, +35/-12, following up on the performance issue for back-to-back calls. Two changes: first, when A fits in one L2 cache partition but B does not, grid grouping lets one group span all rows, so each B block is read from memory only once instead of once per 16-row group; on the q-b shape of DeepSeek V3 prefill (4096 by 24576 by 1536, NT layout), memory reads drop from 168 MB to 94 MB. The split boundary for narrow blocks remains constrained by 16-row groups, and the selection is unchanged. Second, the resident block count on width-uniform dual-cooperative-group grids is changed to a ceiling resident count computed from the block count, total block count, and processor count, rather than using the processor count directly, because the busiest block must run the same number of blocks regardless, while each wave can read fewer distinct panels; ping-pong, narrow-block, and mixed-width grids still use the processor count. Tests run the operator suite on H200, CUDA 13.2, PyTorch 2.13, and TileLang 0.1.12: 212 passed, 24 skipped.

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

3.3 TileFoundry: AMX runtime operators converted to 39 static TIR contracts, a breaking change (#279)

Merged on 10-29 at 15:59, a breaking change (the commit title carries an exclamation mark). The motivation is to give AMX an explicit, typed operator declaration that agents can match and evaluate from the same computation description. The approach describes 24 AMX runtime operator families as 39 statically registered contract versions, with parameter domains, operand and layout patterns, side-effect sets, and target capabilities M1 through M4; it adds a Python evaluator and a raw C simulator as the gold standard, plus 59 representative atomic cases covering all 39 static operators, with high-level and mid-level IR cases running real 2-, 4-, 8-, and 12-thread domains and verifying candidate acceptance and rejection; high-level candidates are discovered from in-operator match declarations, preserving target-specific capability inference. Naming is unified under the AMX namespace, aliases are held by the registration entry point, and registered names are given by a unified query; CPU execution uses thread topology, with private logical AMX register state, and the thread count is not the AMX engine count. Validation is 46 focused tests passing, including one comparison against 57 raw C results and two lifecycle checks.

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

3.4 Adopters: TileKernels and FlashQLA main branches saw no pushes during the window (10-29)

The main branches of both known adopters are stalled at 10-14, with no new pushes during the window: DeepSeek’s operator repo (2,018 stars) and the Qwen-side FlashQLA (742 stars). Today’s ecosystem signals come mainly from community-built repos rather than these two internal users.

Source: https://github.com/deepseek-ai/TileKernels

Source: https://github.com/QwenLM/FlashQLA

3.5 A standalone community operator repo, ops-tilelang, appears on the Ascend side, with pushes during the window (10-29)

The repo was created on 10-20 and had a push on 10-29 at 13:06. It is positioned as a community high-performance operator repo for Ascend NPUs, managing kernel implementations, automated test cases, standard workloads, and performance baselines in a unified way, with the goal of making operator assets reusable, verifiable, and sustainably evolving. It sits under an Ascend ecosystem fork organization rather than the tile-ai organization, indicating that operator supply on the Ascend side is spilling over from the official repo into a layer of community collaboration; its relationship with the official tilelang-ascend is not yet clear, so it is recorded here only as an ecosystem observation, not as an official action.

Source: https://github.com/Atlas-Ascend-Fork/ops-tilelang

3.6 PTO-ISA’s two companion repos active: puzzle repo migrates to PTOAS 0.70; benchmark repo adds SV2 tiled scaling and per-column analysis (10-29)

Both companion repos under the PTO-ISA organization had commits during the window. The puzzle repo migrates PTO’s broadcast load to PTOAS 0.70; the prior week was spent clarifying quantized scaling signs and rewriting the problem statement to “point at your own solution rather than the answer.” The benchmark repo adds two analysis commits: SV2 tiled scaling and per-column analysis, and one piece of evidence (the in-library absolute maximum instruction’s variant count is 31.39, with runtime instruction throughput and roofline analysis, explicitly not a unified buffer metric). These two repos serve as teaching and measurement companions to the official repo and are worth citing as corroboration of Ascend-side methodology.

Source: https://github.com/PTO-ISA/tilelang-puzzles-ascend

Source: https://github.com/PTO-ISA/tilelangGym

3.7 Observation: a tilelang fork repo under the T-HEAD name had pushes during the window, with branch names pointing to PPU matrix instruction layouts (10-29)

Under the T-HEAD organization (whose official site is the developer portal domain) there is a tilelang fork repo, created 2026-09-22, forked from the upstream main repo, with 13 public branches. Besides the main branch and version mirrors v0.1.11 through v0.1.14, there are visible branches for a PPU matrix shared layout fix, a PPU language package, a runtime compilation feature, and several bug fixes; the repo’s main branch last commit is at 10-22, but the repo push time is updated to 10-29 19:27 (Beijing time). The repo has no README and no releases, so it can only serve as an observational lead: a full-fledged team is doing fork development with TileLang as the base of its own hardware stack, in a direction related to PPU matrix instruction layouts. Today’s #3705 bug report was reproduced precisely with a build from this repo’s source.

Source: https://github.com/t-head/tilelang-for-sail

IV. Community, Tutorials, and Events

4.1 Documentation site adds reduction applicability boundaries to the instruction page, synchronized with the redux merge (10-29)

The documentation site saw five automated updates within the window, including a new reduction section on the instruction page aligned with today’s merged hardware reduction changes: CUDA’s cross-thread reduction delegates full warp groups that are warp-aligned with unit thread stride to a helper sharing the same origin as warp reduction; the helper selects hardware reduction instructions for integer sum, max, min, and bitwise operations on SM80 and above, promoting smaller integers to 32 bits; scalar FP32 max/min are used only on sm_100a, sm_100f, sm_103a, and sm_103f, with scalar FP16 and BF16 promoted to FP32; the same rules apply to the final segment of larger reductions and to unpacked batched reductions. Sub-warp, strided, packed, FP64, and unsupported type-operator combinations retain the shuffle implementation, and cross-warp shared-memory exchange and barriers are unchanged. Documenting the fallback boundaries is more useful to users than merely stating “hardware reduction is supported.”

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

4.2 Dev-tools repo completes three follow-ups after extracting the sinking tracker, aligned with the upstream tooling change under review (10-29)

After completing the extraction of the sinking tracker from the main repo on 10-20, the standalone dev-tools repo added three follow-ups within the window: avoiding circular imports when the tracker is active, pinning CI to a version that supports the instrumentation plugin, and initializing TileLang before loading core compatibility aliases. This aligns with the upstream change under review to “extract the sinking tracker into the dev-tools repo,” indicating that tooling is no longer just an upstream branch experiment but has become an independently releasable package.

Source: https://github.com/easy-tilelang/tilelang-devkit

4.3 New repo tilelang-knowledge: four commits in one day — 33 operator design cases, 20 sourced knowledge cards, 7 optimization patterns (10-29)

The repo was created on 10-29, and its four commits that day form an agent-oriented knowledge base: first the repository structure, then 33 operator design cases (each with a fixed source), then 20 sourced programming knowledge cards, and finally documentation of 7 evidence-backed GPU optimization patterns. The repo’s stated purpose is to let an agent facing an operator determine which experiences apply, how the code should be changed, how to verify it, and where to look after a failure. This is the same kind of attempt as TileFoundry’s operator contracts and the Ascend repo’s profiling skill: turning “how to change, how to verify” into searchable assets.

Source: https://github.com/qiuuuuu622/tilelang-knowledge

4.4 Illustrated-course companion repo and distributed-layout experiment repo created the same day (10-29)

Two personal repos were created the same day: an illustrated-course companion repo, described as using original scenarios and animations to explain each knowledge point thoroughly, released alongside videos; and a distributed-layout experiment repo, whose commit message says it adds distributed TileLang layout experiments. Together with the model-porting repo and the Hexagon backend repo that had pushes today (repo push time updated to 10-30 06:10, with the latest commit on its main branch still the 08-24 dependency upgrade), community self-driven investment continues in both education and dialect backends.

Source: https://github.com/zachshi-ai/tilelang-course

Source: https://github.com/Leewin239/tilelang-distributed-layout

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

The six queries covered exact terms plus two days, Chinese plus Ascend, Chinese plus operator, organization name plus two days, English plus GPU, and Chinese plus compiler plus seven days. Only Chinese plus Ascend returned one result — a repost on clinical AI progress unrelated to the operator language — and all others returned nothing. Relevant Hacker News hits over the past four days all fall outside the window (the most recent being a 10-02 repost discussion of the Ascend programming stack). An arXiv search for TileLang found no new papers. With the news side showing zero hits for several consecutive days, the daily report continues to rely on repos and ecosystem repos as its main sources.

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

V. Trend Observations

The rhythm on the Ascend side has shifted from “porting the backend” to “filling in operators.” Yesterday’s entry was the NPUIR backend merging into main, with acceptance criteria of 504 self-tests and 71 versus 72 examples; four of the five commits merged today are operators and examples, arranged according to upstream model requirements: GLM-5.3’s indexer compression, the inner product of block-gated attention, four-bit quantized matrix multiplication, and depthwise convolution for quantized vision models. What these operators share is that quantization and fusion are written into the kernel itself (E4M3 pure-integer encoding, four-bit packing conventions, dual-kernel plus 1×1 fast path) rather than left for host-side remediation. At the same time, two commits establish engineering standards: operator performance profiling enters the process, and unstable test cases caused by concurrent batches move to a checklist-based test channel. This shows that the main contradiction on the Ascend side is shifting from “can it compile” to “can the operator’s numerics and performance be stably accepted.”

The main repo’s three commits today (hardware reduction, concurrent compilation, out-of-bounds copy grouping) correspond to instruction selection, compile latency, and emission quality — all compiler-core issues rather than language-layer features. This is consistent with yesterday’s new-defect direction: the boundaries of static checks such as out-of-bounds, aliasing, and bit width are being opened one by one. Two of today’s three new defects (batched copy with padded buffers, view bit-width overflow) were triggered on real model shapes, indicating that TileLang’s usage has reached a scale where boundary conditions are encountered: NVFP4 weights of 248320 by 5120, and attention head dimensions where K is not a multiple of 16.

The most notable ecosystem development today is “fork-ization.” An independent operator repo has appeared in the Ascend community; T-HEAD has a TileLang branch repo with a PPU branch; PTO-ISA has turned both its problem sets and benchmarks into companion repos; and another dev-tools repo releases the sinking tracker as an independent package. The common premise of these moves is treating TileLang as the foundation of one’s own hardware stack: needing to change layouts for one’s own instruction set, perform acceptance testing with one’s own test system, and accumulate experience in one’s own knowledge format. For a domain-specific language, this is a harder adoption signal than star counts, but it also means upstream must begin to face the long-term cost of “how branches synchronize with upstream”: today’s #3705 defect is exactly what branch users hit first in real builds.

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 / nightly mirror repos) Pushes within window: tilelang, TileOPs, TileFoundry, tilelang-ascend, tilelang-hygon, tilelang.github.io, TileOPs.github.io, TileOPs-nightly; tilelang-metax (10-26), tilelang-musa (09-30), TileRT, tilescale, DeepStack, tilelang-puzzles had no main-branch commits within window
Merged ground-truth search (is:pr is:merged) 12 PRs: tilelang 5, tilelang-ascend 5, TileFoundry 1, TileOPs 1
Open ground-truth search (is:pr is:open updated) 54 PRs, including the Ascend docs review batch, three NPUIR view fixes, and a GEMM scheduling change in TileOPs
New issue check 4 issues: #3704, #3705, #3707 are defects; the other is a daily test-bot failure notification in the Ascend repo
Main-branch commit-level check tilelang 5, tilelang-ascend 5, TileOPs 1, TileFoundry 1, docs site 5; tilelang and the two repos match their merged 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 (10-27); TileFoundry latest is 0.0.4rc1 (10-25); tilelang-ascend latest is v0.1.2.000 (09-09)
Third-party and ecosystem repos 10 repos checked one by one: Ascend community operator repo, the two PTO-ISA repos, devtools repo, knowledge repo, illustrated-course repo, distributed-layout repo, DCU fused-operator repo, model-porting repo, Hexagon backend repo, and T-HEAD fork repo all had pushes within window (Hexagon had no new main-branch commits, only repo push timestamps updated)
Adopters TileKernels and FlashQLA main branches stopped at 10-14, no pushes within window
Google News (six queries, Chinese and English) Only one Chinese repost on clinical AI plus Ascend, unrelated to the topic, not included
Hacker News All relevant hits in the past four days fall outside the window
arXiv No new papers found searching for TileLang

Full Source List