CANN SHMEM RDMA Demo 实战指南：RoCE 环境搭建、编译构建与跨设备 AllGather 通信验证-平芜编程栈

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 个 PE（Processing 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

环境可用时，命令输出应类似下图所示（各网卡配置了有效 IP，net_health返回正常状态）：

Ascend950 平台

Ascend950 平台使用ibv_devinfo命令检查 RDMA 设备信息，根据网卡型号选择对应的过滤关键字：

XSCALE（云脉）网卡，查询xscale关键字：
```
ibv_devinfo | grep xscale
```
正常的输出示例：
HNS 1825 网卡，查询hrn关键字：
```
ibv_devinfo | grep hrn
```
正常的输出示例：

注意：1825 网卡在同一物理端口上多 NPU 通信时，若交换机未开启端口桥（port bridge），RDMA 可能无法正常收发数据。端口桥配置方法详见 Troubleshooting_FAQs - 同端口通信需开启端口桥：多个 NPU 共用同一物理端口时，RDMA 报文会从该端口发出、经交换机后从同一端口返回，交换机会默认丢弃此类同源同宿报文，需要在交换机对应接口上执行port bridge enable并commit提交配置。

IBV_EXTEND_DRIVERS 环境变量

Ascend950 平台运行前必须设置IBV_EXTEND_DRIVERS环境变量，指向对应网卡的用户态 RDMA Verbs provider 插件库：

XSCALE（云脉）网卡：

export IBV_EXTEND_DRIVERS=<path_to_libxscale_nda.so>

HNS 1825 网卡：

export IBV_EXTEND_DRIVERS=<path_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.so"

IBV_EXTEND_DRIVERS是 libibverbs 定义的环境变量，其作用是让 libibverbs 加载位于默认搜索路径之外的 Verbs provider 插件库。A2/A3 平台通常无需设置该变量。

编译构建

在shmem/项目根目录下执行编译命令（RDMA 后端的完整参数说明见 Compilation and Build - RDMA Parameters）。

A2/A3 平台，仅需使能 RDMA 能力：

bash scripts/build.sh -enable_rdma -examples

Ascend950 平台（XSCALE 网卡）：

bash scripts/build.sh -soc_type Ascend950 -enable_rdma -rdma_backend XSCALE -examples

Ascend950 平台（HNS 1825 网卡）：

bash scripts/build.sh -soc_type Ascend950 -enable_rdma -rdma_backend HNS_1825 -examples

RDMA 构建参数与依赖规则

根据 compilation_build_guide_en.md，RDMA 相关构建参数的行为如下：

参数	作用与适用范围
`-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_XSCALE=1，未指定后端时自动添加-DACLSHMEMI_RDMA_K_BACKEND_IN_DIE=1；这些宏无需也不能手动定义。设备侧内部宏ACLSHMEMI_K_RDMA_BACKEND由系统在 shmem_device_rdma.hpp 中根据自动生成的宏统一设置，手动定义会导致编译错误或运行时异常。

运行方式

方式一：使用 run.sh 脚本

在examples/rdma_demo目录下执行：

bash run.sh -pes 4

run.sh通过-pes参数指定启动的 PE 数量，默认值为 2。从 run.sh 的源码可以看到它的实际行为：先自动推导项目根目录并设置LD_LIBRARY_PATH=${PROJECT_ROOT}/build/lib，随后导出SHMEM_UID_SESSION_ID=127.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_ROOT=<shmem-root-directory> export IBV_EXTEND_DRIVERS=<path_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_ROOT=<shmem-root-directory> export IBV_EXTEND_DRIVERS=<path_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_ROOT=<shmem-root-directory> export IBV_EXTEND_DRIVERS=<path_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.so，1825 网卡为libhrn5-rdmav34.so）。跨机测试中tcp://ip1:8765的 IP 必须指向PE0 所在主机。
如需在容器中运行跨机测试，启动容器时指定--net=host模式即可。

命令行参数说明

rdma_demo的命令行格式为：

./rdma_demo <n_pes> <pe_id> <ipport> <g_npus> <f_pe> <f_npu>

参数	含义
`n_pes`	全局 PE 数量
`pe_id`	当前进程的 PE 号
`ipport`	SHMEM 初始化所需的 IP 与端口，格式为`tcp://<IP地址>:<端口号>`；跨机测试时 IP 需设为 PE0 所在 Host 的 IP
`g_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)、aclrtCreateStream；
SHMEM 初始化：通过 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用最简单的方式实现了 AllGather：

extern "C" [[bisheng::core_ratio(0, 1)]] __global__ __aicore__ void device_all_gather_test( GM_ADDR gva, int message_length) { AscendC::TPipe pipe; AscendC::TBuf<AscendC::TPosition::VECOUT> buf; pipe.InitBuffer(buf, UB_ALIGN_SIZE_64 * 2); // 需要用户指定一个长度大于等于128字节的LocalTensor用于RDMA任务下发 AscendC::LocalTensor<uint8_t> ubLocal = buf.GetWithOffset<uint8_t>(UB_ALIGN_SIZE_64 * 2, 0); int64_t my_rank = aclshmem_my_pe(); int64_t pe_size = aclshmem_n_pes(); AscendC::PipeBarrier<PIPE_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 pe=<pe_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 - 同端口通信需开启端口桥。
网络丢包：使能 RDMA（RoCEv2）后出现丢包时，可参考 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 网络上的各类通信原语。

创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考