name: PTO-COMM 通信算子开发指南 description: 基于 PTO-COMM ISA 开发通信算子的完整指南。涵盖 Host-Device 架构、文件结构、通信模式(P2P/集合通信/通算融合)、同步策略、信号矩阵设计、多 Block 调度、远端地址管理、构建系统配置等。触发:需要使用 PTO-COMM 开发通信算子、设计通信 kernel、编写 Host 侧代码、配置 CMakeLists 时。 license: CANN Open Software License Agreement Version 2.0
PTO-COMM 通信算子开发指南
定位
本 Skill 是流程型 Skill,指导从零开发一个基于 PTO-COMM ISA 的通信算子。
架构概述
Host-Device 分离
Host 侧 Device 侧
┌─────────────────┐ ┌─────────────────────────┐
│ main.cpp │ │ comm_kernel.cpp │
│ - MPI 初始化 │ 启动 │ - __global__ AICORE │
│ - ACL 初始化 │──kernel──→ │ - TPUT/TGET/TNOTIFY/... │
│ - HCCL 通信域 │ │ - 信号同步逻辑 │
│ - 内存分配 │ ├─────────────────────────┤
│ - Kernel 启动 │ │ compute_kernel.cpp │
│ - 结果验证 │ 启动 │ - __global__ AICORE │
│ │──kernel──→ │ - TMATMUL/TADD/... │
└─────────────────┘ │ - 计算逻辑 │
└─────────────────────────┘
关键原则:
- Host 侧:负责 MPI/HCCL 通信域初始化、内存分配、远端地址获取、kernel 启动和结果验证
- Device 侧:使用 PTO-COMM 指令执行实际的数据传输和同步
- 计算和通信可以分别编译为独立的
.so文件
编程模型选择
需要 NPU 间通信?
├── 仅需基本 P2P 传输
│ └── 使用 TPUT/TGET → 参考 "开发模式" 之 P2P 模式
│
├── 需要集合通信(AllReduce/AllGather/ReduceScatter 等)
│ ├── 可用内置集合指令完成?
│ │ └── 使用 TGATHER/TSCATTER/TBROADCAST/TREDUCE → 参考 "开发模式" 之集合通信模式
│ └── 需要自定义算法(如 RS+AG 组合 AllReduce)?
│ └── 使用 TPUT<AtomicAdd> + TNOTIFY/TWAIT 组合 → 参考 "开发模式" 之自定义集合通信
│
├── 需要通算融合(计算+通信重叠)
│ └── 使用双 kernel(cube + vec)+ 队列/信号同步 → 参考 "开发模式" 之通算融合模式
│
└── 需要异步大块传输
└── 使用 TPUT_ASYNC/TGET_ASYNC → 参考 pto-comm-isa-reference
文件结构与命名规范
kernels/manual/<platform>/<operator_name>/
├── comm_kernel.cpp # 通信 kernel(Vec 架构)
├── compute_kernel.cpp # 计算 kernel(Cube 架构,如需融合)
├── config.h # Tiling 配置、Block 数量、常量定义
├── kernel_launchers.h # Host 侧 kernel 启动函数声明
├── common.hpp # 远端地址计算等共享工具
├── main.cpp # Host 侧:初始化、启动、验证
├── CMakeLists.txt # 构建配置
├── run.sh # 运行脚本
└── README_zh.md # 算子文档
核心开发模式
四种开发模式的完整代码示例和同步策略详见:
详细指南:开发模式详解
| 模式 | 指令组合 | 适用场景 |
|---|---|---|
| P2P | TPUT/TGET | 两 NPU 间数据传输 |
| 集合通信 | TGATHER/TSCATTER/TBROADCAST/TREDUCE | 标准多 rank 操作 |
| 自定义集合通信 | TPUT<AtomicAdd> + TNOTIFY/TWAIT | RS+AG 组合实现 AllReduce |
| 通算融合 | 双 kernel + 队列 + 信号矩阵 | 计算与通信重叠 |
同步策略与信号设计
信号矩阵布局、DeviceBarrier 实现、流水线同步详见:
详细指南:信号与同步设计
快速参考
| 同步需求 | 推荐方式 |
|---|---|
| 跨 rank barrier | DeviceBarrier(Intra-rank + Cross-rank + 本地广播) |
| 阶段间分隔 | pipe_barrier(PIPE_ALL) |
| 计算→通信通知 | SPSC 就绪队列 + TTEST 轮询 |
| 手动流水线 | set_flag/wait_flag(仅 TLOAD/TSTORE_IMPL 时需要) |
| 多方通知一方 | NotifyOp::AtomicAdd |
| 一方通知多方 | NotifyOp::Set |
远端地址管理与多 Block 调度
远端地址获取方式、地址对齐要求、Block 分配策略、工作均分方法详见:
详细指南:多 Block 调度与地址管理
地址对齐要求
- 所有 GM 地址必须满足 32 字节对齐
- Signal 地址必须 4 字节对齐
- TPUT_ASYNC/TGET_ASYNC 的 workspace 由专用 Manager 管理
多核切分策略
| 切分维度 | 适用场景 | 方法 |
|---|---|---|
| Tile 维度 | 通信量大,Tile 数多 | 均分 Tile 到各 block |
| Row 维度 | 需要精确负载均衡 | 展平为 row-level 分配(推荐) |
| Rank 维度 | 不同 rank 独立传输 | 按 rank 分配给不同 block |
Host 侧与构建系统
Host 侧标准初始化流程、CMakeLists 模板、SOC_VERSION 映射、kernel 启动模式详见:
详细指南:Host 侧与构建系统
SOC_VERSION 与架构映射
| SOC_VERSION | 架构 | Cube Arch | Vec Arch |
|---|---|---|---|
| Ascend910B | A2A3 | dav-c220-cube | dav-c220-vec |
| Ascend910C | A2A3 | dav-c220-cube | dav-c220-vec |
| Ascend950 | A5 | dav-c350-cube | dav-c350-vec |
开发检查清单
开发前
- 确认目标平台(A2A3/A5)和对应的架构编译选项
- 确认通信拓扑(节点内/跨节点)和链路类型
- 确定通信模式(P2P/集合/融合)
- 规划信号矩阵布局
实现中
- TNOTIFY 目标地址为远端,TWAIT/TTEST 监听地址为本地
- 乒乓 Tile 的 UB 偏移不重叠
- 使用
pipe_barrier(PIPE_ALL)分隔不同阶段 - 手动 TLOAD/TSTORE_IMPL 之间有正确的 set_flag/wait_flag
- 所有 rank 使用相同的 rootIdx 构建 ParallelGroup
- 非 root rank 不调用集合通信指令
- 远端地址计算正确(基于通信窗口偏移)
测试前
- 信号矩阵每次运行前清零
- Host 侧
aclrtSynchronizeStream确保 kernel 执行完成 - 内存大小与 Tile 配置一致
- CMakeLists 中 Vec/Cube 架构选择正确
相关 Skills
| Skill | 用途 |
|---|---|
pto-comm-isa-reference | PTO-COMM 指令签名、参数、约束速查 |
pto-comm-testing-debug | 通信算子测试与调试指南 |
pto-comm-performance-optimization | 通信算子性能优化 |
vector-fusion-operator-generate | PTO 向量融合算子开发指南 |