Posted Updated 1-Operator Development2 hours read (About 13820 words)

Ascend 950 RM Target Selection

导言

已经知道组 WQE、敲 doorbell，接下来要弄清：如果同一个本地发送端点要向不同 rank 发送，每条 WQE 的远端目标在哪里设置？ 本文沿 SHMEM 的 pe → remote_mem_info → WQE 追踪源码，再对照 elastic_dispatch.cpp。关键是分清本地队列、远端寻址元数据，以及 put/write 中容易被误读成 Jetty ID 的内存 tokenId。

本文于 2026-09-16 补充 kernel 迁移对照：DeepEP 固定到 19c40e99a622c2778c4c5ce94aa1f6adc33effbc，多 QP 示例核对 SHMEM 73064fa808c8d01bc982f883282d774a44ece09e，并核查 DeepEP 固定的子模块 38a47e5。后附的两卡入门例子及原图仍保留 SHMEM 9ae0bfe 的固定来源。本文只分析 SHMEM 与应用 kernel，不展开 UMDK 底层代码。本次完成了文档、源码和脚本的静态核对，没有在 Ascend 950 上实际编译或运行，也没有测量性能。

先固定问题边界

用户与同事进一步确认，本次关心的 RM/RC 差别是：一个本地 Jetty 能否面向多个远端目标，而不是固定 Jetty 对之间的通信。 下文按这个明确的工程问题分析，不再根据缩写猜需求。所谓“任意目标”，仍以资源建立、远端信息交换及访问权限有效为前提。

更正记录

前版一度将主线改为 Reliability Mechanisms，并撤回端点复用方向；按本次进一步确认，这不符合当前任务。可靠性资料保留在文末作为独立背景。本文也不把下面的通用 UDMA WQE builder 冒充成已经验证开启 RM 的完整实现。

WQE 的远端目标在哪里填

直接看 aclshmemi_udma_fill_sqe_base_ctx()。 它在每次构建 WQE 时，从传入的 remote_mem_info 读取目标字段，而不是等敲 doorbell 时再传目标 rank。完整实现见 shmem_device_udma.hpp:528–555。

每条 WQE 携带什么

下面是函数中目标相关的连续源码摘录。前面的 opcode/flag 构建及后面的函数收尾未列出；dw0、dw1 对应控制字段，dw5 对应归约属性，此处不把它们当作目标 rank 开关。

1
2
3
4
5
6
7
8
9
10
11
uint32_t dw2 = (remote_mem_info->tpn & 0xFFFFFFU) | (1U << 24);
uint32_t dw3 = remote_mem_info->tid & 0xFFFFFU;
uint32_t dw4 = remote_mem_info->rmt_token_value;
uint32_t dw5 = aclshmemi_udma_get_reduce_attr<T, OP_CODE>();
__gm__ uint64_t* rmt_eid = (__gm__ uint64_t*)(remote_mem_info->eid_addr);
sqe64[0] = static_cast<uint64_t>(dw0) | (static_cast<uint64_t>(dw1) << 32);
sqe64[1] = static_cast<uint64_t>(dw2) | (static_cast<uint64_t>(dw3) << 32);
sqe64[2] = rmt_eid[0];
sqe64[3] = rmt_eid[1];
sqe64[4] = static_cast<uint64_t>(dw4) | (static_cast<uint64_t>(dw5) << 32);
sqe64[5] = reinterpret_cast<uint64_t>(remote_addr);

这里 sqe64 是把 WQE 按 64 位整数访问的指针。结合 SQE 结构定义，目标部分对应如下。偏移从 WQE 开头的 0 字节计数。

WQE 位置 字段 SHMEM 填入的来源
字节 8，低 24 位 tp_id remote_mem_info->tpn
字节 12，低 20 位 rmt_jetty_or_seg_id remote_mem_info->tid
字节 16–31 rmt_eid_l/h eid_addr 指向的 128 位远端 EID
字节 32–35 rmt_token_value remote_mem_info->rmt_token_value
字节 40–47 远端地址 当前请求的 remote_addr

因此，目标选择需要一组相互匹配的寻址和访问元数据。target_hint 在当前 Host 路径中为 0，rmt_jetty_type 在 builder 中使用常量；它们不是每次把 rank B 改为 rank C 的位置。

为什么不能把 tid 直接叫 Jetty ID

rmt_jetty_or_seg_id 的名字含有 “or”：这是复用字段。追到 Host 的 FillMemInfo():829–841，可以看到以下赋值。摘录只保留 TP、标识符和访问 token 的来源。

1
2
3
4
const auto& ub = remote_buffer.bufferInfo.rma.protectionInfo.memInfo.ub;
dst_mem.tpn = sq_context.contextInfo.ubJfs.tpID;
dst_mem.tid = ub.tokenId;
dst_mem.rmt_token_value = ub.tokenValue;

对这里的 SHMEM put/write 路径，实际写入的是远端注册内存的 tokenId，不能因为字段名带 Jetty 就将它改成远端接收 Jetty 的编号。 内部头文件也把 tid 注释为来自 segment。远端 EID 则来自 HCOMM 导出的 ubJfs.remoteEID，由 Host 复制到设备 EID 表，再写入 eid_addr。EID 准备代码，第 731–746 行

这不否定 RM 的多目标能力，而是限定当前操作的语义：我们追踪的是 RMA WRITE 的远端内存目标。在所查 put 调用链里，没有发现另一个从 remoteJettyId 取值并写入该字段的步骤；SEND 类操作不能直接套用这一结论。

pe 怎样一路传到 builder

以 direct、多 QP 的 put 为例，pe 是目标 rank，qp_idx 是针对该目标选择的 QP 索引：

  1. 示例 udma_qp_demo_kernel.cpp:52 把 peer 和 qp_idx 传给 aclshmemx_udma_qp_put_nbi()；此示例固定发给环上的下一张卡。
  2. aclshmemx_udma_qp_put_nbi():1472 用 aclshmem_ptr(dst, pe) 得到远端地址，并继续传递 pe、qp_idx。
  3. aclshmemi_udma_post_send():870 在 direct 模式令 slot = pe，按 slot * qp_num + qp_idx 分别查本地 SQ 表和远端元数据表。
  4. builder 将查到的远端元数据写进当前 WQE；提交函数用 st_dev(cur_head, door_bell_addr, 0) 更新本地 SQ doorbell，doorbell 不携带目标 rank。

可把“目标从哪来”这一段压缩成下面的示意代码。它不是完整发送函数：原函数还管理 SQ 信用、WQE 写回、head 和 CQE 计数，不可省去这些步骤直接运行。

1
2
3
4
5
6
// direct 路径的选表关系；对应 post_send() 与两个 get_*()。
uint32_t slot = pe;
auto* local_sq = aclshmemi_udma_get_qp_ctx(udma_info, slot, qp_idx);
auto* remote = aclshmemi_udma_get_mem_info(udma_info, slot, qp_idx);
// 两个表使用相同下标：slot * qp_num + qp_idx。
// fill_post_send_wqe(local_sq, remote, remote_addr, local_addr, ...)

PIPE_MTE3 路径先在 UB scratch 中构造 WQE，再搬到 SQ；它使用相同的远端查表规则及共享 builder，目标语义没有因此改变。MTE3 路径，第 1086–1118 行

这是否已经证明同一 Jetty 跨 rank 发送

尚未。当前代码改变 pe 时，会同时选择另一个本地 SQ 表项和另一个远端元数据表项。 固定 qp_idx=0，分别向 B、C 调 put，并不等于使用同一条本地 SQ。

Host 的 BuildDirectChannels():1236 按 peer × qp_idx 建 channel；CreateChannelsForSlot():1081 设置 remoteEndpoint。这些源码没有直接暴露可确认 RM/RC 的模式选择，既不能证明默认 RC，也不能证明底层多个 channel 共享同一发送队列。

如果要在 950 上验证“B、C 是否共用同一条本地 SQ”，应比较各表项导出的 wqn/jfsID、buf_addr/sqVa、db_addr/dbVa，并核对 Host 的资源模式。确认共享之后，head/tail、SQ 容量和完成计数也必须按实际共享队列统一管理。本地 SQ 导出，第 798–811 行

DeepEP 应当借鉴哪一处

你的 get_sqs_and_cqs_xb_local_ub():1391–1407 把两类信息装进同一个 HcclAiRMAWQ：本地 SQ 地址/门铃，以及远端 EID/TP/token。随后 simt_write_wqe_v4() 全部从 sq 取值。

下面只对比 builder 的元数据依赖；后一段是拟议拆分，不是已经可运行的 RM 补丁。

1
2
3
4
5
6
7
8
9
10
11
12
13
// 现有 DeepEP：本地 SQ 与远端目标捆在 sq 中。
auto sqBaseAddr = sq->sqVA;
auto tpId = sq->tp_id;
auto rmtObjId = sq->rmtObjId;
auto rmtToken = sq->rmtTokenValue;
// EID 同样取自 sq->rmtEid。

// 拟议拆分：localSq 管理实际发送队列；remote 按目标请求选择。
auto sqBaseAddr = localSq->sqVA;
auto tpId = remote->tpn;
auto rmtObjId = remote->tid;          // 保留 WRITE 的 tokenId 语义
auto rmtToken = remote->rmt_token_value;
// EID 改从 remote->eid_addr 读取；remoteAddr 仍来自当前请求。

SHMEM 的 aclshmemi_udma_fill_post_send_wqe():842 已经把 qp_ctx_entry 和 remote_mem_info 作为两个独立参数，这是最值得直接借鉴的设计。若 Host 已提供支持目标集合的 RM 资源，才可以进一步让不同目标请求选择同一 localSq。匹配哪些 TP、EID 和 token 仍由有效资源及路由决定，不能把当前逐 peer 的元数据任意拼到一条队列上。

这次定位到的是改造接口和目标字段来源。完整 RM 改造还要补齐 Host 资源模式、共享队列调度和完成回收，不能只替换一个 rmtObjId。 后文保留原有示例、stage 就绪与接收协议的对照，作为迁移时不能破坏的约束。

代码中的几个术语

  • URMA（Unified Remote Memory Access）：统一远程内存访问的软件接口体系，应用通过它管理资源、提交请求、获取完成结果。
  • RMA（Remote Memory Access）：远端内存访问操作，例如 put/write、get/read；不要只凭相似缩写判断可靠性模式。
  • 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，即核内缓冲区。 两者名称相同，职责不同。

SHMEM 示例里该找什么

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

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

1
attributes.option_attr.data_op_engine_type = ACLSHMEM_DATA_OP_UDMA;

这是 main.cpp 的初始化流程，第 38–96 行中的一个字段，选择数据搬运引擎；它不是可靠性机制的独立配置。

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

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

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

区分后端、协议与端点模式

上述 SHMEM 源码显示了 UDMA 后端与 UBC_CTP 协议选择。这不足以确认 RM/RC 端点模式。 目标字段每条 WQE 都会填写，也不等于当前公共 API 已经跨 peer 复用同一 SQ；模式与共享关系仍需 Host 资源证据。

下面先把 SHMEM 调用与 DeepEP 的裸 WQE 路径对齐，解释哪些部分可以直接复用、哪些还需要配套 Host 资源模式证据。完整两卡入门程序和环境说明保留在后面，方便首次上机时查阅。

从 SHMEM 示例看 kernel 的差别

这里先把能从代码证明的差别讲清楚。检索的 SHMEM master 73064fa 和 DeepEP 所固定的 SHMEM 38a47e5，都没有找到明确标注 RM 与非 RM 的成对 example。 所以这一节比较真实存在的 MTE、UDMA、多 QP kernel；后面再讨论 DeepEP 的目标元数据、完成和接收通知。不能把 RMA 文件名、UDMA 后端或 AiRMAWQ 结构名当作 RM 模式的证据。

普通 MTE 与 UDMA

在 examples/allgather/allgather_kernel.cpp 的小消息路径，数据先从普通 GM 搬到本 PE 的对称区，再由其他 PE 拉到自己的输出。下面按原代码顺序摘取调用，省略分片计算和包装函数；变量名沿用原文件：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
// 1. 本地 input -> 本地对称区。tmp_buff 是业务数据搬运暂存区。
aclshmemx_mte_put_nbi(
    gva_data_gm + gva_offset, input_gm + input_offset,
    tmp_buff, ub_size, num_per_core, my_rank, EVENT_ID0);

// 2. 本地搬运完成、核间同步，然后发布本 PE 的就绪标记。
aclshmem_quiet();
AscendC::SyncAll();
aclshmemx_signal_op(
    gva_sync_gm + flag_offset, magic, ACLSHMEM_SIGNAL_SET, my_rank);
aclshmem_signal_wait_until(
    (__gm__ int32_t*)aclshmem_ptr(gva_sync_gm, x) + flag_offset,
    ACLSHMEM_CMP_EQ, magic);

// 3. 等 PE x 的数据就绪后，拉到本地 output。
aclshmemx_mte_get_nbi(
    output_gm + output_offset, gva_data_gm + gva_offset,
    tmp_buff, ub_size, num_per_core, x, EVENT_ID0);

这里 x 是当前核负责读取的来源 PE，magic 是这轮的同步值。这个例子是 MTE 数据路径，不能直接命名为“URMA 非 RM kernel”。 它的价值是让你看见业务数据的暂存、数据就绪和远端消费是怎样配合的。

UDMA 示例则向硬件 SQ 提交搬运描述，业务 payload 由 UDMA 搬运。带 scratch 的接口虽然也使用 MTE3，但 MTE3 搬的是 WQE，不是整个业务 payload：

1
2
3
4
// 从 udma_demo_kernel.cpp 提取的调用骨架；完整初始化见后附两卡例子。
aclshmemx_udma_put_nbi<uint8_t>(
    symmetric_dst, local_src, wqe_scratch, bytes, peer, sync_id);
aclshmemx_udma_quiet(peer);  // 该默认接口使用 QP 0

因此，不能只看到 PIPE_MTE3 就判断“没有走 UDMA”。应看搬运的对象是业务数据还是通信描述符。

多 QP 示例怎样分工

udma_qp_demo_kernel.cpp提供了与你的 dispatch 更接近的队列使用方式。以下保留 Put 分支的完整核心逻辑，省略 Get、PutSignal 分支及 Host 启动包装：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
const uint32_t qp_idx = AscendC::GetBlockIdx();
const uint32_t qp_num = AscendC::GetBlockNum();
const uint64_t slice_count = elements / qp_num;
const uint64_t slice_offset = qp_idx * slice_count;

AscendC::TPipe pipe;
AscendC::TBuf<AscendC::TPosition::VECOUT> scratch;
pipe.InitBuffer(scratch, ACLSHMEM_UDMA_MTE_STAGING_UB_SIZE);
auto wqe_scratch =
    reinterpret_cast<__ubuf__ uint8_t*>(scratch.Get<uint8_t>().GetPhyAddr());

aclshmemx_udma_qp_put_nbi<uint8_t>(
    symmetric + slice_offset,
    local_buffer + slice_offset,
    wqe_scratch,
    static_cast<uint32_t>(slice_count),
    peer, qp_idx, 0);
aclshmemx_udma_qp_quiet(peer, qp_idx);

读这段代码，抓住三个对应关系：

  1. 一个 AIV 对应一个 qp_idx。 Host 按配置的 QP 数启动相同数量的 block，避免多个核抢同一队列。
  2. 一条 QP 对应一片互不重叠的数据。 Host 保证 elements 能被 QP 数整除，模板是 uint8_t，所以这里元素数就是字节数。
  3. 完成也按 (peer, qp_idx) 等待。 不能发给 QP 1，却只调用使用 QP 0 的默认 aclshmemx_udma_quiet(peer)。

这里的 qp_idx=0 不是所有目标共享的全局 QP 0。SHMEM 直连表按 [peer][qp_idx] 索引；向 PE 1 的 QP 0 与向 PE 2 的 QP 0 提交，访问的是不同的表项。仅此还不能判断底层资源是否存在共享，也不能据此判断事务响应、协议重传或错误处理方式。QP-specific put 与 quiet 实现

对比项 MTE allgather UDMA 多 QP example 当前 DeepEP dispatch
kernel 主要工作 分片、搬运业务数据、发布/等待标记 按 QP 分片并调用 put/get token 打包、专家路由、SIMT 批量写 WQE
核内 UB 用途 业务数据搬运暂存 WQE scratch 打包数据、描述列表、队列元数据；真实 WQE 写入 GM SQ
队列选择 应用不直接操作 UDMA SQ (peer, qp_idx) 先目标 rank，再该 rank 的两条 QP
提交点 MTE 操作及相应同步 普通 put 调用内部提交 显式 st_dev，在 stage 就绪之后
接收就绪 signal/wait 示例结合完成等待和 Host 校验 接收端扫描每个 512 B 块中的标志
能否据此判定 RM 不能 不能 不能；能确认是 UDMA 直驱

elastic_dispatch 现在怎样发送

这一节固定到你给出的 deepep_ccd 分支提交 19c40e9。它已经是 UDMA 发送路径，后续应先核对它的事务完成、错误处理与接收通知契约。 不需要重新实现 token 的专家归属计算。

先确定谁发、发到哪里

Process()把 AIV 分工为打包、远端发送、本地复制、计数等角色：打包核执行 TokenCopyToBuffer() 和 PublishStageDone()，发送核执行 URMASendToken()。当前 rank 的 token 另走 SelfCopyToken()。

buildRemoteWqeDesc() 将 token 到专家的路由转换成一组传输任务。把原函数的关键地址关系单独写出来，就是：

1
2
3
4
5
// 对原函数的等价整理；各变量的含义见下文。
src = sendBufferBase + tokenId * tokenCommSize;
dst = remoteWindowBase[dstRank]
    + (targetExpertOnRank * axisBS + slotInExpert) * tokenCommSize;
bytes = tokenCommSize;
  • tokenId：本 PE 的输入 token 编号；一个 token 可以被路由给多个专家，因此源地址可以重复。
  • targetExpertOnRank：目标 PE 内的专家编号。
  • slotInExpert：该来源 PE 给这个专家的第几个 token。
  • remoteWindowBase：已经包含接收区前缀、本轮 ping-pong 偏移，以及来源 PE 对应专家区的偏移。
  • tokenCommSize：打包后的通信长度，含数据、元信息及分块标志，不能随意替换成原始 hidden dimension 的字节数。

来源：buildRemoteWqeDesc() 第 332–388 行、GetWindAddrByRankId() 第 693 行、目标窗口表初始化第 2230 行。

迁移时先保留这套业务地址布局。 否则接收端即使收到了正确字节，也可能从错误的专家槽取数据。

再选择队列并填写目标

目前 PORT_NUM=1、JETTY_NUM=2。Host 用 aclshmemx_set_qp_num(..., 2) 初始化每个目标 PE 的两条 QP，kernel 读取相同布局。结构体 HcclAiRMAWQ 同时保存 SQ 地址、深度、head、doorbell，以及远端 EID、对象 ID、token、TP 信息。Host 配置

下面摘取真实的索引关系，省略不改变索引的复制语句：

1
2
3
4
5
6
7
8
9
// get_sqs_and_cqs_xb_local_ub：SHMEM GM 表 -> 当前 AIV 的 UB 元数据
const uint32_t qpIdx = portIdx * JETTY_NUM + jettyIdx;
const uint32_t ctxOffset = rankIdx * qpNum + qpIdx;
const uint32_t localOffset =
    (rankIdx - startRankId) * PORT_NUM * JETTY_NUM + qpIdx;

// SQ、CQ、远端内存信息，都从 ctxOffset 对应的条目取得。
// simt_nw_mj_vf：请求按 dstRank 选择该 rank 的队列组。
const uint32_t localSqIdx = rankOffset * jettyPerRank + sqInRank;

sqInRank 在两条 QP 间按照累计 3:1 的请求数比例选择；因为每条请求都是 tokenCommSize 字节，此处请求数比例也就是发送字节比例。这个权重来自现有实现，不能推导成所有 950 的固定带宽比例，也不能把两个 QP 直接理解成两个物理端口。 队列选择，第 485–561 行

simt_write_wqe_v4() 的目标信息来自已经选中的 sq。下面只摘出信息来源；WQE 其余字段没有展示，因此不能把这个片段独立拿去发包：

1
2
3
4
5
6
7
8
9
10
// 当前代码：本地 SQ 和远端目标混放在同一个 sq 对象中。
auto tpId = sq->tp_id & ((1U << 24U) - 1U);
auto rmtObjId = sq->rmtObjId & 0xFFFFFU;
auto rmtEid = (JETTY_POS uint32_t*)(sq->rmtEid);
dw[2] = rmtEid[0] | (static_cast<uint64_t>(rmtEid[1]) << 32);
dw[3] = rmtEid[2] | (static_cast<uint64_t>(rmtEid[3]) << 32);
dw[4] = sq->rmtTokenValue;
dw[5] = remoteAddr;
dw[6] = messageLen;
dw[7] = localAddr;

来源：simt_write_wqe_v4() 第 422–481 行。这说明现有 WQE 的目标信息来自哪里。若要跨 rank 复用本地 SQ，应先解开队列状态与目标元数据的绑定，具体改法见前面的 builder 对照。

最后才敲门铃

现有路径每批最多准备 128 条描述，先按 SQ 计数，原子预留连续位置，再由 SIMT 线程填写 WQE。asc_threadfence()、线程同步和 V→S 事件连接了 WQE 写入与标量提交。流程可概括为：

1
2
3
4
5
6
7
8
9
10
// 控制流示意，不是可独立编译的函数。
buildRemoteWqeDesc(...);  // 列出 src / dst / bytes / dstRank
simt_nw_mj(...);          // 预留 SQ、填 GM WQE，并完成对应发布同步
if (!stageDone) {
    WaitAllStageDone(...); // 等打包核把本地 payload 写好
    stageDone = true;
}
for (each touched SQ) {
    st_dev(newHead, doorbell, 0);
}

这意味着允许提前构造描述符，但不能提前启动对尚未打包完成的数据的读取。发送主循环，第 2250–2304 行

接收端则由 CheckDataArriveWithFlag() 读取各个 512 B 分块在偏移 480 B 处的标志，判断连续到达的 token 数，再由 WaitAndFormatOutput() 拷贝、整理并清理标志。当前算法依赖这些标志与对应 payload 的可见性关系；源码里“检查了标志”并不能独立证明任意新传输模式下这一关系都成立。 可靠性路径改变时，需要验证相应写入/可见性保证，或者改成明确的数据加 signal 协议。接收检查，第 1119 行

从提交到完成的代码对比

先用 SHMEM 的真实 API 建立可核查的发送基线，再检查当前 WQE 的 CQE 与保序配置。这一节只给出完成与通知的对照，不将调用普通 put 等同于启用 RM。

用现有 SHMEM API 替代手写 WQE

适合作为功能对照的是：保留专家槽、打包格式、writer 分工，用 aclshmemx_udma_qp_put_nbi 替代整个裸 WQE 提交路径。以下是迁移片段，放在发送核的标量控制路径；它依赖现有对象和分配好的 128 B scratch，没有在 950 上编译或运行验证。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
// 对比基线：本 AIV 独占自己负责的 (dstRank, qpIdx)。
// 第一次真实提交以前，先等本地打包完成。
WaitAllStageDone(stageDoneLT, stageDoneWorkLT, stageDoneSumLT);

// localWindow：本 PE 的对应对称窗口基址，保留 dataBaseOffset 和本轮偏移。
auto localWindow = GetWindAddrByRankId(epRankId_);

// i 枚举已生成的描述；qpForRequest[i] 沿用现有两 QP 分配结果。
for (uint32_t i = 0; i < wqeCount; ++i) {
    const uint32_t peer = tokenSendWqeDesc->dstRank[i];
    const uint32_t qp = qpForRequest[i];
    const uint64_t src = tokenSendWqeDesc->localAddr[i];

    // 原描述是已转换的远端地址。还原窗口内偏移，再用本地对称表达。
    const uint64_t offset = tokenSendWqeDesc->remoteAddr[i]
        - uint64_t(GetWindAddrByRankId(peer));
    auto symmetricDst = localWindow + offset;

    aclshmemx_udma_qp_put_nbi<uint8_t>(
        symmetricDst, (__gm__ uint8_t*)src, wqeScratch,
        tokenSendWqeDesc->messageLen[i], peer, qp, syncId);

    // 教学基线每条都等：便于限定源生命周期和排查正确性，性能需另行优化。
    aclshmemx_udma_qp_quiet(peer, qp);
}

这里有三个故意写得明确的约束：

  • qpForRequest 是本次迁移需要提供的列表，不是仓库现有变量或新 SHMEM API。 可从现有 wqeSqIdx 对两条队列取模取得，但必须把队列选择从 simt_nw_mj_vf 拆出来，不能先调用原函数填裸 WQE，再补调公共 put。
  • 不能原样传 remoteAddr[i] 给公共 put。 GetWindAddrByRankId(peer) 已调用过 aclshmem_ptr，而公共 put 还会对目标执行一次 aclshmem_ptr(dst, peer)。片段用本地窗口加同偏移重建对称指针，前提是各 PE 的窗口布局、来源区划分和轮次一致。地址转换实现
  • 普通 put 内部就会提交。 原来在 WQE 填好之后的 WaitAllStageDone() 必须提前。若要恢复批量提交，可研究 SHMEM 的 defer/submit 重载，但当前契约要求一批使用相同的 PE 和 QP；不能用它跨 rank 拼一个批次；这项约束本身不定义可靠性模式。批量提交契约

这一步减少了业务 kernel 对 WQE 字段的依赖，但每条都 quiet 会改变原有重叠与吞吐，只适合作为最初的正确性对照。后续应按独占 QP 管理批次和完成，避免把基线性能当成库或可靠性机制的性能上限。

当前 WQE 能看出什么

你的 simt_write_wqe_v4() 固定写入：

1
2
// elastic_dispatch.cpp：WQE Word 0 的 flag 字节。
(0b00100000u << 16)  // flag = 0x20

同版本系列 SHMEM 的 aclshmemi_udma_build_flag() 给出的布局是：

1
2
3
4
5
// SHMEM 实际定义的位域构成，不是协议 RM 开关。
flag = CONFIG.odr
     | (CONFIG.fence << 3)
     | (CONFIG.se << 4)
     | (CONFIG.cqe << 5);

因此按这套布局，0x20 表示 CQE 开启，ODR=NO，fence=0，SE=0。这比“当前有两条 QP”更接近完成与保序的分析入口，但仍不能据此判断具体事务可靠性模式。DeepEP WQE、SHMEM flag 编码

SHMEM 暴露的相关配置可以这样对比：

1
2
3
4
5
6
7
8
// SHMEM 公共头文件中的真实类型及配置值。
// 顺序为 {cqe, se, fence, odr}。
inline constexpr aclshmemx_udma_op_config_t data_ro{
    1U, 0U, 0U, ACLSHMEMX_UDMA_ODR_RO};

// 库内预定义：
// DEFAULT = {1, 0, 0, NO} -> flag 0x20
// SO      = {1, 0, 0, SO} -> flag 0x26

NO/RO/SO 是保序配置，不是“非 RM/RM”的三个取值。 该版本契约规定，后面的 SO 操作等待同一 PE/QP 上前面的 RO/SO 操作完成；前面的 NO 操作不在这项保证内。把最后一个 flag 操作改成 SO、但前面的 payload 仍然使用默认 NO，不能据此推出数据先于 flag 完成。配置契约

对当前业务，更稳妥的核对顺序是：

  1. 保留已有专家槽和每 PE 两 QP 的分工，先确认 Host 创建资源时选择的事务/传输服务及完成策略。
  2. 明确一条 WQE 对应的 CQE 表示什么完成边界，多个事务切片如何汇总，失败状态怎样上报。
  3. 明确接收方用什么条件判断这个 token 完整可读；如果使用独立通知，需要证明 payload 与通知的顺序。
  4. 对齐这些接口后再修改 kernel，不能从白皮书里的抽象机制直接猜测一个 WQE bit。

完成回收与接收通知怎样对照

可以保留原有队列结构，先用 SHMEM 多 QP 示例的公开操作建立可核查的基线：

1
2
3
4
5
6
7
8
9
10
11
12
// 当前发送逻辑的控制流示意；省略描述生成和队列遍历。
fillWqes();
WaitAllStageDone(...);
ringDoorbells();
// 当前 URMASendToken 调用链未见主动 CQ poll；接收端检查分块 flag。

// 对照基线：名称为真实 SHMEM API，参数准备见上一小节。
WaitAllStageDone(...);
aclshmemx_udma_qp_put_nbi<uint8_t>(
    symmetricDst, src, scratch, bytes, peer, qp, syncId);
aclshmemx_udma_qp_quiet(peer, qp);
// 到这里仍需约定接收端何时消费，quiet 本身不替接收方执行协议。

这个对照展示显式完成管理，没有把后一段冒充为“RM 开启后的 kernel”。 公共 API 的完成契约由库负责，裸 WQE 路径则必须自行实现相匹配的记账和检查。当前文件虽然定义了 URMAPollCQ()，但在 URMASendToken() 实际调用链里未见调用；这里需要补齐验证，不能因为接收 flag 成功就忽略发送端 SQ/CQ 的回收。

尤其不要直接在裸 WQE 路径后面补一个库 quiet。SHMEM 的 qp_quiet 按队列的 cqe_cnt 轮询，当前裸写代码没有同步维护这套计数。裸写与公共 API 应统一状态所有权，或使用互不混用的资源。SHMEM 完成计数

队列并发是另外一项约束：atomicAdd(head, count) 只表示预留了槽，不表示前面的生产者已填完 WQE；head % depth 只计算环形地址，不表示还有空闲槽。需要保证 doorbell 发布的是连续、完整、可见的 WQE，并用实际完成结果回收容量。这些规则值得保留，但不能把它们包装成 Reliability Mechanisms 要求跨目标共享 SQ 的证据。

接收协议怎么保留

第一版迁移可以保留 buildRemoteWqeDesc()、SelfCopyToken()、打包格式和专家区布局，把改动集中在已核实的提交、完成和通知层。Host 目前为 transport="udma" 同时启用 MTE 与 UDMA，紧凑元数据仍走 MTE；不要因 payload 可靠性路径改变就顺手关掉 MTE。DeepEP README

保留接收端的前提是新传输仍满足现有分块标志协议。如果没有这项保证，可参考 SHMEM 多 QP example 的真实调用：

1
2
3
4
5
6
7
// 数据和通知由同一个 WRITE_WITH_NOTIFY WQE 承载。
aclshmemx_udma_qp_put_signal_nbi<uint8_t>(
    symmetric + slice_offset, local_buffer + slice_offset,
    static_cast<uint32_t>(slice_count),
    reinterpret_cast<__gm__ uint64_t*>(signal_words) + qp_idx,
    signal, peer, qp_idx, wqe_scratch, 0);
aclshmemx_udma_qp_quiet(peer, qp_idx);

但这是“一个 QP 一片数据、一份 signal”的演示。DeepEP 是多来源、多专家、多个 token 连续消费；接入时要重新定义通知覆盖的槽位/批次、来源和轮次，接收端等待后才能消费，并处理缓冲区复用。不能简单给每条 QP 放一个最终完成标志，就取代所有 token 的到达标志；也不能认为某条 QP 的 signal 覆盖其他 QP 上的数据。

编译与迁移验证顺序

SHMEM 的两卡基础运行见后面的环境与编码说明。多 QP 对照可使用当前固定版本：

1
2
3
4
5
# 在单独的 SHMEM 学习目录中，保留已安装的 CANN 环境。
git checkout 73064fa808c8d01bc982f883282d774a44ece09e
bash scripts/build.sh -examples -soc_type Ascend950
bash examples/udma_qp_demo/run.sh -pes 2 -qp_count 2 -op put -elems 1048576
bash examples/udma_qp_demo/run.sh -pes 2 -qp_count 2 -op put_signal -elems 1048576

要复现你给出的 DeepEP，应保留它的子模块版本：该提交的 3rdparty/shmem 固定为 **38a47e53f37c8814c5b9f4067adfc0be0a3691ba**。本文核对过，它的多 QP example 与上述 master 截面一致；UDMA 实现后者新增两处 CQ 错误路径的 trap()，因此不要把两个安装包当成完全相同。

1
2
3
4
5
6
7
8
git clone -b deepep_ccd https://gitcode.com/ChenDonYY/ascend_deepep.git ascend-deepep-rm-study
cd ascend-deepep-rm-study
git checkout 19c40e99a622c2778c4c5ce94aa1f6adc33effbc
git submodule update --init --recursive
source /usr/local/Ascend/ascend-toolkit/set_env.sh

# 首次构建显式启用 SHMEM 构建。
bash scripts/build.sh -soc_type Ascend950 -s 1

-s 1 来自该版本真实 scripts/build.sh：SHMEM_BUILD_ENABLE 默认是 0。只照 README 运行不带 -s 1 的命令，在还没有 SHMEM 安装目录时可能失败。另需预先准备项目要求的 Python/PyTorch NPU 环境；这个脚本最后会构建并安装扩展。

这些构建参数选择的是 Ascend 950 和 UDMA 可用环境，没有一个已核实的 -enable_rm 开关。 RM 资源怎样创建、是否允许同一发送队列访问多个目标，需要配套 SHMEM/HCOMM 的 Host 契约；不能通过改变 QP 数推导出来。

建议把后续实验分成能逐项归因的四步：

  1. 运行原版。 固定 CANN、驱动、SHMEM、DeepEP 提交，记录原输出、延迟及 QP 数。
  2. 建立公共 API 基线。 保留业务布局，验证 stage 等待、指针转换和完成协议；比较 payload、token 数、专家槽、路由元信息及多轮复用。
  3. 验证完成与接收语义。 核查成功和失败 CQE 的含义、超时与错误处理；检查消费者看到通知时 payload 是否完整，确认源缓冲区能在什么时刻复用。是否需要故障注入及允许的重试方式，依赖实际可靠性接口契约。
  4. 再做并发与性能。 覆盖空专家、自发送、重复路由、不均匀专家负载，累计请求数超过 SQ 深度并执行多轮 ping-pong；确认不覆盖在途数据，再测批量 doorbell、CQ 频率、队列数、尾延迟和吞吐。

在 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 等。编译成功与运行时符号齐全是两次检查。

先做最简单的检查：

1
2
3
4
5
source /usr/local/Ascend/ascend-toolkit/set_env.sh
npu-smi info
command -v bisheng
bisheng -v
cmake --version

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

编译官方示例

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

1
2
3
4
5
6
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。初次验证可以显式缩为同机两卡：

1
2
3
4
5
# 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

参数顺序是：

1
2
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 编译验证；首次运行请使用官方原文件。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
#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 启动它：

1
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。底层有真正的调用：

1
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 UDMA 从 WQE 下发到 quiet 与跨 PE 会合的时序](https://pic.shaojiemike.top/shaojiemike/2026/09/9ba00f74ecdb74eaa6f5e4b00543ff27.png)
自绘时序示意图，依据 SHMEM @9ae0bfe 的 udma_demo_kernel.cpp:35–64 和 shmem_device_udma.hpp:906–931、1182 起的 quiet 实现。反方向的数据传输省略；Host 预先建立资源。图只展示软件提交与等待，不独立证明协议层的事务应答、重传配置或异常恢复语义。

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

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

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

官方示例采用的闭环是：

1
2
3
4
5
各 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 误触发、生产者提前覆盖下一轮数据。

最容易踩的几处坑

这些限制来自当前 SHMEM 头文件和示例，不应扩大成所有可靠性机制的通用规则。

现象或误解 应检查什么
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，并处理跨轮次复用
认为可靠传输天然跨 QP 有序 可靠性、目标执行顺序、完成顺序和消费者可见性要分别处理
编译好了却初始化失败 检查 CANN/HCOMM 版本、动态库路径、端点和内存资源、引导与设备链路

并发与长度约束见 shmem_device_udma.h；多 QP 的具体使用方式见 udma_qp_demo/README.md。多 QP 是这份实现提供的后续并发方案，可靠性机制不会自动解决多个 AIV 并发写同一条 SQ 的问题。

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

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

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

  1. 追清目标字段。 能从 pe 找到 remote_mem_info，说明 EID、TP、tokenId 和远端地址在哪里写进 WQE；知道 WRITE 路径的 tid 不是直接取远端 Jetty ID。
  2. 跑通官方两卡程序。 能解释 16 个 10 和 16 个 11 分别从哪里来，知道每个参数、scratch、quiet 和会合各负责什么，并保存版本、命令、退出码与结果日志。
  3. 能解释迁移边界。 区分“WQE 携带远端信息”和“同一 SQ 支持多个目标”；能指出 DeepEP 需要拆开的本地队列状态和远端元数据，并说明 Host 模式、共享 head/CQ 与接收通知还需要怎样验证。

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

继续阅读时可按 URMA Mental Model → SHMEM Symmetric Memory → SHMEM UDMA Programming 的顺序展开。已有文章保留各自的概念主线；本文连接 WQE 目标选择、950 上手路径与 DeepEP kernel 的队列拆分、完成和通知审查。

补充：可靠性与目标选择分别负责什么

这部分回答的是“发出后怎样确认完成”，与前面“同一发送端点能否选择多个目标”是两个维度。保留它，是因为更换队列映射不能破坏 dispatch 的完成和消费协议。

本文核对的是华为《灵衢基础规范》2.0（封面发布日期 2025 年 9 月）第 6.3–6.4、7.3.1 节，以及《基于灵衢的超节点参考架构白皮书》第 4.2.6 节。来源入口：灵衢官方文档；华为原文公开副本：基础规范 2.0、超节点参考架构白皮书。不把这些层次直接映射为 RM/RC 的模式开关。

先想一个更贴近 dispatch 的场景：卡 0 向卡 1 写一个 token。AIV 已经组好 WQE、敲了 doorbell，但这时还不能直接让卡 1 开始计算。需要继续回答三个问题：数据是否送达、远端事务是否执行完成、接收 kernel 何时可以读取。

传输、事务和应用各负责什么

《基础规范》2.0 的第 6.3 节区分了传输层的 RTP、CTP、UTP；第 7.3.1 节单独定义事务可靠性。它们不能与上层 CQE 或应用 ready flag 混为一谈。

层次 需要解决的问题 已核实的机制 对 kernel 的意义
传输层 包是否正确送达 RTP 使用 PSN、传输应答及重传；CTP 联合下层提供可靠传输 不应自行在 AIV 中仿造协议重传
事务层 远端操作是否正确执行 可靠事务通常返回 TAACK、Read Response 或 Atomic Response；有明确例外 “发出请求”和“远端执行完成”是不同阶段
本地完成队列 发送软件怎样获知完成或错误 SHMEM 提交时维护完成计数，quiet 轮询 CQE 必须核对完成状态及其覆盖范围
接收应用 哪些 token 可以交给计算核 应用定义的 flag、signal 或会合协议 本地 CQE 不会自动变成接收方的消费通知

CTP 不能简单理解成“没有可靠性”。 规范第 6.3.2 节给出的实施例把可靠传输卸载给下层：CTP 本身不执行端到端传输重传，利用链路层点对点重传，在合适的链路条件下为事务层提供可靠服务。相对地，RTP 的第 6.4 节规定了端到端 PSN、应答和重传机制。这两条路线有不同开销和适用边界，不能只凭“可靠性机制”四个字就把现有 CTP 改成 RTP。

TAACK 也不能直接等同于 CQE。 前者是事务响应，后者是软件可观察的完成队列条目。规范第 7.3.1 节还规定了响应聚合、借传输应答携带事务状态，以及特定 CTP + ROL 场景省略正常 TAACK 的例外。必须由配套实现解释这些协议行为最终怎样形成 CQE；不是把 cqe=1 就能独立证明某种可靠性模式已开启。

优势与成本要对应具体机制

可靠性机制的价值，是让错误有检测、响应或恢复路径，并给软件明确的完成边界；它不能保证所有故障都能自动恢复。采用较强的端到端跟踪和应答机制，通常需要相应的状态、响应流量和等待；利用下层可靠性则能减少一部分上层开销，但依赖适用拓扑和错误模型。

对 dispatch 最实用的收益，是让“数据已经可以用”有明确依据。可靠传输不会自动带来跨请求顺序，也不会自动保证接收方已消费数据。 编程时仍需分别处理源缓冲区寿命、数据与通知的先后关系，以及多轮复用。

协议中的自动重传还必须考虑操作语义：规范第 7.3.1 节特别讨论 Atomic 响应失败时盲目重发请求可能导致重复加法。因此，kernel 遇到超时后不能不分情况重发整批 token 或原子计数操作。

Ascend 950 RM Target Selection

http://icarus.shaojiemike.top/2026/09/15/Work/Artificial Intelligence/AIInfra/Ascend-950-URMA-RM/

Author

Shaojie Tan

Posted on

2026-09-15

Updated on

2026-10-30

Licensed under

#

Like this article? Support the author with

Follow

Categories

Subscribe for updates

follow.it

Links

Recents

Archives

Tags

0/1 Adam1
1-bit Adam1
1-bit LAMB1
3FS1
4+1 View Model1
5G1
64bits vs 32bits1
A31
A51
ACLSHMEM1
ADR1
ADXL1
AGENTX2
AGRS1
AI10
AI AGENT1
AI Agent1
AI CODING1
AI Chip1
AI Cluster1
×