TileLang Weekly Report (2026-10-26 ~ 2026-10-30)
Reporting period: Monday, October 26, 2026 to Friday, October 30, 2026, 5 working days in total Sources: 5 issues of this publication’s daily TileLang community update (10-26, 10-27, 10-28, 10-29, 10-30); the combined collection window of the five issues spans 10-23 07:00 to 10-30 07:00 (the Monday issue used a 72-hour window covering the previous weekend; Tuesday through Friday used 24-hour windows); supplemented by a 168-hour window cross-check of repositories in the tile-ai organization, plus verification of each repository’s releases page and tags and the VERSION file (see Appendix)
I. Weekly Highlights
-
TileOPs 0.1.0 officially released, the first formal release with 206 operators, public API frozen as of this release (10-27): The PyPI artifact was uploaded at 07:45 on 10-27 (Beijing time), and the GitHub release page went live at 13:15, marking the repository’s first formal tag. The release notes give three hard metrics: all 206 operators have manifest entries (signatures, test and benchmark workloads, roofline formulas, validated item by item in CI); on H200, 2262 of 2265 nightly benchmark rows are no slower than the fastest reference implementation, with a geometric mean of 1.37x, against references including torch.compile, FlashAttention 2 and 3, FlashInfer, FlashMLA, vLLM, DeepGEMM, flash-linear-attention, FlagGems, and QuACK; of the 206 operators, 195 foundry-tagged merges had kernels generated by TileFoundry and entered the repository through the same CI. The release also freezes the public API and adds an API compatibility check script, so any PR that would break calls written against the previous version fails immediately; from here on, 0.x only adds, never changes. Four known gaps and a three-step roadmap were published the same day, the latter pointing to fused operators, whole models, and compute-communication fusion. Sources: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0, https://pypi.org/project/tileops/
-
The entire NPUIR backend for Ascend 910 A2/A3 merged into the mlir-ascend mainline, with acceptance criteria written as numbers for the first time (merged 10-28): A single change of 55 files, +21490/-1. On the Python side is the tilelang/ascend910 package (backend, adapter, codegen, compiler, execution_backend, jit, pipeline, runtime, target, 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, and three passes: LowerNpuirBlock, NpuLoopVectorize, and RewriteFDivToMul; the runtime adds npu_launcher.h and npu_utils; and the target alias accepts ascend910, ascend-910, a2, a3, and npuir. The author’s acceptance criteria: all 504 pytest testing/npuir items pass; 71 of 72 examples pass with 0 timeouts and 0 errors, and the single failure also fails on the same machine with the old backend, being an environment and tolerance issue. Source: https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
-
The TIRx refactor under review in the main repo replaces the built-in TVM with an external dependency, growing from 1043 files to 1194 files (from 10-25, updated all week): The plan is to delete 3rdparty/tvm and its submodule entries, migrate the frontend, native tensor and scoped IR, lowering passes, backend pipelines, and execution adapters wholesale to a separately built and released TIRx, and pin versions against unpatched upstream source (apache/tvm-tirx 9af5544d, depending on apache-tvm-tirx==0.1.0a2 and apache-tvm-ffi==0.1.15.post0); the integration patches previously carrying a +tilelang.1 suffix are no longer needed. The PR was updated three times from the weekend through midweek, growing from 1043 files, +71428/-104176, to 1194 files, +106648/-104497 — the largest item in the week’s queue, and the one that determines how future TVM-side fixes reach the mainline. As of Friday it was still under review and not merged. Source: https://github.com/tile-ai/tilelang/pull/3666
-
The main repo’s correctness hardening this week centered on the “compiles, runs without error, wrong result” family (10-26 to 10-30): Typical items include the output allocation stream (#3649, where three adapters accepted the caller’s stream but allocated out_idx outputs on the current stream, and the caching allocator reused buffers ordered by allocation stream, corrupting outputs in three of five runs on H100, reduced to zero after the fix); the all-warp reduction switched to redux.sync (#3692, turning a 32-lane FP32 max from 20 shuffles into 4 hardware reductions); static bounds for slice fills and negative indices (T.fill in #3352 writing As[6:12] on a length-8 buffer was previously allowed and actually wrote to 6 through 11; A[-9] in #3347 previously wrapped silently to -1); T.ws(-1) regions running on zero threads (#3354, with all stores discarded and no diagnostic); and seven default math functions on ROCm no longer lowering to fast builtins (#3621, cutting sin/cos error on MI300X from 7.8e-04 to 6.6e-08). Sources: https://github.com/tile-ai/tilelang/pull/3649, https://github.com/tile-ai/tilelang/pull/3692, https://github.com/tile-ai/tilelang/pull/3621
-
On the Ascend side, the operator list is being rolled out one by one against upstream model structures (10-26 to 10-30): Items landed this week include situ_glu aligned with the Kimi K3 Stable LatentMoE gated activation (#1891); the kpool compression operator for the GLM-5.3 sparse attention indexer (#1910, a single-kernel pure-Vector approach where E4M3 encoding is implemented with pure integer bit operations, relying on no FP8 type and doing no approximation); the Chunk Scaled Dot KKT forward pass for chunked gated linear attention (#1803, statically unrolled into four blocks with zero GM workspace, matching at 1.0000 across three cases each at L0 and L1); the root cause and fix for W4A4 int4×int4 GEMM producing entirely wrong results (#1731, where the nibble-packed type was treated as 1 byte, skewing all copy strides and repeat counts by 2x, with symptoms of 99.6% mismatch); and the FP4×FP4 packed input for the main repo’s DeepGEMM example (#3589). Sources: https://github.com/tile-ai/tilelang-ascend/pull/1910, https://github.com/tile-ai/tilelang-ascend/pull/1803, https://github.com/tile-ai/tilelang-ascend/pull/1731
-
Four mainline moves for Hygon this week, shifting from “it runs” to “laid out according to hardware structure” (10-26 to 10-30): The register pipeline plus warp divergence merged (#18, 28 files, +5312/-64, converged over more than two weeks, with the author concluding that example_gemm reaches performance comparable to rocblas); native B16 storage and dynamic WASP GEMM merged (#28, 6 files, +340/-230); the scale buffer split into independent dual banks separate from LDS banks (#29, 12 files, +1389/-314, the largest software-side item in this window); and the LDS swizzle feature branch for the ordinary copy path is under review (#30, 18 files, +2000/-84, including five new tests). Sources: https://github.com/tile-ai/tilelang-hygon/pull/18, https://github.com/tile-ai/tilelang-hygon/pull/29
-
The release image bundles the operator library and TileFoundry into the same image for the first time (10-29): The runner image installs the PyPI releases of TileOPs and the scheduler on top of the development image, tagged tileops0.1.0 plus tilefoundry0.0.3 plus three environment identifiers: cuda13.2, torch2.13, and tilelang0.1.12; an add-on script pulls the image and then runs stack checks, baseline checks, packaging tests, and optional benchmarks on the GPU — since the wheels ship without tests or benchmarks, these are all taken from the tags corresponding to the image’s own versions. The sgl-kernel baseline is removed at the same time. Source: https://github.com/tile-ai/TileOPs/pull/2793
-
The ecosystem is seeing full-fledged forks and community repositories, with TileLang treated as the foundation of proprietary hardware stacks (10-26 to 10-30): The tilelang-for-sail fork repository under T-HEAD has 13 public branches, showing, beyond version mirrors, PPU matrix shared-layout fixes, a PPU language package, and runtime compilation features; the CANN community’s ops-tilelang operator repository, created 10-20, has been pushing continuously and has already strung together a model fragment of DeepSeek V4.1 Flash with seven operators; the Qualcomm Hexagon backend saw 15 commits in a single day, laying out HVX vectorization, the VTCM arena, Crouton layouts, and DMA builtins all the way through to mapping and reduction; and a RISC-V community hub repository, the PTO-ISA problem-set and benchmark repositories, and the tilelang-knowledge repository advanced the same day. Sources: https://github.com/t-head/tilelang-for-sail, https://github.com/Atlas-Ascend-Fork/ops-tilelang
II. Core Project Progress (main repo tilelang)
Versions and Releases
- The main repo’s latest tag and release remain v0.1.15 (released 09-30); no new tags or release-page entries all week, and as of this Friday it has been 30 days. VERSION file check still shows 0.1.15.
- TileOPs v0.1.0 is the only new formal release in the org this week (10-27), and the second new tag since 09-30 after TileFoundry 0.0.3. TileFoundry’s latest remains 0.0.4rc1 (10-25), not yet promoted to formal.
- Adaptation-line tags unchanged all week: tilelang-ascend v0.1.2.000 (09-09), tilelang-metax v0.1.15 (10-16), tilelang-musa v0.1.15+musa.1 (09-30), tilelang-mlir-ascend still on the v0.1.2.020 series.
- Merge volume and cadence: Monday’s issue covered a 72-hour window (10-23 07:00 to 10-26 07:00) with 78 merges (TileOPs 51, tilelang 10, TileFoundry 10, docs site 4, Ascend repo 1, Hygon repo 1, Sunrise 1); it overlaps the previous weekly (10-19 to 10-23) on the single day of 10-23, so it is not summed with the following four issues. The four 24-hour windows from Tuesday to Friday totaled 86 merges: tilelang 47, TileOPs 26, tilelang-ascend 5, TileFoundry 4, tilelang-hygon 3, tilelang-mlir-ascend 1; daily counts were 30, 20, 24, 12, a converging pattern of high-then-low. Open-review updates held between 45 and 76 all week; new issues were 19, 11, 7, 6, 4 respectively.
Compiler, Language Layer, and Execution Paths
- Full-warp reductions switched to hardware redux (#3692, merged 10-29, 4 files +337/-35): previously even when the whole group was one aligned full warp, ordinary cross-thread reductions still expanded into a shuffle tree. On main branch 837486a, four independent 32-lane FP32 max reductions compiled to 20 shuffles and zero redux on sm_103a with CUDA 13.1; after the change it is 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. Integer reductions use hardware redux on SM80 and above; FP32 max/min is only available on sm_100a, sm_100f, sm_103a, sm_103f; narrow floats are promoted to FP32 before entering the helper, and FP64 stays on the shuffle path without precision loss. The author explicitly does not claim end-to-end speedup, only instruction-level comparison.
- Autotuner’s independent JIT configs now compile concurrently (#3476, merged 10-29, 2 files +54/-46): the original implementation held one lock across the whole compile phase, serializing even configs using only default compile options, with frontend, lowering, and device compilation all queued. The change copies an independent JITImpl per override config via dataclasses; default options still take the original lazy path to preserve kernel cache and object identity; only candidates sharing the first-stage key are serialized, and debug-output writes to disk remain serialized. Cross-backend validation data: Ascend 1797 passed (Python 3.12.14, torch_npu 2.10.0, CANN 9.3.0, AscendC examples 50/50 passed, PTO examples 44 passed 6 skipped), ROCm 7.2 on gfx942 3220 passed, Metal 95 passed, CUDA on H100 PCIe cu130 5500 passed, CuTeDSL 75 passed 17 skipped.
- Ascend buffer-view layout and pointer units no longer propagate across aliases (#3623, merged 10-29, 23 files +1139/-438): multiple buffers created by T.view or T.reshape can share the same storage while requiring different layouts; previously layout was resolved uniformly by the shared data variable, potentially selecting the wrong Cube mapping or DMA address. This PR attaches layout to the individual Buffer view, threading through layout inference, fractal storage normalization, multi-buffer materialization, and copy and matmul lowering; raw pointer offsets are remapped by their own element units. On the same versioned allocation, Cube views can now use different MajorK and MajorMN layouts, provided logical strides are dense, base offset is zero, operand types are compatible, and padded storage sizes are equal.
- GM-to-L1 copy and out-of-bounds padding merged into one scheduling task (#3543, merged 10-29): both prepare the same L1 data; previously padding could be split into a separate scheduling task, adding an extra MTE2 pipeline fence on sync insertion. Now consumers wait on the complete data including padding; FP8 MX K-tail padding is likewise merged in and retains in-task conditional padding, and the former MX deferred-padding queue is removed. 47 focused tests passed.
- A set of ROCm and execution-context fixes: seven default functions exp, exp10, log, log2, log10, sin, cos no longer lower to fast builtins on HIP regardless of fast_math setting, and eight underscored fast operators gain minimum lowering rules (#3621); auto target reads the current HIP device instead of hardcoded device 0 (#3594); NVRTC and CuTeDSL take the current stream from the device of the first input tensor (#3359); profiler uses the current device when none is specified, and cache flushing no longer hits device 0 (#3357).
- Minor language-layer and lowering fixes: TMA gather/scatter Python integer row indices converted to int32, so literal forms no longer crash compilation on Blackwell (#3365); pipelined epilogue with runtime trip count now unrolls directly, restoring compilability of paged decode (defect reported downstream by TileOPs#1852) (#3356); VectorizeLoop no longer silently ignores loop step (#3345); T.clamp’s NaN-propagating lowering reverted to min/max semantics (#3628); T.Parallel’s coalesced_width uniformly converted to IntImm (#3374); TMA descriptors record only validated names (#3484); handle suffixes on non-restrict parameters no longer stripped (#3366).
Correctness and Bug Fixes
- Stream and cache identity: after outputs allocated on the caller’s stream (#3649) and NVRTC/CuTeDSL stream capture (#3359), this window brings autotune cache keys incorporating compile_flags (#3338, same kernel with different compile options no longer reuses the old binary), and on the Ascend side Bisheng build identity incorporated into the binary cache key (#3691). Two independent lines doing the same thing in the same week: if the cache is keyed only by source and shape, changing compilers causes false hits.
- Synchronization and memory planning: ThreadSync records address_of as a write and adds fences after opaque calls (#3346; previously T.address_of(B[0]) form had no fence after extern, while T.access_ptr(B, “w”) form did); T.ws(-1) changed to run on all threads (#3354); cross-storage-domain named barrier ID collisions (#3643) and missing CTA barrier on sm_100 TMEM release (#3645 / #3663 / #3664) handled early in the week. Wednesday surfaced three ThreadSync misplaced-fence defects (#3680): fences placed inside divergent branches, fences in branches incorrectly clearing outer pending dependencies, and fence placement in thread-count loops, with two fixes under review in #3682 and #3681.
- Static bounds and types: slice fill validated by region end rather than length, negative indices rejected when still negative after wraparound (#3352 / #3347); SIMT atomic region shape validation added (#3465); sparse MMA K tail block no longer silently skipped (#3367); atomic_addx4 target dtype narrowed to float* (#3369); profiler output index uniformly validated (#3486); scalar-read index coefficient no longer mixed into lane count (#3453, compile crash with row width above 32768 closed); batched all-reduce resolves thread range via scalar branch (#3341, all 2048 outputs wrong when batch>1); int64 atomic max/min and AtomicStore int64, bf16, fp16 crashes (#3307); reduction FP8/FP4 non-finite constants emitting invalid literals (#3189). Thursday moved six tests that assumed cold compilation to without_kernel_cache (#3695 / #3696).
- Direction of newly opened defects: all three opened Thursday fall on static-check boundaries — padded shared buffer via batched copy reports shape mismatch (#3704: global tensor with last dim not 16-aligned copied into padded shared buffer, BK taken as cdiv times 16, lowering reports 112 vs 100), negative-scale iteration graph miscompiles reverse index (#3705, repro environment is v0.1.15 and T-HEAD branch repo source build, a shared upstream path), view bit width multiplied as int32 overflows (#3707: Qwen3.8’s lm_head weight 248320 times 2560 bytes NVFP4 totals 5,085,593,600 bits, wraps to 790,626,304 and then throws literal out of range). Same batch also has 13 example failures on sm_120, exponential IfStmtBinding compilation, without_kernel_cache not forcing recompile (#3701 / #3700 / #3699), and CuTeDSL CUDA graph capture producing an empty graph (#3693). Three fixes entered review Friday: bit width to int64, CPU scalar GEMM clearing only the target region, bypassing two in-process caches when cache is disabled (#3708 / #3709 / #3710).
Runtime, CI, and Toolchain
- CI pins ROCm and Metal jobs back to compatible PyTorch wheels (#3706, merged 10-29): ROCm CI exposes a matching PyTorch ROCm source via an index variable, but the workflow-level “unsafe best match” strategy let a later PyPI candidate win, resolving to PyTorch 2.15.0 (CUDA 13.2) instead of 2.14.1 plus ROCm 7.2 build; the same 2.15.0 rollout also made the macOS 26.6.2 ARM64 runner reject the Apple paravirtual GPU. Three fixes: set backend variables from the configured stable ROCm version, use the first index for nightly installs, and add a pre-test fast-fail validation step (confirm HIP build, not CUDA build, and runner GPU accessible).
- Ascend CI splits AscendC and PTO into two matrix jobs, rerunnable separately (#3687); CI image bumped to the cann9.3.0 weekly build (#3630).
- Four contribution-process governance changes (10-26, all about review cost): CONTRIBUTING.md states PR descriptions should cover problem, approach, and trade-offs, with evidence placed next to the claims it supports (#3671); docs-only PRs skip CUDA, ROCm, Metal, Ascend, and CuTeDSL builds and tests, with a deliberately narrow whitelist (#3672); the commit skill gains a confirmation step that first presents the proposed title to the contributor (#3673); the review bot’s auto-summary and docstring coverage checks are disabled, requiring findings to be anchored to the current diff and actual call paths (#3674).
- Language service ships experimental CUDA debugging (10-26): tilelang-lsp pushed its first debugging commit, adding a tlsp.debug launcher for Python line mapping, compiling with -G and -lineinfo, isolating cache and capturing generated source, wiring Nsight’s cuda-gdb into VS Code, with support for Python breakpoints, pre-launch checks, and browsing generated CUDA source. This is the repo’s first substantive commit since August.
- Docs site continuously auto-synced all week: eight on 10-26, six on 10-27, two on 10-28, plus routine updates to the TileOPs docs site and localization repo.
Review Queue Observations
- The largest item remains the TIRx refactor (#3666, 1194 files +106648/-104497), applying zero patches against a pinned upstream baseline; it is the only change in the queue that will determine “at what cadence subsequent TVM-side fixes enter mainline.”
- The Ascend repo’s a2a3 scheduler migration spans 100 files +18453/-666 (#1951), the largest under review in the dedicated repo, indicating the next phase’s change volume is in the scheduling layer rather than the operator layer.
- The PTO line’s review queue advanced steadily through the week: T.Parallel vectorization, AutoSimVF channel, batched 3D ND2NZ, write-count multi-buffering (#3512 / #3320 / #3689 / #3629), PTO target defaulting to the tvm_ffi execution backend (#3516), plus three view and broadcast fixes (#1962 / #1960 / #1961) and an unsigned byte workspace view fix for bitwise XOR (#1959).
- Open-review updates ranged between 45 and 76 all week, rebounding to 54 on Friday, in contrast with the decline in merges: this was a “clearing week” rather than an “opening week.”
III. Multi-Backend Adaptation (Ascend / Hygon / ROCm / Hexagon)
Ascend: Backend Merged into Mainline, Operators Rolled Out by Model
- Backend level: The NPUIR path for 910 A2/A3 is merged into the mlir-ascend mainline (#207, 55 files +21490/-1); the mode switch TILELANG_ASCEND_MODE takes expert or developer, defaulting to expert. 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 followed by explicit conversion) and copy sinking for dynamic extents.
- Scheduling and addressing: The main repo adds T.annotate_buffer_addresses, allowing precise byte-level addressing of allocations in the UB, L1, L0A, L0B, L0C, and SS pools within the post-scheduling Ascend pipeline. Buffers placed into addressed allocations must first declare a fixed number of versions via T.annotate_buffer_versions; validation covers missing or unsupported version counts, overlaps, misalignment, dynamic sizes, and capacity overflow (#3607, 14 files +982/-26). PerCoreTask candidates now allow internal if/else, so conditional L1 fills are no longer rejected wholesale (#3686); hand-written pipelines preserve configuration register phases (#3624).
- Operator supply: Merged within one week: situ_glu (#1891), kpool compression for GLM-5.3 (#1910), Chunk Scaled Dot KKT forward (#1803), root-cause fix for W4A4 int4×int4 (#1731), and a resubmission of the depthwise_conv_2d example (#1920, now using a manifest-based test channel to avoid concurrent batches). Still under review are two examples, GroupNorm and fla_chunk (#1607 / #1953 / #1952), and three NPUIR view and broadcast PRs (#1962 / #1960 / #1961). What these operators share is that quantization and fusion are written into the kernel itself rather than left for host-side remediation.
- Engineering practices: A new operator performance profiling skill uses msprof op output to determine optimization direction (#1931); the installer’s SHMEM section is fixed on both fronts so empty-shell wheels no longer falsely report success (#1943 / #1946); GM-to-UB copies of more than two dimensions fold leading dimensions by DMA row count (#1805); the Cython wrapper rebuilding its analyzer on every call (about 0.28 ms) was first logged (#1945) and then fixed by the lazy construction now under review (#1948).
- Regression status: The daily full test on 10-25 reported all 2903 items passing; the test job in the early hours of 10-26 reported failure (#1944); 10-27 and 10-28 both passed all 2903 items (#1950). The test suite is now large enough to serve as a gate, with conclusions valid on a weekly timescale.
- Compatibility changes under review: Carver reads the A3 resource budget according to the selected SDK, avoiding placing a 256×256 FP32 accumulator on an L0C with only 128 KiB (#1955); the AI Core Exception Dump dependency is lowered to CANN 9.2.0 and four required fields are filled in (#1956); documentation switches to T.Tensor and deprecates the T.Buffer DSL interface (#1957). The four fact-check items in the Ascend repo’s documentation reached up to 46 rounds of single-comment review (#1628 / #1655 / #1656 / #1659).
Hygon: Four PRs in a Week, Laid Out by Hardware Structure
- The register pipeline can now cooperate with the existing shared-memory pipeline; corresponding commit and barrier instructions are inserted per pipeline stage, and warp divergence is enabled when the warp count exceeds 4 (#18). This PR was opened on 10-08 and converged over more than two weeks; it was merged as a large merge (mainline first, then branch), so its line count far exceeds other changes in the same batch.
- Native B16 storage and dynamic WASP GEMM (#28): preserve MLS Let parameter emission order, propagate the GEMM warp partition to matrix storage, reject packed B16-to-B32 output aliasing, and make WASP dimensions dynamic with alignment inferred per axis.
- Row planning for the scale buffer is changed to an independent dual-bank partition, separate from LDS banks; the same row mapping is handed to allocation, scale copy, and MMAC sinking alike, avoiding three separate layout derivations (#29, 12 files +1389/-314). This PR follows #28; Hygon completed three mainline merges in two days.
- The LDS write phase and swizzle strategy for the ordinary copy path are under review, covering both b8 and fp8 data widths (#30, 18 files +2000/-84, five new tests).
Sources: https://github.com/tile-ai/tilelang-hygon/pull/18, https://github.com/tile-ai/tilelang-hygon/pull/28, https://github.com/tile-ai/tilelang-hygon/pull/29, https://github.com/tile-ai/tilelang-hygon/pull/30
ROCm: pip-Based SDK and Floating-Point Semantics
- Support for the pip-based ROCm 10 SDK merged (#3549, 2 files +106/-19): ROCm 10.x’s rocm/pytorch images ship ROCm as Python wheels, 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 assertions failed outright on MI355X (gfx950), and setting ROCM_PATH had no effect. The fix resolves in the order of the caller’s prefix, find_rocm_path(), torch’s gcnArchName, and a gfx900 fallback, adds a 10-second timeout to each rocminfo probe, and adds a Dockerfile based on rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12. The author measured architecture detection changing from gfx900 to gfx950 on MI355X.
- Precision fixes for exp/sin/cos and context fixes for automatic target, stream capture, and profiler are covered in Chapter II. ROCm-side FP4 conversion and the E8M0 type, along with architecture detection for pip-based images, remain under review (#3665 / related follow-ups to #3549).
MetaX / Moore Threads (MUSA): No Substantive Commits in the Window
- The tilelang-metax organization view has a push timestamp (10-26 15:08), but a branch-by-branch check found no commits within the window: the default branch main still shows the 09-29 version bump, dev’s latest is 10-16, the 0.1.14 and earlier version images are untouched, and the event stream contains no in-window push events. The most recent official release remains 0.1.15 from 10-16.
- tilelang-musa had no in-window main-branch commits all week; the latest tag remains v0.1.15+musa.1 from 09-30. Moore Threads and MetaX appear only in the verification table this week and do not constitute entries.
Sources: https://github.com/tile-ai/tilelang-metax, https://github.com/tile-ai/tilelang-musa
Third Party: Hexagon DSP Backend
- 15 PRs in a single day on 10-27 carried HVX vectorization, the VTCM arena, Crouton layouts, and DMA builtins all the way through to mapping and reduction on Crouton. As of 10-29, the latest main-branch commit still dates to the 08-24 CI dependency upgrade; the repository push time is updated but the main branch has no new code, so this counts only as an observation of community investment intensity, not as upstream progress.
- There is also a zero-modification PPU wrapper (intercepting the import protocol and site mechanism at runtime, neither modifying upstream source nor repackaging, and explicitly erroring rather than silently degrading for CUDA features with no PPU counterpart), a third form of backend migration.
Sources: https://github.com/xwhzz/tilelang-hexagon, https://github.com/flytiger-eco/tilelang-wrapper-for-sail
IV. Ecosystem and Adopters
TileOPs: The Second Week After 0.1.0
- Merge cadence slowed noticeably after the release: 14 merges on Tuesday, 4 on Wednesday, 7 on Thursday, and 1 on Friday, with the focus shifting from “freezing the interface” to wrapping up performance and correctness.
- Four hard numbers 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 from 9 to 32, and 264 shapes with m≤8, using softmin expected log regret for coordinate descent. End-to-end busy-time regret on the held-out manifest drops from 5.79% to 0.48%, and the m≤8 held-out set drops from 14.18% to 1.79% (#2784). Non-NT layouts now go through the transposed persistent kernel coop2_swap: 2112×7168×4096 TN drops from 195.3 µs to 175.3 µs, and on another H200 in the same machine from 174.3 to 158.5 µs, both beating QuACK’s 187.9 and 167.0 µs (#2789). Long-row cumsum switches to a two-stage seed scan: (8,65536) drops from 12.75 µs to 5.18 µs and fixes launch failures above 65535 tiles (#2790). Extrema reduction switches to PTX max.NaN / min.NaN: amax (2048,4096) bf16 goes from 8.352 µs back to 6.976 µs, faster than FlagGems (#2792).
- Back-to-back benchmarking was filed as a separate item: A new issue on Thursday noted that under single-call conditions this implementation leads QuACK, but in back-to-back, power-capped training scenarios six of ten large-row cases lag (#2791). A subsequent GEMM change under review recomputes grid grouping and persistent block counts under the power cap, cutting memory reads on the q-b shape of DeepSeek V3 prefill from 168 MB to 94 MB (#2794).
- Correctness and capability gaps filled: D128 dense forward splits keys, with prefill-rope dropping from 31.04 µs to 17.09 (#2777); D128 persistent forward is rewritten per thread and writes out LSE (#2774); the fused-add LayerNorm row is split across CTAs, with the precision defect filed separately (#2768 / #2773); three classes of offset-store fixes plus SM90 grouped GEMM and the output ABI (#2761 / #2765 / #2762).
- Release form: The release image was merged (#2793), with the operator library and TileFoundry packaged into the same image for the first time. Among known gaps, the FusedAddLayerNormFwdOp precision issue remained unfixed during the week, and no fused operator from the first step of the roadmap has been merged yet.
Sources: https://github.com/tile-ai/TileOPs/pull/2784, https://github.com/tile-ai/TileOPs/pull/2789, https://github.com/tile-ai/TileOPs/pull/2790, https://github.com/tile-ai/TileOPs/pull/2792
TileFoundry: From a Three-Package Split to Static Contracts
- Early in the week it was split into three distribution packages—core / nvidia / apple—with test inputs loaded only once (#268 / #269); sm_89 gained E4M3 warp MMA, and capability validation was extended from “MMA only” to every instruction, since bulk copy, TMA copy, and mbarrier forms previously passed validation and only failed at codegen or ptxas (#261 / #262 / #264).
- Midweek fixes covered cp.async target-layout addressing and MLA backward layout (#263 / #260 / #257), as well as issuing the full window by the operand itself (#273, restoring compilability of sm_89 GEMM) and sizing dynamic shared memory by the kernel’s actual usage (#271; the author did not adopt the suggestion to “allocate the full target capacity,” on the grounds that static and dynamic shared memory share one limit and the per-block request determines how many blocks fit in an SM).
- A breaking change was merged on Thursday (#279): 24 AMX runtime operator families were described as 39 statically registered TIR contracts, with parameter domains, operand and layout patterns, side-effect sets, and target capabilities from M1 to M4; a new Python evaluator and a raw C simulator were added as gold standards, 59 representative atomic test cases were added covering all 39 static operators, and higher- and mid-level IR cases run real 2-, 4-, 8-, and 12-thread domains and verify candidate acceptance and rejection. Validation: 46 focused tests passed.
- One new defect noted on the side: the analysis side overcounts the shared-memory peak and does not propagate it to TIR (#272), the same root cause as the ThreadSync family in the main repo.
Adopters and Community Forks
- Both known adopters were silent all week: DeepSeek’s TileKernels (2,018 stars) and the Tongyi-side FlashQLA (742 stars) both have their main branches stopped at 10-14, with no new pushes in the window. This week’s ecosystem signals came mainly from community-built repos rather than these two internal users.
- T-HEAD’s tilelang-for-sail: Created on 09-22, it is a fork of the upstream main repo with 13 public branches. Besides the main branch and version mirrors from v0.1.11 to v0.1.14, visible branches cover a PPU matrix shared-layout fix, a PPU language package, a runtime compilation feature, and several bug fixes. The main branch’s latest commit is from 10-22, and the repo’s push time is updated to 10-29 19:27; the repo has no README or release records, so it can only serve as an observational lead. Notably, Friday’s defect #3705 was reproduced using a build from this repo’s source—fork users hit real build problems earlier than upstream.
- The CANN community’s ops-tilelang: Created on 10-20, positioned as a community high-performance operator repo for Ascend NPU, centrally managing kernel implementations, automated test cases, standard workloads, and performance baselines. The README explicitly supports only Ascend 950 and above, requires CANN 9.3.0 or later, and uses the community weekly build. Seven operators have been used to string together a “Attention output plus FFN preprocessing” fragment of DeepSeek V4.1 Flash, with HIDDEN=5120 and MHC=4; measured SuperKernel gains for B from 1 to 256 range from 1.17% to 2.75%. Its relationship to the official tilelang-ascend remains unclear, so it is noted only as ecosystem observation.
- PTO-ISA’s two companion repos: The puzzle repo migrated PTO’s broadcast load to PTOAS 0.70, after a prior week of clarifying quantization-scaling notation and rewriting problem statements; the benchmark repo added two commits—SV2 block scaling and per-column analysis—based on a variant count of 31.39 for the library’s absolute-maximum instruction, with runtime instruction throughput and roofline analysis attached, explicitly not a unified buffering basis.
- Distributions: Sunrise’s 0.1.15+sunrise.1.2.0 release PR was merged early in the week (#9), the second downstream distribution form within the organization; if the main repo’s TIRx migration lands, such distributions will need to update their dependency declarations accordingly—a knock-on effect worth tracking next week.
- A community hub repo for RISC-V: TongmingLAIC/tilelang-riscv was created on 10-23 and pushed on 10-26, positioned as the community hub for TileLang on RISC-V, aggregating compiler backends, AI operators, developer tools, and documentation. It is still an empty skeleton, with no runnable operators or backend implementation yet.
Sources: https://github.com/t-head/tilelang-for-sail, https://github.com/Atlas-Ascend-Fork/ops-tilelang, https://github.com/PTO-ISA/tilelang-puzzles-ascend, https://github.com/TongmingLAIC/tilelang-riscv
V. Community, Tutorials, and Events
- Knowledge assets emerged as a new category on the community side this week: tilelang-knowledge was created on 10-29, and four commits that day formed a knowledge base for agents—repo structure, 33 operator design cases (each with a fixed source), 20 programming knowledge cards with sources, and 7 evidence-backed GPU optimization patterns. The repo describes its purpose as letting an agent, when facing an operator, determine which experience applies, 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.
- Continued investment in tutorials and exercises: tilelang-puzzles-ascend builds PTO-to-Ascend quantization kernel tutorials; a Chinese TileLang operator interview practice site has 9 modules and 18 problems; the illustrated-course companion repo describes itself as explaining every knowledge point thoroughly with original scenarios and animations, released alongside videos; a distributed-layout experiment repo was created the same day.
- Docs and localization: tilelang.github.io synced automatically with main-repo merges all week (eight on 10-26, six on 10-27, two on 10-28), the third-party localization repo tilelang-docs-l10n updated as usual, and the Chinese RISC-V community hub repo made its first content commit—indicating the documentation sync pipeline is working normally.
- On the discussion side: A TileScale user asked whether the TileLang version they use can be upgraded (#72). That repo’s last push was on 08-25, and its docs and examples still point to an older version; such questions typically appear after upstream interfaces enter a period of rapid renaming—this week’s TIRx replacement under review in the main repo and TileOPs’ renaming of 26 operator classes will both increase downstream version-tracking costs.
- Zero substantive news hits all week: Six Google News queries in Chinese and English (covering 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) returned only one reposted clinical-AI article across five collection windows, unrelated to operator languages; all relevant Hacker News hits fell outside the window, the most recent being a 10-02 reposted discussion of the Ascend programming stack; an arXiv search for TileLang found no new papers. The daily report therefore continues to rely on repos and ecosystem repos as its main body. On the industry side, there was only one 10-26 report on the narrowing ecosystem gap between Ascend and DeepSeek, which is secondary commentary and contains no new technical data.
VI. Trend Observations
Release engineering has moved from “stuck at the last step” to “image-based delivery.” Monday’s briefing was still reporting that TileOPs’ first release was “docs ready, artifacts not uploaded,” with the latest on PyPI still 0.0.1.dev2 from 2025-10-09; on Tuesday the artifacts were uploaded at 07:45 and the release page went live at 13:15, and by Thursday the release image locked the operator library, TileFoundry, and three environment identifiers (cuda13.2, torch2.13, tilelang0.1.12) together. Within a week, this line went from “can it ship” to “how do we combine and deliver it.” For a library of 200-plus operators, installability determines external adoption speed more than optimizing another percentage point does; this week’s answer was to make the install path itself an artifact.
Defect types are heavily concentrated in “compiles, runs without error, produces wrong results,” and are strongly tied to hardware or context. Output allocated on the wrong stream (corrupted only on non-default streams with queued work), T.ws(-1) running on zero threads, slice fill validated by length rather than end position, negative indices still negative after wraparound, address_of treated as a read—all five produce no compile-time warning and no runtime exception. These defects cost far more than compile errors because what they contaminate is the very premise that “all tests pass”; the six tests that surfaced on Wednesday assuming cold compilation and Thursday’s without_kernel_cache semantic hole show that the test infrastructure sits in the same risk zone. This week’s response was to move boundary validation forward to compile time—the right direction, but three newly opened defects (#3704, #3705, #3707) show that this boundary is far from fully drawn.
Ascend has shifted from “porting the backend” to “filling in operators,” and for the first time offers reproducible acceptance criteria. Wednesday saw NPUIR merged into mainline (504 self-tests, 71/72 examples); four of the five merges from Thursday to Friday were operators and examples, all ordered by 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. The common thread is writing quantization and fusion into the kernel itself (E4M3 pure-integer encoding, four-bit packing conventions, dual-kernel plus 1-multiply-1 fast path) rather than leaving remediation to the host side; at the same time, the profiling skill and the checklist-based test channel were established as process. The main tension is shifting from “can it compile” to “can the operator’s numerics and performance be stably accepted,” and the 100-file preview of the a2a3 scheduler migration signals that the next phase of change is at the scheduling layer.
Subtract from the execution path, add to the backend list. This week the main repo folded the all-warp reduction into redux.sync, split the autotuner’s compile lock more finely, merged the GM-to-L1 copy and fill into a single task, and pulled the buffer view’s layout from a shared variable back into a single view; meanwhile the backend list kept expanding—ROCm added the pip SDK and FP4, Hexagon saw 15 merges in a single day, PPU gained a zero-change wrapper, and Intel AMX became 39 static contracts in TileFoundry. The two directions are not contradictory: upstream converges the execution path while downstream packages and distributes separately, the same approach as TIRx making TVM an external dependency.
“Forking” is the hardest adoption signal this week, and it brings the most concrete synchronization cost. T-HEAD uses TileLang as the foundation of its own hardware stack for the PPU branch, the CANN community builds its own operator repo, PTO-ISA turns both the problem set and benchmarks into companion repos, and a community hub repo appears for RISC-V—none of these are in the tile-ai organization, yet all depend on the stability of the language layer, which is exactly the point of freezing the public API in 0.1.0 and making it a CI gate. The cost is equally clear: on Friday, defect #3705 was first hit by a branch user’s real build, and only then did upstream patch it; the more branches there are, the longer this “downstream finds it first, upstream fixes it later” path becomes, and the main repo needs to start answering the long-term question of how branches and upstream synchronize.
The measurement question is starting to be treated as a product question. On Thursday TileOPs made “leading on single calls, lagging back-to-back” its own item (#2791); Friday’s GEMM change recomputed the tile grouping under a power cap and cut DRAM reads from 168 MB to 94 MB (#2794); and the selector moved from fitting against a pasted manifest to calibration under busy conditions. Compared with “how many microseconds faster a given row is,” these changes acknowledge that the real call patterns of training and inference differ from benchmarks, and write that difference into the scheduling policy. For a language meant to write kernels for large models, this determines long-term adoption more than piling on more operators.
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-mlir-ascend, tilelang-hygon, tilelang.github.io, TileOPs.github.io, TileOPs-nightly; tilelang-metax, tilelang-musa, TileRT, tilescale, DeepStack, tilelang-puzzles had no main-branch commits within the window |
| Merged ground-truth query (is:pr is:merged) | 78 in the Monday 72-hour window; 86 total across the four 24-hour windows Tuesday–Friday (tilelang 47, TileOPs 26, tilelang-ascend 5, TileFoundry 4, tilelang-hygon 3, tilelang-mlir-ascend 1), with 30, 20, 24, and 12 per day respectively |
| Open ground-truth query (is:pr is:open updated) | 76, 49, 45, 48, and 54 across the five windows, including the TIRx refactor, the a2a3 scheduler migration, the Ascend docs verification batch, and TileOPs GEMM scheduling changes |
| New issue check | 19, 11, 7, 6, and 4 across the five windows, each day including failure notifications from the Ascend repo’s test bot |
| Main-branch commit-level check | Merged counts match main-branch commit counts in every window; tilelang had 10, 14, 13, 15, and 5 per day |
| Release pages and tags | tilelang latest remains v0.1.15 (released 09-30), with no new version all week; 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); tilelang-metax is v0.1.15 (10-16) |
| Third-party and ecosystem repos | The Ascend community operator repo, both PTO-ISA repos, the dev tools repo, the knowledge repo, the illustrated-course repo, the distributed layout repo, the DCU fused-operator repo, the model porting repo, the Hexagon backend repo, the T-HEAD fork repo, and the RISC-V community hub repo all had pushes within the window; the Hexagon main branch had no new commits, only a repo push timestamp update |
| Adopters | TileKernels and FlashQLA main branches are stalled at 10-14, with no pushes within the window |
| Google News (six queries, Chinese and English) | Only one Chinese repost on clinical AI plus Ascend and one industry-side commentary, neither containing new technical data, so neither was included as an item |
| Hacker News | All relevant hits in the past four days fall outside the window |
| arXiv | No new papers found when searching for TileLang |
Main References
- [1] TileOPs v0.1.0 release page — https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0
- [2] TileOPs PyPI artifact — https://pypi.org/project/tileops/
- [3] Ascend 910 NPUIR backend merged into main (#207) — https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
- [4] TIRx refactor (#3666) — https://github.com/tile-ai/tilelang/pull/3666
- [5] Output allocation stream fix (#3649) — https://github.com/tile-ai/tilelang/pull/3649
- [6] Full-warp reduction switched to redux.sync (#3692) — https://github.com/tile-ai/tilelang/pull/3692
- [7] autotuner concurrent compilation (#3476) — https://github.com/tile-ai/tilelang/pull/3476
- [8] ROCm math function precision fix (#3621) — https://github.com/tile-ai/tilelang/pull/3621
- [9] Ascend buffer precise addressing interface (#3607) — https://github.com/tile-ai/tilelang/pull/3607
- [10] Ascend kpool compression operator (#1910) — https://github.com/tile-ai/tilelang-ascend/pull/1910
- [11] W4A4 int4×int4 root-cause fix (#1731) — https://github.com/tile-ai/tilelang-ascend/pull/1731
- [12] Hygon scale buffer dual bank (#29) — https://github.com/tile-ai/tilelang-hygon/pull/29
- [13] Hygon register pipeline (#18) — https://github.com/tile-ai/tilelang-hygon/pull/18
- [14] ROCm 10 pip SDK support (#3549) — https://github.com/tile-ai/tilelang/pull/3549
- [15] TileFoundry AMX static contract (#279) — https://github.com/tile-ai/TileFoundry/pull/279
- [16] TileOPs release mirror (#2793) — https://github.com/tile-ai/TileOPs/pull/2793
- [17] TileOPs selector calibrated by busy time (#2784) — https://github.com/tile-ai/TileOPs/pull/2784
- [18] T-HEAD fork repo — https://github.com/t-head/tilelang-for-sail
- [19] CANN community operator repo — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [20] Hexagon backend repo — https://github.com/xwhzz/tilelang-hexagon
- [21] RISC-V community hub repo — https://github.com/TongmingLAIC/tilelang-riscv
- [22] tilelang-knowledge knowledge repo — https://github.com/qiuuuuu622/tilelang-knowledge
Note: All facts, version numbers, and dates in this report are taken from the original text of the five daily reports within the period and their corresponding primary sources, with no speculative additions. The Monday daily report’s collection window was 72 hours (covering the previous weekend), and its merged count includes 10-23, overlapping with the previous weekly report; this report therefore does not add it to the other four. All five weekday daily reports within the period are present, with none missing.