RDNA3 · RCCL · ROCm · measured on 2× RX 7900 XT

ROCr refuses any kernel that declares a hostcall buffer when the root port advertises no AtomicOps,
which is every RCCL collective, and one line of VM configuration puts it back

Traced to hidden_hostcall_buffer on two Radeons; checked against twelve eliminated hypotheses, four shipped RCCL builds, and one line of VM configuration

Two Radeons run --tensor-parallel-size 2 and fail with an error that sends everyone to the wrong layer: the operation cannot be performed in the present state. Twelve of thirteen hypotheses were eliminated along the way. For virtual machines, the eventual fix required one line of configuration rather than a rebuilt library.

2× RX 7900 XT · gfx1100 ROCm 7.14 · RCCL 2.30.4 Proxmox VFIO guest · X399 / Threadripper 1950X filed as ROCm#6520 reported 2026-07-26

Abstract

On a machine without PCIe AtomicOps, every RCCL collective fails at dispatch with hipErrorIllegalState. The cause sits four layers below RCCL. Starting in ROCm 7.2, shipped device kernels declare a hostcall buffer. ROCr allocates this buffer only when the PCIe path provides AtomicOps, refusing dispatch before any collective can execute.

The causal chain reduces to fifty-seven lines of HIP requiring neither RCCL, PyTorch nor vLLM: two kernels share an identical launch path, but the second calls device printf. The plain kernel passes; the one requesting a hostcall fails with the production error.

Twelve competing explanations were tested and eliminated. Two were published by this project before being disproved: first, that RCCL had suffered a version regression; second, that the ROCm runtime was gating kernel submission in guests.

This repository first verified a workaround: rebuilding RCCL with NDEBUG reaching the device compiler. The simpler fix came later. On a passthrough guest, the missing capability is usually self-inflicted by one line of VM configuration. QEMU advertises AtomicOp completer support automatically, declining only when the card is handed over as a multifunction device — the default in Proxmox.

1The failure, and where it sends you

Two GPUs, tensor parallelism, and this at engine start:

AMD_LOG_LEVEL=4 · the failure window, identical on both ranks
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

Because the error surfaces inside RCCL's enqueue.cc, investigators typically debug RCCL, then ROCm, and finally the kernel. Yet the line naming the actual cause sits immediately above the failure. It mentions neither the library nor the collective, and it appears only at log level 4.

2The causal chain

Five links form the chain. Each was established by evidence independent of the link below it.

Figure 1 · from a PCIe capability bit to a refused dispatch derived from root-cause.md
click a link to see the evidence that established it
Figure 1. The chain runs from a capability bit on the root port to a kernel launch that never happens. Two bits are easy to conflate and the distinction decides who is affected: completion is required on the root port above the device, routing only on switch ports in between. Extracted from docs/root-cause.md §1, so this figure cannot drift from the document behind it.

3Fifty-seven lines that take RCCL out of the picture

Link 4 carries the argument because it isolates the failure from all surrounding software. Two kernels share an identical launch configuration; the only distinction is that the second calls device printf, which demands a hostcall buffer.

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

Fifty-seven lines of HIP on one GPU reproduce the identical failure string in half a second. That shifts the issue from a library version story into a platform capability constraint: the failure has nothing to do with collectives. It reproduces for any kernel requesting a hostcall, and RCCL device kernels require one because they contain device-side assert().

4Why "downgrade the library" works, and what it actually shows

The community workaround is to install the ROCm 7.1.1 build of librccl. That works, and it is evidence for this root cause rather than against it. Counting hostcall declarations across shipped device images explains why:

Figure 2 · hostcall declarations in shipped RCCL device images llvm-readelf output not committed
Figure 2. llvm-readelf --notes over each shipped librccl, counting hidden_hostcall_buffer in the ncclDevKernel_Generic* metadata. The platform is identical across all four rows. What changes is whether the library's kernels demand a capability the machine does not have. The raw readelf output is not in the repository; these four rows are transcribed from the write-up.

The newer RCCL build is not broken in general. It simply demands a platform capability these machines never possessed. This reading unifies virtualized reports with bare-metal incidents: if those bare-metal platforms routed AtomicOps, the hostcall requirement would be satisfied and downgrading would not help. The fact that downgrading succeeds is direct evidence that they lack atomics as well.

The version correlation is real and the causation is not Building 2.27.7-b38 from source on the current toolchain fails identically, because that build does not pass NDEBUG to the device stage either. The governing variable is hostcall, not the version number. This is the first of two conclusions this project published before disproving.

5What was ruled out

Thirteen hypotheses were investigated. Twelve proved wrong. We record them so others do not have to repeat the work, and because two of the mistaken hypotheses were our own.

Figure 3 · hypotheses tested derived from root-cause.md
click any card for how it was tested, eliminated or confirmed
Figure 3. Each was tested rather than argued away. Several are eliminated by reading source rather than by measurement — iommu=pt is a string check that warns and does not gate, and both ROCm generations use an ordinary DisableTiming event — which is cheaper and just as conclusive. Extracted from docs/root-cause.md §3.
The second thing we got wrong "ROCm ≥ 7.13's HIP runtime gates kernel submission in VFIO guests" was published here and is false. In diagnose/hipgate.cpp, the exact hipExtLaunchKernel invocation — issued with a completion event, 4 KB kernargs, cross-stream waits, and fine-grained host memory — passes on 7.14. Kernel submission operates normally; only kernels carrying a hostcall buffer are refused.

6The fix that mattered was one line of VM configuration

The remedy verified first was a rebuild: compiling RCCL with NDEBUG reaching the device pass removes the assert() statements that trigger hostcall dependencies. That works on RCCL 2.27.7, where it was verified end to end, and initially appeared complete. However, it is not sufficient on 2.30.4, where the declaration survives a build with no hostcall-calling instructions left in it — §8.

For virtual machines, a rebuild is the wrong first move. QEMU has supported automatic AtomicOp completer advertising on an emulated root port via vfio_pci_enable_rp_atomics() since v8.1.0, declining under a short list of conditions. Primarily, it gives up when the device is passed as multifunction. Proxmox's default hostpciX configuration passes the GPU alongside its companion audio function, taking exactly that branch.

Figure 4 · one line of VM configuration, everything else unchanged check with lspci
Figure 4. Measured on this machine on 2026-08-23, same session, same two cards, same container image, only that line changed. Reverting the line reproduced the failure, so this is not a side effect of reseating the cards. No rebuilt RCCL was involved in the working column: the stock library shipped in the ROCm 7.14 container was used unmodified. The test is all_reduce and all_gather_into_tensor checked elementwise against constructed truth, twelve cases across three dtypes and two sizes.

7Then who is the rebuild for

Rebuilding RCCL remains necessary for systems whose hardware genuinely cannot deliver AtomicOps:

The underlying mechanism is unchanged across these cases, and the NDEBUG rebuild resolves it. What changes is diagnostic ordering: inside a VM, check hypervisor configuration before rebuilding software.

One command decides which case you are in dmesg | grep "PCIE atomic" — output indicates missing atomics; silence is inconclusive rather than reassuring, so verify with lspci. Grep that exact phrase; a bare grep -i atomic also matches unrelated DMA: preallocated ... pool for atomic allocations lines from early boot. If you are affected inside a guest, inspect lspci -vv on the root port above the card before building anything.

8What is not established

Whether the audio function can be kept another way

Passing the GPU as a single function leaves its companion audio function behind. Whether the two can be split across separate emulated root ports so the guest keeps audio alongside atomics was not tested here.

What AMD intends

Whether the hostcall dependency is retained upstream remains an open question in ROCm#6520. On RCCL 2.30.4 the metadata declaration persists even when NDEBUG removes all __ockl_* symbols, so stripping device-side assert() calls is no longer sufficient on current releases.

One machine for the VM half

The configuration A/B reflects one host, one QEMU version, and one motherboard. While the QEMU logic was verified in source, variations across other motherboards were not measured here.

9What has changed since

2026-08-23 · this machine now has atomics The guest was reconfigured to single-function passthrough and has advertised completer support since. Stock unpatched RCCL 2.30.4 now runs 135 of 135 cold communicator initialisations correctly here, so the box that was this repository's counterexample on the atomics axis no longer is one.
2026-08-28 · the severity claim was narrowed This report originally claimed every QEMU root port declines completer support regardless of what sits behind it, concluding that any VFIO passthrough guest is affected, virtualized Instinct included. That conclusion is incorrect for single-function passthrough. Both the upstream issue description and the comment it cites were corrected: a passthrough Instinct is affected when passed as multifunction, not when passed as a single function.
2026-09-05 · the line count was wrong This report described the reproducer as thirty lines. diagnose/hipgate3.cpp is 57 lines (54 non-blank, 39 excluding comments); no reading of the file gives thirty. The text now says fifty-seven, which is what wc -l returns.

Reproduce it yourself. The fifty-seven-line test is diagnose/hipgate3.cpp, requiring neither RCCL nor PyTorch. The platform verification script is diagnose/check-platform.sh. The full evidence chain and elimination matrix are in docs/root-cause.md; the VM analysis is in docs/vfio-atomics.md.

Reported upstream as ROCm#6520, where @adderek independently reproduced the failure on bare metal with IOMMU disabled, and @bkvargyas contributed the switch-versus-atomics trade-off from the other side. Related: ROCm#6074, vllm#38587.