RDNA3 · RCCL · ROCm · 实测于 2× RX 7900 XT

ROCr 在根端口未广告 AtomicOps 时拒绝任何声明 hostcall buffer 的内核，
所有 RCCL 集合通信皆在此列；而一行虚拟机配置就能把这条广告加回来

在两张 Radeon 上溯源至 hidden_hostcall_buffer；经十二个排除的假设、四版 RCCL 发行构建与一行虚拟机配置对照验证

两张 Radeon 启用 --tensor-parallel-size 2 启动运行，随即遭遇了一句把所有人引向错误排查层次的报错：the operation cannot be performed in the present state。调查过程中逐一排除了十三个假设中的十二个。而对于虚拟机，最终的修复只需一行配置，根本无须重新编译库。

2× RX 7900 XT · gfx1100 ROCm 7.14 · RCCL 2.30.4 Proxmox VFIO 直通 · X399 / Threadripper 1950X 已提为 ROCm#6520 上报于 2026-07-26

摘要

在缺失 PCIe AtomicOps 的机器上，所有 RCCL 集合通信都会在 dispatch 阶段报错 hipErrorIllegalState。故障根因位于 RCCL 下方第四层。自 ROCm 7.2 起，官方发布的设备端内核会声明一个 hostcall buffer。ROCr 仅在 PCIe 链路具备 AtomicOps 支持时才分配该缓冲区，否则便在任何集合通信执行前直接拒绝 dispatch。

整条因果链可精简为 五十七行 HIP 代码，完全无须依赖 RCCL、PyTorch 或 vLLM：两个内核共享完全相同的启动路径，区别仅在于第二个内核调用了设备端 printf。普通内核顺利执行，而请求 hostcall 的内核则失败，抛出的正是生产环境下的那句报错。

我们先后测试并排除了十二个竞争性解释。其中两个曾由本项目公开发布，随后才被推翻：其一是误认为 RCCL 出现了版本回退，其二是误认为 ROCm 运行时在客户机内拦截了内核提交。

本仓库最初验证了一种临时变通方案：在编译设备端代码时传入 NDEBUG 并重建 RCCL。更简单的修法是后来才找到的。在直通客户机中，这项缺失的能力通常是 一行虚拟机配置自身所致。QEMU 原本会自动通告 AtomicOp completer 支持，唯独在显卡以多功能设备直通时予以关闭，而这恰恰是 Proxmox 的默认配置。

1故障现象，以及它把人引向哪里

两张 GPU、张量并行，引擎启动时:

AMD_LOG_LEVEL=4 · 失败窗口，两个 rank 完全一致
rocvirtual.cpp:4151  ShaderName : ncclDevKernel_Generic_4(ncclDevKernelArgsStorage<4096ul>)
rocvirtual.cpp:4208  Pcie atomics not enabled, hostcall not supported
rocvirtual.cpp:4636  AQL dispatch failed!
hip_module.cpp:605   hipModuleLaunchKernel: Returned hipErrorIllegalState

报错浮现在 RCCL 的 enqueue.cc 中，因此排查者通常会顺着调试 RCCL、继而调试 ROCm，最后再去翻操作系统内核。然而，直接道明真正诱因的输出恰好紧挨在失败行的上方。那一行对该通信库与集合通信只字未提，且只有把日志级别设为 4 时才会现身。

2因果链

整条因果链由五个环节扣合而成。每一环所依据的证据，均独立于其下游环节。

图 1 · 从一个 PCIe 能力位到一次被拒的 dispatch 由 root-cause.md 提取
点击任一环节查看确立它的证据
图 1. 这条链从根端口上的一个能力位，一路走到一次根本没有发生的 kernel 启动。 有两个位极易混淆，而这个区分决定了谁会中招：completion 要求在设备上方的根端口上， routing 只要求在中间的交换端口上。本图由 docs/root-cause.md §1 提取， 所以它不可能与背后的文档漂移。

3五十七行代码，把 RCCL 移出排查视线

第 4 环是整条论证的关键支柱，因为它彻底将该故障与周遭的所有软件隔离开来。两个内核共享完全相同的启动路径，唯一的差别在于第二个内核调用了设备端 printf，因而需要一个 hostcall 缓冲区。

diagnose/hipgate3.cpp · 无 RCCL、无 PyTorch、无 vLLM
--- dev0 ---
[plain    ] launch:no error   sync:no error
[hostcall ] launch:no error   sync:the operation cannot be performed in the present state

仅凭单张 GPU 上的五十七行 HIP 代码，便可在半秒内完整复现完全相同的故障。这一步将问题从「通信库版本缺陷」的叙事彻底转向了「平台硬件能力限制」：该故障与集合通信毫无干系。任何请求 hostcall 的内核都会触发此错误；而 RCCL 的设备端内核之所以需要 hostcall，纯粹是因为代码中包含了设备端 assert()。

4"降级那个库"为什么有效，以及它真正说明了什么

社区给出的临时解决方案是安装 ROCm 7.1.1 构建的 librccl。此方法确实凑效，而且它恰恰为此处论证的根本原因提供了有力佐证，而非反证。只需统计官方发布的设备镜像中包含多少 hostcall 声明，原因便一目了然：

图 2 · 各版本 RCCL 设备镜像中的 hostcall 声明数 llvm-readelf 输出未入库
图 2. 对每个发行版的 librccl 跑 llvm-readelf --notes， 统计 ncclDevKernel_Generic* 元数据里的 hidden_hostcall_buffer。 四行的平台完全相同。变的是库的 kernel 是否要求一个这台机器没有的能力。 这份 readelf 原始输出不在仓库里；这四行是从写作稿转录的。

由此可见，新版 RCCL 并不是普遍地坏了，它只是要求一项这批机器自始至终都不具备的平台能力。这一解释把虚拟化环境下的报告与裸机上的事故统一了起来：倘若那些裸机平台确实支持 AtomicOps 路由，hostcall 的先决条件本该得到满足，降级也就无从生效。降级库版本之所以奏效，恰恰直接证明了那些机器同样缺失 atomics 支持。

版本存在表面相关性，但不存在因果联系 使用当前工具链从源码编译 2.27.7-b38 同样会失败，原因在于该次构建同样未将 NDEBUG 传递至设备端编译阶段。起决定作用的控制变量始终是 hostcall，而非版本号本身。这是本项目此前公开发表、随后又自行推翻的两个结论中的第一个。

5被排除的假设

排查过程中我们先后调查了十三个假设，其中十二个被证实为错误推断。将它们悉数记录于此，既是为了避免后来者做无谓的重复劳动，也是因为其中两个错误论断正是我们自身早先做出的。

图 3 · 测试过的假设 由 root-cause.md 提取
点击任一条查看它是怎么测的，排除还是确认
图 3. 每一条都是测掉的，不是辩掉的。其中若干条是靠读源码而不是靠测量排除的 —— iommu=pt 只是一个字符串检查，会警告但不会拦截；两代 ROCm 用的都是普通的 DisableTiming 事件 —— 这样更便宜，结论一样确定。 由 docs/root-cause.md §3 提取。
我们搞错的第二个结论 「ROCm ≥ 7.13 的 HIP 运行时会在 VFIO 客户机中阻断内核提交」这一结论曾由本站发表，但已被证明不实。在 diagnose/hipgate.cpp 中，原样调用附带完成事件、4 KB 内核参数（kernarg）、跨流等待及细粒度主机内存的 hipExtLaunchKernel 时，在 7.14 下能够完全顺利通过。内核提交运转一切正常，遭到拒绝的唯独是那些带有 hostcall 缓冲区的内核。

6真正要紧的修复是一行虚拟机配置

最初得到验证的解决方案是重新编译：在编译 RCCL 设备端代码时传入 NDEBUG，以消除引入 hostcall 的 assert() 语句。这一方案在 RCCL 2.27.7 上切实有效（端到端验证即基于该版本完成），并且一度被视作全部答案。然而在 2.30.4 上，该方法已无法奏效：即便构建出的镜像中已不再包含任何调用 hostcall 的指令，该元数据声明依然会残留下来（详见第 8 节）。

然而对于虚拟机而言，重新编译走错了第一步。自 v8.1.0 起，QEMU 便已支持通过 vfio_pci_enable_rp_atomics() 在模拟根端口上自动通告 AtomicOp completer 支持，仅在极少数受限条件下才会放弃。其中最关键的限制是：一旦设备以多功能（multifunction）模式直通，该特性便被关闭。Proxmox 默认的 hostpciX 配置正是将显卡与其配套音频功能捆绑直通，恰好触发了这一分支。

图 4 · 一行虚拟机配置，其余一切不变 用 lspci 自行核对
图 4. 2026-08-23 在这台机器上实测，同一次会话、同样两张卡、同一个容器镜像，只改了那一行。 把这行改回去，故障复现，所以这不是重新插拔显卡的副作用。右列没有用任何重建过的 RCCL： 用的是 ROCm 7.14 容器里自带的库，未加修改。测试是 all_reduce 和 all_gather_into_tensor，逐元素对照构造出来的真值，三种 dtype 两种尺寸共十二个用例。

7那么重建是给谁用的

对于硬件本身确实无法提供 AtomicOps 的系统而言，重新编译 RCCL 依旧不可或缺：

在上述情形中，底层机理并无二致，通过 NDEBUG 重建即可妥善解决。真正改变的是排查诊断的顺序：在虚拟机内部，应当在着手重编软件之前，先行检查虚拟机监视器的配置。

一条命令即可判定所处的状况 dmesg | grep "PCIE atomic" —— 只要有输出即表明缺失 atomics；而无输出并不意味着支持完备，仅代表未明确报错，因而须进一步借助 lspci 进行核实。请务必精确 grep 该完整短语；若仅使用宽泛的 grep -i atomic，则会混入系统引导初期打印的无关信息 DMA: preallocated ... pool for atomic allocations。如果在客户机内受到此问题困扰，在编译任何软件之前，请先检查显卡上方根端口的 lspci -vv。

8没有被确立的部分

音频功能能不能用别的方式保住

将显卡以单功能形式直通，会使其配套的音频功能被阻隔在客户机之外。能否将二者分别挂载到各自独立的模拟根端口上、从而让客户机在享用 atomics 的同时保留音频设备，此方案在本实验中未做验证。

AMD 打算怎么办

上游未来是否仍保留对 hostcall 的依赖，在 ROCm#6520 中仍属悬而未决的问题。在 RCCL 2.30.4 上，即便 NDEBUG 清除了全部 __ockl_* 符号，该元数据声明依然顽固残留；因此在当前的新版本上，单纯剔除设备端的 assert() 已不足以解决问题。

虚拟机那一半只有一台机器

本研究中的配置 A/B 对照仅基于单台宿主机、单一 QEMU 版本与单一主板。尽管 QEMU 的分支逻辑已在源码层面得到核实，但在其他主板上的具体表现，这台机器上并未做实测。

9此后发生的变化

2026-08-23 · 这台机器现在有 atomics 了 客户机已调整为单功能直通模式，此后持续正常通告 completer 支持。未经修改的原版 RCCL 2.30.4 目前在各类测试中，135 次冷启动通信初始化全部顺利通过；因此这台曾因缺失 atomics 而充当反面典型的主机，已不再构成反例。
2026-08-28 · 严重性主张已收窄 本报告最初曾断言：无论下游挂接何种设备，所有 QEMU 根端口均不会通告 completer 支持，并据此得出所有 VFIO 直通客户机（包括虚拟化运行的 Instinct）都会受累的推论。对于单功能直通而言，该结论并不成立。上游 issue 正文及其引用的相关评论均已就地更正；现已明确查实：直通环境下的 Instinct 仅在以多功能形式挂接时才会出现此故障，若以单功能形式直通则不受影响。
2026-09-05 · 行数报错了 本报告此前把复现代码称作三十行。diagnose/hipgate3.cpp 实为 57 行（非空 54 行，除去注释 39 行）， 没有任何一种数法得出三十。正文已改为五十七行，即 wc -l 的结果。

自行复现。用于判别的五十七行独立测试代码位于 diagnose/hipgate3.cpp， 完全无须依赖 RCCL 或 PyTorch。平台环境核对脚本见 diagnose/check-platform.sh。 完整的排查证据链与假设排除矩阵参见 docs/root-cause.md； 虚拟机部分的专门分析参见 docs/vfio-atomics.md。

本问题已作为 ROCm#6520 提交至上游， @adderek 在禁用 IOMMU 的裸机环境下独立完成了复现，@bkvargyas 则从反向补充了「中间交换芯片与 atomics 无法兼得」的制约机制。 相关问题：ROCm#6074、 vllm#38587。