feat(pto): auto-promote persistent fragment loops to full unroll - #1341
Open
jimmychou0 wants to merge 4 commits into
Open
feat(pto): auto-promote persistent fragment loops to full unroll#1341jimmychou0 wants to merge 4 commits into
jimmychou0 wants to merge 4 commits into
Conversation
jimmychou0
force-pushed
the
feat-persistent-unroll-promotion
branch
5 times, most recently
from
August 26, 2026 04:10
85b6fc6 to
cadfca2
Compare
KurrinQu
reviewed
Aug 27, 2026
| @@ -0,0 +1,62 @@ | |||
| // Copyright (c) 2026 Huawei Technologies Co., Ltd. | |||
Collaborator
There was a problem hiding this comment.
整体实现OK，建议做一些端到端验证测试。可以包括：初版寄存器驻留优化上传了几个使用dsl开发的驻留test/vpto/cases/kernels/rmsnorm-fp32-4096-persistent-fragment-ptodsl/kernel.py等，可以尝试将前端语法改为unroll hint进行端到端测试，再进行tilelang codegen修改测试。
jimmychou0
force-pushed
the
feat-persistent-unroll-promotion
branch
3 times, most recently
from
August 29, 2026 02:43
13761b8 to
422212f
Compare
Contributor
Author
|
补充了嵌套 section-wrapping 循环的 lit 用例（Case 6c/6d），固化依赖感知规则在嵌套形态下的判定（
两者均有端到端数值佐证（123 克隆真实 kernel，逐 slot 校验：外层依赖形态每个克隆正确读取各自驻留 slot；内层依赖形态克隆内静态化、外层循环调用正确）。全部 promotion lit（RUN1/UNROLL/NOMETA/invalid/dynamic）与 runtime 回归（rmsnorm persistent 2/2、intrinsics 全量）通过。 |
Persistent fragment materialization requires every loop touching a
persistent buffer to be fully unrolled so each access resolves to a
stable resident slot. Until now this precondition relied on authors
manually writing {pto.unroll = "full"} on those loops; a missing
annotation made materialization fail without a clear cause. The new
pto-promote-persistent-fragment-loops pass (wired into
prepareVPTOForEmission immediately before pto-unroll-loops) automates
the discovery and marking, per the persistent-fragment plan:
- Discovery: every llvm.alloca carrying {pto.persistent} is an explicit
entry point (no structural re-inference). The pass walks the alloca's
use graph and collects every enclosing scf.for of every related op -
loops directly wrapping an access inside a SIMT section, kernel-level
loops wrapping whole pto.section.simt regions, and every nesting
layer.
- Promotion: no-hint and "enable" loops are rewritten to "full"
(overriding enable is deliberate: keeping enable would leave the loop
to the metadata stage and materialization would lose the static access
set it depends on); an existing "full" is kept. Every promoted loop
is additionally marked {pto.persistent_unroll}.
- Fail-fast (hard errors, never a silent fallback): a fixed
pto.unroll_factor on a persistent loop; a persistent access nested
under scf.while; a statically known trip count above the new
max-persistent-unroll-trip-count option (default 128, error names the
function, the trip count and every persistent allocation). Diagnostics
are collected across the whole function before failing once, so the
emitted set stays deterministic under the parallel pass manager.
- The marker makes pto-unroll-loops turn each of its drop-with-remark
fallbacks into a persistent-specific hard error: a dynamic trip count,
an empty body, a non-index induction variable, or a failed unroll.
Silently keeping a persistent loop would break materialization
downstream. The marker disappears together with the unrolled loop.
- The {pto.persistent} attribute name is promoted into PTO.h and shared
with SIMTPersistentFragmentAnalysis instead of being file-private.
Tests: promote_persistent_loops.pto covers no-hint/enable/full
promotion, unrelated loops left untouched, nested loops, kernel-level
loops wrapping sections, iter_args threading, a promotion + unroll end to
end run leaving no attributes behind, and (via a third RUN line) that an
overridden enable hint produces no llvm.loop.unroll metadata;
promote_persistent_loops_invalid.pto covers the fixed-factor, scf.while
and guardrail errors; promote_persistent_dynamic_unroll.pto covers the
three unroll-side hard errors, one module per -split-input-file section
so every diagnostic is observable regardless of scheduling.
Docs: the persistent SIMT fragment plan gains an "automatic promotion"
section (discovery rules, enable-override rationale, fail-fast contract,
and the pass's deliberately narrow responsibilities); the loop unroll
hint design doc records it as revision v4.
Validation (144): full lit suite green except the three
element/predicate/vsts offset-index tests that also fail without this
change (they belong to main's hw-native-sys#1330); CANN simulator dsl-st suite passes
with no failures; compliance check reports errors=0.
- discovery walks the pointer flow only (GEP chains; load/store and any other consumer are terminal accesses) instead of following every result's users, so loops that merely consume a loaded value are no longer promoted or hard-errored - deduplicate the scf.while diagnostic per while op - keep relatedOps in insertion order via SmallSetVector so diagnostics stay stable regardless of the related-op count - skip statically zero-trip loops instead of promoting them into a misleading "no constant trip count" hard error - PTO.h: move the persistent-attr constants below the unroll contract comment they interrupted - docs: move the v5 revision entry to the top of the revision list and sync the pointer-flow wording - tests: cover load-result consumers and zero-trip loops staying untouched, and single emission of the scf.while diagnostic
…d dead loops Address codex review findings: - validate the loop-unroll hint before promotion overwrites it, so a malformed pto.unroll value or attribute type on a persistent loop stays a hard error instead of being silently repaired to "full"; the validation now lives in a shared LoopUnrollUtils.h used by both pto-unroll-loops and the promotion pass - skip the whole enclosing-loop chain when any member is statically zero-trip: the access never executes, so promoting a dynamic outer loop would only produce a spurious "no constant trip count" error - compute the static trip count in uint64_t (ceiling division without span+step-1) so extreme bounds like lb=INT64_MIN/ub=INT64_MAX cannot overflow signed arithmetic; counts beyond int64_t range report as non-constant, which keeps them on the fail-fast paths - add an end-to-end lit test driving an unannotated persistent loop through the full ptoas pipeline into fragment materialization
…apping loops Promoting a kernel-level loop whose body contains a pto.section.simt region made VPTO fatobj emission fail on real persistent kernels (issue hw-native-sys#1379): the full unroll cloned the section once per iteration, SIMT outlining emitted one single-callsite linkonce_odr entry function per clone, and BiSheng folded a (run-to-run varying) subset of them, so their _simt_entry ELF symbols went missing and the VF_SIMT size patch failed fast. The emitted module also ballooned (37k vs 811 lines). The skip is now dependency-aware instead of unconditional. A section-wrapping loop is left alone only when its induction variable never reaches a persistent GEP inside the loop (walked through pure integer arithmetic): that loop cannot contribute to slot resolution - materialization only needs the in-section accesses static - while its unrolling only produces the fragile clone structure. When the induction variable does select the fragment slot, unrolling is the only way to statically resolve the access, so the loop is still promoted. That dependent shape is verified end to end: with no handwritten hints the pass promotes it, compilation succeeds, and every clone reads its own resident slot (per-slot numeric check). A false negative in the dependence walk can only leave a loop unpromoted that materialization then rejects with its own non-constant-index diagnostic, never a wrong promotion. Also verified on the way: materialization rewrites dependent persistent loads into per-clone selections of the resumed resident values, so same-parity clones become machine-code equivalent - the fold BiSheng applies is legal whatever the source-level GEP constants were, which is why keeping clones 'different' cannot avoid the symbol loss. Validation (A5, CANN 9.1.0-beta.3, LLVM 19.1.7 feature-vpto): - lit promote_persistent_loops (RUN1/UNROLL/NOMETA incl. new dependent case 6b), _invalid, _dynamic_unroll (all hard-error contracts), _materialize: PASS - No-hint dependent kernel (123 clones): auto-promoted, compiled, per-slot numerics correct; independent variant untouched as before - TileLang rmsnorm persistent simtvf (both paths): 2/2 numerics PASS, emitted structure back to the two-entry form identical to pre-PR - Full PTO regression (rmsnorm/mhc/intrinsics incl. unroll hints): 26 passed - Review-requested hint variants of the PR#978 case: kernel_handwritten_unroll_hint.py (persistent loops written with unroll="full" hints) and kernel_auto_promoted_unroll.py (no hints; promotion must promote them) both compile through ptoas; for the auto variant, promotion promotes exactly the two in-section loops, the tile loop is left alone, and materialization emits keep/resume (96 sites)
jimmychou0
force-pushed
the
feat-persistent-unroll-promotion
branch
from
August 29, 2026 04:37
422212f to
211635a
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
概述
新增
pto-promote-persistent-fragment-loopspass：自动识别访问 persistent SIMT fragment buffer 的循环并提升为强制 full unroll，落实 persistent fragment 方案 与 unroll-enable 与 persistent promotion 设计 中的 promotion 部分。解决的问题:persistent fragment materialization 要求访问 persistent buffer 的循环被完全展开（否则访问无法归一到稳定的 resident slot)。此前该前置条件依赖调用方手写
{pto.unroll = "full"}——漏标时 materialization 在缺少静态访问集合的情况下失败，且没有清晰诊断。本 pass 把"必须手写"变成"自动识别 + 自动提升 + fail-fast"。实现
识别(discovery):
llvm.alloca {pto.persistent}为入口（不做结构推断），沿 use graph(getelementptr → load/store 及其他 user）收集每个相关 op 的所有外层scf.for:section 内直接包裹 access 的循环、包裹整个pto.section.simt的 kernel 级循环、多层嵌套的所有层级；pto.persistent属性名从SIMTPersistentFragmentAnalysis.cpp的私有常量提升为PTO.h共享常量。提升(promotion):
pto.unroll = "full"pto.unroll = "enable""full"（刻意：保持 enable 会把 loop 留到 metadata 阶段，materialization 失去静态访问集合）pto.unroll = "full"pto.unroll_factor = N命中的循环附加内部 marker
pto.persistent_unroll:marker 让pto-unroll-loops把每一条"丢 hint + remark"的兜底都升级为 persistent 专属硬错误——动态 trip、空 body、非 index 归纳变量、展开失败（这些在 promotion 阶段无法预先检查）;marker 随循环展开一同消失，不残留 IR。Fail-fast（硬错误，绝不静默降级——persistent alloca 不是可安全回退的普通内存）:
unroll_factor;scf.while内的 persistent access（该结构无法承载 full-unroll hint);max-persistent-unroll-trip-count(pass option，默认 128，报错含函数名、trip 值与每个 persistent allocation 的位置）。诊断在整个函数内收集完毕后统一失败，保证并行 pass 调度下发出的诊断集合确定。
管线位置:
prepareVPTOForEmission内、pto-unroll-loops正前方（promotion → native unroll → SCCP/canonicalize/CSE → persistent fragment analysis → materialization)。职责边界：只做识别与标注；slot 分配、keep/resume 生成、SIMT outline、LLVM metadata lowering 仍归各自 pass。
测试与验证
promote_persistent_loops.pto（无 hint/enable/full 三种提升、无关循环不动、嵌套全层级、kernel 级循环、iter_args 穿线、promotion+unroll 端到端无残留 attr、被覆盖的 enable 不产生任何 unroll metadata);promote_persistent_loops_invalid.pto(factor / scf.while / guardrail 报错);promote_persistent_dynamic_unroll.pto（三种 unroll 侧硬错误，每个用-split-input-file独立 module 以保证诊断可观测）。提交列表
bf705b632feat(pto): auto-promote persistent fragment loops to full unroll