From 5df9c444e00299bc374763fba937dd5e79b157e1 Mon Sep 17 00:00:00 2001 From: John Craig Date: Thu, 3 Sep 2026 16:43:34 -0400 Subject: [PATCH 1/2] cuda: probe RCCL with a test all-reduce after init and fall back cleanly (issue #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 Claude-Session: https://claude.ai/code/session_018WHiJ3bdMZsRmkWqMnK1pj --- ggml/src/ggml-cuda/ggml-cuda.cu | 114 ++++++++++++++++++++++++++++++-- 1 file changed, 110 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 1cd0becdab56..fb2a9d05672c 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1498,7 +1498,94 @@ static void ggml_backend_cuda_comm_init_internal(ggml_backend_cuda_comm_context ggml_backend_cuda_comm_init_none(ret); } -static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * ret) { +#ifdef GGML_USE_NCCL +// NCCL/RCCL can report a successful ncclCommInitAll -- P2P channels connected +// and all -- and then fail on the very first collective. Seen on gfx906/MI50 +// with PCIe peer access enabled and RCCL 2.30.4, where the first AllReduce +// dies with "unhandled cuda error"; on kernels without peer access the init +// itself fails instead and the clean fallback below is taken. So run a tiny +// AllReduce here, check it without the aborting NCCL_CHECK macro, and verify +// the sum, while falling back is still possible. GGML_CUDA_RCCL_PROBE=0 skips it. +static bool ggml_backend_cuda_comm_nccl_probe(ggml_backend_cuda_comm_context * ret, const char ** why) { + *why = nullptr; + + const char * env_probe = getenv("GGML_CUDA_RCCL_PROBE"); + if (env_probe && env_probe[0] == '0') { + return true; + } + + const size_t n = ret->comms.size(); + const int ne = 64; + const float expected = (float) (n * (n + 1) / 2); // rank i contributes i+1 + + std::vector bufs(n, nullptr); + std::vector host(ne); + + for (size_t i = 0; i < n && *why == nullptr; ++i) { + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) ret->backends[i]->context; + ggml_cuda_set_device(cuda_ctx->device); + for (int j = 0; j < ne; ++j) { + host[j] = (float) (i + 1); + } + if (cudaMalloc(&bufs[i], ne*sizeof(float)) != cudaSuccess || + cudaMemcpy(bufs[i], host.data(), ne*sizeof(float), cudaMemcpyHostToDevice) != cudaSuccess) { + *why = "probe buffer setup failed"; + } + } + + if (*why == nullptr) { + ncclResult_t rc = ncclGroupStart(); + for (size_t i = 0; i < n && rc == ncclSuccess; ++i) { + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) ret->backends[i]->context; + rc = ncclAllReduce(bufs[i], bufs[i], ne, ncclFloat, ncclSum, ret->comms[i], cuda_ctx->stream()); + } + const ncclResult_t rc_end = ncclGroupEnd(); + rc = rc != ncclSuccess ? rc : rc_end; + if (rc != ncclSuccess) { + *why = ncclGetErrorString(rc); + } + } + + // The MI50 failure is asynchronous - the calls above all return ncclSuccess + // and the error only surfaces on the stream sync and in ncclCommGetAsyncError, + // so check both, then check the arithmetic in case it fails silently. + for (size_t i = 0; i < n && *why == nullptr; ++i) { + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) ret->backends[i]->context; + ggml_cuda_set_device(cuda_ctx->device); + ncclResult_t async_rc = ncclSuccess; + if (cudaStreamSynchronize(cuda_ctx->stream()) != cudaSuccess) { + *why = "probe AllReduce failed on stream synchronize"; + } else if (ncclCommGetAsyncError(ret->comms[i], &async_rc) != ncclSuccess) { + *why = "ncclCommGetAsyncError failed"; + } else if (async_rc != ncclSuccess) { + *why = ncclGetErrorString(async_rc); + } else if (cudaMemcpy(host.data(), bufs[i], ne*sizeof(float), cudaMemcpyDeviceToHost) != cudaSuccess) { + *why = "probe result readback failed"; + } else { + for (int j = 0; j < ne; ++j) { + if (host[j] != expected) { + *why = "probe AllReduce returned a wrong sum"; + break; + } + } + } + } + + for (size_t i = 0; i < n; ++i) { + if (bufs[i] != nullptr) { + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) ret->backends[i]->context; + ggml_cuda_set_device(cuda_ctx->device); + (void) cudaFree(bufs[i]); + } + } + (void) cudaGetLastError(); // clear any sticky error left by the failed probe + return *why == nullptr; +} +#endif // GGML_USE_NCCL + +// forced: GGML_CUDA_ALLREDUCE=nccl was asked for explicitly, so a failing probe +// is an error rather than a reason to fall back. +static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * ret, bool forced) { #ifdef GGML_USE_NCCL // Disabling NCCL path when CUDA virtual devices are in use since NCCL requires one distinct physical GPU per rank. const ggml_cuda_device_info & info = ggml_cuda_info(); @@ -1513,7 +1600,25 @@ static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * re ret->comms.resize(n); ncclResult_t rc = ncclCommInitAll(ret->comms.data(), (int) n, ret->dev_ids.data()); if (rc == ncclSuccess) { - ret->try_allreduce = ggml_backend_cuda_comm_try_allreduce_nccl; + const char * why = nullptr; + if (ggml_backend_cuda_comm_nccl_probe(ret, &why)) { + ret->try_allreduce = ggml_backend_cuda_comm_try_allreduce_nccl; + return; + } + for (ncclComm_t comm : ret->comms) { + (void) ncclCommAbort(comm); // the comms are in an error state, Destroy can hang + } + ret->comms.clear(); + if (forced) { + GGML_ABORT("NCCL init succeeded but the probe AllReduce failed (%s) and " + "GGML_CUDA_ALLREDUCE=nccl was requested; " + "use GGML_CUDA_ALLREDUCE=internal or unset it to fall back\n", why); + } + GGML_LOG_WARN("NCCL initialised but the probe AllReduce failed (%s); " + "falling back to internal AllReduce " + "(GGML_CUDA_ALLREDUCE=nccl|internal|none forces a mode, " + "GGML_CUDA_RCCL_PROBE=0 skips the probe)\n", why); + ggml_backend_cuda_comm_init_internal(ret); return; } @@ -1521,6 +1626,7 @@ static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * re GGML_LOG_WARN("NCCL init failed (%s); falling back to internal AllReduce\n", ncclGetErrorString(rc)); #else // GGML_USE_NCCL + GGML_UNUSED(forced); #ifndef GGML_USE_HIP GGML_LOG_WARN("NCCL not compiled in; falling back to internal AllReduce. " "Recompile with -DGGML_CUDA_NCCL=ON for best multi-GPU performance.\n"); @@ -1562,14 +1668,14 @@ static void * ggml_backend_cuda_comm_init(ggml_backend_t * backends, size_t n_ba if (!env) { // Platform default: Linux uses NCCL, otherwise (generally Windows) internal #if defined(__linux__) - ggml_backend_cuda_comm_init_nccl(ret); + ggml_backend_cuda_comm_init_nccl(ret, /*forced =*/ false); #else ggml_backend_cuda_comm_init_internal(ret); #endif // defined(__linux__) } else { std::string env_str(env); if (env_str == "nccl") { - ggml_backend_cuda_comm_init_nccl(ret); + ggml_backend_cuda_comm_init_nccl(ret, /*forced =*/ true); } else if (env_str == "internal") { ggml_backend_cuda_comm_init_internal(ret); } else if (env_str == "none") { From 70207b8c724f50d7d8535309111689b95bbe9525 Mon Sep 17 00:00:00 2001 From: John Craig Date: Thu, 3 Sep 2026 17:21:36 -0400 Subject: [PATCH 2/2] cuda: after a failed RCCL probe, initialise the internal AllReduce the 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 Claude-Session: https://claude.ai/code/session_018WHiJ3bdMZsRmkWqMnK1pj --- ggml/src/ggml-cuda/ggml-cuda.cu | 29 +++++++++++++++++++++++++---- 1 file changed, 25 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index fb2a9d05672c..b281cf1aa9a5 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1484,6 +1484,20 @@ static void ggml_backend_cuda_comm_init_none(ggml_backend_cuda_comm_context * re ret->try_allreduce = ggml_backend_cuda_comm_try_allreduce_butterfly; } +// The internal AllReduce is a CUDA-only feature: in HIP and MUSA builds +// ggml_cuda_ar_pipeline_init() is a stub that always returns nullptr (see the +// #else branch at the end of ggml-cuda/allreduce.cu), so "internal" resolves to +// the meta-backend butterfly on those backends no matter what. Knowing that at +// compile time lets the messages below say what actually happens instead of +// reporting an init failure that was never avoidable. +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) +static constexpr bool ggml_cuda_has_internal_allreduce = false; +static constexpr const char * ggml_cuda_internal_allreduce_name = "meta-backend butterfly AllReduce"; +#else +static constexpr bool ggml_cuda_has_internal_allreduce = true; +static constexpr const char * ggml_cuda_internal_allreduce_name = "internal AllReduce"; +#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + static void ggml_backend_cuda_comm_init_internal(ggml_backend_cuda_comm_context * ret) { ret->ar_pipeline = ggml_cuda_ar_pipeline_init(ret->dev_ids.data(), ret->dev_ids.size()); if (ret->ar_pipeline) { @@ -1493,8 +1507,14 @@ static void ggml_backend_cuda_comm_init_internal(ggml_backend_cuda_comm_context // Clear sticky CUDA error from the failed init. (void) cudaGetLastError(); - GGML_LOG_WARN("internal AllReduce init failed (n_devices != 2?); " - "falling back to meta-backend butterfly\n"); + if (ggml_cuda_has_internal_allreduce) { + GGML_LOG_WARN("internal AllReduce init failed (n_devices != 2?); " + "falling back to meta-backend butterfly\n"); + } else { + // Nothing failed -- this backend has no internal AllReduce to begin with. + GGML_LOG_DEBUG("%s: internal AllReduce is not built for this backend; " + "using the meta-backend butterfly\n", __func__); + } ggml_backend_cuda_comm_init_none(ret); } @@ -1615,9 +1635,10 @@ static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * re "use GGML_CUDA_ALLREDUCE=internal or unset it to fall back\n", why); } GGML_LOG_WARN("NCCL initialised but the probe AllReduce failed (%s); " - "falling back to internal AllReduce " + "falling back to the %s " "(GGML_CUDA_ALLREDUCE=nccl|internal|none forces a mode, " - "GGML_CUDA_RCCL_PROBE=0 skips the probe)\n", why); + "GGML_CUDA_RCCL_PROBE=0 skips the probe)\n", + why, ggml_cuda_internal_allreduce_name); ggml_backend_cuda_comm_init_internal(ret); return; }