English | 中文版

附录 G：Tile API 与 Buffer API 的对比——与 FlashTile/PTO 的横向比较

核心发现：对于 attention 类内核，基于 tile 的内核 API（ascend_std::tile）的可用性远优于基于 buffer 的 API（ascend_std 缓冲区操作）——将一个 50 行的 softmax 缩减为 5 行，同时彻底消除显式 pipe barrier 管理。ptoas 汇编器已在 910c 服务器上确认可用；mlir_to_pto.rs 已重写为生成真实 PTO-MLIR 方言格式，向量操作（add/mul/exp）和 softmax 规约操作的端到端路径均已通过 ptoas 验证。translate_matmul() 现已正确生成 cube unit tile 序列（loc=mat CBUF 暂存 → loc=left L0A / loc=right L0B → pto.tmatmul → loc=acc L0C），通过 ptoas 验证可正确输出带 __ca__/__cb__/__cc__ 限定符的 AscendC C++。当前阻塞点：ptoas 生成的 C++ 使用 pto/pto-inst.hpp，与 CANN 8.5.0 的 bisheng（Clang 15）不兼容，待升级至 CANN 9.x 解决。

G.1 背景：ascend-rs 的两条代码生成路径

ascend-rs 目前为 Rust NPU 内核提供两条代码生成路径：

Tile API 路径实现了 PTO/FlashTile 集成方案的第三阶段。PTO（可编程 Tile 操作，Programmable Tile Operations）是面向昇腾 NPU 的虚拟指令集，ptoas 是其汇编器。FlashTile 指的是 PTO ISA 所暴露的 tile 级编程模型——tile 加载、存储和 tile.softmax 等融合操作——有别于 AscendC 更底层的 buffer/DMA 模型。

G.2 可用性差距：以 Softmax 为例

相同的逐行 softmax 计算在两种 API 下所需代码量截然不同：

Buffer API（mha/kernels/src/lib.rs，约 50 行内核代码）：

#![allow(unused)]
fn main() {
#[ascend_std::aiv_kernel]
pub unsafe fn softmax_rows_f16(
    input: *const u16, output: *mut u16,
    row_len: *const u32, num_rows: *const u32,
) {
    let cols = *row_len;
    let rows = *num_rows;
    let buf_in    = ascend_std::ascend_buf_alloc(cols);
    let buf_out   = ascend_std::ascend_buf_alloc(cols);
    let buf_work  = ascend_std::ascend_buf_alloc(cols);
    let buf_rwork = ascend_std::ascend_buf_alloc(cols);

    let mut row = 0u32;
    loop {
        if row >= rows { break; }
        let in_ptr  = input.wrapping_add((row * cols) as usize);
        let out_ptr = output.wrapping_add((row * cols) as usize);

        ascend_std::ascend_buf_load_f16(buf_in, in_ptr, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 1

        let max_val = ascend_std::ascend_reduce_max_f16(buf_rwork, buf_in, buf_work, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 2

        ascend_std::ascend_adds_f16(buf_out, buf_in, -max_val, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 3

        ascend_std::ascend_exp_f16(buf_out, buf_out, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 4

        let sum_val = ascend_std::ascend_reduce_sum_f16(buf_rwork, buf_out, buf_work, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 5

        ascend_std::ascend_muls_f16(buf_out, buf_out, 1.0f32 / sum_val, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 6

        ascend_std::ascend_buf_store_f16(out_ptr, buf_out, cols);
        ascend_std::ascend_pipe_barrier();                          // 屏障 7

        row = row + 1;
    }
}
}

Tile API（tile_softmax/kernels/src/lib.rs，5 行内核逻辑）：

#![allow(unused)]
fn main() {
#[ascend_std::aiv_kernel]
pub unsafe fn tile_softmax<const ROWS: usize, const COLS: usize>(
    input: *const f32,
    output: *mut f32,
) {
    let block_idx = ascend_std::get_block_idx() as usize;
    let offset = block_idx * ROWS * COLS;
    let t_in  = tile_load_f32::<ROWS, COLS>(input.add(offset));
    let t_out = tile_softmax_f32::<ROWS, COLS>(t_in);
    tile_store_f32::<ROWS, COLS>(output.add(offset), t_out);
}
}

差距一目了然：Buffer API 需要 7 个显式 pipe_barrier() 调用、4 个命名缓冲区分配和一个手动行循环；而 Tile API 零屏障、零显式缓冲区、无循环。mlir_to_pto 代码生成路径自动抑制 pipe_barrier 调用，因为 PTO 隐式管理流水线同步。

G.3 PTO 格式与端到端验证

ptoas 所消费的实际 .pto 格式是带有 pto 方言的 MLIR。mlir_to_pto.rs 已完成重写，现可生成正确的 PTO-MLIR 方言格式。

端到端已验证的格式（以 vec_add 为例）：

module {
  func.func @vec_add(%arg0: !pto.ptr<f32>, %arg1: !pto.ptr<f32>, %arg2: !pto.ptr<f32>) {
    %c0 = arith.constant 0 : index
    %c1 = arith.constant 1 : index
    %c32 = arith.constant 32 : index
    %0 = pto.make_tensor_view %arg0, shape = [%c32, %c32] strides = [%c32, %c1] : !pto.tensor_view<32x32xf32>
    %2 = pto.partition_view %0, offsets = [%c0, %c0], sizes = [%c32, %c32] : !pto.tensor_view<32x32xf32> -> !pto.partition_tensor_view<32x32xf32>
    %5 = pto.alloc_tile : !pto.tile_buf<loc=vec, dtype=f32, rows=32, cols=32, v_row=32, v_col=32, blayout=row_major, slayout=none_box, fractal=512, pad=0>
    pto.tload ins(%2 : !pto.partition_tensor_view<32x32xf32>) outs(%5 : !pto.tile_buf<...>)
    pto.tadd  ins(%5, %6 : ...) outs(%7 : ...)
    pto.tstore ins(%7 : ...) outs(%8 : ...)
    return
  }
}

ptoas 通过若干 MLIR 变换趟（PTO Infer Mem Scope → PTO plan Mem → PTOToEmitC）降级，最终生成带自动同步屏障的 AscendC C++：

__global__ AICORE void vec_add(__gm__ float* v1, __gm__ float* v2, __gm__ float* v3) {
  // TLOAD → TADD → TSTORE，带 set_flag/wait_flag 自动插入
}

完整流水线为：

Rust 内核 → MLIR → PTO-MLIR (.pto) → ptoas --enable-insert-sync → AscendC C++ → bisheng → NPU 二进制

ptoas 流水线状态：ptoas（LLVM 19.1.7）能够正确解析并编译完整的 softmax 分解流程——pto.trowmax、pto.trowexpandsub、pto.texp、pto.trowsum、pto.trowexpanddiv 全部通过验证，输出带 TROWMAX/TROWEXPANDSUB/TEXP/TROWSUM/TROWEXPANDDIV 调用的 AscendC C++。

当前阻塞点：pto-inst.hpp 与 CANN 8.5.0 bisheng 不兼容。 ptoas 生成的 C++ 包含 #include "pto/pto-inst.hpp"，但 CANN 8.5.0 附带的 bisheng（基于 Clang 15）无法编译该头文件，错误包括：缺少 MrgSortExecutedNumList 类型、copy_gm_to_ubuf_align_b32 内置函数不支持目标特性、bfloat16_t 类型未定义。pto-inst.hpp 是为更新版 bisheng 设计的，需要 CANN 9.x 或更高版本方能完整支持。

G.4 与 PTO Tile Library（pto-isa）的横向比较

PTO Tile Library（pto-isa，2025-12-27 在 GitCode 开源，网址 https://pto-isa.gitcode.com）是 Huawei 发布的 tile 级 C++ 头文件库，提供约 90 条 tile 操作，包含 FlashAttention 的参考实现（kernels/manual/a2a3/flash_atten/）。

pto-isa 提供的 softmax（FlashAttention 流式 softmax 的核心）使用如下 C++ 模板：

#include <pto/pto-inst.hpp>
// 流式 FA softmax：初始 tile
TROWMAX(new_global_max, input_x, tmp_float);   // 规约求行最大值
pipe_barrier(PIPE_V);
TROWEXPANDSUB(p_tile_f32, input_x, new_global_max); // x - max（广播）
TMULS(p_tile_f32, p_tile_f32, scale);
TEXP(p_tile_f32, p_tile_f32);
pipe_barrier(PIPE_V);
TROWSUM(new_global_sum, p_tile_f32, tmp_float); // 规约求和

对应 PTO-MLIR 方言中的 pto.trowmax、pto.trowexpandsub、pto.trowsum 算子。

Rust 方案的核心结构优势在于编译期形状系统与内存安全：Tile<16, 1024, f32> 与 Tile<1, 1024, f32> 是不同的类型，形状不匹配在 rustc 编译期即可检出。pto-isa 的 C++ 模板在形状维度上同样有编译期检查，但设备端内存安全无 Rust 级保障。

互补关系：pto-isa 是 ascend-rs tile API 的理想验证参照——ascend_std::tile 中的 tile_softmax_f32 最终应生成与 pto-isa 的 TROWMAX/TROWEXPANDSUB/TROWSUM/TEXP/TROWEXPANDDIV 链等价的 PTO-MLIR，经 ptoas 编译为相同的 AscendC。

G.5 量化对比

V-pipe 工作负载（softmax）— 易用性

V-pipe 工作负载（softmax）— 昇腾 910B2 实测性能

M-pipe 工作负载（矩阵乘法/GEMM）

G.6 当前状态与后续步骤

mlir_to_pto.rs 已完成重写，现在生成正确的 PTO-MLIR 方言格式，向量操作和规约操作（tadd、tmul、texp、trowmax、trowexpandsub、trowsum、trowexpanddiv、tmatmul）的端到端路径已通过 ptoas 验证并输出 AscendC C++。

已完成：

• mlir_to_pto.rs：从虚构文本汇编重写为真实 PTO-MLIR 方言（现已超 950 行）
• 10 项单元测试全部通过
• softmax 分解（trowmax → trowexpandsub → texp → trowsum → trowexpanddiv）经 ptoas 完整验证，输出正确的 TROWMAX/TROWEXPANDSUB/TEXP/TROWSUM/TROWEXPANDDIV AscendC C++
• ptoas 已接入 KernelBuilder：ACLRS_CODEGEN_PATH=pto 触发完整 MLIR → PTO-MLIR → ptoas → .pto.cpp → bisheng 流水线
• translate_matmul() 现已正确生成 cube unit tile 类型。修正后的 mlir_to_pto.rs 现在生成完整的 cube unit 流水线：pto.alloc_tile（含 loc=mat CBUF 暂存、loc=left L0A、loc=right L0B、loc=acc L0C，fractal=1024），随后是 pto.tload GM→mat、pto.tmov mat→left/right（MTE1 流水线），以及 pto.tmatmul left×right→acc

待完成：

1. pto-inst.hpp 与 bisheng 的版本兼容性。 ptoas 生成的 C++ 使用了 pto/pto-inst.hpp，而 CANN 8.5.0 附带的 bisheng（Clang 15）无法编译该头文件。解决路径：升级至 CANN 9.x，或联系 pto-isa 维护者提供 Clang 15 兼容的兼容层。

2. 在 910B2 上对比 buffer API 与 tile API softmax 性能。 完整端到端路径为：

Rust 内核 → MLIR → PTO-MLIR (.pto) → ptoas --enable-insert-sync → AscendC C++ → bisheng → NPU 二进制 → 硬件

对比 910B2 上的内核执行时间，实证回答 PTO 生成的 AscendC 是否能避免 buffer 路径中 pipe_barrier(PIPE_ALL) 带来的流水线停顿，以及与 pto-isa FlashAttention 参考内核的性能差距。