CANN PTO-ISA 通信算子 Host 侧开发与构建系统指南从标准初始化到双架构 Kernel 编译【免费下载链接】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本指南是 CANN PTO-ISA 仓库中 PTO-COMM 通信算子开发指南 的配套深度解读聚焦Host 侧与构建系统这一主题如何编写通信算子的 Host 侧入口MPI/ACL/HCCL 初始化、内存分配、kernel 启动、结果验证与清理如何设计 kernel 启动函数以及如何编写同时承载 Vec通信与 Cube计算双架构 kernel 的 CMakeLists.txt 并配置编译选项。读完本文你将能够独立搭建一个基于 PTO-COMM ISA 的通信算子工程并在 A2A3Ascend910B/910C与 A5Ascend950平台上完成构建与多机多卡运行。一、Host 侧与 Device 侧的职责划分PTO-COMM 通信算子的整体架构采用 Host-Device 分离模型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 侧main.cpp负责 MPI/HCCL 通信域初始化、内存分配、远端地址获取、kernel 启动和结果验证Device 侧comm_kernel.cpp/compute_kernel.cpp使用 PTO-COMM 指令执行实际的数据传输与同步如 TPUT/TGET/TNOTIFY 等或使用 PTO 计算指令如 TMATMUL/TADD 等执行计算计算 kernel 与通信 kernel 可以分别编译为独立的.so文件再由 Host 可执行文件统一链接启动这一模式在仓库的 allgather_gemm 示例A5 中体现得十分典型。二、Host 侧标准初始化流程9 步骨架通信算子 Host 侧入口遵循一个固定的九步流程这是所有多 rank 通信 demo 的通用骨架int main(int argc, char **argv) { // 1. MPI 初始化 MPI_Init(argc, argv); int rank, nranks; MPI_Comm_rank(MPI_COMM_WORLD, rank); MPI_Comm_size(MPI_COMM_WORLD, nranks); // 2. ACL 初始化 aclInit(nullptr); aclrtSetDevice(rank % device_count); aclrtStream computeStream, commStream; aclrtCreateStream(computeStream); aclrtCreateStream(commStream); // 3. HCCL 通信域创建 HcclRootInfo rootInfo; if (rank 0) HcclGetRootInfo(rootInfo); MPI_Bcast(rootInfo, sizeof(rootInfo), MPI_BYTE, 0, MPI_COMM_WORLD); HcclComm hcclComm; HcclCommInitRootInfo(nranks, rootInfo, rank, hcclComm); // 4. 获取通信上下文远端地址 // 5. 内存分配 uint8_t *buffer; aclrtMalloc((void**)buffer, size, ACL_MEM_MALLOC_HUGE_FIRST); // 6. 信号矩阵初始化清零 aclrtMemset(signal_matrix, 0, signal_size); // 7. 启动 kernel launchCommKernel(buffer, ..., commStream); aclrtSynchronizeStream(commStream); // 8. 验证结果 // 9. 清理 HcclCommDestroy(hcclComm); aclrtDestroyStream(computeStream); aclrtDestroyStream(commStream); aclrtResetDevice(rank % device_count); aclFinalize(); MPI_Finalize(); }各步骤的工程要点如下MPI 初始化确定当前进程的rank与总进程数nranks。进程数与设备数通常一一对应aclrtSetDevice(rank)多机场景下由 mpirun 负责将进程分布到各节点ACL 初始化aclInit建立 runtime 上下文随后为每个 rank 创建独立的 compute/comm 双流。仓库示例 allgather_gemm/main.cpp 中通过aclrtCreateStream分别创建 computeStream 与 commStream为通算融合计算与通信并行执行提供流级基础HCCL 通信域创建rank 0 通过HcclGetRootInfo生成 root info再经MPI_Bcast广播给所有 rank最后由各 rank 调用HcclCommInitRootInfo共同创建通信域。仓库实际实现会额外校验 MPI world size 与N_RANKS一致见 allgather_gemm/main.cpp获取通信上下文从 HCCL 域中得到包含各 rank 通信窗口基址的上下文windowsInDevice 侧据此计算远端 GM 地址内存分配使用aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)分配大块 GM 内存。注意在 HCCL 通信窗口场景下输入缓冲区需从窗口内偏移分配详见下文远端地址获取信号矩阵清零见下一节是每次 kernel 执行前必须完成的同步前置条件启动 kernel通过自定义 launch 函数见第四节在指定流上启动并以aclrtSynchronizeStream等待完成验证结果将 Device 输出拷回 Host 与 golden 数据比对数值容差比较仓库中PtoTestCommon::ResultCmp承担该职责见 allgather_gemm/main.cpp清理按通信域 → 流 → 设备 → runtime → MPI的逆序释放资源。工程提示仓库示例在aclrtSetDevice之外还会调用rtSetDevice并容忍kAclRepeatInit 100002的重复初始化错误码同时用HcclHostBarrier在 rank 之间做阶段同步这些细节在多 rank 调试中非常实用可参考 allgather_gemm/main.cpp。三、信号矩阵清零每次 kernel 执行前的强制步骤关键每次 kernel 执行前必须清零信号矩阵否则上次的残留值会导致同步错误例如 TWAIT/TTEST 误判为已就绪从而提前消费尚未到达的数据。aclrtMemset(signal_matrix, signal_size, 0, signal_size); aclrtSynchronizeStream(stream);清零动作本身是异步提交的因此必须用aclrtSynchronizeStream确保清零在 kernel 启动前真正落盘。仓库示例中PrepareStreamingState在每轮迭代开始时都会重置 chunk 标志矩阵ChunkFlagMatrixResetChunkFlagMatrixSummaryInit并通过aclrtMemcpy同步到 Device随后才启动通信/计算 kernel其语义与信号矩阵清零完全一致见 allgather_gemm/main.cpp。四、Kernel 启动函数模式launchers 声明 实现分离Host 侧不直接调用 Device kernel而是通过一层启动函数launcher封装典型做法是声明与实现分离// kernel_launchers.h 声明 void launchCommKernel(uint8_t *data, uint8_t *signal, uint8_t *ctx, int rank, int nranks, void *stream); // comm_kernel.cpp 实现 void launchCommKernel(uint8_t *data, uint8_t *signal, uint8_t *ctx, int rank, int nranks, void *stream) { CommKernelEntryCOMM_BLOCK_NUM, nullptr, stream( data, signal, ctx, rank, nranks, COMM_BLOCK_NUM); }要点启动语法KernelNameblockNum, nullptr, stream(args...)其中blockNum为 AICore 核数nullptr表示无 L2 缓存配置stream决定 kernel 在哪个流上排队执行void* streamlauncher 签名统一使用void*承载aclrtStream避免 Host 头文件与 runtime 头文件的强耦合仓库实证A5 的 allgather_gemm 示例中launcher 声明位于 kernel_launch.hpp实现位于 allgather_gemm_comm_kernel.cpp。实际实现还会按每个远端 rank 分到一份 block 预算进行二次拆分blocksPerDest COMM_BLOCK_NUM / (nRanks - 1)再将总 block 数传给 kernel实现多 block 并行的按目的 rank 分片传输双 kernel 协同通算融合场景下launchCommKernel与launchComputeKernel分别在不同流上启动通信与计算通过 Device 侧的信号/就绪队列完成流水线衔接Host 侧只负责分别提交与等待。五、CMakeLists.txt 模板与关键编译选项5.1 模板总览通信算子工程的 CMakeLists.txt 需要同时处理三类目标通信 kernelVec 架构 .so、计算 kernelCube 架构 .so与Host 可执行文件。模板如下cmake_minimum_required(VERSION 3.16) project(my_comm_operator) set(CMAKE_CXX_COMPILER bisheng) set(CMAKE_CXX_STANDARD 17) # PTO 头文件路径优先使用仓库内版本 set(PTO_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../../../../include) include_directories(BEFORE ${PTO_INCLUDE_DIR}) # CANN 环境 if(DEFINED ENV{ASCEND_HOME_PATH}) set(ASCEND_HOME $ENV{ASCEND_HOME_PATH}) else() set(ASCEND_HOME /usr/local/Ascend/ascend-toolkit/latest) endif() include_directories(${ASCEND_HOME}/include) link_directories(${ASCEND_HOME}/lib64) # 通信 KernelVec 架构 add_library(comm_kernel SHARED comm_kernel.cpp) target_compile_options(comm_kernel PRIVATE --cce-aicore-archdav-c220-vec -DMEMORY_BASE -D_GLIBCXX_USE_CXX11_ABI0) target_link_options(comm_kernel PRIVATE --cce-fatobj-link) target_link_libraries(comm_kernel runtime) # 计算 KernelCube 架构如需通算融合 add_library(compute_kernel SHARED compute_kernel.cpp) target_compile_options(compute_kernel PRIVATE --cce-aicore-archdav-c220-cube -DMEMORY_BASE -D_GLIBCXX_USE_CXX11_ABI0) target_link_options(compute_kernel PRIVATE --cce-fatobj-link) target_link_libraries(compute_kernel runtime) # Host 可执行文件 add_executable(my_operator main.cpp) target_link_libraries(my_operator comm_kernel compute_kernel ascendcl hccl tiling_api platform)5.2 关键配置项速查表配置说明--cce-aicore-archdav-c220-vec通信 kernel 使用 Vec 架构--cce-aicore-archdav-c220-cube计算 kernel 使用 Cube 架构-DMEMORY_BASE启用远端地址计算宏--cce-fatobj-link启用 fat object 链接-D_GLIBCXX_USE_CXX11_ABI0关闭新版 libstdc ABI与 CANN 预编译库二进制兼容--cce-pto-enable开启 PTOParallel Tile Operation编译通道A5 示例中置于公共 CCE 编译选项内5.3 SOC_VERSION 与架构映射--cce-aicore-arch的具体取值取决于目标平台。仓库 Skill 文档给出了权威映射表SOC_VERSION架构Cube ArchVec ArchAscend910BA2A3dav-c220-cubedav-c220-vecAscend910CA2A3dav-c220-cubedav-c220-vecAscend950A5dav-c350-cubedav-c350-vec从仓库源码还可以进一步确认两点A2A3 平台的 dav-c220 双架构用法A2A3 目录下的 allgather_gemm 与 conv2d_forward 工程分别使用dav-c220-cube计算与dav-c220-vec通信/向量见 kernels/manual/a2a3/allgather_gemm/CMakeLists.txtA5 平台实际为 dav-c310 双架构A5 示例工程中计算 kernel 使用dav-c310-cube、通信 kernel 使用dav-c310-vec见 kernels/manual/a5/allgather_gemm/CMakeLists.txt比 Skill 表格中的 dav-c350 命名更细粒度。从源码结构可以推断dav-c310 与 dav-c350 是 A5 代际内不同型号的架构代号实际使用时应以目标机型对应的 CANN 工具链支持列表为准。5.4 仓库工程中的完整配置进阶参考模板是精简骨架实际仓库工程如 kernels/manual/a5/allgather_gemm/CMakeLists.txt还包含几类重要增强可直接借鉴CCE 公共编译选项-xcce -Xhost-start -Xhost-end及一系列-mllvm -cce-aicore-*选项栈大小、溢出记录、地址变换等并统一追加--cce-pto-enable开启 PTO 通道尺寸与 Block 数可配置化通过-DG_M/-DG_K/-DG_N/-DG_BASE_M/-DG_BASE_N以及-DCOMPUTE_BLOCK_NUM/-DCOMM_BLOCK_NUM在编译期注入配置gemm_config.hpp 中再用static_assert强制整除约束如G_M % G_BASE_M 0见 gemm_config.hpp链接库差异A5 示例的通信 kernel 额外链接hcommHCCL 通信底层与nnopbaseHost 可执行文件则根据RUN_MODE在runtimeNPU 实机与runtime_camodel模拟器之间选择并按平台自动探测ASCEND_HOME_PATH下的x86_64-linux/aarch64-linux内核头文件路径。六、MPI 多机多卡运行构建完成后通过 mpirun 指定进程数与 NPU 卡数一致运行# 单机多卡 mpirun -np 8 ./my_operator # 多机 mpirun -np 16 -H host1:8,host2:8 ./my_operator仓库示例的 run.sh 提供了可直接复用的运行流水线自动探测并 source CANN 的set_env.sh可用ASCEND_CANN_PATH覆盖、在常见路径中搜索 mpich 的mpirun并注入 PATH/LD_LIBRARY_PATH、支持--run-modenpu/sim、--soc-version、--n-ranks、--gm/--gk/--gn、--base-m/--base-n、--compute-blocks/--comm-blocks等参数最终执行python3 scripts/gen_data.py --n-ranks ${N_RANKS} --m ${G_M} --k ${G_K} --n ${G_N} --output-dir ./out cmake -DRUN_MODE${RUN_MODE} -DSOC_VERSION${SOC_VERSION} -DG_M${G_M} ... .. make -j16 N_RANKS${N_RANKS} ALLGATHER_GEMM_DATA_DIR... mpirun -n ${N_RANKS} ./allgather_gemm运行层面的注意事项均来自仓库脚本与源码的实际行为进程数约束mpirun -n的进程数必须等于N_RANKS否则程序直接报错退出同时N_RANKS不得超过MAX_RING_RANKS环形拓扑上限见 allgather_gemm/main.cpp环境变量N_RANKS与ALLGATHER_GEMM_DATA_DIR分别控制 rank 数与数据目录Host 侧通过std::getenv读取main.cpp残留状态清理多轮运行前建议清理/dev/shm/sem.hccl*残留信号量与 IPC 资源避免上次异常退出影响本次 HCCL 初始化模拟器支持RUN_MODEsim时链接runtime_camodel并依赖目标 SOC 对应的模拟器库${ASCEND_HOME_PATH}/tools/simulator/${SOC_VERSION}/lib适合无实机的开发环境先行验证。七、与远端地址管理的衔接Host 侧构建好的通信上下文最终要服务于 Device 侧的数据传输。Host 侧在第 4 步取得的通信窗口windowsIn[rank]为 Device 侧提供了两种远端地址计算手段详见 多 Block 调度与地址管理方式 1HCCL 通信窗口在 Device 侧用本地指针相对ctx-windowsIn[myRank]的偏移加上ctx-windowsIn[remote_rank]基址得到远端地址方式 2ParallelGroup将各 rank 远端地址打包为comm::ParallelGroupGlobalData供 TGATHER/TSCATTER 等集合指令使用。因此 Host 侧内存分配有两个硬性对齐要求构建与运行前务必核对所有 GM 地址必须满足32 字节对齐Signal 地址必须4 字节对齐TPUT_ASYNC/TGET_ASYNC 的 workspace 由专用 Manager 管理无需额外对齐。八、开发检查清单Host 侧与构建相关在动手写 Host 侧代码与 CMakeLists 之前对照 Skill 文档的检查清单逐项确认确认目标平台A2A3/A5和对应的架构编译选项dav-c220-cube/vec 或 dav-c310/dav-c350确认通信拓扑节点内/跨节点和链路类型据此规划进程数与 mpirun 参数确定通信模式P2P/集合/融合决定是单个 comm kernel 还是 comm compute 双 kernel规划信号矩阵布局并保证每次 kernel 执行前清零远端地址计算正确基于通信窗口偏移内存大小与 Tile 配置一致GM 地址 32 字节对齐、Signal 4 字节对齐Host 侧aclrtSynchronizeStream确保 kernel 执行完成后再验证结果CMakeLists 中 Vec/Cube 架构选择正确-DMEMORY_BASE、--cce-fatobj-link、-D_GLIBCXX_USE_CXX11_ABI0均已配置mpirun 进程数与 N_RANKS 一致且不超过 MAX_RING_RANKS九、相关资源导航Skill 主文档PTO-COMM 通信算子开发指南含 SOC_VERSION 映射表与完整检查清单配套参考多 Block 调度与地址管理、信号与同步设计、开发模式详解A5 完整工程示例kernels/manual/a5/allgather_gemm/main.cpp、kernel_launch.hpp、双 kernel 源文件、CMakeLists.txt、run.sh 一应俱全A2A3 双架构工程示例kernels/manual/a2a3/allgather_gemm/CMakeLists.txt通信指令速查pto-comm-isa-reference Skill调试与测试pto-comm-testing-debug Skill【免费下载链接】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），仅供参考