English | 中文版

13. 下一步:路线图与展望

当前状态

ascend-rs 已经远远跨过了前面几章覆盖领域中的 alpha 阶段。本章的路线图只关注剩下的事情——那些第 2–7、9、10、11、12 章还没有演示过的工作。已经演示过的东西都按“已交付“处理,这里不再重复。

• 宿主机 API:alpha 完成。ACL、内存、stream、event、HCCL、DVPP、profiling、BLAS 都有安全的 Rust 封装。
• ascend_compile crate:独立编译库,提供 Rust API、C ABI、CLI、Python 绑定——AscendC C++ 到 NPU 二进制的唯一统一路径,服务栈中所有前端(架构见 §7.1.1)。
• 设备运行时:1565 个 Rust NPU 内核(489 个编译测试 + 16 个可部署),413 个在 Ascend 910B3 上通过 NPU 正确性验证,覆盖 MultiKernelBench 全部 17 个类别。
• PyPTO / PTO-MLIR 路径:已集成。emitter(mlir_to_pto)→ ptoas 0.26 → AscendC → bisheng。通过这条路径,DeepSeek-R1-Distill-Qwen-1.5B 在 910B2 上端到端 decode 达到 114–187 tok/s(第 10 章)。
• 安全卫士:已交付(第 11、12 章)。在 ptoas 产出的 stage-2 plan 上跑六个 check_* pass;Path A + Path C 两种 ingress 方案覆盖来自第三方前端的 linalg 内核;能捕获 ptoas 自己 rc=0 通过、但 PlanMemoryPass 放置有 bug 的情况。

下面是三条方向——不是任务清单。每一条都吸收了之前分散追踪的多条线索。

方向一:闭合内核编写回路——双缓冲、迭代器、调试信息

核心 MLIR 后端在第 3–5 章所覆盖的运算上已经功能完备:算术、归约、一元数学、标量-向量、16 种激活、17 个组合算子、cube engine 矩阵乘(含硬件 L1→L0B 转置)、以及编译期防止 UB/L1/L0 混用的类型安全 buffer newtype。剩下的是表达力和开发体验问题,不是新运算。

• 基于 queue 的流水线(TQue)。当前 codegen 发射直线型内核并自动推断 pipe_barrier(§7 的 BufDepTracker)。切换到基于 TQue 的双/三缓冲能让 DMA 与 compute 重叠,这是 memory-bound 内核下一步的性能空间。DeepSeek decode 的 lm_head 已经通过手工分块利用了这一点(§9 chunk sweep);这件事应该由编译器自动完成。
• 内核代码里的迭代器组合子。map、filter、fold、zip、enumerate——内核作者期望能写的形状。这些需要 mlir_to_cpp / mlir_to_pto 里对应的 codegen 支持;运算本身已经存在。
• 调试信息。MLIR 后端当前不发射任何 DWARF。加上之后,就能用 gdb/lldb 在生成的 AscendC 里单步调试,这正是“能跑但结果错“且通过 oracle 的内核当前缺失的那块。

这些是工程而不是研究。每一条都是有边界的工作,有清晰的验收标准:cargo run 一个用 .iter().map().sum() 的内核,看到生成的 .cpp 正确使用 ReduceSum;在调试器里单步一个内核的 UB 访问。

方向二:ascend_compile 作为通用编译后端

第 7 章的架构已经命名了我们自己的 mlir_to_cpp / mlir_to_pto 之外的四个前端:TileLang、Triton-Ascend、torch.compile、PyPTO。每一个都产出 AscendC C++;每一个当前都直接调用 bisheng,标志不一致,也没有验证。这条方向的路线图更多是接线而不是写新代码:

• TileLang 今天驱动一个无验证的 subprocess.run(bisheng, ...);通过 Python wrapper 替换为 ascend_compile,TileLang 就能自动获得目标检测、三项验证 pass、以及与我们自己内核一致的标志路径。
• Triton-Ascend 把它的 IR 下沉到 AscendC;最后一公里对任何 C++ 前端都是相同的。
• 带 Ascend 后端的 torch.compile 可以通过 ctypes 调用 libascend_compile.so,完全绕开 Python-to-Rust 依赖。
• PyPTO,当它与 CANN 一起发布时,是最自然的用户:它那 ~90 条的虚拟 ISA 已经下沉到 AscendC,通过 ascend_compile 跑一遍就意味着安全卫士能看到同一批 plan。

交付物不是更多的后端——而是昇腾生态里更少的定制编译流水线。LLVM 的图景适用于这里:多前端,一个经过验证的后端。

这条方向也为方向三铺路,因为它为安全卫士建立了一个公共的拦截点。

方向三:扩大安全卫士的覆盖面

第 11 章和第 12 章分别把 oracle 交付到了 PTO-MLIR 和 ingested linalg 上。自然的下一步是保留检查 pass、替换parser:

• PTX(NVIDIA)。第 12 章 §12.3.3 描述了这条路径:六个 check_* pass 在逻辑上并非 Ascend-specific——它们作用于一个由 (space, offset, rows, cols, dtype, blayout, slayout) 元组组成的 stage-2 plan。写一个 parse_ptx_stage2 就能让它们跑在 mlir_to_gpu 发射的 PTX 上、跑在 OxiCUDA 这样的 runtime-PTX 项目上、或者跑在任何其他源头上。shared memory bank 冲突、__shared__ 数组别名、对 48 KB 或 100 KB per-SM 限额的超限,都能映射到现有的 check 上。
• 通过共享 tile IR 扩到其他厂商。crates/rustc_codegen_mlir/src/ 今天有 15 个 mlir_to_* 后端(aie、bang、cpp、csl、gaudi、gpu、hexagon、linalg、msl、musa、nki、pto、spirv……)。每一个都从同一个 ascend_tile_* dialect 下沉。直接读这个 dialect 的 parser——在任何 vendor-specific lowering 之前——给 oracle 对每个目标最早、最干净的切入点。
• 向 Rust 上游贡献。upstream-tier3/ 下已经准备好了一份 Tier-3 目标规格(davinci-huawei-none):目标三元组、ABI、platform-support 文档、mod.rs/platform-support.md/bootstrap/sanity.rs 的补丁、以及社区材料(Zulip 帖子、可选 MCP 草案、PR 描述)。参与计划:(1) 在 Zulip #t-compiler/help 上就 triplet 名称征求反馈,(2) 如果新颖的 MLIR codegen 需要编译器团队共识就提交 MCP,(3) 向 rust-lang/rust 提交 draft PR。Tier-3 门槛最低——不需要 RFC,不需要 CI,单个 reviewer 批准即可——且我们的 in-tree 改动不含任何专有代码。

三条小节背后潜伏着一个长期问题:ascend_std 里的 #![no_core] 重实现,最终能否在上游 target 之上被 -Zbuild-std=core 替代?那将切掉这个项目当前最大的一笔维护税。

社区参与

ascend-rs 正在等待开源决定。一旦公开,贡献入口包括:

1. 新增 ascend_std intrinsic——遵循 extern "C" stub + mlir_to_cpp handler 模式。
2. 内核语料——写真实的内核,反馈 codegen 的缺口。
3. 宿主机 API 覆盖——CANN 的 API 比我们封装的要多。
4. 前端集成——如果你在做 TileLang、Triton、PyPTO、或 torch.compile 的 Ascend 路径,尝试把你的编译步骤换成 ascend_compile 并反馈问题。
5. Oracle 的 parser——为另一种 IR(PTX、SPIR-V、LLVM NVPTX)写一个 stage-2 parser,六个 check_* pass 就白送给你。