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

Sources: GitHub tile-ai organization per-repo verification (tilelang / TileOPs / TileOPs-nightly / TileFoundry / tilelang-hygon / tilelang-metax / tilelang-ascend / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / docs site and nightly benchmark repo) + merged and in-review ground-truth search (20 merged, 45 in-review updates) + full check of newly added issues (7) + commit-level diff verification + release page and PyPI artifact verification + external backend repo verification (Hexagon / CANN operator repo / PPU wrapper / model snippet repo) + Google News multi-query search in Chinese and English (via proxy) + Hacker News + arXiv.

Index

  • Today’s Highlights: TileOPs 0.1.0 officially released — 206 operators, and 2262 of 2265 benchmark rows are no slower than the fastest implementation; 20 merges in 24 hours, 45 updates under review, 7 new issues; Hygon scale buffer dual-bank merged with 1389 lines; third-party Hexagon backend 15 commits in one day
  • I. Core Project Progress
    • 1.1 Main repo: ROCm exp/sin/cos no longer default to fast builtins; sin/cos error drops from 7.8e-04 to 6.6e-08 (#3621) (10-27)
    • 1.2 Main repo: auto target now reads the current HIP device; multi-GPU machines no longer pin to device 0’s architecture (#3594) (10-27)
    • 1.3 Main repo: constant row indices in TMA gather/scatter converted to int32; literal forms no longer crash compilation (#3365) (10-27)
    • 1.4 Main repo: NVRTC and CuTeDSL take the current stream from the device of the first input tensor (#3359) (10-27)
    • 1.5 Main repo: autotune cache key now includes compile_flags; changing compile options for the same kernel no longer reuses the old binary (#3338) (10-27)
    • 1.6 Main repo: ThreadSync records address_of as a write; fences after opaque calls are restored (#3346) (10-27)
    • 1.7 Main repo: pipeline epilogue with runtime trip count is unrolled directly; paged decode compiles again (#3356) (10-27)
    • 1.8 Main repo: two fixes for slice fill and static bounds of negative indices; out-of-bounds writes and silent negative addresses closed (#3352 / #3347) (10-27)
    • 1.9 Main repo: T.ws(-1) regions run on all threads; standalone forms no longer lose stores (#3354) (10-27)
    • 1.10 Main repo: profiler uses the current device when none is specified; cache flushes no longer hit device 0 (#3357) (10-27)
    • 1.11 New defects: kernel cache hits break six tests that assume cold compilation; CuTeDSL CUDA graph capture yields an empty graph (#3694 / #3693) (10-27)
  • II. Multi-Backend Adaptation (Ascend / Hygon / ROCm / Hexagon)
    • 2.1 Hygon: HCU scale buffer partitioned into dual banks; blockscaled GEMM shares one row mapping under manual warp specialization (#29) (10-27)
    • 2.2 Ascend main repo: PerCoreTask candidate bodies allow internal if/else; conditional L1 fills are no longer rejected wholesale (#3686) (10-27)
    • 2.3 Ascend main repo: CI splits AscendC and PTO into two matrix jobs that can be rerun separately (#3687) (10-27)
    • 2.4 Ascend repo under review: Cython wrapper lazily constructs the analyzer, targeting the per-call rebuild issue noted yesterday (#1948) (10-27)
    • 2.5 Ascend repo under review: GLM-5.3 kpool compression operator, tensor shift amounts on PTO (#1910 / #1949) (10-27)
    • 2.6 Ascend main repo under review: PTO T.Parallel vectorization, AutoSimVF channel, batched 3D ND2NZ, write-count multi-buffering, Bisheng build identity in cache key (#3512 / #3320 / #3689 / #3629 / #3691) (10-27)
    • 2.7 ROCm under review: architecture detection and images for pip ROCm 10, FP4 conversion and E8M0 type (#3549 / #3665) (10-27)
    • 2.8 Ascend repo daily tests: 2903 items all pass, two consecutive days with no failures (#1950) (10-28)
    • 2.9 MetaX: org view has push timestamps; none of the six branches has new commits in the window (10-27)
    • 2.10 Third-party: Hexagon DSP backend 15 commits in one day; HVX vectorization, VTCM arena, Crouton layouts and DMA builtins all in place (10-27)
  • III. Ecosystem and Adopters (TileOPs / TileFoundry)
    • 3.1 TileOPs 0.1.0 released: first official version with 206 operators; public API now frozen (10-27)
    • 3.2 TileOPs 0.1.0 benchmarks: on H200, 2262 of 2265 rows are no slower than the fastest implementation, geometric mean 1.37x (10-27)
    • 3.3 TileOPs 0.1.0 known gaps and next steps: fused operators, whole models, compute-communication fusion (10-27)
    • 3.4 TileOPs merges four: where and masked_fill tunable, D128 causal diagonal scored over 64 keys, sparse long GDN prefill split, changelog dated by release day (#2782 / #2783 / #2779 / #2781) (10-27)
    • 3.5 TileFoundry merges two: full window dispatched by operand itself, dynamic shared memory sized to actual kernel usage (#273 / #271) (10-27)
    • 3.6 External ecosystem: CANN community TileLang operator repo, zero-change wrapper for PPU, DeepSeek V4.1 Flash snippet example (10-27)
  • IV. Community, Tutorials and Events
    • 4.1 Documentation site: six automated updates in one day (10-27)
    • 4.2 Ascend repo: four documentation fact-checks, one thread with up to 46 comments (#1628 / #1655 / #1656 / #1659) (10-27)
    • 4.3 Under review: extract lower trace into a standalone developer toolkit (#3583) (10-27)
    • 4.4 News side: no new coverage from Google News or HN within 24 hours for this window (10-27)
  • V. Trend Observations
  • Appendix: Sources and Verification Notes

Today’s Highlights

TileOPs 0.1.0 landed within today’s window, ending a multi-day state of “stuck on the last square.” The 0.1.0 artifact on PyPI was uploaded at 10-27 07:45 (Beijing time), 45 minutes after the cutoff of the previous daily report; the GitHub release page went up at 13:15, and the repository got its first official tag. The release notes cite three hard metrics: all 206 operators have manifest entries (signature, test and benchmark workloads, roofline formula, each validated 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 3, FlashInfer, vLLM, DeepGEMM and FlagGems; and 195 of the 206 operators came from foundry-tagged merges where TileFoundry generated the kernel and it entered the repo through the same CI. The release also freezes the public API — from here on 0.x only adds, never changes — and adds an API compatibility check script, so any PR that calls the API the way the previous version did will fail outright.

In 24 hours the tile-ai org merged 20 PRs (tilelang 13, TileOPs 4, TileFoundry 2, tilelang-hygon 1), with 45 updates under review and 7 new issues (2 of them from the Ascend repo’s daily test bot). Merges concentrated on two fronts: one is floating-point semantics and execution context — ROCm switched seven default math functions from fast builtins back to precise implementations (sin/cos error on MI300X drops from 7.8e-04 to 6.6e-08), auto target reads the current device instead of pinning device 0, NVRTC and CuTeDSL take the stream from the device of the input tensor, and the autotune cache key now includes compile_flags; the other is static bounds and thread domains — slice fill is validated against the region end rather than the length, out-of-range negative indices are rejected at compile time, and T.ws(-1) regions now run on all threads. Hygon merged 12 files and 1389 lines, splitting the scale buffer into an independent dual bank separate from LDS banks. On the third-party side, the Hexagon DSP backend saw 15 commits in one day, carrying HVX vectorization, the VTCM arena, Crouton layouts and DMA builtins all the way through to mappings and reductions on Crouton.

I. Core Project Progress

1.1 Main repo: ROCm exp/sin/cos no longer lower to fast builtins by default; sin/cos error drops from 7.8e-04 to 6.6e-08 (#3621)

Merged 10-27 22:39, 5 files +107/-17, fixes #3618. On HIP, seven default functions — T.exp, T.exp10, T.log, T.log2, T.log10, T.sin, T.cos — lowered to fast builtins like __expf regardless of the tl.enable_fast_math setting, while the CUDA side follows the #875 convention of using precise implementations; conversely, explicit fast operators such as T.__exp did not lower at all on HIP and raised AttributeError directly in tilelang.rocm.language. This PR makes HIP adopt the same policy as CUDA: the seven default operators are redirected to HIPMath, eight underscore-prefixed fast operators get a new minimum-lowering rule pointing to HIPFastMath, and all are moved into the FAST_MATH_INTRINSICS list exported by the ROCm language layer. Precision comparison uses the #3618 reproduction: on MI300X, exp and exp10 are measured by relative error, sin and cos by absolute error (x up to 1e4); the four columns are T.exp, T.exp10, T.sin, T.cos in order, with mainline at 3.7e-06 / 4.1e-06 / 7.8e-04 / 7.8e-04 and this PR at 8.3e-08 / 8.1e-08 / 6.6e-08 / 6.4e-08.

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

1.2 Main repo: Auto target now reads the current HIP device; multi-GPU machines no longer always take device 0’s architecture (#3594)

Merged 10-27 22:37, 2 files +24/-1. determine_target(“auto”) lands on the registered HIP detector on ROCm, and that detector unconditionally reads torch.cuda.get_device_properties(0), so when the current HIP device is not device 0, auto-detection returns device 0’s architecture; the CUDA side has long used torch.cuda.current_device(), and ROCm had not caught up. The fix takes the current device index first, then queries properties, and strips the feature suffix with normalize_rocm_arch. The new test stubs current_device() to return 1, records every probed index, and asserts that gfx942 is detected and only device 1 was touched; before the fix that assertion yielded probed == [0]. The author notes there is no ROCm device on the local machine, so the test only validates the detection helper with a stubbed torch and does not run the real HIP runtime; it also touches a different function from the other PR merged today, #3549 (#3549 is in tilelang/contrib/rocm.py, this one in tilelang/rocm/target.py).

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

1.3 Main repo: Constant row indices in TMA gather/scatter converted to int32; literal forms no longer crash compilation (#3365)

Merged 10-27 09:15, 2 files +44/-2, fixes #2601. The two frontend helpers tma_gather4 and tma_scatter4 passed Python integer row indices straight through, leaving them with no definite integer width in TIR, which crashed compilation on Blackwell; other row expressions are left as-is. This PR only converts Python integers to int32 TIR constants and adds a Blackwell-gated compile regression that compiles gather and scatter with literal row indices and then checks that the corresponding instructions actually appear in the generated CUDA source. The author notes the test only verifies that compilation succeeds, not runtime numerical correctness; also, on the submitter’s local Windows Python 3.13 environment, Home Assistant’s pytest plugin depends on fcntl, so pytest fails at collection and the TileLang tests in this PR were not actually run locally.

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

1.4 Main repo: NVRTC and CuTeDSL take the current stream from the device of the first input tensor (#3359)

Merged 10-27 08:11, 2 files +5/-4, follows F149 of #3301 and #3408. Both adapters previously determined the target type by string prefix and defaulted to the current stream of the caller’s current device, so they picked the wrong stream when the caller’s current GPU differed from the tensor’s GPU. This PR switches to target-kind checks and looks up the current stream by the first input tensor’s device; an explicitly passed stream handle still has highest priority, and NVRTC calls without tensors keep their original semantics. This is the second stream-semantics cleanup after yesterday’s out_idx output-allocation stream (#3649). The author draws the boundary: this PR only fixes stream selection, and cross-device default-stream launch context remains unsolved — such callers must set the tensor’s device as current before launching themselves, as the patch adds no per-call device guard; validation covered same-device and cross-device non-default streams with real Torch streams and GPU launches, but did not establish correctness for cross-device default-stream context.

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

1.5 Main repo: Autotuning cache key includes compile_flags; same kernel with different compile options no longer reuses the old binary (#3338)

Merged 10-27 16:57, 3 files +28/-1, follows F131 of #3301. The autotuner persists compilation results in a cache, but the cache identity did not include compile_flags, so the same kernel with a different set of device compiler options still hit the old binary, and measured performance did not match the actual compilation result. This PR makes CompileArgs actually store the flags and hash them in their original order, and threads the flags from the JIT implementation through AutoTuner.set_compile_args() all the way to candidate compilation, per-config pass merging, grouped device compilation, and cached kernel loading. Validation covers the three existing cache-key tests plus a targeted CUDA comparison using the native runtime: both list and string flag forms survive per-config overrides; the author notes the previous merge logic dropped compiler macros, which is exactly the gap being closed here.

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

1.6 Main repo: ThreadSync records address_of as a write; fences after opaque calls restored (#3346)

Merged 10-27 05:49, 2 files +62/-1, follows F305 of #3301 and continues yesterday’s fence-semantics thread. address_of itself carries no access information — it only hands a pointer downstream — but the access classifier hardcoded its type as a read, so a shared buffer written through an opaque call never entered the writes set, subsequent whole-block reads found no conflicting write and planned no fence, and threads raced past the write to read pre-write contents. The measured CUDA output is telling: the T.address_of(B[0]) form has no fence after the extern, while the T.access_ptr(B, “w”) form does; the latter explicitly declares the access type, which is why this hole went unnoticed. loop_unswitching.cc already treats address_of conservatively as a write, and ThreadSync was the sole exception. This PR records address_of as a write, a conservative direction: an unknown pointer use may be written, and fence planning only needs to err on the safe side; two new cases assert that after lowering, the kernel has __syncthreads() following the extern.

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

1.7 Main repo: Pipeline epilogue with runtime trip count unrolled directly; paged decode compiles again (#3356)

Merged 10-27 07:02, 1 file +5/-1. When a T.Pipelined loop’s trip count is a runtime value, CUDA codegen reports TypeError: Downcast from tirx.Sub to ir.IntImm failed as soon as the epilogue unrolls two or more steps with async copies. The root cause is in EmitImpl: prologue and epilogue are only unrolled when their extent is constant, ensuring each iteration’s wait count is constant; the epilogue extent is end - start, i.e. (n + max_stage) - n, which folds at construction time when n is constant but stays a Sub when n is runtime, so as_const_int returns null and the epilogue is emitted as a T.unroll loop whose wait is written as T.ptx_wait_group(1 - k), reading the loop variable, while ptx_wait_group only accepts immediates. The fix inspects the simplified extent before the check; constant-bound cases simplify back to themselves, matching the original behavior. The defect was reported downstream in TileOPs#1852 and bisected to after afcebed1, breaking paged decode, whose trip count comes from the KV length. The author compared both versions compiled on sm_89 with CUDA 13.0; the new test fails on mainline and passes with this PR.

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

1.8 Main repo: Two fixes for static bounds in slice fill and negative indexing; out-of-bounds writes and silent negative addresses closed (#3352 / #3347)

Merged 10-27 05:37 and 05:14, 4 files +100/-3 combined. The former targets the Fill constructor’s static bounds check: it previously validated only min >= 0 and extent <= shape per dimension, and extent <= shape holds only for regions starting at the buffer origin; a slice starting further in passes both checks yet lowers to an out-of-bounds write — T.fill(As[6:12], 7) compiles on a buffer of length 8 and emits As[threadIdx.x + 6] = 7, writing to 6 through 11. The fix also validates the region end against the target shape, with the new condition requiring both start and length to be static; zero-anchored regions and symbolic-start regions are unaffected; T.clear shares the same node and had the same defect, now closed as well. The latter targets LegalizeNegativeIndex: it previously just added the buffer length once to a negative index with no range check, so on a buffer of length 8, A[-9] yielded -1 and A[-10] yielded -2, escaping the allocation origin with no diagnostic, whereas NumPy and PyTorch raise IndexError for the same indices. The fix rejects only when the analyzer can prove the wrapped index is still negative, reporting the index and length; all in-range constants and all dynamic indices keep their original behavior.

Source: https://github.com/tile-ai/tilelang/pull/3352, https://github.com/tile-ai/tilelang/pull/3347

1.9 Main repo: T.ws(-1) regions run on all threads; standalone form no longer drops stores (#3354)

Merged 10-27 04:52, 2 files +40/-3, follows #2978. WarpSpecialize’s documentation defines -1 as “skip threadIdx.x binding”, i.e. the region is not constrained by threadIdx; the implementation instead treated -1 as an ordinary warp-group index, constructing the interval [-warp_group_size, 0), which holds for no thread, so the region ran on zero threads and every store inside was silently dropped. The minimal reproduction is clean: with 256 threads, first write 0 everywhere, then write 5 inside T.ws(-1), and all 256 threads still output 0; the same store without the ws wrapper has all 256 threads writing 5. The fix emits no range constraint when the id contains -1, real group indices still use their own thread constraints, and the docstring gains the -1 example — the original documentation’s examples happened to omit it. The author checked the whole repo and found no ws(-1) usage, so no existing code depends on that empty region.

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

1.10 Main repo: Profiler uses the current device when none is specified; cache flush no longer hits device 0 (#3357)

Merged 10-27 07:26, 1 file +1/-1. When do_bench is called without an explicit device and the current CUDA device is nonzero, _cache_device(None) fell back to the module-level default cuda:0, so allocations made to flush the cache could land on a different GPU than the measured workload, and the numbers measured were not on the same card as the real load. The fix changes the default string to “cuda” without an index, letting PyTorch resolve the current device when allocating the buffer; explicit devices and the NPU and MPS defaults are unchanged. Upstream CI (run 37915443170) passed on Quick Lint, CUDA, ROCm, Metal, Ascend, and CuTeDSL examples; the author notes that subsequent PR-triggered runs require maintainer approval and have not been executed, and that the local machine has no CUDA hardware or timing benchmark.

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

1.11 New defects: Kernel cache hits break six tests that assume cold compilation; CuTeDSL CUDA graph capture yields an empty graph (#3694 / #3693)

Both issues were opened by lucifer1004 on 10-27. The former notes that six tests pass with a cold cache but fail on a second run with the same cache directory, because they assume compilation always happens: for example, the sm75 boundary case in test_tilelang_kernel_gemm_sm75.py reports AttributeError: ‘NoneType’ object … under four parameters. The latter is more severe: kernel launches emitted from the CuTeDSL backend are not recorded into the CUDA graph, so torch.cuda.CUDAGraph capture yields an empty graph and replay() runs nothing, while the same kernel compiled with the cuda backend captures and replays correctly; the report includes a minimal reproduction. Both are extensions of the “cache and execution context” family of defects, pointing at the same area as the merges in 1.4 and 1.5.

Source: https://github.com/tile-ai/tilelang/issues/3694, https://github.com/tile-ai/tilelang/issues/3693

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

2.1 Hygon: HCU scale buffer partitioned into dual banks; blockscaled GEMM shares a single row mapping under manual warp specialization (#29)

Merged 10-27 08:26, 12 files +1389/-314 — the largest software-side change in this window. The change plans the rows of the scale buffer as an independent dual-bank partition, separate from the LDS banks, and feeds the same row mapping to all three of allocation, scale copy, and MMAC sinking, so the layout is not re-derived three times. Details include: retaining the physically-intervalized lifetime and m0-aligned paging; reducing the number of scale-copy issues; spreading the work of idle waves across the two scale banks; and using this layout for blockscaled GEMM under manual warp specialization. This lands on top of yesterday’s native B16 storage and dynamic WASP GEMM (#28); Hygon has completed three mainline merges in two days.

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

2.2 Ascend main repo: PerCoreTask candidates may contain internal if/else; conditional L1 fill no longer rejected wholesale (#3686)

Merged 10-27 07:24, 2 files +85/-24. A T.Task inside T.PerCoreTask was rejected whenever its body contained if/else, even if the branch only controlled internal work — for example, a conditional L1 fill after a copy. Reproduced at d85a70a9, where the materialization stage reported “control flow selecting a candidate must be outside the task scope.” This change removes that blanket rejection and instead validates the task body uniformly: a sync barrier wraps the entire candidate, while each branch still checks pipeline and core compatibility, nested tasks, and unsupported cross-core synchronization item by item. Regression tests accept internal if and if/else, preserve the candidate body, and retain the original rejection cases; they also check that the completion flag follows the entire candidate under the outer guard. All 113 tests across task, scheduling unit, and latency pass, and kernels with conditional L1 fill now complete full Ascend source sinking; device execution was not tested.

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

2.3 Ascend main repo: CI splits AscendC and PTO into two matrix jobs that can be re-run independently (#3687)

Merged 10-27 07:05, 7 files +56/-27. Previously AscendC and PTO were two steps in the same Test for Ascend job, so any backend-specific failure turned the whole check red, and a re-run necessarily dragged in the other suite of examples. This change splits examples and unit tests into two matrix jobs: Test for AscendC runs examples/ascend (target ascend, 4-way concurrency) and testing/ascend -m “not pto” (32-way concurrency), while Test for PTO runs examples/ascend (target pto) and testing/ascend -m pto (32-way concurrency). The split relies on the existing pto marker, with markers added to PTO-only tests and to the PTO parameters of dual-backend tests; the two marker expressions are mutually exclusive and together cover the full suite without duplicate execution. fail-fast is set to false so one job’s failure does not cancel the other; both retain the “still run unit tests when examples fail but the build succeeds” behavior and each publishes its own build and test summary. Shared tests (including contract checks that run the same frontend IR through both code generators) run only once on the AscendC side, as documented in testing/ascend/README.md.

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

2.4 Ascend repo under review: Cython wrapper lazily constructs the analyzer, targeting the per-call rebuild issue noted yesterday (#1948)

Updated 10-27 07:42, author zwh1025. The new defect #1945 recorded in yesterday’s briefing noted that the Ascend repo’s Cython wrapper creates a new TVM analyzer on every call, measured at roughly 0.28 ms of fixed overhead; this PR, titled “lazily construct TVM analyzer,” targets exactly that overhead. Such small costs accumulate into observable latency across per-layer calls in large-model inference, making this low-hanging fruit for Ascend-side performance cleanup. Still under review.

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

2.5 Ascend repo under review: GLM-5.3 kpool compression operator, tensor shift amounts on PTO (#1910 / #1949)

Both updated 10-27. #1910 adds a GLM-5.3 kpool compression operator example for the Ascend NPU — moving a concrete model’s structural component into the examples repo, consistent with the TileOPs approach of “operators grow from model requirements”; #1949 lets T.tile.bitwise_lshift accept tensor-form shift amounts on PTO, which previously accepted only constants. PTO is a relatively independent code-generation path on the Ascend side, and tensorizing shift amounts moves that path one step toward general expressiveness for mask and bit-packing operators.

Sources: https://github.com/tile-ai/tilelang-ascend/pull/1910, https://github.com/tile-ai/tilelang-ascend/pull/1949

2.6 Ascend main repo under review: PTO T.Parallel vectorization, AutoSimVF channel, batched 3D ND2NZ, write-count multi-buffering, Bisheng build identity in cache key (#3512 / #3320 / #3689 / #3629 / #3691)

These five under-review PRs form the Ascend-side queue for this window, all updated 10-27. #3512 is the pipeline framework and elementwise support for vectorizing T.Parallel on PTO — a structural change; #3320 adds an AutoSimtVF channel registered into the Ascend pipeline, aimed at automatically choosing between SIMT and vector paths; #3689 supports batched 3D GM-to-L1 ND2NZ copies, directly addressing the transfer needs of batched attention and tiled GEMM; #3629 adds write-count multi-buffering and reorders DeepGEMM loading; #3691 incorporates the Bisheng compiler build identity into the binary cache key, the same idea as 1.5’s compile_flags in the cache key — the Ascend side likewise suffers from stale binary hits after a compiler change.

Sources: https://github.com/tile-ai/tilelang/pull/3512, https://github.com/tile-ai/tilelang/pull/3320, https://github.com/tile-ai/tilelang/pull/3689, https://github.com/tile-ai/tilelang/pull/3629, https://github.com/tile-ai/tilelang/pull/3691

2.7 ROCm under review: architecture detection and image for pip-installed ROCm 10, FP4 conversion and E8M0 type (#3549 / #3665)

#3549 supports the pip-installed ROCm 10 SDK, with changes including architecture detection and a Dockerfile; this relates to the auto-target fix in 1.2 but lands in a different function — this PR touches the /opt/rocm-based get_rocm_arch detection, while 1.2 touches the detection helper in tilelang/rocm/target.py. #3665 fixes float-to-FP4 conversion in the HIP template and adds E8M0 type support, pointing at the availability of block scaling and low-precision quantization on ROCm. Both under review.

Sources: https://github.com/tile-ai/tilelang/pull/3549, https://github.com/tile-ai/tilelang/pull/3665

2.8 Ascend repo daily test: all 2903 items pass, two consecutive days with no failures (#1950)

The Ascend repo’s daily test bot reported at 10-28 06:49 a total of 2903 items, 2903 passed, 100% pass rate, workflow run 40590261193; the previous day’s report (#1947) was likewise 2903 items all passing. For an NPU backend still evolving rapidly, two days of failure-free daily tests indicate stable mainline availability, and this also provides the baseline for the CI split in 2.3 — after the split, both jobs should still maintain this number.

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

2.9 MetaX: organization view has a push timestamp; none of the six branches has new commits in the window (10-27)

The tilelang-metax repo’s pushed_at shows 10-26 07:08, before the start of this window; checking all six branches individually, there are no new commits in the window, so it is treated as having no substantive changes and no item-level content is listed.

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

2.10 Third party: Hexagon DSP backend lands 15 commits in one day — HVX vectorization, VTCM arena, Crouton layout, and DMA builtins all in place (10-27)

The hexagon-backend branch of xwhzz/tilelang-hexagon landed 15 commits on 10-27 alone, the densest external backend activity in this window. The day’s mainline falls into three parts: the low level rebuilds the code generator from CodeGenTileLangCBase, makes DMA an IR builtin with commit groups, sinks the launch grid and worker thread pool into TIR, and plans the VTCM arena in TIR; the middle layer lays down HVX’s single vectorization path (arithmetic mapping, vector math plus an elementwise path), uses an HVX transpose tree for multi-row reductions, software-pipelines HVX mapping, and supports fragments and T.reduce_*; the upper layer introduces Crouton layouts and region copies, runs HVX mapping and reductions directly on Crouton tiles, votes on DSP frequency per session, and finally adds build caching and parallel device tests. The commits carry a breaking-change marker (feat(hexagon)!), indicating that Crouton layouts and region copies changed existing interfaces. Hexagon is the DSP for phones and embedded devices; if this line works out, TileLang’s hardware coverage extends from GPU and NPU to DSP.

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

III. Ecosystem and Adopters (TileOPs / TileFoundry)

3.1 TileOPs 0.1.0 released: the first official version with 206 operators, public API now frozen (10-27)

v0.1.0 was posted at 13:15 on 10-27, the repo’s first GitHub release; the 0.1.0 artifact on PyPI was uploaded earlier, at 07:45 on 10-27 (Beijing time, i.e. 23:45 UTC on 10-26), 45 minutes after the cutoff of the previous daily brief—exactly the change that happened right after yesterday’s brief wrote “the first release is still in the last slot.” The release notes are clear about positioning: an operator library rebuilt for the agent era, where each operator first gets a YAML spec (declaring the signature, the workloads used for tests and benchmarks, the roofline formula), coding agents develop and tune the kernels, and tests and benchmarks derived from the spec decide what makes it into the library; 195 merges tagged foundry entered this way. This version freezes the public API: operators are imported as tileops.., docs/stability.md lists the public modules, and scripts/ci/check_api_compat.py makes any PR that breaks the previous version's usage fail outright; subsequent 0.x versions only add, never change. Several capabilities also land in this version: torch.compile(fullgraph=True) works for all manifest entries that contain call-time tensor inputs and no composition; CUDA Graph capture of operators in the graph test group was verified on H200; automatic tuning after op.request_tune(), with compiled artifacts and tuning results cached under TILELANG_CACHE_DIR and reused across processes; tileops.backend lets backend distribution builds plug their own device's whole operators in via register_kernel_builder; tileops.trace records in-kernel timelines. The release acknowledgments note that @FeathBow alone contributed 98 merges.

Source: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0, https://pypi.org/project/tileops/

3.2 TileOPs 0.1.0 benchmark: on H200, 2262 of 2265 rows no slower than the fastest implementation, geometric mean 1.37x (10-27)

The nightly benchmark compares each TileOPs kernel against every other implementation of the same operator on the same GPU; each row’s ratio is the fastest comparator’s time divided by TileOPs’ time, with a value above 1 meaning faster. Across all 2265 rows and 16 operator families, the geometric mean is 1.37x, with 2262 rows at or above 0.99. By family the differences are pronounced: sequence_modeling 4.00x, sampling 2.95x, convolution 1.86x, mamba 1.61x, linear_attention 1.54x, pool 1.51x, quantization 1.45x, attention 1.42x; the largest family, elementwise at 527 rows, is only 1.09x, and gemm at 235 rows is 1.28x. Comparators include torch, torch.compile, FlashAttention 2 and 3, FlashInfer, FlashMLA, vLLM, DeepGEMM, flash-linear-attention, FlagGems, QuACK, Mamba, Marlin, and cuBLASLt. The three rows below 0.99 are Where and ClampScalar on 16M float32 elements, and an FP8 GEMM on the DeepSeek-V3 prefill output projection, with gaps of 1.1% to 1.7%. This distribution shows the gains come mainly from long-tail operator families rather than mainstream GEMM and elementwise—the latter being exactly where every implementation competes hardest.

Source: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.3 TileOPs 0.1.0’s known gaps and next steps: fused operators, whole models, compute-communication fusion (10-27)

The release notes list four known issues: FusedAddLayerNormFwdOp loses precision on rows far from zero and overflows on rows near 1e36 (#2773, already recorded in yesterday’s brief); if a single host has GPUs of different architectures, a kernel compiled for one GPU may run on another (#2302); Bwd-class operators must be called directly and do not hook into torch.autograd; a single operator instance supports only one thread and one CUDA stream at a time. Validation covers only H200, Python 3.12, and 3.13; some kernels have additional SM80, SM86, and SM89 paths, which run on the corresponding devices but are outside the release validation scope. The roadmap has three steps: first, complex fusion across operator boundaries (attention together with its preceding and following projections and normalization, whole-layer MoE, quantization inside the GEMM epilogue); then whole-model prefill and decode compared end-to-end against serving engines; finally compute-communication fusion (tensor-parallel GEMM with all-reduce, expert-parallel MoE with all-to-all). If these three steps land, TileOPs’ competitor shifts from single-operator libraries to inference frameworks.

Source: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.4 Four TileOPs merges: where and masked_fill become tunable, D128 causal diagonal scores by 64 keys, sparse long GDN prefill split, changelog dated by release day (#2782 / #2783 / #2779 / #2781)

Four merges landed on 10-27, exactly the wrap-up before the release. #2782 fixes request_tune() throwing ValueError on WhereFwdOp, MaskedFillTensorFwdOp, and MaskedFillScalarFwdOp: automatic tuning rejects randomly generated integer inputs, while these three kernels pass bool conditions or masks as uint8; the random selector only changes which operand each element takes and does not change the candidates’ workload, so declaring random integer inputs is safe; three test nodes were added, and all three fail on main. #2783 optimizes the causal diagonal of the D128 dense program: when a causal call has only one key segment and both lengths are on 128-row boundaries, the diagonal key tile of the first consuming warpgroup masks out the upper 64 keys for all rows, so this step now runs as an m64n64 QK with a 64-key softmax; the choice is a compile-time flag (diag_aligned, part of the kernel key), because branching at runtime between 64- and 128-wide QK would make ptxas serialize every WGMMA on some specializations (C7511); with the flag, all D128 prefill specializations in the manifest no longer emit C75xx. #2779 splits sparse long-sequence GDN prefill on SM89 into a state scan plus a separate output program, applicable to dense, key-major D128 inputs with at least 16384 tokens and a state-scan CTA count no more than a quarter of the SM count; in the benchmark, (1, 16384, 4, 128) bf16 drops from about 0.612 ms to about 0.470 ms, which is 0.939x against flash-linear-attention’s 0.441 ms. #2781 sets the date of 0.1.0 in the CHANGELOG to the day it entered PyPI (2026-10-27), and removes the one-time release configuration and Actions click path from the contributor docs.

Source: https://github.com/tile-ai/TileOPs/pull/2782, https://github.com/tile-ai/TileOPs/pull/2783, https://github.com/tile-ai/TileOPs/pull/2779, https://github.com/tile-ai/TileOPs/pull/2781

3.5 Two TileFoundry merges: whole windows lowered as the operand itself, dynamic shared memory sized to actual kernel usage (#273 / #271)

Both merged on 10-27. #273 fixes sm_89 GEMM lowering: what schedule finalize emits does not compile, because lowering wraps the entire accumulator in a window that starts at zero and spans itself, T.ptr_of(acc[0:0+16, 0:0+8]), and CUDA codegen rejects pointer slicing of sharded register tiles with “expected a plain source layout”; MMA also reads register fragments as bare Slices like lhs_frag_dst[0:0+16, 0:0+32], which codegen likewise has no handler for. The fix is on the lowering side: StoragePlan.operand already treats “a window that starts at zero and spans the full operand” as the operand itself, but materialization and instruction lowering still construct a Slice; the change makes the determination once—applying the same substitution to MLIR’s TrivialSubViewOpFolder. #271 fixes dynamic shared memory: the synthesized host entry launches every TIR kernel with dynamic_smem=0, so kernels that access shared memory through integer TensorView offsets get no space allocated, and above 48 KiB cudaFuncAttributeMaxDynamicSharedMemorySize must be set, which the launcher never did; the author did not adopt the suggestion in #267 to “allocate the full target capacity,” instead sizing to the kernel’s actual usage, on the grounds that static and dynamic shared memory count against the same limit and how much each block requests determines how many blocks fit on an SM—on sm_89, filling 99 KiB would let only one block fit even for a kernel using 512 bytes. CUTLASS and TileLang both size the same way.

Source: https://github.com/tile-ai/TileFoundry/pull/273, https://github.com/tile-ai/TileFoundry/pull/271

3.6 External ecosystem: CANN community’s TileLang operator repo, a zero-change wrapper for the PPU build, a DeepSeek V4.1 Flash fragment example (10-27)

Three external repos were active within the window, and together they trace the formation of “Ascend-side operator assets.” Atlas-Ascend-Fork/ops-tilelang is the CANN community’s TileLang operator repo for Ascend, created 10-20, with a 10-27 commit “auto-discover incremental tests and remove the operator mapping table”; the README states it supports only Ascend 950 and above, requires CANN 9.3.0 or later with the community weekly build, and its design principles are separation of interface and scheduling, aggregation by business domain, correctness before performance, separation of benchmark and tests, and baselines saved sharded by test file and chip cluster. kevin3113/models-tilelang uses seven operators from this repo to chain together a DeepSeek V4.1 Flash “attention output plus pre-FFN processing” fragment (mhc_post_forward, GEMM with sum of squares, reduction plus RMSNorm, mix split, sinkhorn, apply_mix, ffn_norm), with HIDDEN=5120, MHC=4, default decode shape SEQ_LEN=1, supporting SuperKernel and npugraph_ex’s ACL graph; measured across B from 1 to 256, SuperKernel yields 1.17% to 2.75%. flytiger-eco/tilelang-wrapper-for-sail addresses migration in the other direction: running CUDA-style TileLang code unchanged on tilelang-sail (the PPU backend distribution build), by intercepting the import protocol and site mechanism purely at runtime, without modifying upstream source or repackaging, and explicitly erroring rather than silently degrading for CUDA features with no PPU counterpart. The three correspond to the three links of operator supply, model deployment, and backend migration.

Source: https://github.com/Atlas-Ascend-Fork/ops-tilelang, https://github.com/kevin3113/models-tilelang, https://github.com/flytiger-eco/tilelang-wrapper-for-sail

IV. Community, Tutorials, and Events

4.1 Documentation site receives six automated updates in a single day (10-27)

tilelang.github.io had six Update docs commits within the window (05:20, 06:22, 07:19, 09:28, 22:42, 22:43), all pushed by a bot, with content coming from documentation regeneration after merges into the main repo. On release day, the documentation site’s update frequency kept pace with the main branch merge cadence—a normal linkage, so the content is not itemized separately.

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

4.2 Four documentation fact-check items in the Ascend repo, with one thread reaching 46 comments (#1628 / #1655 / #1656 / #1659)

The Ascend repo had four documentation PR updates titled “fact check” on 10-27: T.tile.clear (#1628), T.tile.bitwise_or (#1655), T.tile.bitwise_not (#1656), and T.tile.bitwise_lshift (#1659), with 10, 46, 18, and 20 comments respectively; three of them carry a bot label. A single round of checking rolling up to 46 comments shows the discussion is not just about wording—this kind of per-API behavioral verification is effectively building an archive of semantic conventions for the Ascend side. For a new backend line, the cost of aligning documentation with implementation is often underestimated; this set of PRs is visible investment.

Sources: https://github.com/tile-ai/tilelang-ascend/pull/1655, https://github.com/tile-ai/tilelang-ascend/pull/1656

4.3 Under review: extracting lower trace into a standalone development toolkit (#3583)

This PR under review, updated on 10-27, extracts lower trace from the main repo into tilelang-devkit. Trace information from the lowering process is very useful for debugging backend adaptation, but keeping it in the main repo adds compile-path and dependency burden; once extracted into a standalone toolkit, the main repo can stay lightweight and attach it when needed. Changes of this kind—”externalizing debugging capability”—usually appear when a project starts to have multiple downstream backends; within this window, the Hexagon, PPU, and Ascend lines are all moving at once, which is exactly that kind of moment.

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

4.4 News side: no new coverage of this window on Google News or HN within 24 hours (10-27)

Four query groups in Chinese and English (TileLang, TileOPs, TileLang operators, tile-ai intelligent computing kernel) returned items all falling between late September and mid-October: one batch covers DeepSeek and Huawei open-sourcing the Ascend version of TileLang on September 29–30 (SCMP, Reuters, Tom’s Hardware, Leiphone, InfoQ, etc.), and the latest is an October 22 SemiAnalysis analysis of the symbiotic relationship between Huawei and DeepSeek—all outside the window. The exact-match query for TileOPs returned zero hits. Within the window, Hacker News had only one item related to a GameMaker reimplementation, which was filtered out as noise. On arXiv, the most recent paper with the TileLang keyword is dated 2026-07-24. Today’s content focus is entirely on the GitHub side, which matches the timing of release day—the official release has not yet entered the media coverage cycle.

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

V. Trend Observations

First, TileOPs 0.1.0 changes the cost structure of “operator libraries” to a different algorithm. The traditional approach is that the number of operators grows linearly with hardware and model requirements, with manpower as the bottleneck; here, each operator first gets a YAML spec that clearly declares its signature, payload, and roofline, then an agent generates and tunes the kernel, and finally spec-derived tests and benchmarks serve as the admission gate. The 206 operators and 195 foundry-labeled merges are the first full delivery of this process. Whether it can be sustained depends on two things: whether the spec can cover complex fused operators (the first step of the roadmap is exactly this), and whether the benchmarks can block kernels that are “tuned to the benchmark but not actually robust”—the known issue of FusedAddLayerNormFwdOp losing precision on rows far from zero is precisely a sample of this kind of risk.

Second, the merges in this window are almost all “silent errors.” ROCm’s math functions defaulting to fast built-ins can make sin/cos differ by three orders of magnitude; address_of being treated as a read can cause shared memory after an opaque call to read pre-write contents; a T.ws(-1) region runs on zero threads, silently discarding all stores without reporting an error; slice filling validates by length rather than by end point, writing directly out of bounds; and a negative index remains negative after wrapping, landing before the allocation start. What these five share is that they compile, run without errors, and produce wrong results—and all were discovered in a concentrated burst within the past two weeks. For a language meant for writing kernels for large models, the cost of such defects is far higher than compile-time errors, and they point to the same engineering requirement: reject at compile time the patterns that “don’t error but behave incorrectly.” The two boundary-check fixes and the ws(-1) fix in this window are doing exactly that.

Third, backend coverage is shifting from “vendor-led” to “community-spontaneous.” Beyond the four officially maintained lines—Ascend, Hygon, MetaX, and Moore Threads—the Hexagon DSP backend landed 15 commits in one day and extended linear algebra primitives to Crouton tiles, a zero-change wrapper appeared for the PPU version, and the CANN community built its own operator repo and used it to string together a model fragment for DeepSeek V4.1 Flash. None of these actions are in tile-ai organization repositories, yet all depend on the stability of TileLang’s language layer. This is also the significance of 0.1.0 freezing the public API and making API compatibility checks a CI gate—the more downstream users there are, the more valuable backward compatibility becomes.

Fourth, cache identity has been targeted by two backends at the same time. The main repo incorporates compile_flags into the autotuning cache key (#3338), and Ascend incorporates the Bisheng build identity into the binary cache key (#3691). The same kernel produces different binaries under different compile options or different compilers; if the cache keys only on source and shape, it will produce false hits, and the measured performance will not match the actual compilation result. Two independent lines doing the same thing on the same day shows that this kind of problem is already painful enough in real use.

Appendix: Sources and Verification Notes

Source Verification Table

Source Verification Result
GitHub tile-ai organization (29 repos) 7 repos pushed within the window: tilelang, TileOPs, TileOPs-nightly, TileOPs.github.io, tilelang-hygon, TileFoundry, tilelang.github.io
Merged ground truth (is:merged) 20 PRs, each checked individually to confirm merge time falls within the window (tilelang 13, TileOPs 4, TileFoundry 2, tilelang-hygon 1)
Open ground truth (is:open updates) 45 PRs, all checked after sorting by update time
New issues 7, each read in full; 2 of them are daily test bot reports from the Ascend repo
Releases page and PyPI TileOPs v0.1.0 posted 10-27 13:15, PyPI artifact uploaded 10-27 07:45; tilelang latest remains 0.1.15 (09-30), tilefoundry latest 0.0.4rc1 (10-25)
tilelang-ascend repo No commits within the window, but dense PR and daily test issue activity
tilelang-metax repo Organization view shows a push timestamp, but none of the six branches has commits within the window; treated as no substantive change
tilelang-musa repo Last push 09-30, no commits within the window
Adopters (TileKernels / FlashQLA) No pushes within the window (10-14)
External backend repos Hexagon backend 15 commits; CANN operator repo, PPU wrapper, and model snippet repo all have activity within the window, see 2.10 and 3.6
Google News multi-query in Chinese and English Exact queries returned zero hits; all returned items fall between late September and mid-October, all outside the window
Hacker News One GameMaker reimplementation item within the window, unrelated to TileLang, excluded
arXiv Keyword search found no new papers within the window; most recent is 2026-07-24

Full Source List