[Bug] PR #1341 导致 persistent fragment SIMT kernel 编译失败：VF_SIMT size patch 缺失 ELF callee symbol（rmsnorm_persistent_simtvf）

[jimmychou0]

摘要

test_rmsnorm_persistent_simtvf.py（ub-to-fragment 与 gm-to-fragment 两条路径）在 feat-persistent-unroll-promotion 分支（PR #1341）上编译失败：

Error: VF_SIMT size patch: missing ELF callee symbol 'main_kernel_simt_1_simt_entry'

同一份 TileLang 生成的 kernel 源码（经逐字节 source-diff 验证一致，且该 kernel 不使用 T.unroll——循环是 T.SimtVF 内的 T.Parallel lane 循环）在当前 main（PTOAS 0ee47471）上编译、运行均正常，因此这是 PR #1341 分支引入的回归。

环境

• PTOAS：340650086（feat-persistent-unroll-promotion，PR feat(pto): auto-promote persistent fragment loops to full unroll #1341 代码）
• 对照：PTOAS main-latest 0ee47471 构建——同一 kernel 通过
• TileLang：asc@505acc0f；kernel 源码在两个 codegen 变体间逐字节一致
• 机器：A5 板卡（CANN 9.1.0-beta.3），LLVM 19.1.7 feature-vpto

复现

TileLang 侧（tilelang 仓库）：

pytest examples/ascend/test_rmsnorm_persistent_simtvf.py
# main PTOAS 构建：2 passed
# PR #1341 PTOAS 构建：2 failed，报上述 VF_SIMT size patch 错误

kernel 是带 persistent weight fragment 的 RMSNorm：T.alloc_buffer(..., scope="local.fragment") 式 weight 定义在 SIMT region 之外，在一个 section 初始化、另一个 section 消费，with T.SimtVF(threads=256) 内是 T.Parallel(TILE) 循环。

两种 codegen 下生成的 PTODSL 均保持 SIMT body 循环为 trace 期 pto.static_range（source-diff 逐字节一致），因此差异完全在 PTOAS 对 persistent fragment + SIMT size patch 路径的处理上。

分析

PR #1341 新增的 pto-promote-persistent-fragment-loops 会把 touching llvm.alloca {pto.persistent} 的循环改写为 pto.unroll = "full" 并加 pto.persistent_unroll 标记。对这个 kernel，promotion 把包裹整个 pto.section.simt 的 kernel 级 tile 循环（trip 64）也展开了：

1. 展开把 SIMT section 按 trip 克隆 64 份，SIMT outlining 为每个克隆生成独立的 entry 函数（加上 init entry 共 65 个），每个都是单调用点的 linkonce_odr simt_entry 函数（noinline 已设置——并非普通的内联候选）；
2. BiSheng 对这些单调用点 ODR 克隆做折叠/内联，其中一部分的 _simt_entry ELF symbol 消失。两次相同 IR 的编译缺失的 symbol 还不同（一次缺 main_kernel_simt_1、一次缺 main_kernel_simt_3），说明折叠行为非确定；
3. VFSIMTSizePatcher 的 manifest 来自 bisheng 编译前的 LLVM module，要求每个 callee symbol 存在，于是 fail-fast 报 missing ELF callee symbol；发射模块同时膨胀 45 倍（37k 行 vs 811 行）。

关键事实：物化只要求 section 内 的 persistent access 解析到静态 slot。失败 kernel 中所有 persistent GEP 索引都是常量（从不引用 kernel 级循环变量），展开包裹 section 的循环对 slot 静态化没有任何贡献——它只制造了脆弱的克隆结构。而 section 内部（或不包裹 section 的）循环的展开与提升不受影响，仍然必要且保留。

修复（分支本地提交，见后续评论）

pto-promote-persistent-fragment-loops 在 enclosing-loop 收集时跳过体内包含 pto.section.simt region 的循环，pass 头注释、lit case 6 与设计文档同步更新。验证（A5，CANN 9.1.0-beta.3）：

• lit：promote_persistent_loops（RUN1/UNROLL/NOMETA）、_invalid、_dynamic_unroll（全部硬错误契约）、_materialize 全部 PASS；
• TileLang persistent rmsnorm 两路径数值 2/2 PASS，发射结构恢复为 PR 前的双 entry 形态（811 行）；
• 全量 PTO 回归（rmsnorm/mhc/intrinsics 含 unroll hint 用例）26 passed；
• review 要求的 hint 变体（kernel_handwritten_unroll_hint.py / kernel_auto_promoted_unroll.py）：均编译通过；无 hint 变体上 promotion 恰好提升 section 内两个循环、tile 循环不被触碰、物化发出 96 处 keep/resume。

后续评论补充了"手写 unroll="full" 在 kernel 级循环上同样失败"的对照证据：该形态在本工具链上从未可用，属于 SIMT outlining（单调用点 ODR 克隆）+ BiSheng（折叠）的上游能力缺口，需要 outlining/bisheng 侧保持 entry symbol 稳定（如对 simt entry 克隆禁用 ODR 折叠）才能真正支持。

[jimmychou0]

Root cause identified and fixed on the feat-persistent-unroll-promotion branch (local; will be pushed to the PR)

Root cause chain (confirmed on the failing RMSNorm persistent kernel, traced PTO IR captured):

1. pto-promote-persistent-fragment-loops promoted the kernel-level tile loop wrapping the whole pto.section.simt (trip 64) to full unroll, exactly as the pass's discovery rules intended.
2. The unroll cloned the SIMT section 64 times; SIMT outlining then emitted 64 outlined entry functions plus the init entry (65 total), each a single-callsite linkonce_odr simt_entry function (noinline was already set — these were not plain inlining candidates).
3. BiSheng folded/inlined a subset of those single-callsite ODR clones into their callers, so their _simt_entry ELF symbols went missing. Which clones disappear varies run to run (we observed main_kernel_simt_1 and main_kernel_simt_3 missing on two identical-IR compiles), which is why the failing symbol name was not stable.
4. VFSIMTSizePatcher's manifest is built from the pre-bisheng LLVM module, requires every callee symbol, and fails fast: missing ELF callee symbol. The emitted module also ballooned 45x (37k vs 811 lines).

Key insight: materialization only requires the persistent accesses inside the section to resolve to static slots. In the failing kernel every persistent GEP index is constant (they never reference the kernel-level induction variable), so unrolling the section-wrapping loop contributed nothing to slot resolution — it only produced the fragile clone structure. Loops inside the section (or wrapping accesses without wrapping a section) still need promotion and are unaffected.

Fix (branch-local commit): during enclosing-loop collection, skip any scf.for whose body contains a pto.section.simt region, with the rationale documented in the pass header and the design doc. Lit case 6 (promote_kernel_level_loop) is updated to assert the loop is now left alone (kept, no pto.unroll/pto.persistent_unroll attrs, survives the unroll run).

Validation (A5 board, CANN 9.1.0-beta.3, LLVM 19.1.7 feature-vpto):

• test/lit/vpto/promote_persistent_loops.pto RUN1/RUN2(UNROLL)/RUN3(NOMETA): PASS with the updated case 6
• promote_persistent_loops_invalid.pto, promote_persistent_dynamic_unroll.pto (all hard-error contracts), promote_persistent_loops_materialize.pto (MATERIALIZE): PASS
• TileLang examples/ascend/test_rmsnorm_persistent_simtvf.py (both paths): 2/2 PASS — emitted structure back to the two-entry form (811 lines, identical to pre-PR), numerics correct
• Full PTO regression (rmsnorm pto, mhc_copy pto, pto_intrinsics incl. unroll-hint cases): 26 passed

[jimmychou0]

Additional evidence that the skip rule is the right call, and that this was never a supported shape: putting a handwritten unroll="full" on the section-wrapping tile loop (bypassing the promotion pass entirely, on the fixed build) fails with the exact same VF_SIMT size patch: missing ELF callee symbol error. Fully unrolling a loop that wraps a pto.section.simt has never been compilable end-to-end on this toolchain — the SIMT outlining emits one single-callsite linkonce_odr entry per clone and BiSheng folds a subset of them, so the symbol goes missing regardless of whether the unroll came from a handwritten hint or from promotion. Supporting that shape would need the outlining/bisheng side to keep the entry symbols stable (e.g. disable ODR folding for simt entry clones), which is an upstream capability gap rather than something promotion can fix.

[jimmychou0]

修复已升级为依赖感知版（PR 分支 13761b83f），先前的一律跳过过于保守：

• 跳过条件收窄：仅当包裹 section 的循环的 induction variable 不出现在该循环体内任何 persistent GEP index（沿纯整数算术链追踪）时才跳过——即循环只搬数据、fragment 按 lane 常量索引的形态（rmsnorm 的 tile 循环）。
• 依赖形态保留提升：induction variable 参与 fragment slot 选择时循环仍被提升——展开是这类 access 静态化的唯一途径。端到端验证：无手写 hint 的依赖形态 kernel（123 克隆）由 pass 自动提升、编译通过、逐 slot 数值校验每个克隆正确读取各自驻留值。
• 中途关键发现：materialization 会把依赖形态的 persistent load 改写为各克隆对 resume 驻留值的选择，GEP 常量差异在进入 LLVM 前即被抹平——同奇偶克隆机器码等价，BiSheng 的折叠对源码层常量差异"免疫"（依赖/非依赖两种形态的展开同样失败），因此依赖感知是正确的精确边界而非规避手段。

验证矩阵（A5，CANN 9.1.0-beta.3）：全部 promotion lit（含新增依赖用例 6b）PASS；无 hint 依赖 kernel 自动提升+数值正确；rmsnorm persistent 2/2、全 PTO 回归 26 passed。