Repository navigation
Eval bug: 2x MI50 tensor split aborts in RCCL on the first AllReduce once peer access works, no fallback to butterfly #7
Description
Name and Version
mx-llama.cpp tag b10811, run from the Docker image mixa3607/llama.cpp-gfx906:b10811-rocm-7.14-mxxm (ROCm 7.14, RCCL 2.30.4). llama-server --version inside the image prints version: 0.3.0-dev (build 1, commit 4e6a72b).
Host: Debian 13 with the Ubuntu mainline 6.12.98 kernel. I moved to that kernel because the stock Debian 6.12.94 kernel is built without CONFIG_PCI_P2PDMA and CONFIG_HSA_AMD_P2P, so peer access between the two cards never comes up on it.
Operating systems
Linux
GGML backends
HIP
Hardware
Two AMD MI50 32 GB (gfx906) on a B650E board with a Ryzen 5 9600X. Both cards run at PCIe Gen4 x16. There is no bridge or NVLink style link, just the PCIe bus.
Models
Qwen3.8-27B Q8_0 GGUF with the MTP draft head (--spec-type draft-mtp). The problem is not model specific as far as I can tell. It happens on the first request of any size.
Problem description & steps to reproduce
I run the model split across both cards with -sm tensor. On the stock kernel, where peer access is not available, RCCL init fails, the fork logs that it is falling back to the butterfly path, and everything works at about 35 tokens per second.
On the 6.12.98 kernel, peer access works. RCCL init now succeeds. The log shows Connected all rings and Connected all trees. The first prompt-sized AllReduce then fails inside RCCL and the server aborts. Every request does this, including a one-line prompt. The server never answers.
The RCCL warning is enqueue.cc:2061 NCCL WARN HIP failure 'the operation cannot be performed in the present state'. That is hipErrorIllegalState, which HIP returns when you enqueue work while a graph capture is active on the stream. The abort lands in ggml_backend_cuda_comm_allreduce_nccl at ggml-cuda.cu:1366, which is the ncclGroupEnd() of the large-tensor BF16 branch. So the RCCL collective is being issued while the fork's HIP graph is being captured.
Two things in the code make this worse than a slow path:
ggml_backend_cuda_comm_allreduce_ncclcan only return true. Every RCCL call in it is wrapped inNCCL_CHECK, which ends inGGML_ABORT. The comment on the comm context atggml-cuda.cu:1038says per-call failures return false and the meta backend handles them. That is not what happens for the NCCL variant. There is no way to get back to the butterfly path once RCCL has initialised.- The size gate in front of the custom AllReduce checks the compile-time
s_have_nccl, not whether RCCL actually initialised. So withGGML_CUDA_ALLREDUCE=internalornone, the large tensors are still declined by the custom AllReduce and go to the butterfly instead of the two-shot path.
Things I tried that did not change the abort:
--shm-size=1g(this only removed a separate/dev/shmwarning)NCCL_PROTO=SimpleNCCL_P2P_DISABLE=1--ipc=host- running without
HSA_FORCE_FINE_GRAIN_PCIE=1
Workaround that works: GGML_CUDA_ALLREDUCE=internal. The internal AllReduce is compiled out on HIP, so in practice this means the custom peer-write AllReduce takes the decode-sized tensors and the butterfly takes the rest. With that set I get about 45 tokens per second on the P2P kernel against 35 on the no-P2P kernel, so the P2P path is worth keeping.
What I would like: either keep RCCL calls out of the graph capture (or capture around them), or let the NCCL AllReduce return false on failure so the fallback the comment describes actually happens. I am happy to run a patched build on this box and send back logs. I have a full NCCL_DEBUG=INFO trace of the crash if that helps.
Steps to reproduce:
- Two MI50s on a kernel where
cudaDeviceCanAccessPeerreports true between them. - Run the image with the default AllReduce selection, for example:
docker run --rm --device /dev/kfd --device /dev/dri --group-add video --security-opt seccomp=unconfined \
-e GGML_ENABLE_CUSTOM_AR=1 -e HSA_FORCE_FINE_GRAIN_PCIE=1 -e HIP_VISIBLE_DEVICES=0,1 \
-v /path/to/models:/models -p 8080:8080 mixa3607/llama.cpp-gfx906:b10811-rocm-7.14-mxxm \
--server --host 0.0.0.0 --port 8080 -m /models/Qwen3.8-27B-Q8_0.gguf -c 262144 \
-ngl 99 -np 1 -lm dio -sm tensor --flash-attn on --jinja --spec-type draft-mtp --spec-draft-n-max 3
- Send any chat completion request. The server aborts on the first AllReduce.
- Add
-e GGML_CUDA_ALLREDUCE=internaland the same request works.
First Bad Commit
Unknown. b10811 is the only tag I have run on the P2P kernel. Earlier tags were only run on the stock kernel, where RCCL never initialises, so they would not have shown it.
Relevant log output
NCCL INFO Channel 00/0 : 0[3000] -> 1[9000] via P2P/direct pointer comm 0x5bf07deecf90 nRanks 02
NCCL INFO Channel 00/0 : 1[9000] -> 0[3000] via P2P/direct pointer comm 0x5bf07e9696d0 nRanks 02
NCCL INFO Connected all rings, use ring PXN 0 GDR 1
NCCL INFO Connected all trees
NCCL INFO ncclCommInitAll_impl comm 0x5bf090ca9630 rank 0 nranks 2 cudaDev 0 nvmlDev 0 busId 3000 - Init COMPLETE
NCCL INFO AllReduce: opCount 0 sendbuff 0x7b773b98d000 recvbuff 0x7b773b98d000 count 215040 datatype 9 op 0 root 0 comm 0x5bf07deecf90 [nranks=2]
NCCL INFO AllReduce: opCount 0 sendbuff 0x7b773b18d000 recvbuff 0x7b773b18d000 count 215040 datatype 9 op 0 root 0 comm 0x5bf07e9696d0 [nranks=2]
.../rccl/build/hipify/src/enqueue.cc:2061 NCCL WARN HIP failure 'the operation cannot be performed in the present state'
NCCL INFO .../rccl/build/hipify/src/group.cc:408 -> 1
NCCL INFO .../rccl/build/hipify/src/group.cc:791 -> 1
NCCL INFO .../rccl/build/hipify/src/group.cc:932 -> 1
/build/llamacpp/ggml/src/ggml-cuda/ggml-cuda.cu:109: ROCm error
E ROCm error: unhandled cuda error (run with NCCL_DEBUG=INFO for details)
E current device: 1, in function ggml_backend_cuda_comm_allreduce_nccl at /build/llamacpp/ggml/src/ggml-cuda/ggml-cuda.cu:1366
libggml-base.so.0(ggml_print_backtrace+0x21c)
libggml-base.so.0(ggml_abort+0x15b)
One more detail from the same trace. With --spec-type draft-mtp the fork creates a second llama_context for the draft, and each context runs its own ggml_backend_cuda_comm_init. In some runs one context's init walked the fallback chain and logged internal AllReduce init failed ... falling back to meta-backend butterfly while the other context held a live RCCL comm and aborted. So seeing the fallback line in the log does not mean RCCL is out of the picture.
Activity
model: Muse Glimmer Support (ggml-org#26841)
mixa3607 commentedon Sep 2, 2026
try run
all_reduce_perf -b 8 -e 1G -f 2 -g <gpus count>in llama containerand environment
rocm-host-report-llama-cpp-amd-llamacpp-6cf97999cb-k5s4t-20260902T183422Z.tar.gz
JCraigWasTaken commentedon Sep 2, 2026
Thanks. I ran both.
all_reduce_perf -b 8 -e 1G -f 2 -g 2inside theb10811-rocm-7.14-mxxmimage fails on the first collective, before it prints a single row. Init is fine: both ranks come up, both channels connect overP2P/direct pointer,Connected all ringsandConnected all treesprint. The first AllReduce (8 bytes, 2 floats) then dies with the same error the fork hits:Exit code 3, same result on two runs. So this reproduces without llama.cpp: RCCL 2.30.4 in this image cannot complete an all-reduce between these two MI50s once peer access is available. On the stock Debian kernel, where peer access is off, RCCL init fails instead and the fork falls back to the butterfly, which is why it never showed before.
RCCL warns at init that
iommu=ptis missing from my kernel command line. I rebooted with it added and ran the same test: same failure at the first AllReduce, exit 3. Both cards' IOMMU groups reportidentitywith or without the flag, so that is ruled out. Log from that run is in the gist too.One difference from your setup that may matter: my cards are the Apple / V420 vBIOS MI50s, on a consumer B650E board, 2 cards, PCIe Gen4 x16 each, no bridge. Your run was 4 cards on a different board.
The
NCCL_DEBUG=INFOoutput from the failing run, the same withiommu=pt, and the host report from your script are here: https://gist.github.com/JCraigWasTaken/134bbe2f32ab870404485bd0e3c14d1amixa3607 commentedon Sep 3, 2026
Try run all_reduce on 6.3.3 🤔
To be honest, I don't really understand how it all works. I just built packages and they worked. I don't have enough knowledge to debug rccl.
JCraigWasTaken commentedon Sep 4, 2026
I built the test tool from source inside the 6.3.3 image and ran the same two-card all-reduce test. Same result as on 7.14: setup completes, then the first all-reduce fails with "unhandled cuda error", exit code 3. So it is not something new in RCCL 2.30; both versions fail on this pair of cards once peer access is on.
I also tried RCCL's transport switches (peer-to-peer off, shared memory off, the copy engines off, fine-grained memory, both protocols): every one fails the same way.
Logs are in the same gist as before: https://gist.github.com/JCraigWasTaken/134bbe2f32ab870404485bd0e3c14d1a (the 6.3.3 run, the switch summary, and one switch run with debug output).
I have a small change for this fork that tests RCCL right after it starts up and falls back to the fork's own exchange instead of crashing, so the peer-access kernel works without setting
GGML_CUDA_ALLREDUCEby hand. PR coming separately.cuda: probe RCCL with a test all-reduce after init and fall back clea…
cuda: test RCCL with a small all-reduce after init and fall back when…