cuda: test RCCL with a small all-reduce after startup and fall back cleanly (#7) - #9
Open
JCraigWasTaken wants to merge 2 commits into
Open
cuda: test RCCL with a small all-reduce after startup and fall back cleanly (#7)#9JCraigWasTaken wants to merge 2 commits into
JCraigWasTaken wants to merge 2 commits into
Conversation
…nly (issue mxxm-t#7) On a 2x MI50 (gfx906) box running a kernel that supports PCIe peer access, ncclCommInitAll() returns ncclSuccess and RCCL reports its P2P channels as connected, but the very first real all-reduce dies with "unhandled cuda error" and the aborting NCCL_CHECK macro takes the process down. The fault is in RCCL 2.30.4 itself, not in llama.cpp: rccl-tests all_reduce_perf reproduces it standalone on the same machine. On kernels without peer access the init call fails instead, so the existing init-failure path fires and the fork falls back to its own exchange - which is why this was never seen before. The only workaround so far was setting GGML_CUDA_ALLREDUCE=internal by hand. Run a tiny all-reduce right after ncclCommInitAll and take the same fallback the init-failure path takes if it does not work. The probe checks the ncclResult_t values directly instead of using NCCL_CHECK, and because the MI50 failure is asynchronous it also checks cudaStreamSynchronize and ncclCommGetAsyncError on every rank, then verifies the sum numerically so a silently wrong result also triggers the fallback. On failure the comms are aborted and cleared, one GGML_LOG_WARN line explains what happened and names GGML_CUDA_ALLREDUCE, and init continues into the internal path. If GGML_CUDA_ALLREDUCE=nccl was set explicitly the probe still runs but a failure aborts with the same explanation, since the user asked for NCCL. GGML_CUDA_RCCL_PROBE=0 skips the probe; it is on by default. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_018WHiJ3bdMZsRmkWqMnK1pj
…e same way GGML_CUDA_ALLREDUCE=internal does It already did. The post-probe fallback calls ggml_backend_cuda_comm_init_internal(), which is byte-for-byte the same function GGML_CUDA_ALLREDUCE=internal calls, with the same context, the same dev_ids and no state consumed by the NCCL branch in between. What made it look otherwise is that on this box (2x MI50, HIP) *both* paths end up on the meta-backend butterfly, because the internal AllReduce is compiled out: allreduce.cu guards the whole implementation with "#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)" and HIP builds get a stub ggml_cuda_ar_pipeline_init() that returns nullptr always. init_internal reads that nullptr as an init failure and logs "internal AllReduce init failed (n_devices != 2?)", which is wrong on two counts: nothing was tried, and n_devices was fine. The production GGML_CUDA_ALLREDUCE=internal runs log the same line and have all along. So there is no fallback to repair, only messages that misdescribe what happened. Give the file a compile-time pair -- whether the internal AllReduce exists in this build, and the name of whatever "internal" resolves to -- and use it in both places: on HIP/MUSA the init step now says at debug level that there is no internal AllReduce in this build rather than warning about a failure, and the probe-failure warning names the path actually being taken. On CUDA every message is unchanged. Behaviour is unchanged everywhere; only log text moves. On this box the no-env fallback run already reached 46.29 t/s decode and 317 t/s prefill, matching the GGML_CUDA_ALLREDUCE=internal champion, which is the same thing measured twice. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_018WHiJ3bdMZsRmkWqMnK1pj
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
On two MI50s with PCIe peer access enabled, RCCL starts up fine, connects both cards, and then fails on the first real all-reduce. The test tool that ships with RCCL shows the same failure without llama.cpp, on both the 7.14 and 6.3.3 builds, and none of RCCL's transport switches change it, so this is RCCL itself, not the fork. Today the fork only falls back when RCCL fails to start, so on a peer-access kernel it crashes on the first token.
This change runs a tiny all-reduce right after RCCL starts: 64 numbers per card, each card contributing its own number, checked against the expected total. It checks the return codes, the stream, and RCCL's asynchronous error state, because on these cards the failure only shows up after the calls return. If anything is wrong it shuts the RCCL connections down, prints one warning saying which path it is using instead, and carries on with the same fallback the startup-failure case already uses. If
GGML_CUDA_ALLREDUCE=ncclwas set on purpose, a failed test stops with a message pointing at that setting.GGML_CUDA_RCCL_PROBE=0skips the test.Found on the way: the "internal" all-reduce is not compiled into AMD builds at all, so
GGML_CUDA_ALLREDUCE=internalhas always meant the fork's butterfly exchange on AMD. The messages now say that instead of reporting a failure that could never have succeeded. Nothing about the behaviour changes.Measured on two MI50s with Qwen3.8-27B at 8-bit split across both cards, production settings: with nothing set, the server starts, prints the warning once per model load, and runs at 46.4 tokens per second decode and 315 prefill on an 8,148-token prompt, the same as with the setting forced by hand. With
GGML_CUDA_ALLREDUCE=ncclset, it stops at startup with the clear message.Fixes #7
🤖 Generated with Claude Code
https://claude.ai/code/session_018WHiJ3bdMZsRmkWqMnK1pj