T.Parallel 对接 PTO Roadmap #3

1. 背景与目标

现状：

• target="ascend" 后端已合入 tile-ai 主线（tile-ai/tilelang#3308），PTO driver/codegen 已合入主线（tile-ai/tilelang#3310）。
• VMI 语言面（tl.vmi.*、VMI codegen）还没有合入主线，只在 fork 的 asc 分支（deepseek-ai/tilelang#250）。
• T.Parallel 对接 PTO 的完整实现目前只在 fork 侧：PR269（T.Parallel→PTO VMI 自动向量化）、PR262（控制流 + reduce_max/min/sum），均为 OPEN。
• ST 验收套件已移植到 PTO-ISA/tilelang 的 pto-dev 分支：PR #2，examples/ascend/t_parallel_st_suite/（21 个 SimtVF T.Parallel kernel + 21 个手写 PTODSL 对照版）。

目标：用户在 T.SimdVF(lanes=L) 内写标量并行的 T.Parallel 程序（element-wise、控制流、reduce），编译器自动向量化为 PTO VMI，经 PTOAS 编译在 Ascend 950 上正确运行。全部能力分若干 PR 合入 tile-ai 主线，每个 PR 的发起是一个路标。

分工原则：

2. Roadmap 总览

• M3 的主线 numeric 验收在 PTOAS P1（mem_bar，见 §5.1）合入后闭环；P1 需要修复后合入 PTOAS 主线。

3. 里程碑表

PR-1 拆分方式（待定）：当前倾向把 VMI 语言面并入 PR-1 一次合入；备选是 VMI 语言面单独先行一个 PR（review 工作量更小）。视社区反馈决定。

4. TileLang 侧工作分解

5. PTOAS 侧工作分解

5.1 P1 问题说明（供 PTOAS 对接）

最小复现（TileLang 源码）：

@T.prim_func
def main(A: T.Tensor((128,), "float32"), B: T.Tensor((1,), "float32")):
    with T.Kernel(1):
        a = T.alloc_shared((128,), "float32")
        best = T.alloc_shared((1,), "float32")
        T.copy(A, a)
        with T.SimdVF(lanes=64):
            for i in T.Parallel(128):
                a[i] = a[i] * 2                       # Parallel 写 UB（in-place）
            T.reduce_max(a, best, dim=0, clear=True)  # 同 region direct reduce 读同一 UB
        T.copy(best, B)

历史复现中 TileLang 生成的 PTODSL（关键片段，保留当时输出，不代表当前目标写回形态）：

with pto.vecscope():
    mask = pto.vmi.create_mask(64, size=64)
    for q in range(0, 2, 1):                       # unit 循环：vload → vmul → vstore 写 UB
        v = pto.vmi.vload(pto.addptr(a_ptr, q * 64), 0, size=64)
        pto.vmi.vstore(pto.vmi.vmul(v, two, mask), pto.addptr(a_ptr, q * 64), 0, mask)
    # ← unit 的 vstore 与 reduce 的 vload 之间没有任何保序
    r0 = pto.vmi.vload(a_ptr, 0, size=64)          # reduce：读同一 UB 地址
    r1 = pto.vmi.vload(a_ptr, 64, size=64)
    red = pto.vmi.vcmax(pto.vmi.vsel(mask, pto.vmi.vmax(r0, r1, mask), r0), mask)
    pto.vmi.vstore(pto.vmi.vbrc(red, size=64), best_ptr, 0, pto.vmi.create_mask(1, size=64))

问题：同一 SIMD VF 内部、同一 UB 地址的 vstore（Parallel 写）→ vload（reduce 读）RAW 依赖没有被保序，设备上 reduce 读到写入前的旧值。确定性复现（非竞态）；9 用例数值矩阵 6 失败 3 通过；region 之后 copy-out 能读到新值（store 本身已落盘）。

需要 PTOAS 做的工作：在 VPTO 后端的 vecscope 内做 load/store 依赖分析，对真实依赖边按需插入 pto.mem_bar（向量 store/load 之间 RAW 用 VST_VLD、WAR 用 VLD_VST）。pto.mem_bar 是 vecscope 内合法原语（16 种 kind，TileOps a5 生产模板已有 vsts→vlds 同 tile 插 VST_VLD 的同构先例）；注入单行 pto.mem_bar "VST_VLD" 后 9 用例矩阵 9/9 全部通过（手工注入的实验性验证，不代表 PTOAS 自动插入已落地）。同一机制还需按真实指令队列覆盖 direct Reduce 写回后的读取：vector store → vector load 使用 VST_VLD，vector store → auxiliary scalar load 使用 VST_LD，不能将两者混用。TileLang 侧不实现该分析。mem_bar 有运行时开销，应按需而非无条件插入。

6. 验收用例

• ST 套件（PTO-ISA/tilelang#2，examples/ascend/t_parallel_st_suite/）：kernels/ 21 个 kernel 全部纳入本期验收，验收路径为 SimdVF→PTO 向量化，暂不设置 opsim 性能门槛（kernels_ptodsl/ 21 个手写 PTODSL 版作正确性与性能参照）。
• 产品用例：TopK（非整除、tie-break、-inf 边界）、RMSNorm（compile + numeric）；回归：vecadd/vintlv/PTO dtype/PTO SIMT/手写 VMI、AscendC 不回归。

6.1 逐 kernel 特性需求矩阵（21 个 kernel 逐个分析的结论）

逐 kernel 需求（"读/写"标注间接访存方向）：

路标解锁汇总：

7. 风险与依赖

1. P1（mem_bar）需在 PR-3 发起前合入 PTOAS 主线，正由 PTOAS 团队开发；若滑期则 M3 顺延。
2. PR-1 拆分方式需社区反馈：VMI 语言面并入首个 T.Parallel PR（当前倾向）vs 单独先行（备选）。
3. TileLang 与 PTOAS 的版本配套：TileLang 跟随 PTOAS/CANN 主线演进，上游合入时需声明依赖的 PTOAS 最低版本（P1 落地后 M3 起有硬性版本依赖）。
4. Fragment/slot 是新子系统，无现成迁移来源，工作量不确定性最高；且 M4 是 ST 用例解锁数最多（10 个）的路标，其延期对验收面影响最大。
5. ST 套件验收需经 SimdVF→PTO 向量化路径：现有 kernels/ 为 SimtVF 形态，验收时需以 SimdVF 形态编译通过并数值正确。

8. 待定事项

• PR-1 拆分方式：VMI 并入首个 T.Parallel PR（倾向）还是单独先行（备选）——视社区反馈。
• CF6 内嵌串行循环的路线选择：cf6 的 Parallel 体内含常数界 T.serial(3) Newton 迭代。cf6 已纳入本期验收范围，以下路线均以覆盖 cf6 为前提：(1) Legalize 展开常数界串行循环，并合并分支臂内同位置写链；(2) Vectorize 直接支持 Parallel 体内 uniform 串行循环；(3) 验收层改写用例形态。待讨论确定。