CANN SHMEM RDMA Demo 实战指南RoCE 环境搭建、编译构建与跨设备 AllGather 通信验证【免费下载链接】shmemCANN SHMEM 是面向昇腾平台的多机多卡内存通信库基于OpenSHMEM 标准协议实现跨设备的高效内存访问与数据同步。项目地址: https://gitcode.com/cann/shmem本篇技术指南以 CANN SHMEM 仓库中的rdma_demo样例为核心系统讲解基于 RDMA/RoCE 的跨设备集合通信示例从环境检查、编译构建到单机/跨机运行验证的完整链路。读完本文后你可以在 A2/A3 与 Ascend950 平台上正确配置 RDMA 运行环境理解IBV_EXTEND_DRIVERS等关键环境变量的作用掌握rdma_demo的六个命令行参数并能够结合 main.cpp 与 rdma_demo_kernel.cpp 的源码理解 AllGather 在设备侧通过 RoCE Put 与 Barrier 实现的基本原理。样例定位用一次 AllGather 验证 RDMA 通路rdma_demo是 CANN SHMEM 提供的最小化 RDMA 验证样例其功能非常聚焦在 N 个 PEProcessing Element之间执行一次 AllGather 集合通信并在主机侧对结果做逐元素校验最终打印check transport result success与[SUCCESS] demo run success。它不依赖任何外部数据集也不需要额外的输入文件因此特别适合在部署完 RDMA 网卡与驱动后作为验证 RDMA 数据面是否打通的第一道关卡。从代码结构看该样例由三部分组成主机侧入口 main.cpp完成 ACL 初始化、SHMEM 初始化、数据准备、kernel 下发与结果校验设备侧 kernel rdma_demo_kernel.cpp使用aclshmemx_roce_put_nbi与aclshmemx_roce_barrier_all实现 AllGather构建与启动脚本CMakeLists.txt 通过aclshmem_add_collective_example(rdma_demo)注册构建目标run.sh 负责一键拉起多个 PE 进程。环境要求运行rdma_demo前必须保证机器具备可用的 RDMA 环境即 RDMA 网卡及对应驱动已正确安装并配置。在此基础上不同平台还有各自的额外约束。Ascend950 平台的 CANN 版本要求Ascend950 平台上的 RDMA 样例要求安装CANN 9.1.0版本的 CANN 包其他版本不在当前样例的支持范围内。请从 CANN 官方下载渠道获取对应版本的安装包后再进行编译运行。检查 RDMA 环境A2/A3 平台A2/A3 平台可直接使用hccn_tool检查网卡 IP 配置与网络健康状态{0..7}中的7需根据实际要检查的卡数量修改for i in {0..7}; do hccn_tool -i $i -ip -g; done for i in {0..7}; do hccn_tool -i $i -net_health -g; done环境可用时命令输出应类似下图所示各网卡配置了有效 IPnet_health返回正常状态Ascend950 平台Ascend950 平台使用ibv_devinfo命令检查 RDMA 设备信息根据网卡型号选择对应的过滤关键字XSCALE云脉网卡查询xscale关键字ibv_devinfo | grep xscale正常的输出示例HNS 1825 网卡查询hrn关键字ibv_devinfo | grep hrn正常的输出示例注意1825 网卡在同一物理端口上多 NPU 通信时若交换机未开启端口桥port bridgeRDMA 可能无法正常收发数据。端口桥配置方法详见 Troubleshooting_FAQs - 同端口通信需开启端口桥多个 NPU 共用同一物理端口时RDMA 报文会从该端口发出、经交换机后从同一端口返回交换机会默认丢弃此类同源同宿报文需要在交换机对应接口上执行port bridge enable并commit提交配置。IBV_EXTEND_DRIVERS 环境变量Ascend950 平台运行前必须设置IBV_EXTEND_DRIVERS环境变量指向对应网卡的用户态 RDMA Verbs provider 插件库XSCALE云脉网卡export IBV_EXTEND_DRIVERSpath_to_libxscale_nda.soHNS 1825 网卡export IBV_EXTEND_DRIVERSpath_to_libhrn5-rdmav34.so需要特别强调的是libxscale_nda.so与libhrn5-rdmav34.so均为对应网卡驱动安装包自带的用户态 Verbs provider 库不是 SHMEM 项目的编译产物。其中libxscale_nda.so随 XSCALE 网卡驱动安装libhrn5-rdmav34.so随 HNS 1825 网卡驱动安装。安装网卡驱动后可以通过以下命令定位库的实际路径并将其作为IBV_EXTEND_DRIVERS的值find / -name libhrn5-rdmav34.so # 或 find / -name libxscale_nda.soIBV_EXTEND_DRIVERS是 libibverbs 定义的环境变量其作用是让 libibverbs 加载位于默认搜索路径之外的 Verbs provider 插件库。A2/A3 平台通常无需设置该变量。编译构建在shmem/项目根目录下执行编译命令RDMA 后端的完整参数说明见 Compilation and Build - RDMA Parameters。A2/A3 平台仅需使能 RDMA 能力bash scripts/build.sh -enable_rdma -examplesAscend950 平台XSCALE 网卡bash scripts/build.sh -soc_type Ascend950 -enable_rdma -rdma_backend XSCALE -examplesAscend950 平台HNS 1825 网卡bash scripts/build.sh -soc_type Ascend950 -enable_rdma -rdma_backend HNS_1825 -examplesRDMA 构建参数与依赖规则根据 compilation_build_guide_en.mdRDMA 相关构建参数的行为如下参数作用与适用范围-enable_rdma编译并使能 RDMA 能力。Ascend 910B/C 上默认使用内置 RDMA 后端无需额外指定后端类型-soc_type Ascend950显式声明目标 SoC 为 Ascend950配合-rdma_backend使用-rdma_backend XSCALE指定使用 XSCALE云脉网卡后端仅在-soc_type Ascend950下合法-rdma_backend HNS_1825指定使用 HNS 1825 网卡后端仅在-soc_type Ascend950下合法参数依赖规则使用-rdma_backend时必须同时指定-enable_rdma否则构建失败并提示Error: -rdma_backend requires -enable_rdma to be specified.-rdma_backend仅在-soc_type Ascend950下有效在其他 SoC 上指定会报错Error: -rdma_backend can only be specified when SOC_TYPE is Ascend950.三个参数-rdma_backend、-enable_rdma、-soc_type xxx的书写顺序不限但所有依赖参数必须齐全。需要留意的是CMake 会根据-rdma_backend的值自动生成编译宏例如指定XSCALE时自动添加-DACLSHMEMI_RDMA_K_BACKEND_XSCALE1未指定后端时自动添加-DACLSHMEMI_RDMA_K_BACKEND_IN_DIE1这些宏无需也不能手动定义。设备侧内部宏ACLSHMEMI_K_RDMA_BACKEND由系统在 shmem_device_rdma.hpp 中根据自动生成的宏统一设置手动定义会导致编译错误或运行时异常。运行方式方式一使用 run.sh 脚本在examples/rdma_demo目录下执行bash run.sh -pes 4run.sh通过-pes参数指定启动的 PE 数量默认值为 2。从 run.sh 的源码可以看到它的实际行为先自动推导项目根目录并设置LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib随后导出SHMEM_UID_SESSION_ID127.0.0.1:8899最后循环拉起num_pes个./build/bin/rdma_demo ${num_pes} ${i} tcp://127.0.0.1:8899 ${num_pes} 0 0进程并等待所有进程结束后以首个非零返回值作为脚本退出码。因此该方式只适用于单机多卡场景。注意Ascend950 平台在运行前必须设置IBV_EXTEND_DRIVERS环境变量详见上文 IBV_EXTEND_DRIVERS 环境变量 一节。方式二在 shmem/ 目录手动执行命令单机双卡执行shmem-root-directory为 SHMEM 项目根目录export PROJECT_ROOTshmem-root-directory export IBV_EXTEND_DRIVERSpath_to_plugin.so # 仅 Ascend950 平台需要按网卡类型设置 export LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib:$LD_LIBRARY_PATH ./build/bin/rdma_demo 2 0 tcp://127.0.0.1:8765 2 0 0 # PE 0 ./build/bin/rdma_demo 2 1 tcp://127.0.0.1:8765 2 0 0 # PE 1跨机双卡执行假设服务器 A 的 IP 为ip1服务器 B 的 IP 为ip2在服务器 A 上执行export PROJECT_ROOTshmem-root-directory export IBV_EXTEND_DRIVERSpath_to_plugin.so # 仅 Ascend950 平台需要按网卡类型设置 export LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib:$LD_LIBRARY_PATH ./build/bin/rdma_demo 2 0 tcp://ip1:8765 1 0 0 # PE 0与此同时在服务器 B 上执行export PROJECT_ROOTshmem-root-directory export IBV_EXTEND_DRIVERSpath_to_plugin.so # 仅 Ascend950 平台需要按网卡类型设置 export LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib:$LD_LIBRARY_PATH ./build/bin/rdma_demo 2 1 tcp://ip1:8765 1 1 0 # PE 1注path_to_plugin.so为按网卡类型确定的插件库路径XSCALE 网卡为libxscale_nda.so1825 网卡为libhrn5-rdmav34.so。跨机测试中tcp://ip1:8765的 IP 必须指向PE0 所在主机。如需在容器中运行跨机测试启动容器时指定--nethost模式即可。命令行参数说明rdma_demo的命令行格式为./rdma_demo n_pes pe_id ipport g_npus f_pe f_npu参数含义n_pes全局 PE 数量pe_id当前进程的 PE 号ipportSHMEM 初始化所需的 IP 与端口格式为tcp://IP地址:端口号跨机测试时 IP 需设为 PE0 所在 Host 的 IPg_npus当前服务器上启动的 NPU 卡数量f_pe当前服务器上使用的第一个 PE 号f_npu当前服务器上执行本样例使用的第一张 NPU 卡的卡号这些参数在 main.cpp 中按顺序解析其中设备号由device_id pe_id % g_npus f_npu计算得出local_mem_size在样例中固定为 1 GiB。深入理解主机侧与设备侧的协作流程主机侧从 ACL 初始化到结果校验main.cpp 中的test_aclshmem_team_all_gather展示了完整的调用链ACL 环境初始化依次调用aclInit、aclrtSetDevice(device_id)、aclrtCreateStreamSHMEM 初始化通过 utils.h 中的test_set_attr填充aclshmemx_init_attr_t设置my_pe、n_pes、ip_port、local_mem_size等字段随后将attributes.option_attr.data_op_engine_type显式指定为ACLSHMEM_DATA_OP_ROCE再调用aclshmemx_init_attr(ACLSHMEMX_INIT_WITH_DEFAULT, attributes)完成初始化——这一步是样例走 RDMA 数据面的关键异常快照使能调用aclshmemx_enable_exception_report(nullptr, ACLSHMEMX_EXCEPTION_REPORT_DEBUG)开启 Runtime 异常快照与 RDMA 队列诊断若返回ACLSHMEM_NOT_SUPPORTED则跳过老版本 Runtime 不支持该能力对称内存分配与数据准备aclshmem_malloc(1024)分配 1024 字节对称内存每个 PE 将自己的数据pe_id 10通过aclrtMemcpy写入ptr aclshmem_my_pe() * trans_size * sizeof(int32_t)对应的分片kernel 下发调用allgather_demo(1, stream, ptr, trans_size * sizeof(int32_t))以单 block 下发 AllGather kernel随后aclrtSynchronizeStream等待完成异常上报与结果校验aclshmemx_report_exception()上报设备侧 trap流同步后即为安全的上报点随后将对称内存拷回主机逐元素校验y_host[trans_size * i trans_size / block_size * j]是否等于10 i全部通过则打印check transport result success资源释放依次aclshmem_finalize、销毁流、复位设备、aclFinalize。设备侧AllGather 的 RoCE 实现rdma_demo_kernel.cpp 中的device_all_gather_test用最简单的方式实现了 AllGatherextern C [[bisheng::core_ratio(0, 1)]] __global__ __aicore__ void device_all_gather_test( GM_ADDR gva, int message_length) { AscendC::TPipe pipe; AscendC::TBufAscendC::TPosition::VECOUT buf; pipe.InitBuffer(buf, UB_ALIGN_SIZE_64 * 2); // 需要用户指定一个长度大于等于128字节的LocalTensor用于RDMA任务下发 AscendC::LocalTensoruint8_t ubLocal buf.GetWithOffsetuint8_t(UB_ALIGN_SIZE_64 * 2, 0); int64_t my_rank aclshmem_my_pe(); int64_t pe_size aclshmem_n_pes(); AscendC::PipeBarrierPIPE_ALL(); aclshmemx_roce_barrier_all(); for (int i 0; i pe_size; i) { if (i my_rank) { continue; } aclshmemx_roce_put_nbi( gva message_length * my_rank, gva message_length * my_rank, (__ubuf__ uint8_t*)ubLocal.GetPhyAddr(), message_length, i, 0); } aclshmemx_roce_barrier_all(); }其核心逻辑分为三步分配一块不小于 128 字节的 UB LocalTensor 作为 RDMA 任务下发缓冲区这是aclshmemx_roce_*系列接口的硬性要求见 shmem_device_rdma.h 中aclshmemx_roce_barrier_all的注释先执行aclshmemx_roce_barrier_all()保证所有 PE 的数据已写入各自对称内存分片对除自身外的每个 PEi以gva message_length * my_rank为源地址和目的地址源目标指向自身对称内存中的分片通过aclshmemx_roce_put_nbi(dst, src, buf, elem_size, pe, 0)非阻塞 Put 到对端循环结束后再执行一次aclshmemx_roce_barrier_all()确保所有写操作完成。aclshmemx_roce_put_nbi的设备侧声明位于 shmem_device_rdma.h其语义约束值得注意dst会被翻译为对端 PE 上的对应地址src为本地 RDMA 操作数两个操作数都必须指向对称内存且每次传输的完整范围必须落在对应分配区间内此外RDMA 作为底层传输时同一 PE 上的并发 RMA/AMO 操作不受支持需要使用device_state.rdma_config中的sync_id做流水线同步即aclshmemx_roce_put_nbi带sync_id的重载形式。输出与判据每个 PE 成功通过校验后会打印check transport result success, relative pepe_id [SUCCESS] demo run success in relative pe pe_id若校验失败main.cpp 会打印具体的数值不匹配项如xx ! 10 i并返回-1此时脚本退出码为非零说明 RDMA 数据面存在问题需要回到环境检查环节排查。常见问题与后续排查1825 网卡同端口通信失败单机多 NPU 共用同一物理端口且未在交换机开启端口桥时RDMA 收发包异常。排查与配置步骤含 NPU 与网卡端口对应关系、ibv_devinfo/hiroce5 gids查询、display mac-address定位交换机端口、port bridge enable配置等见 Troubleshooting_FAQs - 同端口通信需开启端口桥。网络丢包使能 RDMARoCEv2后出现丢包时可参考 Troubleshooting_FAQs - 通信丢包 一节从网卡与流控配置角度排查。会话建立失败若初始化阶段 UID 会话无法建立可检查SHMEM_UID_SESSION_ID/SHMEM_UID_SOCK_IFNAME等环境变量配置相关说明见 env_vars_intro.md 与 Troubleshooting_FAQs.md。rdma_demo是了解 CANN SHMEM RDMA 数据面的最小完整闭环从编译期-enable_rdma/-rdma_backend的后端选型到运行期IBV_EXTEND_DRIVERS的 provider 加载再到设备侧aclshmemx_roce_put_nbi与 Barrier 的对称内存操作语义全链路均可在该样例中直接验证。将rdma_demo跑通后再进一步阅读 rdma_aggregate_demo、rdma_atomic_demo 与 rdma_perftest_demo 等样例即可系统掌握 CANN SHMEM 在 RoCE 网络上的各类通信原语。【免费下载链接】shmemCANN SHMEM 是面向昇腾平台的多机多卡内存通信库基于OpenSHMEM 标准协议实现跨设备的高效内存访问与数据同步。项目地址: https://gitcode.com/cann/shmem创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN SHMEM RDMA Demo 实战指南RoCE 环境搭建、编译构建与跨设备 AllGather 通信验证【免费下载链接】shmemCANN SHMEM 是面向昇腾平台的多机多卡内存通信库基于OpenSHMEM 标准协议实现跨设备的高效内存访问与数据同步。项目地址: https://gitcode.com/cann/shmem本篇技术指南以 CANN SHMEM 仓库中的rdma_demo样例为核心系统讲解基于 RDMA/RoCE 的跨设备集合通信示例从环境检查、编译构建到单机/跨机运行验证的完整链路。读完本文后你可以在 A2/A3 与 Ascend950 平台上正确配置 RDMA 运行环境理解IBV_EXTEND_DRIVERS等关键环境变量的作用掌握rdma_demo的六个命令行参数并能够结合 main.cpp 与 rdma_demo_kernel.cpp 的源码理解 AllGather 在设备侧通过 RoCE Put 与 Barrier 实现的基本原理。样例定位用一次 AllGather 验证 RDMA 通路rdma_demo是 CANN SHMEM 提供的最小化 RDMA 验证样例其功能非常聚焦在 N 个 PEProcessing Element之间执行一次 AllGather 集合通信并在主机侧对结果做逐元素校验最终打印check transport result success与[SUCCESS] demo run success。它不依赖任何外部数据集也不需要额外的输入文件因此特别适合在部署完 RDMA 网卡与驱动后作为验证 RDMA 数据面是否打通的第一道关卡。从代码结构看该样例由三部分组成主机侧入口 main.cpp完成 ACL 初始化、SHMEM 初始化、数据准备、kernel 下发与结果校验设备侧 kernel rdma_demo_kernel.cpp使用aclshmemx_roce_put_nbi与aclshmemx_roce_barrier_all实现 AllGather构建与启动脚本CMakeLists.txt 通过aclshmem_add_collective_example(rdma_demo)注册构建目标run.sh 负责一键拉起多个 PE 进程。环境要求运行rdma_demo前必须保证机器具备可用的 RDMA 环境即 RDMA 网卡及对应驱动已正确安装并配置。在此基础上不同平台还有各自的额外约束。Ascend950 平台的 CANN 版本要求Ascend950 平台上的 RDMA 样例要求安装CANN 9.1.0版本的 CANN 包其他版本不在当前样例的支持范围内。请从 CANN 官方下载渠道获取对应版本的安装包后再进行编译运行。检查 RDMA 环境A2/A3 平台A2/A3 平台可直接使用hccn_tool检查网卡 IP 配置与网络健康状态{0..7}中的7需根据实际要检查的卡数量修改for i in {0..7}; do hccn_tool -i $i -ip -g; done for i in {0..7}; do hccn_tool -i $i -net_health -g; done环境可用时命令输出应类似下图所示各网卡配置了有效 IPnet_health返回正常状态Ascend950 平台Ascend950 平台使用ibv_devinfo命令检查 RDMA 设备信息根据网卡型号选择对应的过滤关键字XSCALE云脉网卡查询xscale关键字ibv_devinfo | grep xscale正常的输出示例HNS 1825 网卡查询hrn关键字ibv_devinfo | grep hrn正常的输出示例注意1825 网卡在同一物理端口上多 NPU 通信时若交换机未开启端口桥port bridgeRDMA 可能无法正常收发数据。端口桥配置方法详见 Troubleshooting_FAQs - 同端口通信需开启端口桥多个 NPU 共用同一物理端口时RDMA 报文会从该端口发出、经交换机后从同一端口返回交换机会默认丢弃此类同源同宿报文需要在交换机对应接口上执行port bridge enable并commit提交配置。IBV_EXTEND_DRIVERS 环境变量Ascend950 平台运行前必须设置IBV_EXTEND_DRIVERS环境变量指向对应网卡的用户态 RDMA Verbs provider 插件库XSCALE云脉网卡export IBV_EXTEND_DRIVERSpath_to_libxscale_nda.soHNS 1825 网卡export IBV_EXTEND_DRIVERSpath_to_libhrn5-rdmav34.so需要特别强调的是libxscale_nda.so与libhrn5-rdmav34.so均为对应网卡驱动安装包自带的用户态 Verbs provider 库不是 SHMEM 项目的编译产物。其中libxscale_nda.so随 XSCALE 网卡驱动安装libhrn5-rdmav34.so随 HNS 1825 网卡驱动安装。安装网卡驱动后可以通过以下命令定位库的实际路径并将其作为IBV_EXTEND_DRIVERS的值find / -name libhrn5-rdmav34.so # 或 find / -name libxscale_nda.soIBV_EXTEND_DRIVERS是 libibverbs 定义的环境变量其作用是让 libibverbs 加载位于默认搜索路径之外的 Verbs provider 插件库。A2/A3 平台通常无需设置该变量。编译构建在shmem/项目根目录下执行编译命令RDMA 后端的完整参数说明见 Compilation and Build - RDMA Parameters。A2/A3 平台仅需使能 RDMA 能力bash scripts/build.sh -enable_rdma -examplesAscend950 平台XSCALE 网卡bash scripts/build.sh -soc_type Ascend950 -enable_rdma -rdma_backend XSCALE -examplesAscend950 平台HNS 1825 网卡bash scripts/build.sh -soc_type Ascend950 -enable_rdma -rdma_backend HNS_1825 -examplesRDMA 构建参数与依赖规则根据 compilation_build_guide_en.mdRDMA 相关构建参数的行为如下参数作用与适用范围-enable_rdma编译并使能 RDMA 能力。Ascend 910B/C 上默认使用内置 RDMA 后端无需额外指定后端类型-soc_type Ascend950显式声明目标 SoC 为 Ascend950配合-rdma_backend使用-rdma_backend XSCALE指定使用 XSCALE云脉网卡后端仅在-soc_type Ascend950下合法-rdma_backend HNS_1825指定使用 HNS 1825 网卡后端仅在-soc_type Ascend950下合法参数依赖规则使用-rdma_backend时必须同时指定-enable_rdma否则构建失败并提示Error: -rdma_backend requires -enable_rdma to be specified.-rdma_backend仅在-soc_type Ascend950下有效在其他 SoC 上指定会报错Error: -rdma_backend can only be specified when SOC_TYPE is Ascend950.三个参数-rdma_backend、-enable_rdma、-soc_type xxx的书写顺序不限但所有依赖参数必须齐全。需要留意的是CMake 会根据-rdma_backend的值自动生成编译宏例如指定XSCALE时自动添加-DACLSHMEMI_RDMA_K_BACKEND_XSCALE1未指定后端时自动添加-DACLSHMEMI_RDMA_K_BACKEND_IN_DIE1这些宏无需也不能手动定义。设备侧内部宏ACLSHMEMI_K_RDMA_BACKEND由系统在 shmem_device_rdma.hpp 中根据自动生成的宏统一设置手动定义会导致编译错误或运行时异常。运行方式方式一使用 run.sh 脚本在examples/rdma_demo目录下执行bash run.sh -pes 4run.sh通过-pes参数指定启动的 PE 数量默认值为 2。从 run.sh 的源码可以看到它的实际行为先自动推导项目根目录并设置LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib随后导出SHMEM_UID_SESSION_ID127.0.0.1:8899最后循环拉起num_pes个./build/bin/rdma_demo ${num_pes} ${i} tcp://127.0.0.1:8899 ${num_pes} 0 0进程并等待所有进程结束后以首个非零返回值作为脚本退出码。因此该方式只适用于单机多卡场景。注意Ascend950 平台在运行前必须设置IBV_EXTEND_DRIVERS环境变量详见上文 IBV_EXTEND_DRIVERS 环境变量 一节。方式二在 shmem/ 目录手动执行命令单机双卡执行shmem-root-directory为 SHMEM 项目根目录export PROJECT_ROOTshmem-root-directory export IBV_EXTEND_DRIVERSpath_to_plugin.so # 仅 Ascend950 平台需要按网卡类型设置 export LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib:$LD_LIBRARY_PATH ./build/bin/rdma_demo 2 0 tcp://127.0.0.1:8765 2 0 0 # PE 0 ./build/bin/rdma_demo 2 1 tcp://127.0.0.1:8765 2 0 0 # PE 1跨机双卡执行假设服务器 A 的 IP 为ip1服务器 B 的 IP 为ip2在服务器 A 上执行export PROJECT_ROOTshmem-root-directory export IBV_EXTEND_DRIVERSpath_to_plugin.so # 仅 Ascend950 平台需要按网卡类型设置 export LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib:$LD_LIBRARY_PATH ./build/bin/rdma_demo 2 0 tcp://ip1:8765 1 0 0 # PE 0与此同时在服务器 B 上执行export PROJECT_ROOTshmem-root-directory export IBV_EXTEND_DRIVERSpath_to_plugin.so # 仅 Ascend950 平台需要按网卡类型设置 export LD_LIBRARY_PATH${PROJECT_ROOT}/build/lib:$LD_LIBRARY_PATH ./build/bin/rdma_demo 2 1 tcp://ip1:8765 1 1 0 # PE 1注path_to_plugin.so为按网卡类型确定的插件库路径XSCALE 网卡为libxscale_nda.so1825 网卡为libhrn5-rdmav34.so。跨机测试中tcp://ip1:8765的 IP 必须指向PE0 所在主机。如需在容器中运行跨机测试启动容器时指定--nethost模式即可。命令行参数说明rdma_demo的命令行格式为./rdma_demo n_pes pe_id ipport g_npus f_pe f_npu参数含义n_pes全局 PE 数量pe_id当前进程的 PE 号ipportSHMEM 初始化所需的 IP 与端口格式为tcp://IP地址:端口号跨机测试时 IP 需设为 PE0 所在 Host 的 IPg_npus当前服务器上启动的 NPU 卡数量f_pe当前服务器上使用的第一个 PE 号f_npu当前服务器上执行本样例使用的第一张 NPU 卡的卡号这些参数在 main.cpp 中按顺序解析其中设备号由device_id pe_id % g_npus f_npu计算得出local_mem_size在样例中固定为 1 GiB。深入理解主机侧与设备侧的协作流程主机侧从 ACL 初始化到结果校验main.cpp 中的test_aclshmem_team_all_gather展示了完整的调用链ACL 环境初始化依次调用aclInit、aclrtSetDevice(device_id)、aclrtCreateStreamSHMEM 初始化通过 utils.h 中的test_set_attr填充aclshmemx_init_attr_t设置my_pe、n_pes、ip_port、local_mem_size等字段随后将attributes.option_attr.data_op_engine_type显式指定为ACLSHMEM_DATA_OP_ROCE再调用aclshmemx_init_attr(ACLSHMEMX_INIT_WITH_DEFAULT, attributes)完成初始化——这一步是样例走 RDMA 数据面的关键异常快照使能调用aclshmemx_enable_exception_report(nullptr, ACLSHMEMX_EXCEPTION_REPORT_DEBUG)开启 Runtime 异常快照与 RDMA 队列诊断若返回ACLSHMEM_NOT_SUPPORTED则跳过老版本 Runtime 不支持该能力对称内存分配与数据准备aclshmem_malloc(1024)分配 1024 字节对称内存每个 PE 将自己的数据pe_id 10通过aclrtMemcpy写入ptr aclshmem_my_pe() * trans_size * sizeof(int32_t)对应的分片kernel 下发调用allgather_demo(1, stream, ptr, trans_size * sizeof(int32_t))以单 block 下发 AllGather kernel随后aclrtSynchronizeStream等待完成异常上报与结果校验aclshmemx_report_exception()上报设备侧 trap流同步后即为安全的上报点随后将对称内存拷回主机逐元素校验y_host[trans_size * i trans_size / block_size * j]是否等于10 i全部通过则打印check transport result success资源释放依次aclshmem_finalize、销毁流、复位设备、aclFinalize。设备侧AllGather 的 RoCE 实现rdma_demo_kernel.cpp 中的device_all_gather_test用最简单的方式实现了 AllGatherextern C [[bisheng::core_ratio(0, 1)]] __global__ __aicore__ void device_all_gather_test( GM_ADDR gva, int message_length) { AscendC::TPipe pipe; AscendC::TBufAscendC::TPosition::VECOUT buf; pipe.InitBuffer(buf, UB_ALIGN_SIZE_64 * 2); // 需要用户指定一个长度大于等于128字节的LocalTensor用于RDMA任务下发 AscendC::LocalTensoruint8_t ubLocal buf.GetWithOffsetuint8_t(UB_ALIGN_SIZE_64 * 2, 0); int64_t my_rank aclshmem_my_pe(); int64_t pe_size aclshmem_n_pes(); AscendC::PipeBarrierPIPE_ALL(); aclshmemx_roce_barrier_all(); for (int i 0; i pe_size; i) { if (i my_rank) { continue; } aclshmemx_roce_put_nbi( gva message_length * my_rank, gva message_length * my_rank, (__ubuf__ uint8_t*)ubLocal.GetPhyAddr(), message_length, i, 0); } aclshmemx_roce_barrier_all(); }其核心逻辑分为三步分配一块不小于 128 字节的 UB LocalTensor 作为 RDMA 任务下发缓冲区这是aclshmemx_roce_*系列接口的硬性要求见 shmem_device_rdma.h 中aclshmemx_roce_barrier_all的注释先执行aclshmemx_roce_barrier_all()保证所有 PE 的数据已写入各自对称内存分片对除自身外的每个 PEi以gva message_length * my_rank为源地址和目的地址源目标指向自身对称内存中的分片通过aclshmemx_roce_put_nbi(dst, src, buf, elem_size, pe, 0)非阻塞 Put 到对端循环结束后再执行一次aclshmemx_roce_barrier_all()确保所有写操作完成。aclshmemx_roce_put_nbi的设备侧声明位于 shmem_device_rdma.h其语义约束值得注意dst会被翻译为对端 PE 上的对应地址src为本地 RDMA 操作数两个操作数都必须指向对称内存且每次传输的完整范围必须落在对应分配区间内此外RDMA 作为底层传输时同一 PE 上的并发 RMA/AMO 操作不受支持需要使用device_state.rdma_config中的sync_id做流水线同步即aclshmemx_roce_put_nbi带sync_id的重载形式。输出与判据每个 PE 成功通过校验后会打印check transport result success, relative pepe_id [SUCCESS] demo run success in relative pe pe_id若校验失败main.cpp 会打印具体的数值不匹配项如xx ! 10 i并返回-1此时脚本退出码为非零说明 RDMA 数据面存在问题需要回到环境检查环节排查。常见问题与后续排查1825 网卡同端口通信失败单机多 NPU 共用同一物理端口且未在交换机开启端口桥时RDMA 收发包异常。排查与配置步骤含 NPU 与网卡端口对应关系、ibv_devinfo/hiroce5 gids查询、display mac-address定位交换机端口、port bridge enable配置等见 Troubleshooting_FAQs - 同端口通信需开启端口桥。网络丢包使能 RDMARoCEv2后出现丢包时可参考 Troubleshooting_FAQs - 通信丢包 一节从网卡与流控配置角度排查。会话建立失败若初始化阶段 UID 会话无法建立可检查SHMEM_UID_SESSION_ID/SHMEM_UID_SOCK_IFNAME等环境变量配置相关说明见 env_vars_intro.md 与 Troubleshooting_FAQs.md。rdma_demo是了解 CANN SHMEM RDMA 数据面的最小完整闭环从编译期-enable_rdma/-rdma_backend的后端选型到运行期IBV_EXTEND_DRIVERS的 provider 加载再到设备侧aclshmemx_roce_put_nbi与 Barrier 的对称内存操作语义全链路均可在该样例中直接验证。将rdma_demo跑通后再进一步阅读 rdma_aggregate_demo、rdma_atomic_demo 与 rdma_perftest_demo 等样例即可系统掌握 CANN SHMEM 在 RoCE 网络上的各类通信原语。【免费下载链接】shmemCANN SHMEM 是面向昇腾平台的多机多卡内存通信库基于OpenSHMEM 标准协议实现跨设备的高效内存访问与数据同步。项目地址: https://gitcode.com/cann/shmem创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
