Eval bug: 2x MI50 tensor split aborts in RCCL on the first AllReduce once peer access works, no fallback to butterfly

[JCraigWasTaken]

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:

1. ggml_backend_cuda_comm_allreduce_nccl can only return true. Every RCCL call in it is wrapped in NCCL_CHECK, which ends in GGML_ABORT. The comment on the comm context at ggml-cuda.cu:1038 says 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.
2. The size gate in front of the custom AllReduce checks the compile-time s_have_nccl, not whether RCCL actually initialised. So with GGML_CUDA_ALLREDUCE=internal or none, 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/shm warning)
• NCCL_PROTO=Simple
• NCCL_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:

1. Two MI50s on a kernel where cudaDeviceCanAccessPeer reports true between them.
2. 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

3. Send any chat completion request. The server aborts on the first AllReduce.
4. Add -e GGML_CUDA_ALLREDUCE=internal and 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.

[mixa3607]

try run all_reduce_perf -b 8 -e 1G -f 2 -g <gpus count> in llama container

root@llama-cpp-amd-llamacpp-6cf97999cb-k5s4t:/app# apt update && apt install -y amdrocm-rccl-test7.14-gfx906 && all_reduce_perf -b 8 -e 1G -f 2 -g 4 # <= 4 is GPUs count
...
# rccl-tests version 2.18.3-HEAD:418cd5f rccl-headers=23004 rccl-library=23004
# Collective test starting: all_reduce_perf
# nThread 1 nGpus 4 minBytes 8 maxBytes 1073741824 step: 2(factor) warmup iters: 1 iters: 20 agg iters: 1 validation: 1 graph: 0 unalign: 0
#
# Using devices
#  Rank  0 Group  0 Pid   1409 on llama-cpp-amd-llamacpp-6cf97999cb-k5s4t device  0 [0000:36:00] AMD Instinct MI60 / MI50
#  Rank  1 Group  0 Pid   1409 on llama-cpp-amd-llamacpp-6cf97999cb-k5s4t device  1 [0000:33:00] AMD Instinct MI60 / MI50
#  Rank  2 Group  0 Pid   1409 on llama-cpp-amd-llamacpp-6cf97999cb-k5s4t device  2 [0000:50:00] AMD Instinct MI60 / MI50
#  Rank  3 Group  0 Pid   1409 on llama-cpp-amd-llamacpp-6cf97999cb-k5s4t device  3 [0000:4d:00] AMD Instinct MI60 / MI50
#
#                                                              out-of-place                       in-place
#       size         count      type   redop    root     time   algbw   busbw  #wrong     time   algbw   busbw  #wrong
#        (B)    (elements)                               (us)  (GB/s)  (GB/s)             (us)  (GB/s)  (GB/s)
           8             2     float     sum      -1    30.90    0.00    0.00       0    26.55    0.00    0.00       0
          16             4     float     sum      -1    24.51    0.00    0.00       0    24.02    0.00    0.00       0
          32             8     float     sum      -1    25.63    0.00    0.00       0    25.10    0.00    0.00       0
          64            16     float     sum      -1    30.02    0.00    0.00       0    29.95    0.00    0.00       0
         128            32     float     sum      -1    29.47    0.00    0.01       0    29.73    0.00    0.01       0
         256            64     float     sum      -1    29.64    0.01    0.01       0    29.53    0.01    0.01       0
         512           128     float     sum      -1    29.35    0.02    0.03       0    29.04    0.02    0.03       0
        1024           256     float     sum      -1    29.02    0.04    0.05       0    29.32    0.03    0.05       0
        2048           512     float     sum      -1    29.33    0.07    0.10       0    29.65    0.07    0.10       0
        4096          1024     float     sum      -1    29.54    0.14    0.21       0    29.78    0.14    0.21       0
        8192          2048     float     sum      -1    60.21    0.14    0.20       0    58.26    0.14    0.21       0
       16384          4096     float     sum      -1    56.92    0.29    0.43       0    57.57    0.28    0.43       0
       32768          8192     float     sum      -1    53.68    0.61    0.92       0    53.07    0.62    0.93       0
       65536         16384     float     sum      -1    55.92    1.17    1.76       0    55.99    1.17    1.76       0
      131072         32768     float     sum      -1    68.86    1.90    2.86       0    67.34    1.95    2.92       0
      262144         65536     float     sum      -1    88.83    2.95    4.43       0    89.88    2.92    4.38       0
      524288        131072     float     sum      -1   138.10    3.80    5.69       0   137.69    3.81    5.71       0
     1048576        262144     float     sum      -1   226.25    4.63    6.95       0   225.95    4.64    6.96       0
     2097152        524288     float     sum      -1   398.54    5.26    7.89       0   398.64    5.26    7.89       0
     4194304       1048576     float     sum      -1   771.23    5.44    8.16       0   775.61    5.41    8.11       0
     8388608       2097152     float     sum      -1  1610.47    5.21    7.81       0  1607.30    5.22    7.83       0
    16777216       4194304     float     sum      -1  3570.16    4.70    7.05       0  3577.83    4.69    7.03       0
    33554432       8388608     float     sum      -1  7139.72    4.70    7.05       0  7131.44    4.71    7.06       0
    67108864      16777216     float     sum      -1  14050.1    4.78    7.16       0  14026.4    4.78    7.18       0
   134217728      33554432     float     sum      -1  28149.3    4.77    7.15       0  28182.1    4.76    7.14       0
   268435456      67108864     float     sum      -1  56131.8    4.78    7.17       0  56919.0    4.72    7.07       0
   536870912     134217728     float     sum      -1   112470    4.77    7.16       0   112153    4.79    7.18       0
  1073741824     268435456     float     sum      -1   225836    4.75    7.13       0   224934    4.77    7.16       0
# Out of bounds values : 0 OK
# Avg bus bandwidth    : 3.47786
#
# Collective test concluded: all_reduce_perf
#

and environment

# curl -fsSL https://gist.githubusercontent.com/mixa3607/2ed3494b14d401ddd8bfcb03be8abeaf/raw/rocm-host-report.sh | bash
Created report directory: rocm-host-report-llama-cpp-amd-llamacpp-6cf97999cb-k5s4t-20260902T183422Z
Created archive: rocm-host-report-llama-cpp-amd-llamacpp-6cf97999cb-k5s4t-20260902T183422Z.tar.gz

rocm-host-report-llama-cpp-amd-llamacpp-6cf97999cb-k5s4t-20260902T183422Z.tar.gz

[JCraigWasTaken]

Thanks. I ran both.

all_reduce_perf -b 8 -e 1G -f 2 -g 2 inside the b10811-rocm-7.14-mxxm image fails on the first collective, before it prints a single row. Init is fine: both ranks come up, both channels connect over P2P/direct pointer, Connected all rings and Connected all trees print. The first AllReduce (8 bytes, 2 floats) then dies with the same error the fork hits:

# rccl-tests version 2.18.3-HEAD:418cd5f rccl-headers=23004 rccl-library=23004
#  Rank  0 Group  0 Pid      1 on e20a3f6a8c03 device  0 [0000:03:00] AMD Radeon Graphics
#  Rank  1 Group  0 Pid      1 on e20a3f6a8c03 device  1 [0000:09:00] AMD Radeon Graphics
NCCL INFO Channel 00/0 : 0[3000] -> 1[9000] via P2P/direct pointer comm 0x5a3ea68dbe40 nRanks 02
NCCL INFO Channel 01/0 : 1[9000] -> 0[3000] via P2P/direct pointer comm 0x5a3ea69f50d0 nRanks 02
NCCL INFO Connected all rings, use ring PXN 0 GDR 1
NCCL INFO Connected all trees
NCCL INFO ncclCommInitRankConfig_impl comm 0x5a3ea68dbe40 rank 0 nranks 2 cudaDev 0 busId 3000 - Init COMPLETE
NCCL INFO AllReduce: opCount 0 sendbuff 0x7cb27a400000 recvbuff 0x7cb23a200000 count 2 datatype 7 op 0 root 0 comm 0x5a3ea68dbe40 [nranks=2]
e20a3f6a8c03: Test NCCL failure .../rccl-tests/build/src/hipify/common.cu.cpp:775 'unhandled cuda error (run with NCCL_DEBUG=INFO for details) / '

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=pt is 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 report identity with 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=INFO output from the failing run, the same with iommu=pt, and the host report from your script are here: https://gist.github.com/JCraigWasTaken/134bbe2f32ab870404485bd0e3c14d1a

[mixa3607]

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]

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_ALLREDUCE by hand. PR coming separately.