CANN pto-isa 异步 RDMA 通信测试tput/tget/tput_notify实战指南【免费下载链接】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导读本文面向在 Ascend A5HNS1825 目标网卡平台上验证 pto-isa 异步 RDMA 通信能力的开发者围绕tput_async_rdma、tget_async_rdma、tput_async_notify_rdma三个 STSystem Test目标展开。阅读本文后你将掌握如何配置PTO_RDMA_BACKENDHNS_1825并驱动run_st.py完成构建与多 rank 运行、RDMA 端点物理设备 id 与 IPv4的自动发现机制与兜底变量、TPUT_ASYNC/TGET_ASYNC/TPUT_ASYNC_NOTIFY三类异步指令在 Kernel 中的真实调用方式以及常见问题的定位手段。背景三个目标复用同一套 RDMA 测试 Kernel本目录tests/npu/a5/comm/st/testcase/tput_async_rdma存放 RDMA 异步 ST 的共享实现由三个测试 target 复用同一套 RDMA Kernel 实现测试目标验证内容对应异步指令tput_async_rdma远程 WRITETPUT_ASYNCDmaEngine::RDMAtget_async_rdma远程 READTGET_ASYNCDmaEngine::RDMAtput_async_notify_rdma带Set通知的远程 WRITETPUT_ASYNC_NOTIFYDmaEngine::RDMA三个 target 的目录结构保持一致例如 tget_async_rdma 与 tput_async_notify_rdma 均只包含CMakeLists.txt、main.cpp、gen_data.py与 README真正的共享实现位于tput_async_rdma/下的main.cpp、tput_async_rdma_kernel.cpp/.h以及backends/目录。main.cpp通过#include tput_async_rdma_kernel.h调用同一组 host 侧入口函数RunPutAsyncRdmaRootPut、RunPutAsyncRdmaRootPutPlan、RunGetAsyncRdmaRootGetPlan、RunPutAsyncNotifyRdmaSet从源码结构看三个 target 的区别仅在于编译期宏开关PTO_RDMA_GET_TEST控制 GET 入口的实例化与 gtest 过滤条件。说明当前 RDMA 后端不支持AtomicAdd因此没有对应的 RDMA 用例详见后文 Kernel 分析。前置条件运行 RDMA 异步 ST 需要满足以下环境要求硬件配备 HNS1825 网卡及匹配驱动、HCOMM 组件的 A5 环境软件按顶层 测试说明 配置 MPI 和 HCCL网络各参与 rank 使用的 RDMA 网卡 IPv4 互相可达。构建与运行RDMA 后端在CMake 配置阶段选定配置前设置PTO_RDMA_BACKEND再运行所需的 RDMA ST targetexport PTO_RDMA_BACKENDHNS_1825 python3 tests/script/run_st.py -r npu -v a5 -t comm/tput_async_rdma -d -n 2 python3 tests/script/run_st.py -r npu -v a5 -t comm/tget_async_rdma -d -n 2 python3 tests/script/run_st.py -r npu -v a5 -t comm/tput_async_notify_rdma -d -n 2各参数含义-r npu指定运行平台为 NPU-v a5指定 A5 版本-t comm/xxx指定测试 target-d进入调试/详细模式-n 2指定 2 个 rankMPI 进程数。首次验证 TPUT_ASYNC_NOTIFY首次验证TPUT_ASYNC_NOTIFY时建议先只运行 2 个 rank 的定向用例缩小排查范围export PTO_RDMA_BACKENDHNS_1825 python3 tests/script/run_st.py -r npu -v a5 -t comm/tput_async_notify_rdma \ -g TPutAsyncNotifyRdma.Int32SetAndCanaries -d -n 2-g用于按 gtest filter 只跑指定用例。该用例的校验逻辑对应 tput_async_rdma_kernel.cpp 中TPutAsyncNotifyRdmaKernelImpl检查远端 payload接收端非 root rank轮询本地 signalTTEST(signal, kRdmaNotifySignalValue, WaitCmp::EQ)轮询上限kRdmaNotifyPollLimit 10000000观察到 signal 后先通过dcci维护数据缓存再逐元素校验recvBuf内容期望值为index rootRank * 10000检查远端 signal及其两侧 canary发送端以NotifyOp::Set将kRdmaNotifySignalValue 37写入远端 signal测试同时校验 canary 前后值0x13572468与0x24681357未被破坏等待接口返回的AsyncEvent完成通过CompleteRdmaEvent走PUBLIC_EVENT_WAIT_TEST完成模式先event.Test再event.Wait最后再Test。关于 PTO_RDMA_BACKEND 的构建期语义PTO_RDMA_BACKEND仅在 CMake 配置该 ST 构建时读取。未设置、空值或不支持的值会构建出不含 RDMA 支持的测试二进制此时运行用例会走SkipUnsupportedRdmaKernel路径直接 SKIP日志打印built without PTO_RDMA_SUPPORTED。注意run_st.py默认会重新构建修改PTO_RDMA_BACKEND后不要使用-w/--without-build复用已有二进制否则会出现构建与运行不一致的假象。构建期宏与 CMake 接线从源码可以确认 RDMA 支持由编译期宏控制backends/rdma_test_backend.hpp依据PTO_RDMA_BACKEND_HNS_1825_SUPPORTED宏引入 HNS1825 后端头文件与pto/comm/async/rdma/backends/hns_1825/hns_1825_types.hpp并在未配置任何后端时触发#error The RDMA ST has no adapter for the configured backend.tput_async_rdma_kernel.cpp以#ifdef PTO_RDMA_SUPPORTED包裹数据面实现未定义时 Kernel 只做pipe_barrier同步并置零状态GET 入口额外受PTO_RDMA_GET_TEST控制target 的 CMake 配置极简见 tput_async_rdma/CMakeLists.txt仅一行pto_comm_st(tput_async_rdma)后端使能逻辑由公共测试框架统一处理。端点发现Bootstrap发现流程Bootstrap 先解析每个 rank 的物理设备 idphyId和RDMA IPv4再通过MPIAllgather交换端点peer IP、peer phyId与注册内存信息对称通信缓冲区的基地址symAddr。该流程实现在 hns_1825_bootstrap.hpp 的BootstrapConfig::Init中本地 IPv4 的查找顺序为固定/etc/hccl_rootinfo.json解析与物理设备匹配的 rank 块取net_type:CLOS的 IPv4 地址ResolveLocalRdmaIp。该解析器不依赖额外库按device_id出现的次数切分 rank 块后扫描addr字段HCOMM topology 组件解析固定/var/run/ascend-topologyd/virtualTopology.xml通过dlsym(RTLD_DEFAULT, GetRoceIpFromXml)动态解析 HCOMM 符号缺失时尝试dlopen(libtopoaddrinfo.so)得到该 phyId 对应的 RoCE IPv4ResolveLocalRdmaIpFromVirtualTopology使用测试专用 IP 变量兜底PTO_ROCE_LOCAL_IP/PTO_ROCE_IPS见下表。ST 不会生成或修改这两个拓扑文件也不提供路径覆盖变量。路径以源码常量形式固定kDefaultRootInfoPath /etc/hccl_rootinfo.jsonkDefaultVirtualTopologyPath /var/run/ascend-topologyd/virtualTopology.xml。phyId 的解析优先调用aclrtGetPhyDevIdByUserDevId其次为已废弃的ByLogicDevId与rtGetDevicePhyIdByIndex均通过 dlsym 获取全部失败时回退到PTO_ROCE_PHYIDS或 ACL device id。环境变量一览变量说明PTO_RDMA_BACKEND配置阶段选择项当前唯一支持值为HNS_1825。PTO_ROCE_PHYIDS可选按 MPI rank 索引、逗号分隔的物理设备 id。PTO_ROCE_LOCAL_IP当前 MPI 进程使用的最终兜底 IPv4必要时需为各 rank 分别设置。PTO_ROCE_IPS最终兜底列表按 MPI rank 排序且 IPv4 数量必须等于 rank 数。PTO_ROCE_BASE_PORT各 rank 一致的 channel base port默认60032。PTO_ROCE_VERBOSE设为1打印端点、MR、channel 和释放进度。HCCL_RDMA_TCHCOMM traffic class默认132。HCCL_RDMA_SLHCOMM service level默认4。优先级与一致性约束PTO_ROCE_LOCAL_IP的优先级高于PTO_ROCE_IPSroot-info 或 virtual topology 解析成功时两者均被忽略对应ResolveBootstrapLocalIp的短路逻辑所有 rank 必须使用相同的 base portResolveAndAgreeBasePort会通过MPI_Allgather汇总各 rank 的端口并校验一致性不一致直接报错使用PTO_ROCE_IPS时所有 rank 还必须使用相同的按 rank 排序列表任一 rank 无法解析本地 IP 时会通过AgreeOnLocalIp做集合级 skipcollective skip而不是产生假通过false pass——这是BootstrapConfig中skipped标志的用途host 侧据此返回RdmaTestResult::SKIPPED。Kernel 数据面实现解析对称通信缓冲区布局PUT/GET/NOTIFY 共用与 URMA 测试一致的对称缓冲区布局见 tput_async_rdma_kernel.cpp 中的常量[64 x int32 头部区][sendBuf: count x T][recvBuf: count x T]头部区前 4 字节为deviceStatusKernel 向 host 回传状态kRdmaNotifySignalOffset 2 * sizeof(uint32_t)处为 notify signal其前后各 4 字节为 canary远端目标地址由peer 的注册 MR 基地址PeerMrBaseAddr(rdmaWorkspace, targetPeer) recv 区域偏移计算得出实现真正的单侧one-sided远程写。三种异步指令的调用形态TPUT_ASYNC远程 WRITETPUT_ASYNCDmaEngine::RDMA(remoteRecvGlobal, sendGlobal, session, targetPeer)root rank 将本地 send buffer 写入每个 peer 的 recv bufferTGET_ASYNC远程 READTGET_ASYNCDmaEngine::RDMA(localRecvGlobal, remoteSendGlobal, session, sourcePeer)root rank 从每个 peer 的 send 区域读取数据到本地仅PTO_RDMA_GET_TEST编译TPUT_ASYNC_NOTIFY带通知的远程 WRITETPUT_ASYNC_NOTIFYDmaEngine::RDMA(destination, source, remoteSignal, kRdmaNotifySignalValue, NotifyOp::Set, session, targetPeer)数据写入完成后将 signal 置为37。数据面从 AIV 核发起Kernel 以[[bisheng::core_ratio(0, 1)]]标注、__global__ AICORE修饰每次 PUT/GET 前通过BuildAsyncSessionDmaEngine::RDMA构建AsyncSession失败时写入kRdmaSessionBuildError并pipe_barrier(PIPE_ALL)。三种完成模式测试通过RdmaCompletionMode枚举覆盖三种完成语义见 tput_async_rdma_kernel.h完成模式语义覆盖的测试STATUS_WAIT_EACH每个 WQE 提交后立即WaitEventStatus排空Vec_Int32_MultiWqe_WaitEachSTATUS_WAIT_LAST批量提交多个 WQE 后只等最后一个事件Vec_Int32_MultiWqe_WaitLastPUBLIC_EVENT_WAIT_TEST走公开AsyncEvent::Test/Wait接口先Test后Wait再TestVec_Float_PublicEventWaitTest、NOTIFY 用例多 WQE 用例operationCount 16刻意在同一个 QP 上反复推进 SQ/CQ而不是每次 WQE 后重建队列用于验证队列推进与事件排空的正确性。用例清单与校验main.cpp中注册的 PUT 用例每个用例都有SKIP_IF_RANKS_LT(2)守卫rank 不足自动跳过用例类型/规模验证点Vec_FloatSmallfloat × 256基础小包 PUTVec_Int32Largeint32 × 4096大包 PUTVec_Uint8Smalluint8 × 512字节类型 PUTVec_Uint8_SingleChunkuint8 × 64单 chunk 传输Vec_Float_ExactChunkfloat × 64恰好一个 chunkVec_Uint8_Offset_63B偏移 17、63 元素、1 次操作非零源/目的偏移 两侧 canaryVec_Int32_MultiWqe_WaitEach/WaitLast16 次 × 256 元素单 QP 多 WQE 两种排空模式Vec_Float_PublicEventWaitTestfloat × 256公开 AsyncEvent 分发Test before/after WaitVec_Float_MR_4MBfloat × 524288约 4 MiB 注册 MR对齐 URMA 大 MR 档位Vec_FloatSmall_4Ranks4 rank4 rank 多进程场景需mpirun -n 4匹配数据校验采用确定性构造输入值RdmaInputValue(index, rankId)int 类型为index rankId * 10000uint8 类型为位运算散列未写入区域预填 sentinelint 类型为-1uint8 类型为0xa5派生值。host 侧VerifyPutResult/VerifyGetResult逐元素比对并把校验结果通过AllRanksReady做集合归约确保任一 rank 失败即整体失败。生命周期管理每个用例拥有一个完整的后端 channel 生命周期RdmaTestContextaclrtSetDevice→rtStreamCreate→aclrtMalloc(ACL_MEM_MALLOC_HUGE_FIRST)分配对称缓冲区 →SetupBootstrap解析 phyId/IP/端口并MPI_Allgather交换→ConnectRdma填充WorkspaceConfigrankId、rankCount、phyId、localIp、basePort、peerIps、peerPhyIds、peerSymAddrs、symmetricAddr/Size 后调用rdmaMgr.Init→ 启动 Kernel →aclrtSynchronizeStream→Verify→CleanuprdmaMgr.Finalize、aclrtFree、rtStreamDestroy。复用同一端口跑完整个用例集本身也验证了 channel 销毁后重连的能力。问题定位现象排查手段CMake 提示 RDMA 未使能设置PTO_RDMA_BACKENDHNS_1825并在不使用-w的情况下重新配置构建。端点发现失败检查物理设备映射以及 root-info 或 virtual topology 中的 CLOS IPv4必要时使用测试专用 IP 变量PTO_ROCE_PHYIDS/PTO_ROCE_LOCAL_IP/PTO_ROCE_IPS兜底。HCOMM 无法从默认路径加载 HNS1825 verbs provider将IBV_EXTEND_DRIVERS指向驱动提供的libhrn5-rdmav34.so。无法区分失败阶段设置PTO_ROCE_VERBOSE1日志会按[RDMA][HNS_1825][case N][rank M]前缀区分端点发现、MR 注册、channel 建链和释放阶段的错误host 侧PrintRdmaDeviceStatus还会把 Kernel 回传的DevStatus翻译为可读名称session_build_fail、public_event_wait_failed、poll_cq_timeout、cqe_error_without_syndrome等映射见 rdma_test_backend.hpp 的DescribeBackendCompletionStatus。另外运行前会做集合级前置检查AgreeOnRdmaPreflight所有 rank 都编译进 RDMA 后端才继续READY全部未使能则整体 SKIP打印no RDMA backend was compiled into this binary各 rank 选择不一致则直接 ERROR——因此当部分 rank 二进制不一致时需要回到构建阶段统一PTO_RDMA_BACKEND重新构建。延伸阅读顶层 ST 运行入口与参数说明tests/script/run_st.py、tests/README_zh.mdRDMA 异步指令与工作空间管理include/pto/comm/async/rdma/下的rdma_async_intrin.hpp、rdma_workspace_manager.hpp以及backends/hns_1825/后端实现相关 URMA 对照实现tput_async_urma、tget_async_urma、tput_async_notify_urma缓冲区布局与完成模式与 RDMA 版对齐【免费下载链接】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），仅供参考