一个 RDMA visibility 问题的讨论

发表于 |

这次在 NVIDIA B300 和 PPU 上做 RDMA visibility 实验，我想建立的对照是：不加额外处理时读错，加入处理后恢复正确。但截至 2026 年 10 月 14 日的实验记录，两组一直都通过。值得整理的是，每次未复现以后改了什么条件，以及查到的 OWNER 保证怎样改变了我对结果的理解。

实验目标：RDMA Write 的 GPU visibility

我关心的是网卡直接写入 GPU memory 之后，GPU kernel 读到的究竟是哪一轮数据。这条绕过接收端 host memory 中转的路径，在 NVIDIA 上由 GPUDirect RDMA 支持。本文的 visibility 是指：在约定的传输完成和同步条件成立后，指定 GPU 消费者能够从目标地址读到本轮写入结果。

NVIDIA 文档给出了一个具体风险场景：已经运行的 kernel 与外部 RDMA Write 并发访问同一块 GPU memory，缺少相应的同步关系，即使网卡已经完成 RDMA Write，kernel 仍会面临旧数据、部分更新或乱序更新的问题。[1] 因而，RDMA Write completion 与 GPU visibility 不能脱离平台保证直接画等号。我想验证的是这个边界，而不是在本轮尚未写完时读到半成品。

候选处理是在写完后，额外做一次 loopback RDMA Read，再等待它完成。这里的 loopback 是让接收端网卡向本机目标发起请求，读的是刚接收数据的 GPU memory。我把这次额外访问称为 readback。实验选择读一个字节到临时 host buffer，目的是插入一笔需要完成确认的访问，不承担整份数据的校验。GPU 通信库 NCCL 的公开 RDMA 后端也采用了这种 readback，源码将请求长度设为 1。[2]

我把不做 readback 记为 OFF，增加 readback 记为 ON。这笔 RDMA Read 使用独立 QP，不能拿同一 QP 内的 ordering guarantee，直接推导它是跨 QP、跨设备的 memory barrier。它是否建立了所需的 visibility，正是要检验的假设。

基础实验：同机 RDMA 传输与设备端校验

起步时先确认网络和 memory registration 能够工作，也在两台 PPU 主机之间用 host memory 跑通过 RDMA Write。这个检查只覆盖跨机网络，没有验证跨机 GPU memory 的 visibility。

GPU memory 测试从同机 loopback 开始，随后拆成独立的发送、接收进程，让完成确认和消费分别发生在通信两端。同机没有绕过网卡，TCP 只传控制通知，数据仍由 RDMA 设备直接写入注册过的 GPU memory。每轮的流程是：

1
2
3
4
5
6
接收端准备好结果区，完成初始化后才允许本轮 RDMA Write
    → 发送端提交 RDMA Write，确认本轮 completion
    → 通过 TCP 通知接收端
    → OFF 直接继续，ON 额外做 readback 并等待完成
    → 接收端 CPU 发起 kernel launch，提交验证 kernel
    → GPU 校验完成，返回结果后才开始下一轮

一轮普通 RDMA Write 在同一条 RC（Reliable Connected）QP 上顺序提交。每次提交给 RDMA 接口的操作请求称为 WR（Work Request），只有批末 WR 设置 signaled，要求网卡为该请求生成 CQE（Completion Queue Entry）。[6] 发送端核对这条 CQE 的成功状态、操作类型和请求编号后才发通知。这个完成条件针对本轮、这条 QP 上的 RDMA Write 批次，不能拿其他 QP 或上一轮的 CQE 来代替。

发送内容同时随轮次和位置变化，kernel 直接读取原本注册给网卡的 GPU memory，逐位置比较。CPU 只取回小块结果。因此，PASS 不是 host 副本正确，而是设备确实执行了读取和校验。等校验结果返回后才进入下一轮，则把“数据还没用完就被覆盖”排除在这个协议之外。

两种模式都不调用 CUDA 提供的显式 flush API，唯一的对照变量是额外 readback。基础版本的 OFF 和 ON 都通过了，于是继续改变 RDMA Write 的条件。

PPU：单连接负载与 WR 拆分

10 月 12 日的 PPU 单连接记录覆盖了下面三种负载。数据量和轮数是按检查目的与测试预算选的，不是从硬件规格推出来的阈值。

每轮数据量 每种配置的轮数 用途
4 KiB 2 基础检查
1 MiB 1000 反复写入和校验
64 MiB 10 扩大单轮访问范围

除了增加数据量，我还改变了 WR 拆分：同一份数据分别用 1 个和 16 个 WR 发送。例如 1 MiB 拆成 16 个 RDMA Write，每个就是 64 KiB。这样改变请求批次，而不同时改变总数据量，才能分清测试增加的到底是什么。

另一个变量是 CU_POINTER_ATTRIBUTE_SYNC_MEMOPS。按 CUDA 的定义，它是作用于单个 allocation 的 pointer attribute。开启后，该区域上的 synchronous memory operation 必须满足接口规定的同步完成条件。[3] 具体例子是，CPU 调用 cuMemcpyDtoD() 复制 GPU memory 后，让网卡读取目的地址。该复制原本允许排入设备队列后就向 CPU 返回，开启这个属性会禁止这里的提前返回，维持“复制完成以后再供网卡读取”的依赖。[1] 它不是缓存开关，也不会等待未来的外部 RDMA Write。

PPU 测试同样比较了这个属性的 0/1 设置，但引用 CUDA 定义不等于确认 PPU 后端具有相同实现。每种负载都覆盖 WR 数 1/16、属性 0/1 和 OFF/ON，因此每行有 2 × 2 × 2 = 8 种配置，全部通过。这些结果没有区分出哪项额外处理是必需的，我也开始检查设备已经提供的保证，而不是继续只从负载大小找原因。

B300：OWNER memory ordering

我查询了本次测试中 B300 的 CU_DEVICE_ATTRIBUTE_GPU_DIRECT_RDMA_WRITES_ORDERING，返回 100，在 CUDA 枚举中对应 OWNER。这个属性描述的是设备对 GPUDirect RDMA writes 原生提供的 memory ordering，范围按哪些 GPU 能一致地读取写入结果来区分：[4]

返回值 名称 原生保证
0 NONE 不提供原生的 GPUDirect RDMA writes ordering
100 OWNER 拥有目标内存的 GPU 能一致地消费 GPUDirect RDMA writes
200 ALL_DEVICES 保证范围扩展到系统内其他 CUDA 设备

把接收卡记为 GPU A：网卡写入 A 的 GPU memory，再由 A 校验，就在 OWNER 范围内。即使源数据来自 GPU B，最终仍由 A 读取 A 的内存，也没有改变这个范围。只有改成由 B 直接读取 A 的目标内存，才需要考虑其他消费者的 visibility。这是消费者范围的例子，不是说这些访问自动具备所需的映射和权限。

CUDA 的显式接口 cuFlushGPUDirectRDMAWrites() 会阻塞到外部 GPUDirect RDMA writes 满足指定 scope 内的 visibility 要求。scope 指需要覆盖的消费者范围，TO_OWNER 只要求内存所属设备读到结果，而不是“清空 GPU 缓存”。当这个 scope 已经被设备报告的 OWNER 覆盖时，文档规定调用是 no-op，可以省略。[5]

这让我重新理解了“不加处理”：在本次 B300 环境中，由目标内存的 owner GPU 消费时，去掉 TO_OWNER flush 并没有撤掉一项额外保证。这里的 no-op 指原生 CUDA API，不是说 ON 里的 RDMA Read 没有执行。后者是否有额外作用，仍然需要对照。

OWNER 也没有替应用确认整轮 RDMA Write 完成，更不保证整块数据原子更新。文档还提醒，不同硬件写入路径之间的 ordering 需要额外协调，不能把设备属性当成跨路径的 memory barrier。[5] 这些是本次 B300 结果的契约依据，不能直接搬过去解释 PPU。

B300：显式 flush 与传输路径对照

另一组独立实验记录在 10 月 13 日。它检查的是，移除显式 CUDA flush，并减少完成通知到消费之间的附加操作后，是否还能读对。kernel 和结果区提前准备、预热，收到 TCP 完成通知就提交校验，中间不再插入整块数据的 readback、清零、分配或显式 CUDA synchronization。

这组还改变了数据来源和设备路径：既有 H2D（host-to-device），即 host memory 经网卡写入 B300 GPU memory，也有 D2D（device-to-device），即两张 B300 的 GPU memory 之间经网卡传输，覆盖同网卡、不同网卡和交换传输方向。这里 H2D、D2D 都指 RDMA 数据路径。记录中的最大场景，单轮有效数据超过 1 GiB，用来扩大单轮访问范围。

结果仍然通过。不过，它没有做额外 readback 的 OFF/ON 对照，而且保留了 SYNC_MEMOPS=1，不能与前面的组合混算。结合 OWNER 再看这组记录，源设备和网卡路径虽然变了，消费者始终是目标内存的 owner GPU，并且仍在收到通知后才发起 kernel launch。换成 D2D 传输，没有自动换成“其他 GPU 读取目标内存”的实验。

并发负载与验证器检查

单连接没有暴露问题，接下来把负载改为多个发送者汇聚到一个接收进程。10 月 14 日的 B300 并发记录覆盖 64 和 128 个发送者，每个发送者每轮传 1 MiB、拆成 16 个 RDMA Write，分别跑 100 和 1000 轮，并比较同步属性开关和 OFF/ON。这些数值是选定的并发规模和重复次数，所有组合仍然通过。

针对 WR 拆分，又在 B300 上补跑四个发送者、每个十轮的短测，每轮仍是 1 MiB，但拆成 256 个 RDMA Write，即每个 4 KiB。这组关闭 SYNC_MEMOPS，分别跑 OFF 和 ON，仍然没有错误。

PPU 上也验证了 64 和 128 个发送者，但每个发送者每轮只传 4 KiB，运行三轮，用于检查并发实现能否在该环境运行。同步属性开关和 OFF/ON 都通过了，但这不是 B300 那组每轮 1 MiB、1000 轮长测的重复。

并发没有改变每块数据的消费条件：每个发送者都有独立 buffer，自己的 RDMA Write 完成后才提交校验。即便出现“GPU 读已经写完的 buffer A，网卡同时写 buffer B”，也不等于已经运行的 kernel 在观察同一个 buffer 的外部更新。日志中的 64 或 128 个并发窗口，从允许写入计数到结果返回，只能证明轮次重叠，不能证明同一瞬间有相同数量的 kernel 在执行，也没有测定具体的传输与 kernel 执行重叠。

验证器另外做过反例检查。在 B300 的 H2D、D2D 测试中，我故意跳过某轮 RDMA Write，仍按新一轮内容校验，设备识别出了旧数据，并把失败传回发送端。并发版也检查了设备校验失败、kernel launch error 和错误 CQE 的传播。这些检查验证了检测链路，没有替代始终缺少的 OFF 自然失败。

实验局限：GPU 消费时序

两种设备上的测试都保留着同一个顺序：RDMA Write 完成，CPU 收到通知，再发起新的 kernel launch。 GPUDirect RDMA 文档要求先确认外部 RDMA Write 相对于提交线程已经完成，再通过 CPU 发起的相关 CUDA work submission 或 synchronization API 建立依赖。[1] 对本次 B300，这与查询到的 OWNER 保证相吻合。所以，去掉显式 flush 不等于消除了消费前的 ordering guarantee。

对 PPU，只能确认当前时序下全部通过。我没有能区分设备写入路径与运行时提交行为的证据，不能照搬 NVIDIA 的机制解释。类似地，记录既没有证明新 kernel 的缓存为空，也没有测出完成通知到实际读取的间隔，不能用“TCP 恰好多等了一会儿”作为原因。

尚未覆盖的一个条件是 persistent kernel：它启动后持续等待后续轮次，而不是每轮由 CPU 重新提交。这样的消费者不经过每轮的新 kernel launch，但这只是需要验证的差异，不是已经确认的故障根因。现有材料没有完整还原原始失败场景，不能断言改成 persistent kernel 就一定复现。

继续实验时，需要同时说清本轮完成通知、数据发布和 buffer 复用的规则。尤其是发布标志与数据的顺序，必须有平台支持的外部 RDMA Write synchronization 依据，不能直接套用 CPU 线程之间的 release/acquire 关系。否则，读到半成品、上一轮残留或下一轮覆盖，都能表现为校验失败，却不一定是在检验同一个 visibility 问题。

这次能得出的结论是：传输和设备校验跑通了，在记录覆盖的条件下，额外 readback 前后都正确。实验没有证明 readback 必需，也没有证明它足以解决原本想抓的错误。要进一步判断它的作用，缺少的是一个完成条件明确、能够失败的 baseline，而不是更多通过记录。

参考资料

实验参数与结果来自本次记录，下面的公开资料用于核对 API 语义，不代替实验结果。

  1. GPUDirect RDMA: Synchronization and Memory Ordering。
  2. NCCL v2.27.7-1 RDMA implementation：loopback readback、请求长度和 ncclIbIflush()。版本用于固定源码依据，不代表实验环境。
  3. CUDA 13.0.1: Unified Addressing。
  4. CUDA 13.0.1: Data Types。
  5. CUDA 13.0.1: Device Management：flush scope 与 no-op 条件。
  6. rdma-core: ibv_post_send：WR 提交与 signaled 设置。