| 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 |
开发检查清单
开发前
实现中
测试前
相关 Skills
| Skill | 用途 |
|---|
pto-comm-isa-reference | PTO-COMM 指令签名、参数、约束速查 |
pto-comm-testing-debug | 通信算子测试与调试指南 |
pto-comm-performance-optimization | 通信算子性能优化 |
vector-fusion-operator-generate | PTO 向量融合算子开发指南 |
