mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-20 06:30:07 +08:00
24a023c636
* [Transform] Never classify side-effecting binds as replayable IsReplayableScalarBind (introduced in #2245) decides whether a scalar bind may be dissolved and its value re-evaluated at every use site (IfStmtBinding, default-on) or replayed across pipeline stages (PipelinePlanning, InjectSoftwarePipeline). The predicate checked the dtype and read/write disjointness, but never the value's effect kind, so a bind like pos = T.atomic_add(counter[0], 1, return_prev=True) used at two sites (bound check + store index) was inlined into both use sites, executing the atomic once per use site instead of exactly once and silently corrupting compaction/histogram/top-k kernels since v0.1.11. The atomic intrinsic is correctly registered kOpaque; its pointer argument enters the expression through tvm_access_ptr, which the access collector records as a read, so the write-disjointness check passed and the bind was declared replayable. Reject any bind whose value has SideEffect > kReadState, mirroring the effect guard StmtSimplifier already applies to let inlining. Read-only values remain replayable; read consistency is enforced by the existing write-set disjointness check. The shipped examples/deepseek_v32/topk_selector.py uses exactly this pattern and is miscompiled on v0.1.11/v0.1.12; its test only printed recall and could not fail, so it also gains a recall assertion. Regression tests: - transform: IfStmtBinding must keep a bind whose value is an atomic RMW returning the previous value. - language (CUDA): compaction kernel using atomic_add(return_prev=True) must contain exactly one AtomicAddRet and produce a permutation with counter == N. * [Transform] Validate scalar bind replay safety Track access-pointer write masks during bind and pipeline analysis, and reject unsafe precomputed replay annotations during pipeline injection. Add regressions for atomic target reads and side-effecting replay masks. * Remove recall calculation from top-k selector Removed recall calculation and assertion for top-k selection. --------- Co-authored-by: LeiWang1999 <leiwang1999@outlook.com> Co-authored-by: Lei Wang <34334180+LeiWang1999@users.noreply.github.com>