Ascend 950 URMA RM

导言

如果目前只知道“URMA 是组 WQE、敲 doorbell”，下一步需要补上三个问题：请求可以发给谁、资源由谁准备、什么时候才算完成。 RM 主要回答第一个问题；SHMEM 的 Ascend 950 UDMA 示例把后两个问题串成了可阅读、可编译的 Ascend C 程序。本文从定义走到两卡示例，并给出原生 URMA 显式选择 RM 的办法。

本文核对于 2026-09-15：SHMEM 固定到 9ae0bfe55b621185e0f89dd1db748d430e346eb3，UMDK 固定到 b429422a183d7379bfc1fcdb6617b78c218f80a1。下面的命令和签名均按这两个源码截面解释。本次完成了文档、源码和脚本的静态核对，没有在 Ascend 950 上实际编译或运行，也没有测量性能。

RM 到底是什么

先想一个三卡场景：卡 0 先给卡 1 写入一块数据，接着给卡 2 写入另一块数据。我们希望发起端能复用一个通信端点，在每次请求里指定不同目标，同时由传输系统提供可靠传输。

RM 是 Reliable Message，可靠消息模式。它允许一个本地 Jetty/JFS 面向多个远端目标通信。 Jetty 可以先理解成“持有通信队列等资源的端点”，JFS 是发送侧队列资源。RM 的名字带有 Message，但它支持的不只是 send/receive，也包括访问远端已注册内存的 WRITE、READ 等操作。URMA 用户指南，5.3 数据面

“无连接可靠”是资料里常见的另一种表述。这里的“无连接”描述的是应用无需把本地端点永久绑定到唯一远端；底层仍然可能建立传输连接、维护状态、确认与重试。因此，不能把它理解成“无需初始化、无需交换远端信息”。

对照起来就容易记了：

模式	应用看到的目标关系	可靠性与使用边界
RC：Reliable Connection	一个本地 Jetty 绑定一个目标 Jetty	可靠；需要多个独立目标时要安排相应端点关系
RM：Reliable Message	一个本地端点可以面向多个远端	可靠；顺序能力还要看配置、操作和后端
UM：Unreliable Message	可以面向多个目标	不保证可靠；官方指南中的 UM 不支持单边语义

这里的可靠是传输语义，不是保证机器永远不故障，更不是保证业务已经消费数据。资源失效、访问权限错误或重试耗尽，仍然需要通过错误返回、完成状态或超时处理。

RM、RMA、UDMA 分别回答什么

这几个词很像，却位于不同层次：

• URMA（Unified Remote Memory Access）：统一远程内存访问的软件接口体系，应用通过它管理资源、提交请求、获取完成结果。
• RM：URMA 的一种传输模式，关心端点与目标的关系以及可靠性。
• RMA（Remote Memory Access）：远端内存访问这类操作，例如 put/write、get/read。RM 中可以执行 RMA。
• UDMA（Unified DMA）：本文 Ascend 950 示例使用的数据搬运后端；SHMEM 用 ACLSHMEM_DATA_OP_UDMA 选择它。
• SHMEM：在底层通信能力之上提供对称内存、PE 编号、put/get 和同步接口的通信库。

还有一个特别容易混淆的缩写：UB 网络中的 UB 指 Unified Bus；Ascend C 的 __ubuf__ / UB scratch 中的 UB 指 Unified Buffer，是核内缓冲区。 本文会分别说“UB 网络”和“核内 UB”。

RM 的优势来自哪里

对于需要和许多目标交换数据的程序，RM 提供了更灵活的端点复用方式。例如卡 0 的同一个本地端点，第一条请求选择卡 1 的目标句柄，第二条请求选择卡 2 的目标句柄。应用不需要为这个端点维持唯一对端的绑定关系。UMDK 的 sample_import_jetty 与发送请求

由此可以推导出三点价值，但也要把收益的来源分清：

1. 多目标通信更容易组织。 对多对多交换、动态选择目的端的程序，端点复用更自然。这是 RM 直接带来的编程模型优势。
2. 有机会减少应用端点管理开销。 复用本地端点可能降低应用管理的资源数量；但底层仍有连接、目标导入记录和队列，不能推导为“整个系统只剩一个 QP”或“内存从平方规模必然降到线性规模”。
3. 可与单边访问及设备直驱结合。 单边 WRITE 的目标 CPU 不必为每块数据执行一次接收拷贝；Ascend C 直驱又能让 AIV 发出传输请求。这分别来自操作语义和设备执行路径，不能全算成 RM 独有的优势。

RM 本身不等于低延迟开关。 排队、消息大小、保序要求、路径配置、同步频率和实际拓扑都会影响结果。本文没有可用于声称“比 RC 快多少”的对照实验。

SHMEM 示例里该找什么

领导提示的方向是有用的：docs/ 解释环境与 API，examples/udma_demo/ 展示 Ascend C 的完整使用路径。它们能教会你如何把 WQE 与 doorbell 放进一个可运行程序。

但在固定版本里，应用明确设置的是：

attributes.option_attr.data_op_engine_type = ACLSHMEM_DATA_OP_UDMA;

这是 main.cpp 的初始化流程，第 38–96 行中的一个字段，不是 URMA_TM_RM。

继续向下追，SHMEM 的 UdmaTransportManager 通过 HCOMM（CANN 的底层通信资源接口）创建端点、注册内存、创建通道，再把队列上下文提供给设备端。端点配置中可以看到：

endpoint_desc.protocol = COMM_PROTOCOL_UBC_CTP;
endpoint_desc.commAddr.type = COMM_ADDR_TYPE_EID;

来源是 device_udma_transport_manager.cpp，第 1530–1545 行。EID 是端点地址标识；CTP 是传输路径类型所在的另一层配置。

不要用 UDMA 或 CTP 证明 RM

trans_mode 和 tp_type 是不同参数。 原生 URMA 示例同时列出了 RM + CTP 和 RC + CTP 组合。因此，SHMEM 示例成功只能直接证明该环境下 UDMA 数据路径可用；仅凭上述字段，不能确认配套 HCOMM 最终创建资源时选择了 RM。需要继续核对机器上所用 HCOMM/Provider 的资源配置、源码或可观测信息。UMDK 示例组合表

这决定了学习顺序：先用 SHMEM 跑通 Ascend C 的数据路径，再用原生 URMA 或配套运行时证据确认 RM 配置。 如果验收项明确写着“验证 RM”，后一步不能省。

在 950 上编译和运行

先确认配套环境

这些操作在 Ascend 950 Linux 服务器上执行。Mac 可以阅读与编辑源码，不能据此完成 NPU 验证。

按当前 udma_demo/README.md，准备：

• Ascend 950 与可用的 UB 通信拓扑。两张卡都在服务器里，并不自动证明所需链路已经可用；跨机更需要确认设备侧路径。
• CANN 9.1.0 对应的 950 toolkit、ops 及配套驱动/固件。该示例明确不把低于 9.1.0 的 CANN 纳入支持范围；后续版本仍需按实际配套验证。
• bisheng 编译器、CMake 和源码依赖。详细系统要求见快速开始。
• 完整的 HCOMM 运行时符号，包括 HcommEndpointCreate、HcommMemReg、HcommChannelCreate、HcommChannelGetStatus 等。编译成功与运行时符号齐全是两次检查。

先做最简单的检查：

source /usr/local/Ascend/ascend-toolkit/set_env.sh
npu-smi info
command -v bisheng
bisheng -v
cmake --version

自定义安装路径时，使用实际的 set_env.sh。不要同时混用不同 CANN 安装目录的编译器、头文件和动态库。

编译官方示例

在全新目录中取固定源码，方便后续和本文对照：

git clone https://gitcode.com/cann/shmem.git shmem-rm-study
cd shmem-rm-study
git checkout 9ae0bfe55b621185e0f89dd1db748d430e346eb3

source /usr/local/Ascend/ascend-toolkit/set_env.sh
bash scripts/build.sh -examples -soc_type Ascend950

这里 -examples 要求构建示例，-soc_type Ascend950 选择 950 平台。当前普通 Ascend C 构建路径使用 bisheng、-xcce 和 --cce-aicore-arch=dav-c310；工程还负责设备代码链接与库依赖。因此，入门时直接复用官方 CMake，比手写一条漏参数的编译命令更容易定位问题。顶层 CMake 第 180–198 行

离线机器应提前准备依赖。当前项目说明 Ascend 950 构建需要 nlohmann/json，UT 构建涉及 googletest；以固定版本的构建脚本和编译指南为准。出现依赖下载失败时，先排查构建环境，不要直接归因于 RM。

从两卡开始

官方脚本默认拉起 8 个 PE。初次验证可以显式缩为同机两卡：

# all-gather：2 个 PE，使用本机卡 0、1
bash examples/udma_demo/run.sh 0 2 2 127.0.0.1:8899 0 2 0

# 上一轮退出后，再运行 put-with-signal
bash examples/udma_demo/run.sh 1 2 2 127.0.0.1:8899 0 2 0

参数顺序是：

test_type  n_pes  g_npus  ipport  f_pe  local_pes  f_npu
测试类型    总PE数  本机卡数  引导地址  起始PE  本机进程数  起始卡号

PE（Processing Element）是通信参与者编号。该示例一个进程对应一个 PE、一张 NPU。PE 编号与 AIV block 编号是两套东西。

脚本会设置 PROJECT_ROOT、LD_LIBRARY_PATH 和 SHMEM_UID_SESSION_ID。直接执行二进制时不能忘记这些准备。run.sh，第 13–98 行

all-gather 的结果很直观：每个 PE 提供 16 个 int32_t。PE 0 填入 10，PE 1 填入 11，最后两个 PE 都应读到“16 个 10 + 16 个 11”。检查每个进程的退出码和 check transport result success / [SUCCESS] 日志；put-with-signal 还检查来自其他 PE 的 signal 值为 1000。这里描述的是源码定义的预期结果，不是本次运行记录。

跨机时，127.0.0.1 必须换成所有节点可达的 node0 地址，同时安排不重复的全局 PE 范围。这个 TCP 地址用于引导和信息交换，ping 通它不等于 NPU 的 UDMA 路径已经通了。

Ascend C 代码怎样写

Host 准备什么

Host 可以理解为 CPU 侧程序，Device 是 NPU 侧 kernel。AIV 是 NPU 的向量计算核；GM（Global Memory）是设备全局内存，通常落在 HBM 上。先沿着官方 main.cpp 读这条顺序：

1. aclInit、aclrtSetDevice、aclrtCreateStream：初始化运行时，选择设备，创建执行流。
2. test_set_attr：准备 PE 总数、当前 PE、引导地址和堆大小。这是示例辅助函数，不是 SHMEM 公共 API。
3. 设置 ACLSHMEM_DATA_OP_UDMA，调用 aclshmemx_init_attr：库准备通信资源。
4. 所有 PE 按相同顺序、相同大小调用 aclshmem_malloc，得到对称内存。
5. 初始化各自数据，launch kernel，调用 aclrtSynchronizeStream，检查通信异常并校验结果。
6. 完成通信后释放对称内存，再 aclshmem_finalize，销毁流、复位设备、结束 ACL。

“对称内存”不是“所有卡内容一样”。它表示各 PE 的分配布局满足对应关系，库可以根据本地对称指针 + 目标 PE找到远端对应位置。底层会进行地址转换，不能把本地普通指针直接当作远端实际地址。

在两卡例子里，令 gva 为本地对称缓冲区起点：

• gva + 0 对应 PE 0 的 64 字节数据槽。
• gva + 64 对应 PE 1 的 64 字节数据槽。
• PE 0 把自己的第一个槽复制到 PE 1 的第一个槽；PE 1 反向复制第二个槽。

双方都有完整输出缓冲区，传输复制数据，不会自动合并或求和。 数据源要一直保留到相应通信完成；结果使用完、所有远端访问结束后才能释放整个堆对象。

Device 发出什么

核心调用位于 udma_demo_kernel.cpp，第 35–64 行。下面是按这个示例整理的教学 kernel，保留缓冲区初始化、发送、等待和全局会合；删去了 dump 支路，使用固定的两卡数据长度。它是源码改写示意，未在 950 编译验证；首次运行请使用官方原文件。

#include "kernel_operator.h"
#include "shmem.h"

extern "C" [[bisheng::core_ratio(0, 1)]]
__global__ __aicore__ void study_allgather(GM_ADDR gva)
{
    AscendC::TPipe pipe;
    AscendC::TBuf<AscendC::TPosition::VECOUT> scratch;
    constexpr uint32_t scratch_bytes = ACLSHMEM_UDMA_MTE_STAGING_UB_SIZE;
    constexpr uint32_t sync_id = 0;
    constexpr uint32_t bytes_per_pe = 16 * sizeof(int32_t);

    pipe.InitBuffer(scratch, scratch_bytes);
    auto local = scratch.GetWithOffset<uint8_t>(scratch_bytes, 0);
    auto ub = (__ubuf__ uint8_t*)local.GetPhyAddr();
    auto words = reinterpret_cast<__ubuf__ uint64_t*>(ub);
    for (uint32_t k = 0; k < scratch_bytes / sizeof(uint64_t); ++k) {
        words[k] = 0;
    }

    const int me = aclshmem_my_pe();
    const int count = aclshmem_n_pes();
    for (int peer = 0; peer < count; ++peer) {
        if (peer == me) {
            continue;
        }
        aclshmemx_udma_put_nbi<uint8_t>(
            gva + bytes_per_pe * me, // 远端同偏移槽的对称地址表达
            gva + bytes_per_pe * me, // 本地数据源
            ub,                     // 核内 WQE staging scratch
            bytes_per_pe,           // uint8_t 元素数，恰好等于字节数
            peer,
            sync_id);
        aclshmemx_udma_quiet(peer);
    }
    aclshmemx_sync_vec_all();
}

Host 以一个 AIV block 启动它：

study_allgather<<<1, nullptr, stream>>>(gva);

这里假设 Host 已完成上一节所有准备，gva 至少容纳 64 × PE 数量 字节，并在启动前写好本 PE 的数据槽；stream 是 Host 创建的设备流。这个调用本身不会替你执行初始化。

最值得看清的是三个参数：

• dst 与 src 数值表达相同，所属设备不同。 第一个通过 peer 解释为远端目标，第二个是本地源。
• elem_size 是元素数。 这里使用 uint8_t 才能直接传字节数；改成 int32_t 时，64 字节应传 16 个元素。
• ub 暂存 WQE，不暂存整个 64 字节业务数据。 当前示例预留并清零 128 字节，用于容纳完整的 WQE staging block；大数据本体仍由 UDMA 在设备内存间搬运。

最后这点能把你熟悉的 WQE 接回代码：AIV 在核内 UB 准备 WQE，通过 MTE3（核内 UB 向 GM 搬运的流水线）把它搬到 SQ（Send Queue，发送队列），做好流水线同步，再敲 doorbell。底层有真正的调用：

st_dev(cur_head, door_bell_addr, 0);

shmem_device_udma.hpp 第 906–931 行同时给出了 doorbell 更新及 S_MTE3 / MTE3_S 同步。cur_head 是新的队列生产位置，doorbell 告诉硬件“有新任务可取”，不是把业务数据直接写进门铃。

这张图只展开 PE 0 到 PE 1 的一笔传输，用来分开“发请求”“搬数据”和“获知完成”。


自绘时序示意图，依据 SHMEM @9ae0bfe 的 udma_demo_kernel.cpp:35–64 和 shmem_device_udma.hpp:906–931、1182 起的 quiet 实现。反方向的数据传输省略；Host 预先建立资源。图不判定 HCOMM 内部所选的 RM/RC 模式。

要看的是：数据箭头与完成箭头走向不同。 发起方通过完成队列获知传输完成，接收方仍需要约定何时读取目标内存。

完成与接收方读取怎样衔接

nbi 表示非阻塞、隐式完成：调用返回后，传输可能仍在进行。CQE（Completion Queue Entry）是硬件完成队列条目；quiet 在这个实现中会轮询相关完成并推进队列状态。

官方示例采用的闭环是：

各 PE 完成自己的发送提交
→ 各 PE quiet 自己发出的 UDMA 操作
→ 各 PE 参加 sync_vec_all 会合
→ kernel 返回
→ Host 同步流并校验结果

其中 quiet(peer) 等待本 PE 发起的相应传输，不能替接收方完成整个业务协议。示例在每个 PE 都 quiet 后再会合，才让全体进入读取结果阶段。sync_vec_all 也不应脱离前面的 quiet，被当作所有异步搬运的通用完成操作。同步接口说明

需要更细粒度的点对点消费时，可以学习同目录的 udma_put_signal_kernel：把数据写入与远端 signal 更新组合成 aclshmemx_udma_put_signal_nbi，接收方按库支持的 signal/wait 协议等待，再读取数据。官方这个 demo 仍然用了 quiet 和全局会合，并在 Host 检查 signal；不能把它说成已经演示了设备端边收到边消费的完整流水线。 多轮复用缓冲区还需要轮次号或确认协议，避免旧 signal 误触发、生产者提前覆盖下一轮数据。

怎样显式选择 RM

如果现在要回答“RM 的开关究竟在哪”，最直接的证据来自原生 URMA。它属于 Host C 程序，用于理解和验证 RM 配置，不是把 urma_* API 直接搬进 Ascend C kernel。

编译并运行原生示例

在已经安装匹配版本 UMDK 开发头文件、liburma、Provider 与驱动的 Linux 环境中，可按官方 CMake 的依赖关系直接编译示例：

# 在与已安装 UMDK 配套的源码根目录中执行
cc -std=gnu99 -O2 src/urma/examples/urma_sample.c \
   -I/usr/include/ub/umdk/urma -L/usr/lib64 \
   -o urma_sample -lurma -pthread

urma_admin show

这个命令使用固定版本 CMake 的默认安装目录：头文件在 /usr/include/ub/umdk/urma，库在 /usr/lib64；其他发行包或自定义安装时，替换成实际 -I / -L 与运行时库路径。若尚未安装开发环境，应按该版本 UMDK README准备配套组件，不能只复制一个头文件凑编译。安装路径见URMA core CMake。官方示例 CMake的链接依赖就是 urma 和 pthread。

也可使用官方源码构建方式：

cmake -S src -B src/build -DBUILD_ALL=disable -DBUILD_URMA=enable
cmake --build src/build --target urma_sample -j 8
# 生成的示例位于 src/build/urma/examples/urma_sample

取 urma_admin show 实际显示的设备名，不要默认两端都叫 udma2。下面沿用官方 README 中的名字和示例 IP，运行前必须替换成自己的设备名与服务端地址：

# 服务端：RM + CTP
./urma_sample -m 0 -t 1 -d udma2

# 客户端：另一台具有可用 URMA 设备和相通链路的机器
./urma_sample -m 0 -t 1 -d udma2 -i 192.168.100.100

-m 0 是这个示例定义的 RM 选项，-t 1 是 CTP 选项。命令行里的 0 不是 URMA_TM_RM 枚举值；当前 urma_types.h 第 332–336 行中 URMA_TM_RM = 0x1，示例会做转换。不能在自己的资源配置里用 trans_mode = 0 代替 URMA_TM_RM。模式转换，第 149–177 行

原生示例使用 Host 分配、注册的内存，并包含 write/read、消息交互等路径。因此它的成功证明的是该 Host URMA 路径的 RM 能力；不会自动证明 SHMEM 的 NPU 内存与 AIV 直驱路径也正确。

编码时改哪些对象

在这个固定示例中，RM 的关键设置分布在创建和导入阶段，而非临时敲 doorbell 时。下面仅列模式字段与提交字段；完整资源初始化、检查和释放请保留官方示例：

/* 创建本地接收资源、发送资源；其余字段已按示例初始化。 */
jfr_cfg.trans_mode = URMA_TM_RM;
jfs_cfg.trans_mode = URMA_TM_RM;

/* 导入远端端点时，使用匹配的模式与受支持的路径类型。 */
remote_jetty.trans_mode = URMA_TM_RM;
remote_jetty.tp_type = URMA_CTP;

/* 为一次 WRITE 请求选择目标。 */
wr.opcode = URMA_OPC_WRITE;
wr.tjetty = target_jetty;
wr.flag.bs.complete_enable = 1;

变量的含义是：jfr_cfg / jfs_cfg 描述本地收发资源，remote_jetty 描述待导入远端，target_jetty 是 urma_import_jetty 返回的远端句柄，wr 是软件工作请求。真实样例用的目标变量名是 ctx->c.t_jetty。本地资源，第 234–273 行；远端导入，第 481–509 行；WRITE，第 834–912 行

把完整顺序记为：

1. 初始化 URMA、查询设备、创建 Context、完成队列及本地端点。
2. 注册本地内存；通过引导通道交换远端端点 ID、内存描述和所需权限信息。
3. 导入远端 Segment 与 Jetty。RM 路径不执行 RC 那种一对一 urma_bind_jetty。 官方代码只有相应 RC/RS 分支才 bind。
4. 用 SGE（Scatter/Gather Entry）描述源和目标的地址、长度、Segment 句柄，填入 WR（Work Request）。
5. urma_post_jetty_send_wr 提交 WR，Provider 将其编码为 WQE 并通知硬件。
6. 轮询 JFC（完成队列资源），检查 CR（Completion Record）的状态及 user_ctx，按约定通知消费者；请求完成前保持被引用内存有效。
7. 全部访问结束后，按依赖逆序反导入、注销和释放资源。

要发给第二个目标，先为它准备有效的远端导入句柄与内存信息，再把那一条 WR 的目标和地址换过去。只换 wr.tjetty、却保留第一个远端的地址与 Segment，是错误的。

最容易踩的几处坑

这些限制来自当前 SHMEM 头文件和示例，不应扩大成“所有 RM 实现都如此”。

现象或误解	应检查什么
put_nbi 返回后立刻改写源数据	等待对应 quiet 或明确的完成协议；Get 目标同样不能提前读
多个 AIV 同时向同一 PE 使用默认接口	默认队列路径有并发限制；先一个 AIV 跑通，再参考 udma_qp_demo 为不同执行者分配独立 QP 并匹配 qp_quiet
payload 偶尔对、偶尔错	先检查缓冲区生命周期、不同流水线的数据可见性、接收方同步，再查传输顺序
使用 128 字节 scratch 却仍出问题	大小之外还要检查完整初始化、scratch 所有权，以及 sync_id 是否与业务流水线冲突
以为所有接口自动选择同一 WQE 路径	当前带 scratch 的默认重载走 PIPE_MTE3；某些无 scratch 兼容重载走直接路径，签名与语义需逐一核对
把长度都当字节	模板接口通常传元素数；一次 UDMA 请求上限为 256 MiB 字节，大传输要拆分
同一个 signal 被多个发送者写	为来源或消息槽分配独立 signal，并处理跨轮次复用
认为 RM 天然跨 QP 有序	可靠性、目标执行顺序、完成顺序和消费者可见性要分别处理
编译好了却初始化失败	检查 CANN/HCOMM 版本、动态库路径、端点和内存资源、引导与设备链路

并发与长度约束见 shmem_device_udma.h；多 QP 的具体使用方式见 udma_qp_demo/README.md。多 QP 是这份实现提供的后续并发方案，不是看到 RM 三个字就可以无锁抢写同一条 SQ。

高阶 UDMA RMA 接口还有隐式 scratch 配置：当前默认位于核内 UB 的 189 * 1024 偏移，大小 128 字节、事件 ID 为 0，可通过 aclshmemx_set_udma_config 调整。本文 kernel 使用自己通过 TPipe 分配并显式传入的 scratch，不要把两套工作区管理方式混在一起。示例说明

学到什么程度算第一步完成

我的建议是把第一次验收缩成三个能复述、能核查的结果：

1. 讲清定义。 RM 让一个本地端点面向多个远端可靠通信；它与 RMA 操作、UDMA 后端、Ascend C 执行位置是不同维度。
2. 跑通官方两卡程序。 能解释 16 个 10 和 16 个 11 分别从哪里来，知道每个参数、scratch、quiet 和会合各负责什么，并保存版本、命令、退出码与结果日志。
3. 给出模式证据。 原生示例用 -m 0 显式选择 RM；如果任务要求 SHMEM 底层必须为 RM，则补齐实际 HCOMM/Provider 的模式证据。不要用“UDMA demo 成功”替代这个结论。

回到“组 WQE、敲 doorbell”：你已经知道的是发送链条中间的两步。补齐前面的资源与寻址、后面的完成与消费协议，才有一个可靠可用的程序。

继续阅读时可按 URMA Mental Model → SHMEM Symmetric Memory → SHMEM UDMA Programming 的顺序展开。已有文章保留各自的概念主线；本文专门承担 RM 定义与 950 上手路径的入口。