CANN/PTO-ISA通信算子开发指南

【免费下载链接】pto-isaParallel Tile Operation (PTO) is a virtual instruction set architecture designed by Ascend CANN, focusing on tile-level operations. This repository offers high-performance, cross-platform tile operations across Ascend platforms.项目地址: https://gitcode.com/cann/pto-isaname: PTO-COMM 通信算子开发指南 description: 基于 PTO-COMM ISA 开发通信算子的完整指南。涵盖 Host-Device 架构、文件结构、通信模式P2P/集合通信/通算融合、同步策略、信号矩阵设计、多 Block 调度、远端地址管理、构建系统配置等。触发需要使用 PTO-COMM 开发通信算子、设计通信 kernel、编写 Host 侧代码、配置 CMakeLists 时。 license: CANN Open Software License Agreement Version 2.0PTO-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 → 参考 开发模式 之集合通信模式 │ └── 需要自定义算法如 RSAG 组合 AllReduce │ └── 使用 TPUTAtomicAdd TNOTIFY/TWAIT 组合 → 参考 开发模式 之自定义集合通信 │ ├── 需要通算融合计算通信重叠 │ └── 使用双 kernelcube vec 队列/信号同步 → 参考 开发模式 之通算融合模式 │ └── 需要异步大块传输 └── 使用 TPUT_ASYNC/TGET_ASYNC → 参考 pto-comm-isa-reference文件结构与命名规范kernels/manual/platform/operator_name/ ├── comm_kernel.cpp # 通信 kernelVec 架构 ├── compute_kernel.cpp # 计算 kernelCube 架构如需融合 ├── config.h # Tiling 配置、Block 数量、常量定义 ├── kernel_launchers.h # Host 侧 kernel 启动函数声明 ├── common.hpp # 远端地址计算等共享工具 ├── main.cpp # Host 侧初始化、启动、验证 ├── CMakeLists.txt # 构建配置 ├── run.sh # 运行脚本 └── README_zh.md # 算子文档核心开发模式四种开发模式的完整代码示例和同步策略详见详细指南开发模式详解模式指令组合适用场景P2PTPUT/TGET两 NPU 间数据传输集合通信TGATHER/TSCATTER/TBROADCAST/TREDUCE标准多 rank 操作自定义集合通信TPUTAtomicAdd TNOTIFY/TWAITRSAG 组合实现 AllReduce通算融合双 kernel 队列 信号矩阵计算与通信重叠同步策略与信号设计信号矩阵布局、DeviceBarrier 实现、流水线同步详见详细指南信号与同步设计快速参考同步需求推荐方式跨 rank barrierDeviceBarrierIntra-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 到各 blockRow 维度需要精确负载均衡展平为 row-level 分配推荐Rank 维度不同 rank 独立传输按 rank 分配给不同 blockHost 侧与构建系统Host 侧标准初始化流程、CMakeLists 模板、SOC_VERSION 映射、kernel 启动模式详见详细指南Host 侧与构建系统SOC_VERSION 与架构映射SOC_VERSION架构Cube ArchVec ArchAscend910BA2A3dav-c220-cubedav-c220-vecAscend910CA2A3dav-c220-cubedav-c220-vecAscend950A5dav-c350-cubedav-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 架构选择正确相关 SkillsSkill用途pto-comm-isa-referencePTO-COMM 指令签名、参数、约束速查pto-comm-testing-debug通信算子测试与调试指南pto-comm-performance-optimization通信算子性能优化vector-fusion-operator-generatePTO 向量融合算子开发指南【免费下载链接】pto-isaParallel Tile Operation (PTO) is a virtual instruction set architecture designed by Ascend CANN, focusing on tile-level operations. This repository offers high-performance, cross-platform tile operations across Ascend platforms.项目地址: https://gitcode.com/cann/pto-isa创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考