From 799bad73489370646a4b12885be7ed29d8c6ccd4 Mon Sep 17 00:00:00 2001 From: piggidragon Date: Fri, 16 Oct 2026 16:23:02 +0200 Subject: [PATCH 1/4] cuda : AllReduce paths for two GPUs with peer access over PCIe With peer access llama.cpp left every AllReduce to NCCL. Over a PCIe link NCCL's small reductions have more latency than needed and its kernels hold an SM on each GPU, which slows down the compute of the other ubatch with --split-overlap; a large reduction is bound by the slower link either way. - small reductions (below 512 KiB): a peer-to-peer kernel (allreduce-p2p.cu). Each block writes its slice of the own BF16 contribution into the memory of the other GPU and then its token into that GPU's flag, and polls only its own flag; the data crosses the link once, as posted writes. Two slots alternate; the stream order of the kernels and the growing tokens make that safe without host synchronization. - large reductions: the copy engines copy GPU to GPU (cudaMemcpyPeerAsync) instead of through pinned host memory. Both are on by default with peer access without native peer atomics (PCIe); NVLink keeps NCCL, and so does HIP. GGML_CUDA_AR_P2P_SMALL and GGML_CUDA_AR_CE_P2P override the choice. RTX 4070 + RTX 3060 (PCIe 3.0 x4) with BAR1 P2P, Qwen3.8-27B, -ts 0.625/0.375 -as 1/3, q5_1, -ub 256 --split-overlap, against the same GPUs without peer access: tg32 36.95 -> 37.73 t/s, pp3 97.4 -> 99.6, pp2048 1179 -> 1188, at 16k 35.77 -> 36.48, 77.6 -> 79.2, 1046 -> 1067. Against NCCL with peer access: tg32 34.2 -> 37.7, pp2048 with --split-overlap 865 -> 1188. A small reduction takes 10 us on the 3060 (internal pipeline 17 us, NCCL without P2P 41 us). Co-Authored-By: Claude Opus 5.5 --- docs/multi-gpu.md | 26 +++-- ggml/src/ggml-cuda/allreduce-ce.cu | 64 +++++++++- ggml/src/ggml-cuda/allreduce-ce.cuh | 4 +- ggml/src/ggml-cuda/allreduce-p2p.cu | 168 +++++++++++++++++++++++++++ ggml/src/ggml-cuda/allreduce-p2p.cuh | 23 ++++ ggml/src/ggml-cuda/ggml-cuda.cu | 74 +++++++++--- 6 files changed, 327 insertions(+), 32 deletions(-) create mode 100644 ggml/src/ggml-cuda/allreduce-p2p.cu create mode 100644 ggml/src/ggml-cuda/allreduce-p2p.cuh diff --git a/docs/multi-gpu.md b/docs/multi-gpu.md index c6514373315e..bd1cc79aafba 100644 --- a/docs/multi-gpu.md +++ b/docs/multi-gpu.md @@ -113,28 +113,32 @@ GGML_CUDA_P2P=1 llama-cli -m model.gguf -sm tensor P2P requires driver support (usually restricted to workstation/datacenter GPUs) and **may cause crashes or corrupted outputs on some motherboards or BIOS configurations** (e.g. when IOMMU is enabled). If you see instability after enabling it, unset the variable. -### 7. Two GPUs without peer access +### 7. Two GPUs over PCIe -With `--split-mode tensor`, two GPUs that cannot access each other's memory (the usual case for consumer cards) exchange their partial results through host memory. With NCCL compiled in, each reduction takes one of three paths: +With `--split-mode tensor`, two GPUs exchange their partial results after every layer. With NCCL compiled in, each reduction takes one of these paths: -| Reduction size | Path | Why | -|---|---|---| -| below 512 KiB (decode, small batches) | internal pipeline, a single kernel with a handshake through pinned memory | lower latency than NCCL | -| from 512 KiB (prompt processing) | copy engines: each GPU copies its part to pinned memory in chunks and the other one copies them back as they arrive | no SM is busy, both directions of a link are used at once | -| anything the two above do not handle | NCCL | | +| Link | Reduction size | Path | Why | +|---|---|---|---| +| no peer access (the usual case for consumer cards) | below 512 KiB (decode, small batches) | internal pipeline: a single kernel with a handshake through pinned memory | lower latency than NCCL | +| no peer access | from 512 KiB (prompt processing) | copy engines: each GPU copies its part to pinned memory in chunks, the other copies them back as they arrive | no SM is busy, both directions of a link are used at once | +| peer access over PCIe | below 512 KiB | peer-to-peer kernel: each GPU writes its part and a flag into the memory of the other and polls only its own | lower latency than NCCL and the internal pipeline | +| peer access over PCIe | from 512 KiB | copy engines from GPU to GPU | as above, without the copy through host memory | +| NVLink (native peer atomics), or anything the paths above do not handle | | NCCL | | -With peer access, NCCL handles every reduction. +The paths matter most with `--split-overlap` (recipe 8): NCCL keeps an SM on each GPU busy, which slows down the compute of the other ubatch. Over PCIe, the bandwidth of the slower link limits a large reduction with or without peer access, so peer access mostly lowers the latency of the small ones. Environment variables: | Variable | Effect | |---|---| | `GGML_CUDA_ALLREDUCE=nccl` | NCCL for every reduction | -| `GGML_CUDA_AR_NCCL_MIN_BYTES=N` | size from which the internal pipeline is no longer used, and the default of `GGML_CUDA_AR_CE_MIN_BYTES` (0: NCCL for every reduction). Setting it also enables the internal pipeline with peer access; the copy engines stay off there | -| `GGML_CUDA_AR_CE=0` | NCCL instead of the copy engines for large reductions | +| `GGML_CUDA_AR_NCCL_MIN_BYTES=N` | size below which the small-reduction path is used, and the default of `GGML_CUDA_AR_CE_MIN_BYTES` (0: NCCL for every reduction). Setting it also enables the internal pipeline with peer access when the peer-to-peer kernel is off | +| `GGML_CUDA_AR_CE=0` | without peer access: NCCL instead of the copy engines for large reductions | +| `GGML_CUDA_AR_CE_P2P=0` or `1` | with peer access: copy engines from GPU to GPU off or on (default: on over PCIe) | +| `GGML_CUDA_AR_P2P_SMALL=0` or `1` | with peer access: peer-to-peer kernel off or on (default: on over PCIe) | | `GGML_CUDA_AR_CE_MIN_BYTES=N` | size from which the copy engines are used | -The copy engines need 16 MiB of VRAM and 8 MiB of pinned memory per GPU and per context that has large reductions. They are allocated with the first large reduction. An MTP draft context processes the prompt batches as well, so it has its own buffers: 32 MiB of VRAM per GPU in total. +The copy engines need 16 MiB of VRAM per GPU, plus 8 MiB of pinned memory without peer access, per context that has large reductions. They are allocated with the first large reduction. An MTP draft context processes the prompt batches as well, so it has its own buffers: 32 MiB of VRAM per GPU in total. The peer-to-peer kernel needs 512 KiB of VRAM per GPU and context. ### 8. Overlapping the AllReduce with compute (`--split-overlap`) diff --git a/ggml/src/ggml-cuda/allreduce-ce.cu b/ggml/src/ggml-cuda/allreduce-ce.cu index 0106b2c16e2e..4f02b993ae65 100644 --- a/ggml/src/ggml-cuda/allreduce-ce.cu +++ b/ggml/src/ggml-cuda/allreduce-ce.cu @@ -20,6 +20,9 @@ static constexpr int64_t GGML_CUDA_CE_AR_CHUNK_ELEMS = GGML_CUDA_CE_AR_CHUNK_BYT struct ggml_cuda_ce_ar { int devices[2] = { -1, -1 }; + // copy from GPU to GPU with peer access instead of through host memory + bool p2p = false; + // the buffers are set up with the first large reduction, a context that never has one allocates nothing (an // MTP draft does: it processes the prompt batches as well); if that fails, every call is left to the caller bool allocated = false; @@ -105,7 +108,7 @@ void ggml_cuda_ce_ar_free(ggml_cuda_ce_ar * ar) { delete ar; } -ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int * devices, size_t n_devices) { +ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int * devices, size_t n_devices, bool peer) { if (n_devices != 2) { return nullptr; } @@ -113,6 +116,7 @@ ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int * devices, size_t n_devices) { auto * ar = new ggml_cuda_ce_ar; ar->devices[0] = devices[0]; ar->devices[1] = devices[1]; + ar->p2p = peer; return ar; } @@ -127,7 +131,17 @@ static bool ggml_cuda_ce_ar_alloc(ggml_cuda_ce_ar * ar) { ok = ok && cudaStreamCreateWithFlags(&ar->h2d[i], cudaStreamNonBlocking) == cudaSuccess; ok = ok && cudaMalloc((void **) &ar->own[i], bytes) == cudaSuccess; ok = ok && cudaMalloc((void **) &ar->peer[i], bytes) == cudaSuccess; - ok = ok && cudaHostAlloc((void **) &ar->host[i], bytes, cudaHostAllocPortable) == cudaSuccess; + if (ar->p2p) { + // a direct copy needs peer access, without it cudaMemcpyPeerAsync goes through host memory + const cudaError_t err = cudaDeviceEnablePeerAccess(ar->devices[1 - i], 0); + if (err == cudaErrorPeerAccessAlreadyEnabled) { + (void) cudaGetLastError(); + } else { + ok = ok && err == cudaSuccess; + } + } else { + ok = ok && cudaHostAlloc((void **) &ar->host[i], bytes, cudaHostAllocPortable) == cudaSuccess; + } for (int s = 0; s < GGML_CUDA_CE_AR_N_SLOTS && ok; ++s) { for (cudaEvent_t * ev : { &ar->ev_conv[i][s], &ar->ev_d2h_done[i][s], &ar->ev_h2d[i][s], &ar->ev_add[i][s] }) { @@ -146,8 +160,8 @@ static bool ggml_cuda_ce_ar_alloc(ggml_cuda_ce_ar * ar) { return false; } - GGML_LOG_INFO("%s: copy-engine AllReduce: %d slots of %zu KiB per GPU (%zu MiB VRAM and %zu MiB pinned memory per GPU)\n", __func__, - GGML_CUDA_CE_AR_N_SLOTS, GGML_CUDA_CE_AR_SLOT_BYTES >> 10, (2*bytes) >> 20, bytes >> 20); + GGML_LOG_INFO("%s: copy-engine AllReduce%s: %d slots of %zu KiB per GPU (%zu MiB VRAM and %zu MiB pinned memory per GPU)\n", __func__, + ar->p2p ? " from GPU to GPU" : "", GGML_CUDA_CE_AR_N_SLOTS, GGML_CUDA_CE_AR_SLOT_BYTES >> 10, (2*bytes) >> 20, ar->p2p ? (size_t) 0 : bytes >> 20); return true; } @@ -193,6 +207,46 @@ bool ggml_cuda_ce_ar_allreduce(ggml_cuda_ce_ar * ar, ggml_backend_t * backends, const int64_t s_off = s*GGML_CUDA_CE_AR_SLOT_ELEMS; + if (ar->p2p) { + // each GPU copies its contribution into the peer buffer of the other one + for (int i = 0; i < 2; ++i) { + const int p = 1 - i; + ggml_cuda_set_device(ar->devices[i]); + cudaStream_t stream = cuda_ctx[i]->stream(); + + nv_bfloat16 * own = ar->own[i] + s_off; + + // own is read by the copy of the previous use of the slot + CUDA_CHECK(cudaStreamWaitEvent(stream, ar->ev_d2h_done[i][s])); + if (compute[i]) { + to_bf16((const float *) tensors[i]->data + off, own, n, stream); + } else { + CUDA_CHECK(cudaMemsetAsync(own, 0, n*sizeof(nv_bfloat16), stream)); + } + CUDA_CHECK(cudaEventRecord(ar->ev_conv[i][s], stream)); + + // the peer buffer of the other GPU is read by its sum of the previous use of the slot + CUDA_CHECK(cudaStreamWaitEvent(ar->d2h[i], ar->ev_conv[i][s])); + CUDA_CHECK(cudaStreamWaitEvent(ar->d2h[i], ar->ev_add[p][s])); + CUDA_CHECK(cudaMemcpyPeerAsync(ar->peer[p] + s_off, ar->devices[p], own, ar->devices[i], + n*sizeof(nv_bfloat16), ar->d2h[i])); + CUDA_CHECK(cudaEventRecord(ar->ev_d2h_done[i][s], ar->d2h[i])); + } + for (int i = 0; i < 2; ++i) { + const int p = 1 - i; + ggml_cuda_set_device(ar->devices[i]); + cudaStream_t stream = cuda_ctx[i]->stream(); + + CUDA_CHECK(cudaStreamWaitEvent(stream, ar->ev_d2h_done[p][s])); + const int block_size = 256; + const int64_t n_blocks = (n + block_size - 1) / block_size; + ggml_cuda_ce_ar_add<<>>((float *) tensors[i]->data + off, ar->own[i] + s_off, ar->peer[i] + s_off, n); + CUDA_CHECK(cudaGetLastError()); + CUDA_CHECK(cudaEventRecord(ar->ev_add[i][s], stream)); + } + continue; + } + // the own contribution to host memory for (int i = 0; i < 2; ++i) { ggml_cuda_set_device(ar->devices[i]); @@ -257,7 +311,7 @@ bool ggml_cuda_ce_ar_allreduce(ggml_cuda_ce_ar * ar, ggml_backend_t * backends, #else // defined(GGML_USE_MUSA) // MUSA lacks cudaHostAlloc / cudaHostAllocPortable (see allreduce.cu): the caller keeps its other AllReduce -ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int *, size_t) { +ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int *, size_t, bool) { return nullptr; } void ggml_cuda_ce_ar_free(ggml_cuda_ce_ar *) { diff --git a/ggml/src/ggml-cuda/allreduce-ce.cuh b/ggml/src/ggml-cuda/allreduce-ce.cuh index 38302dcd25e9..baaa88724b2b 100644 --- a/ggml/src/ggml-cuda/allreduce-ce.cuh +++ b/ggml/src/ggml-cuda/allreduce-ce.cuh @@ -14,8 +14,10 @@ struct ggml_cuda_ce_ar; // the buffers (16 MiB VRAM and 8 MiB pinned memory per GPU) are allocated with the first reduction +// peer: the GPUs have peer access, each copies its contribution straight into the other's memory instead of +// through host memory (no pinned memory then) // returns nullptr on MUSA, which lacks the pinned-memory API used here -ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int * devices, size_t n_devices); +ggml_cuda_ce_ar * ggml_cuda_ce_ar_init(const int * devices, size_t n_devices, bool peer = false); void ggml_cuda_ce_ar_free(ggml_cuda_ce_ar * ar); diff --git a/ggml/src/ggml-cuda/allreduce-p2p.cu b/ggml/src/ggml-cuda/allreduce-p2p.cu new file mode 100644 index 000000000000..13ea55fd90cb --- /dev/null +++ b/ggml/src/ggml-cuda/allreduce-p2p.cu @@ -0,0 +1,168 @@ +#include "allreduce-p2p.cuh" + +#if !defined(GGML_USE_MUSA) + +#include + +// Each AllReduce is one kernel per GPU. Block b of GPU i writes its slice of the own contribution (BF16) into the +// receive buffer of the other GPU and then its token into the flag of the other GPU, both through peer access; +// then it waits for its own flag, which the other GPU writes the same way, and adds. Only local memory is polled, +// and the data crosses the link once, as posted writes. +// +// Two slots alternate. GPU i writes slot k%2 of call k only after its kernel of call k-1 saw the arrival of the +// other GPU for call k-1, which that GPU signals after its kernel of call k-2 has read slot k%2 (stream order). +// The tokens grow, so a flag that already holds a later token also counts as arrived. + +static constexpr int GGML_CUDA_P2P_AR_BLOCKS = 8; +static constexpr int GGML_CUDA_P2P_AR_THREADS = 256; +static constexpr int64_t GGML_CUDA_P2P_AR_MAX_NE = 512*1024/sizeof(float); // the small reductions only + +struct ggml_cuda_p2p_ar { + int devices[2] = { -1, -1 }; + + nv_bfloat16 * recv [2] = {}; // device i: [slot][GGML_CUDA_P2P_AR_MAX_NE], written by the other GPU + int * flags[2] = {}; // device i: [block], written by the other GPU + + unsigned int token = 0; +}; + +static __global__ void ggml_cuda_p2p_ar_kernel( + float * data, const bool compute, nv_bfloat16 * peer_recv, int * peer_flags, + const nv_bfloat16 * own_recv, const int * own_flags, const int64_t n, const int token) { + const int64_t per = (n + gridDim.x - 1) / gridDim.x; + const int64_t i0 = blockIdx.x*per; + const int64_t i1 = i0 + per < n ? i0 + per : n; + + for (int64_t i = i0 + threadIdx.x; i < i1; i += blockDim.x) { + peer_recv[i] = compute ? __float2bfloat16(data[i]) : __float2bfloat16(0.0f); + } + __syncthreads(); + + if (threadIdx.x == 0) { + // the data before the token, for the other GPU + __threadfence_system(); + *(volatile int *) &peer_flags[blockIdx.x] = token; + + while ((int) (*(const volatile int *) &own_flags[blockIdx.x] - token) < 0) { + } + __threadfence(); + } + __syncthreads(); + + for (int64_t i = i0 + threadIdx.x; i < i1; i += blockDim.x) { + const float own = compute ? __bfloat162float(__float2bfloat16(data[i])) : 0.0f; + data[i] = own + __bfloat162float(own_recv[i]); + } +} + +void ggml_cuda_p2p_ar_free(ggml_cuda_p2p_ar * ar) { + if (ar == nullptr) { + return; + } + for (int i = 0; i < 2; ++i) { + if (ar->devices[i] < 0) { + continue; + } + ggml_cuda_set_device(ar->devices[i]); + if (ar->recv[i]) { + CUDA_CHECK(cudaFree(ar->recv[i])); + } + if (ar->flags[i]) { + CUDA_CHECK(cudaFree(ar->flags[i])); + } + } + delete ar; +} + +size_t ggml_cuda_p2p_ar_max_bytes() { + return GGML_CUDA_P2P_AR_MAX_NE*sizeof(float); +} + +ggml_cuda_p2p_ar * ggml_cuda_p2p_ar_init(const int * devices, size_t n_devices) { + if (n_devices != 2) { + return nullptr; + } + + auto * ar = new ggml_cuda_p2p_ar; + ar->devices[0] = devices[0]; + ar->devices[1] = devices[1]; + + bool ok = true; + for (int i = 0; i < 2 && ok; ++i) { + int can_access = 0; + ok = cudaDeviceCanAccessPeer(&can_access, devices[i], devices[1 - i]) == cudaSuccess && can_access; + if (!ok) { + break; + } + ggml_cuda_set_device(devices[i]); + const cudaError_t err = cudaDeviceEnablePeerAccess(devices[1 - i], 0); + if (err == cudaErrorPeerAccessAlreadyEnabled) { + (void) cudaGetLastError(); + } else { + ok = err == cudaSuccess; + } + ok = ok && cudaMalloc((void **) &ar->recv[i], 2*GGML_CUDA_P2P_AR_MAX_NE*sizeof(nv_bfloat16)) == cudaSuccess; + ok = ok && cudaMalloc((void **) &ar->flags[i], GGML_CUDA_P2P_AR_BLOCKS*sizeof(int)) == cudaSuccess; + ok = ok && cudaMemset(ar->flags[i], 0, GGML_CUDA_P2P_AR_BLOCKS*sizeof(int)) == cudaSuccess; + } + + if (!ok) { + (void) cudaGetLastError(); + GGML_LOG_WARN("%s: cannot set up the peer-to-peer AllReduce\n", __func__); + ggml_cuda_p2p_ar_free(ar); + return nullptr; + } + + GGML_LOG_INFO("%s: peer-to-peer AllReduce up to %zu bytes (%zu KiB VRAM per GPU)\n", __func__, + ggml_cuda_p2p_ar_max_bytes(), (2*GGML_CUDA_P2P_AR_MAX_NE*sizeof(nv_bfloat16)) >> 10); + return ar; +} + +bool ggml_cuda_p2p_ar_allreduce(ggml_cuda_p2p_ar * ar, ggml_backend_t * backends, ggml_tensor ** tensors) { + const int64_t ne = ggml_nelements(tensors[0]); + for (int i = 0; i < 2; ++i) { + if (tensors[i]->type != GGML_TYPE_F32 || ggml_nelements(tensors[i]) != ne || !ggml_is_contiguously_allocated(tensors[i])) { + return false; + } + } + if (ne > GGML_CUDA_P2P_AR_MAX_NE) { + return false; + } + if (ne == 0) { + return true; + } + + ar->token++; + const int token = (int) ar->token; + const int64_t s_off = (ar->token % 2)*GGML_CUDA_P2P_AR_MAX_NE; + + for (int i = 0; i < 2; ++i) { + const int p = 1 - i; + auto * cuda_ctx = (ggml_backend_cuda_context *) backends[i]->context; + GGML_ASSERT(cuda_ctx->device == ar->devices[i]); + ggml_cuda_set_device(ar->devices[i]); + + const bool compute = (tensors[i]->flags & GGML_TENSOR_FLAG_COMPUTE) != 0; + ggml_cuda_p2p_ar_kernel<<stream()>>>( + (float *) tensors[i]->data, compute, ar->recv[p] + s_off, ar->flags[p], ar->recv[i] + s_off, ar->flags[i], ne, token); + CUDA_CHECK(cudaGetLastError()); + } + + return true; +} + +#else // defined(GGML_USE_MUSA) + +ggml_cuda_p2p_ar * ggml_cuda_p2p_ar_init(const int *, size_t) { + return nullptr; +} +void ggml_cuda_p2p_ar_free(ggml_cuda_p2p_ar *) { +} +size_t ggml_cuda_p2p_ar_max_bytes() { + return 0; +} +bool ggml_cuda_p2p_ar_allreduce(ggml_cuda_p2p_ar *, ggml_backend_t *, ggml_tensor **) { + return false; +} + +#endif // !defined(GGML_USE_MUSA) diff --git a/ggml/src/ggml-cuda/allreduce-p2p.cuh b/ggml/src/ggml-cuda/allreduce-p2p.cuh new file mode 100644 index 000000000000..0b8a7047b476 --- /dev/null +++ b/ggml/src/ggml-cuda/allreduce-p2p.cuh @@ -0,0 +1,23 @@ +#pragma once + +#include "common.cuh" +#include "ggml-backend-impl.h" + +#include + +// AllReduce of two GPUs with peer access for small reductions (decode): one kernel per GPU writes its contribution +// and then a flag straight into the memory of the other GPU and waits for the flag the other one writes into its +// own memory. Only local memory is polled and the data crosses the link once, as posted writes. +struct ggml_cuda_p2p_ar; + +// enables peer access between the two devices; nullptr if that is not possible +ggml_cuda_p2p_ar * ggml_cuda_p2p_ar_init(const int * devices, size_t n_devices); + +void ggml_cuda_p2p_ar_free(ggml_cuda_p2p_ar * ar); + +// largest reduction in bytes (F32) the kernel takes +size_t ggml_cuda_p2p_ar_max_bytes(); + +// in-place sum of tensors[0] and tensors[1], contiguous F32 of the same size on the devices of backends[0] and [1] +// returns false without enqueuing anything if the tensors are not supported +bool ggml_cuda_p2p_ar_allreduce(ggml_cuda_p2p_ar * ar, ggml_backend_t * backends, ggml_tensor ** tensors); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index e9d25ec5fe15..67979b316fc1 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4,6 +4,7 @@ #include "ggml-cuda/allreduce.cuh" #include "ggml-cuda/allreduce-ce.cuh" +#include "ggml-cuda/allreduce-p2p.cuh" #include "ggml-cuda/common.cuh" #include "ggml-cuda/acc.cuh" #include "ggml-cuda/add-id.cuh" @@ -1031,6 +1032,9 @@ struct ggml_backend_cuda_comm_context { ggml_cuda_ce_ar * ce_ar = nullptr; size_t ce_min_bytes = 0; + // with NCCL and peer access: reductions smaller than nccl_min_bytes go through the peer-to-peer kernel + ggml_cuda_p2p_ar * p2p_ar = nullptr; + #ifdef GGML_USE_NCCL std::vector comms; #endif // GGML_USE_NCCL @@ -1043,6 +1047,7 @@ struct ggml_backend_cuda_comm_context { #endif // GGML_USE_NCCL ggml_cuda_ar_pipeline_free(ar_pipeline); ggml_cuda_ce_ar_free(ce_ar); + ggml_cuda_p2p_ar_free(p2p_ar); } }; @@ -1191,6 +1196,10 @@ static bool ggml_backend_cuda_comm_try_allreduce_nccl( ggml_cuda_ce_ar_allreduce(comm_ctx->ce_ar, comm_ctx->backends.data(), tensors)) { return true; } + if (comm_ctx->p2p_ar != nullptr && ggml_nbytes(tensors[0]) < comm_ctx->nccl_min_bytes && + ggml_cuda_p2p_ar_allreduce(comm_ctx->p2p_ar, comm_ctx->backends.data(), tensors)) { + return true; + } if (comm_ctx->ar_pipeline != nullptr && ggml_nbytes(tensors[0]) < comm_ctx->nccl_min_bytes && ggml_backend_cuda_comm_allreduce_internal(comm_ctx, tensors)) { return true; @@ -1257,38 +1266,73 @@ static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * int can_access = 0; CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access, ret->dev_ids[0], ret->dev_ids[1])); - size_t min_bytes = 512*1024; - if (const char * env = getenv("GGML_CUDA_AR_NCCL_MIN_BYTES")) { - min_bytes = std::strtoull(env, nullptr, 10); - } else if (can_access) { - // NCCL takes the peer path, which has not been measured against the internal pipeline + // With peer access over PCIe the same holds: NCCL's kernels hold an SM, its small reductions have more latency + // than a kernel that writes into the other GPU and polls only its own memory, and a large one is bound by the + // link either way. The same GPUs with BAR1 P2P: tg 36.9 -> 37.6 t/s and pp3 97.5 -> 99.7 against the internal + // pipeline, with --split-overlap prompt processing 865 -> 1183 t/s against NCCL. A link with native peer atomics + // (NVLink) keeps NCCL. GGML_CUDA_AR_P2P_SMALL and GGML_CUDA_AR_CE_P2P (0 or 1) override the choice. + int native_atomics = 0; +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if (can_access) { + CUDA_CHECK(cudaDeviceGetP2PAttribute(&native_atomics, cudaDevP2PAttrNativeAtomicSupported, ret->dev_ids[0], ret->dev_ids[1])); + } + const bool p2p_pcie = can_access && !native_atomics; +#else + const bool p2p_pcie = false; +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + + auto env_flag = [](const char * name, bool def) { + const char * env = getenv(name); + return env != nullptr ? atoi(env) != 0 : def; + }; + + const char * env_min = getenv("GGML_CUDA_AR_NCCL_MIN_BYTES"); + const bool small_p2p = can_access && env_flag("GGML_CUDA_AR_P2P_SMALL", p2p_pcie); + const bool small_internal = !small_p2p && (!can_access || env_min != nullptr); + const bool large_ce = can_access ? env_flag("GGML_CUDA_AR_CE_P2P", p2p_pcie) : env_flag("GGML_CUDA_AR_CE", true); + if (!small_p2p && !small_internal && !large_ce) { return; } + size_t min_bytes = 512*1024; + if (env_min != nullptr) { + min_bytes = std::strtoull(env_min, nullptr, 10); + } + if (min_bytes == 0) { return; } // the reductions it gets stay below the copy-engine threshold, its 32 MiB staging per GPU is not needed // it needs Volta or newer, the copy engines do not - ret->ar_pipeline = ggml_cuda_ar_pipeline_init(ret->dev_ids.data(), ret->dev_ids.size(), /*small_only =*/ true); - if (ret->ar_pipeline == nullptr) { - (void) cudaGetLastError(); - } else { - ret->nccl_min_bytes = min_bytes; + if (small_p2p) { + ret->p2p_ar = ggml_cuda_p2p_ar_init(ret->dev_ids.data(), ret->dev_ids.size()); + if (ret->p2p_ar != nullptr) { + ret->nccl_min_bytes = std::min(min_bytes, ggml_cuda_p2p_ar_max_bytes()); + } + } + if (small_internal) { + ret->ar_pipeline = ggml_cuda_ar_pipeline_init(ret->dev_ids.data(), ret->dev_ids.size(), /*small_only =*/ true); + if (ret->ar_pipeline == nullptr) { + (void) cudaGetLastError(); + } else { + ret->nccl_min_bytes = min_bytes; + } } - // with peer access NCCL copies from GPU to GPU, the copy engines would go through host memory - const char * env_ce = getenv("GGML_CUDA_AR_CE"); - if (!can_access && (env_ce == nullptr || atoi(env_ce) != 0)) { - ret->ce_ar = ggml_cuda_ce_ar_init(ret->dev_ids.data(), ret->dev_ids.size()); + if (large_ce) { + ret->ce_ar = ggml_cuda_ce_ar_init(ret->dev_ids.data(), ret->dev_ids.size(), /*peer =*/ can_access != 0); ret->ce_min_bytes = min_bytes; if (const char * env = getenv("GGML_CUDA_AR_CE_MIN_BYTES")) { ret->ce_min_bytes = std::strtoull(env, nullptr, 10); } } - if (ret->ar_pipeline != nullptr && ret->ce_ar != nullptr) { + if (ret->p2p_ar != nullptr) { + GGML_LOG_INFO("%s: AllReduce below %zu bytes through the peer-to-peer kernel, from %zu bytes through %s\n", + __func__, ret->nccl_min_bytes, ret->ce_ar != nullptr ? ret->ce_min_bytes : ret->nccl_min_bytes, + ret->ce_ar != nullptr ? "the copy engines" : "NCCL"); + } else if (ret->ar_pipeline != nullptr && ret->ce_ar != nullptr) { GGML_LOG_INFO("%s: AllReduce below %zu bytes through the internal pipeline, from %zu bytes through the copy engines\n", __func__, min_bytes, ret->ce_min_bytes); } else if (ret->ar_pipeline != nullptr) { From 75ba452e5f1d38228aaa25e57c5f2f44824cd38e Mon Sep 17 00:00:00 2001 From: piggidragon Date: Fri, 16 Oct 2026 17:01:21 +0200 Subject: [PATCH 2/4] cuda : P2P AllReduce review fixes - compare the flag with the token in unsigned arithmetic: the signed difference overflowed at token 2^31, and the first GPU to arrive then read a stale slot - peer access over PCIe only with GGML_CUDA_P2P (as in the rest of the backend), not with NCCL_P2P_DISABLE - the copy engines start where the peer-to-peer kernel ends (512 KiB) when GGML_CUDA_AR_NCCL_MIN_BYTES is larger Co-Authored-By: Claude Opus 5.5 --- docs/multi-gpu.md | 16 +++++++++------- ggml/src/ggml-cuda/allreduce-p2p.cu | 19 ++++++++++--------- ggml/src/ggml-cuda/ggml-cuda.cu | 18 ++++++++++++------ 3 files changed, 31 insertions(+), 22 deletions(-) diff --git a/docs/multi-gpu.md b/docs/multi-gpu.md index bd1cc79aafba..44467ed74940 100644 --- a/docs/multi-gpu.md +++ b/docs/multi-gpu.md @@ -113,16 +113,18 @@ GGML_CUDA_P2P=1 llama-cli -m model.gguf -sm tensor P2P requires driver support (usually restricted to workstation/datacenter GPUs) and **may cause crashes or corrupted outputs on some motherboards or BIOS configurations** (e.g. when IOMMU is enabled). If you see instability after enabling it, unset the variable. +With two GPUs over PCIe and `--split-mode tensor`, it also switches the AllReduce to the peer-to-peer paths of recipe 7. `NCCL_P2P_DISABLE=1` keeps those off. + ### 7. Two GPUs over PCIe With `--split-mode tensor`, two GPUs exchange their partial results after every layer. With NCCL compiled in, each reduction takes one of these paths: | Link | Reduction size | Path | Why | |---|---|---|---| -| no peer access (the usual case for consumer cards) | below 512 KiB (decode, small batches) | internal pipeline: a single kernel with a handshake through pinned memory | lower latency than NCCL | -| no peer access | from 512 KiB (prompt processing) | copy engines: each GPU copies its part to pinned memory in chunks, the other copies them back as they arrive | no SM is busy, both directions of a link are used at once | -| peer access over PCIe | below 512 KiB | peer-to-peer kernel: each GPU writes its part and a flag into the memory of the other and polls only its own | lower latency than NCCL and the internal pipeline | -| peer access over PCIe | from 512 KiB | copy engines from GPU to GPU | as above, without the copy through host memory | +| no peer access (the usual case for consumer cards), or `GGML_CUDA_P2P` not set | below 512 KiB (decode, small batches) | internal pipeline: a single kernel with a handshake through pinned memory | lower latency than NCCL | +| no peer access, or `GGML_CUDA_P2P` not set | from 512 KiB (prompt processing) | copy engines: each GPU copies its part to pinned memory in chunks, the other copies them back as they arrive | no SM is busy, both directions of a link are used at once | +| peer access over PCIe and `GGML_CUDA_P2P` set (recipe 6) | below 512 KiB | peer-to-peer kernel: each GPU writes its part and a flag into the memory of the other and polls only its own | lower latency than NCCL and the internal pipeline | +| peer access over PCIe and `GGML_CUDA_P2P` set | from 512 KiB | copy engines from GPU to GPU | as above, without the copy through host memory | | NVLink (native peer atomics), or anything the paths above do not handle | | NCCL | | The paths matter most with `--split-overlap` (recipe 8): NCCL keeps an SM on each GPU busy, which slows down the compute of the other ubatch. Over PCIe, the bandwidth of the slower link limits a large reduction with or without peer access, so peer access mostly lowers the latency of the small ones. @@ -132,10 +134,10 @@ Environment variables: | Variable | Effect | |---|---| | `GGML_CUDA_ALLREDUCE=nccl` | NCCL for every reduction | -| `GGML_CUDA_AR_NCCL_MIN_BYTES=N` | size below which the small-reduction path is used, and the default of `GGML_CUDA_AR_CE_MIN_BYTES` (0: NCCL for every reduction). Setting it also enables the internal pipeline with peer access when the peer-to-peer kernel is off | +| `GGML_CUDA_AR_NCCL_MIN_BYTES=N` | size below which the small-reduction path is used, and the default of `GGML_CUDA_AR_CE_MIN_BYTES` (0: NCCL for every reduction). The peer-to-peer kernel handles at most 512 KiB; the copy engines then start there. Setting it also enables the internal pipeline with peer access when the peer-to-peer kernel is off | | `GGML_CUDA_AR_CE=0` | without peer access: NCCL instead of the copy engines for large reductions | -| `GGML_CUDA_AR_CE_P2P=0` or `1` | with peer access: copy engines from GPU to GPU off or on (default: on over PCIe) | -| `GGML_CUDA_AR_P2P_SMALL=0` or `1` | with peer access: peer-to-peer kernel off or on (default: on over PCIe) | +| `GGML_CUDA_AR_CE_P2P=0` or `1` | with peer access in use: copy engines from GPU to GPU off or on (default: on over PCIe) | +| `GGML_CUDA_AR_P2P_SMALL=0` or `1` | with peer access in use: peer-to-peer kernel off or on (default: on over PCIe) | | `GGML_CUDA_AR_CE_MIN_BYTES=N` | size from which the copy engines are used | The copy engines need 16 MiB of VRAM per GPU, plus 8 MiB of pinned memory without peer access, per context that has large reductions. They are allocated with the first large reduction. An MTP draft context processes the prompt batches as well, so it has its own buffers: 32 MiB of VRAM per GPU in total. The peer-to-peer kernel needs 512 KiB of VRAM per GPU and context. diff --git a/ggml/src/ggml-cuda/allreduce-p2p.cu b/ggml/src/ggml-cuda/allreduce-p2p.cu index 13ea55fd90cb..3dec00ce86d6 100644 --- a/ggml/src/ggml-cuda/allreduce-p2p.cu +++ b/ggml/src/ggml-cuda/allreduce-p2p.cu @@ -11,7 +11,8 @@ // // Two slots alternate. GPU i writes slot k%2 of call k only after its kernel of call k-1 saw the arrival of the // other GPU for call k-1, which that GPU signals after its kernel of call k-2 has read slot k%2 (stream order). -// The tokens grow, so a flag that already holds a later token also counts as arrived. +// The tokens grow, so a flag that already holds a later token also counts as arrived. They wrap around at 2^32: +// a flag counts as arrived when it is at most 2^31 - 1 ahead of the token, in unsigned arithmetic. static constexpr int GGML_CUDA_P2P_AR_BLOCKS = 8; static constexpr int GGML_CUDA_P2P_AR_THREADS = 256; @@ -21,14 +22,14 @@ struct ggml_cuda_p2p_ar { int devices[2] = { -1, -1 }; nv_bfloat16 * recv [2] = {}; // device i: [slot][GGML_CUDA_P2P_AR_MAX_NE], written by the other GPU - int * flags[2] = {}; // device i: [block], written by the other GPU + unsigned int * flags[2] = {}; // device i: [block], written by the other GPU unsigned int token = 0; }; static __global__ void ggml_cuda_p2p_ar_kernel( - float * data, const bool compute, nv_bfloat16 * peer_recv, int * peer_flags, - const nv_bfloat16 * own_recv, const int * own_flags, const int64_t n, const int token) { + float * data, const bool compute, nv_bfloat16 * peer_recv, unsigned int * peer_flags, + const nv_bfloat16 * own_recv, const unsigned int * own_flags, const int64_t n, const unsigned int token) { const int64_t per = (n + gridDim.x - 1) / gridDim.x; const int64_t i0 = blockIdx.x*per; const int64_t i1 = i0 + per < n ? i0 + per : n; @@ -41,9 +42,9 @@ static __global__ void ggml_cuda_p2p_ar_kernel( if (threadIdx.x == 0) { // the data before the token, for the other GPU __threadfence_system(); - *(volatile int *) &peer_flags[blockIdx.x] = token; + *(volatile unsigned int *) &peer_flags[blockIdx.x] = token; - while ((int) (*(const volatile int *) &own_flags[blockIdx.x] - token) < 0) { + while (*(const volatile unsigned int *) &own_flags[blockIdx.x] - token > 0x7fffffffu) { } __threadfence(); } @@ -102,8 +103,8 @@ ggml_cuda_p2p_ar * ggml_cuda_p2p_ar_init(const int * devices, size_t n_devices) ok = err == cudaSuccess; } ok = ok && cudaMalloc((void **) &ar->recv[i], 2*GGML_CUDA_P2P_AR_MAX_NE*sizeof(nv_bfloat16)) == cudaSuccess; - ok = ok && cudaMalloc((void **) &ar->flags[i], GGML_CUDA_P2P_AR_BLOCKS*sizeof(int)) == cudaSuccess; - ok = ok && cudaMemset(ar->flags[i], 0, GGML_CUDA_P2P_AR_BLOCKS*sizeof(int)) == cudaSuccess; + ok = ok && cudaMalloc((void **) &ar->flags[i], GGML_CUDA_P2P_AR_BLOCKS*sizeof(unsigned int)) == cudaSuccess; + ok = ok && cudaMemset(ar->flags[i], 0, GGML_CUDA_P2P_AR_BLOCKS*sizeof(unsigned int)) == cudaSuccess; } if (!ok) { @@ -133,7 +134,7 @@ bool ggml_cuda_p2p_ar_allreduce(ggml_cuda_p2p_ar * ar, ggml_backend_t * backends } ar->token++; - const int token = (int) ar->token; + const unsigned int token = ar->token; const int64_t s_off = (ar->token % 2)*GGML_CUDA_P2P_AR_MAX_NE; for (int i = 0; i < 2; ++i) { diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 67979b316fc1..cddeb0c0f8a6 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1271,14 +1271,19 @@ static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * // link either way. The same GPUs with BAR1 P2P: tg 36.9 -> 37.6 t/s and pp3 97.5 -> 99.7 against the internal // pipeline, with --split-overlap prompt processing 865 -> 1183 t/s against NCCL. A link with native peer atomics // (NVLink) keeps NCCL. GGML_CUDA_AR_P2P_SMALL and GGML_CUDA_AR_CE_P2P (0 or 1) override the choice. + // Peer access over PCIe is used only on request (GGML_CUDA_P2P, docs/multi-gpu.md), as in the rest of this + // backend: some platforms report it but corrupt or hang the transfers (IOMMU). NCCL_P2P_DISABLE also turns it + // off. Without it the GPUs are treated like GPUs without peer access. int native_atomics = 0; #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) if (can_access) { CUDA_CHECK(cudaDeviceGetP2PAttribute(&native_atomics, cudaDevP2PAttrNativeAtomicSupported, ret->dev_ids[0], ret->dev_ids[1])); } - const bool p2p_pcie = can_access && !native_atomics; + const bool p2p_pcie = can_access && !native_atomics && ggml_cuda_p2p_enabled() && !ggml_cuda_env_flag("NCCL_P2P_DISABLE"); + const bool peer = p2p_pcie || (can_access && native_atomics); #else const bool p2p_pcie = false; + const bool peer = can_access != 0; #endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) auto env_flag = [](const char * name, bool def) { @@ -1287,9 +1292,9 @@ static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * }; const char * env_min = getenv("GGML_CUDA_AR_NCCL_MIN_BYTES"); - const bool small_p2p = can_access && env_flag("GGML_CUDA_AR_P2P_SMALL", p2p_pcie); - const bool small_internal = !small_p2p && (!can_access || env_min != nullptr); - const bool large_ce = can_access ? env_flag("GGML_CUDA_AR_CE_P2P", p2p_pcie) : env_flag("GGML_CUDA_AR_CE", true); + const bool small_p2p = peer && env_flag("GGML_CUDA_AR_P2P_SMALL", p2p_pcie); + const bool small_internal = !small_p2p && (!peer || env_min != nullptr); + const bool large_ce = peer ? env_flag("GGML_CUDA_AR_CE_P2P", p2p_pcie) : env_flag("GGML_CUDA_AR_CE", true); if (!small_p2p && !small_internal && !large_ce) { return; } @@ -1321,8 +1326,9 @@ static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * } if (large_ce) { - ret->ce_ar = ggml_cuda_ce_ar_init(ret->dev_ids.data(), ret->dev_ids.size(), /*peer =*/ can_access != 0); - ret->ce_min_bytes = min_bytes; + ret->ce_ar = ggml_cuda_ce_ar_init(ret->dev_ids.data(), ret->dev_ids.size(), /*peer =*/ peer); + // the copy engines take over where the small path ends: the peer-to-peer kernel handles at most 512 KiB + ret->ce_min_bytes = small_p2p && ret->p2p_ar != nullptr ? ret->nccl_min_bytes : min_bytes; if (const char * env = getenv("GGML_CUDA_AR_CE_MIN_BYTES")) { ret->ce_min_bytes = std::strtoull(env, nullptr, 10); } From 933605b22db7b09836441d94eadda5e9045b22dd Mon Sep 17 00:00:00 2001 From: piggidragon Date: Fri, 16 Oct 2026 17:24:30 +0200 Subject: [PATCH 3/4] llama, cuda : opt in to the AllReduce paths with --allreduce host|p2p The paths of #14, #16 and this PR were chosen by the hardware and switched by environment variables. They are now opt-in with one flag, like the other features of this fork, and the default is upstream's NCCL: - --allreduce host: small reductions through the internal pipeline, large ones through the copy engines, through pinned host memory (#14, #16) - --allreduce p2p: the peer-to-peer kernel and the copy engines from GPU to GPU; host without peer access. The paths enable peer access themselves, GGML_CUDA_P2P is not needed - llama_context_params.allreduce_type, passed to the backends through the new reg proc ggml_backend_comm_set_mode before the meta backends of a context are created - llama-bench -ar, test-llama-archs configs "Meta -ar host" and "Meta -ar p2p -so" - GGML_CUDA_AR_CE, GGML_CUDA_AR_CE_P2P and GGML_CUDA_AR_P2P_SMALL are gone; the two size thresholds stay as tuning variables, GGML_CUDA_ALLREDUCE still overrides - README: the features of this fork as one list (what, benefit, how to enable), and a table upstream vs this fork Co-Authored-By: Claude Opus 5.5 --- README.md | 85 +++++++++------------ common/arg.cpp | 21 ++++++ common/common.cpp | 1 + common/common.h | 1 + docs/multi-gpu.md | 40 +++++----- ggml/include/ggml-backend.h | 3 + ggml/src/ggml-cuda/ggml-cuda.cu | 121 ++++++++++++++++-------------- include/llama.h | 13 ++++ src/llama-context.cpp | 13 ++++ src/llama.cpp | 12 +++ tests/test-llama-archs.cpp | 18 ++++- tools/llama-bench/llama-bench.cpp | 40 +++++++++- 12 files changed, 240 insertions(+), 128 deletions(-) diff --git a/README.md b/README.md index a8f38ffa9ca1..31017dbab6e4 100644 --- a/README.md +++ b/README.md @@ -1,55 +1,40 @@ # What this fork adds -Branch `master` mirrors [ggml-org/llama.cpp](https://github.com/ggml-org/llama.cpp). Branch `main` is upstream plus the two features below. The rest of this file is the upstream README. - -## Quantized-native CUDA FlashAttention - -The CUDA MMA FlashAttention kernel reads a quantized K/V cache in place, instead of casting the visible attention window to F16 first. - -- It is opt-in with `--flash-attn-native-quants` (`-fanq 1` in `llama-bench`). Without the flag everything runs as upstream. -- Cache types: `q4_0`, `q4_1`, `q5_0`, `q5_1` and `q8_0`, the same type for K and V. Head sizes 64, 128, 256 and 512, at every query batch size, on Turing (`sm_75`) and newer. -- Decode and the small batches of speculative decoding take tiles of 8 Q heads, so each K/V row is read once per KV head instead of once per Q head. -- It is faster at long context. Qwen3.8-27B at 16k context: - - token generation +6-11% on an RTX 3060 or 4070; - - 2-token speculative steps +14-31%. - - On both GPUs at 64k context the gain is up to +39%. On D=128 models, 4-token batches run up to twice as fast. -- The F16 copy of the window no longer takes compute buffer: 1.6 GiB less for Qwen3.8-27B at 262k context. -- `q4_0` and `q8_0` are always compiled. `q4_1`, `q5_0` and `q5_1` are compiled with their pair in `GGML_CUDA_FA_QUANTS` (for example `-DGGML_CUDA_FA_QUANTS="q5_1-q5_1"`) or with `all`. -- Results differ from the F16 path in rounding only. - -Details, validation and all measurements: [docs/quantized-native-flash-attention.md](docs/quantized-native-flash-attention.md). - -## Prompt cache hierarchy: slots, RAM, disk - -`--cache-disk PATH` adds a disk level below the RAM prompt cache (`--cache-ram`). - -- Prompts that RAM evicts are written to disk in the background, and are loaded back when a later request shares their prefix. -- On exit and before sleeping, the server also writes its slots and RAM entries. Conversations survive restarts and model swaps. - -How prompts are stored: - -- A prompt is stored as KV blocks of `--cache-block-size` tokens (default 2048), keyed by a hash of their prefix. -- Prompts that share a prefix share its blocks, in RAM and on disk. -- A growing conversation writes only its new blocks. In one 25k-token conversation with Qwen3.8-27B, the next turn wrote 0.16 GB instead of 2.3 GB. - -The new libllama functions `llama_state_seq_get_data_range()` and `llama_state_seq_set_data_range()` provide the blocks. All model types are supported: - -- Attention, sliding-window, recurrent and hybrid models work. -- Models that roll back only to context checkpoints (hybrid, recurrent, sliding-window) are stored at their newest checkpoint. -- Caches without state ranges store whole states. - -Size limit and sharing: - -- `--cache-disk-size` caps the directory (default 32 GiB). The least recently used prompt is removed first, and 5% of the file system always stays free. -- File names carry a fingerprint of the model files and the KV settings, so several models can share one directory. - -Two more server changes come with the disk level: - -- Slot selection: if a cached prompt shares at least 256 more tokens with a request than the most similar slot does, the server restores the cached prompt instead of processing those tokens again. One case is a slot that holds another conversation with the same system prompt. -- SIGINT and SIGTERM are handled on their own thread. A sleeping server now stops cleanly and writes its cache. - -Details: [tools/server/README.md](tools/server/README.md#prompt-cache-on-disk). +Branch `master` mirrors [ggml-org/llama.cpp](https://github.com/ggml-org/llama.cpp). Branch `main` is upstream plus the features below. Each one is off until you enable it; without the flags everything runs as upstream. The rest of this file is the upstream README. + +## Features + +- **Quantized-native FlashAttention (CUDA).** The FlashAttention kernel reads a quantized K/V cache in place instead of first copying the attention window to F16. + - Benefit: faster generation at long context (Qwen3.8-27B: +6-11% at 16k, up to +39% at 64k), and no F16 copy in the compute buffer (1.6 GiB less at 262k context). + - Enable: `--flash-attn-native-quants` (`-fanq 1` in llama-bench) with a quantized cache, the same type for K and V (`-ctk q5_1 -ctv q5_1`). `q4_1`, `q5_0` and `q5_1` need `-DGGML_CUDA_FA_QUANTS=` or `all` at build time. [Details](docs/quantized-native-flash-attention.md) +- **Prompt cache on disk.** A disk level below the RAM prompt cache. Prompts are stored as KV blocks keyed by their prefix, so conversations with a shared prefix share blocks, and a growing conversation writes only its new blocks. + - Benefit: conversations survive RAM eviction, restarts and model swaps, and resume without processing the prompt again. The next turn of a 25k-token conversation writes 0.16 GB instead of 2.3 GB. + - Enable: `--cache-disk PATH` (server), sized with `--cache-disk-size` (MiB) and `--cache-block-size` (tokens). [Details](tools/server/README.md#prompt-cache-on-disk) +- **Attention split** (`--split-mode tensor`). The attention heads and their KV cache get their own split across the GPUs, separate from `--tensor-split`. + - Benefit: with GPUs of different speed, the faster one takes more of the weights while the other one holds more of the KV cache. + - Enable: `--attn-split 1,3` (`-as`), next to e.g. `--tensor-split 0.625,0.375`. [Details](docs/multi-gpu.md) +- **AllReduce by size** (`--split-mode tensor`, two CUDA GPUs). Small reductions go through a kernel with a handshake instead of NCCL, and large ones through the copy engines. With `p2p`, the GPUs write directly into each other's memory. + - Benefit: lower latency per decode step, and prompt processing that keeps every SM free for compute. + - Enable: `--allreduce host` (no peer access needed) or `--allreduce p2p` (needs peer access between the GPUs; falls back to `host` without it). [Details](docs/multi-gpu.md#7-two-gpus-over-pcie---allreduce) +- **Split overlap** (`--split-mode tensor`). Large ubatches alternate between two schedulers, so the AllReduce of one ubatch overlaps with the compute of the next. + - Benefit: faster prompt processing when the link between the GPUs is slow (Qwen3.8-27B on two PCIe GPUs: 733 -> 947 t/s at 10k context). + - Enable: `--split-overlap`, with half the usual `--ubatch-size` (each scheduler has its own compute buffers). Best together with `--allreduce host` or `p2p`. [Details](docs/multi-gpu.md#8-overlapping-the-allreduce-with-compute---split-overlap) + +## Upstream vs this fork + +Setup: RTX 4070 (12 GB, PCIe 4.0 x16) + RTX 3060 (12 GB, PCIe 3.0 x4), Qwen3.8-27B UD-Q4_K_S with its MTP head (2 draft tokens), `llama-server` with 2 slots of 131k context, q5_1 K/V cache. + +- Upstream: `-sm tensor -fa on -ctk q5_1 -ctv q5_1`, the best split that fits. +- This fork: additionally `--flash-attn-native-quants -ts 0.625,0.375 --attn-split 1,3 --allreduce p2p --split-overlap -ub 256 --cache-disk ...`. + +| | upstream | this fork | +|---|--:|--:| +| prompt processing, 10k context (t/s) | | | +| prompt processing, 95k context (t/s) | | | +| generation, 10k context (t/s) | | | +| generation, 95k context (t/s) | | | +| resume a 25k-token conversation after a restart (s) | | | +| VRAM after a 95k request, 4070 / 3060 (MiB) | | | --- diff --git a/common/arg.cpp b/common/arg.cpp index 1eead7c308b6..606cc136ffce 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -1779,6 +1779,27 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.split_overlap = value; } ).set_env("LLAMA_ARG_SPLIT_OVERLAP")); + add_opt(common_arg( + {"--allreduce"}, "{default,host,p2p}", + "with --split-mode tensor on two CUDA GPUs: how the GPUs exchange their partial results (docs/multi-gpu.md)\n" + "- default: NCCL\n" + "- host: by size, through host memory: small reductions through a kernel with a handshake in pinned memory, " + "large ones through the copy engines (no SM busy, both directions of the link at once)\n" + "- p2p: by size, with CUDA peer access between the GPUs: small reductions through a kernel that writes into " + "the other GPU, large ones through the copy engines from GPU to GPU; host without peer access. Some " + "platforms report peer access but corrupt the transfers (IOMMU)", + [](common_params & params, const std::string & value) { + if (value == "default") { + params.allreduce_type = LLAMA_ALLREDUCE_TYPE_DEFAULT; + } else if (value == "host") { + params.allreduce_type = LLAMA_ALLREDUCE_TYPE_HOST; + } else if (value == "p2p") { + params.allreduce_type = LLAMA_ALLREDUCE_TYPE_P2P; + } else { + throw std::invalid_argument("invalid value for --allreduce: " + value); + } + } + ).set_env("LLAMA_ARG_ALLREDUCE")); add_opt(common_arg( {"--cache-idle-slots"}, {"--no-cache-idle-slots"}, diff --git a/common/common.cpp b/common/common.cpp index 8d58f1d40d8c..193d1904d647 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1734,6 +1734,7 @@ struct llama_context_params common_context_params_to_llama(const common_params & cparams.pooling_type = params.pooling_type; cparams.attention_type = params.attention_type; cparams.flash_attn_type = params.flash_attn_type; + cparams.allreduce_type = params.allreduce_type; cparams.cb_eval = params.cb_eval; cparams.cb_eval_user_data = params.cb_eval_user_data; cparams.offload_kqv = !params.no_kv_offload; diff --git a/common/common.h b/common/common.h index 0b98923331a0..063f8a205fe0 100644 --- a/common/common.h +++ b/common/common.h @@ -505,6 +505,7 @@ struct common_params { enum llama_pooling_type pooling_type = LLAMA_POOLING_TYPE_UNSPECIFIED; // pooling type for embeddings enum llama_attention_type attention_type = LLAMA_ATTENTION_TYPE_UNSPECIFIED; // attention type for embeddings enum llama_flash_attn_type flash_attn_type = LLAMA_FLASH_ATTN_TYPE_AUTO; // whether to use Flash Attention + enum llama_allreduce_type allreduce_type = LLAMA_ALLREDUCE_TYPE_DEFAULT; // split mode tensor: AllReduce paths bool flash_attn_native_quants = false; // read a quantized K/V cache in place instead of casting it to F16 bool split_overlap = false; // split mode tensor: overlap the AllReduce of a ubatch with the compute of the next diff --git a/docs/multi-gpu.md b/docs/multi-gpu.md index 44467ed74940..cb9fe78b97da 100644 --- a/docs/multi-gpu.md +++ b/docs/multi-gpu.md @@ -113,32 +113,36 @@ GGML_CUDA_P2P=1 llama-cli -m model.gguf -sm tensor P2P requires driver support (usually restricted to workstation/datacenter GPUs) and **may cause crashes or corrupted outputs on some motherboards or BIOS configurations** (e.g. when IOMMU is enabled). If you see instability after enabling it, unset the variable. -With two GPUs over PCIe and `--split-mode tensor`, it also switches the AllReduce to the peer-to-peer paths of recipe 7. `NCCL_P2P_DISABLE=1` keeps those off. +`--allreduce p2p` (recipe 7) uses peer access for the AllReduce of `--split-mode tensor` without this variable: it enables peer access between the two GPUs itself, with the same caveat. -### 7. Two GPUs over PCIe +### 7. Two GPUs over PCIe (`--allreduce`) -With `--split-mode tensor`, two GPUs exchange their partial results after every layer. With NCCL compiled in, each reduction takes one of these paths: +With `--split-mode tensor`, two GPUs exchange their partial results after every layer. By default NCCL does all of it. With NCCL compiled in, `--allreduce` picks a path for each reduction by its size instead: -| Link | Reduction size | Path | Why | +| `--allreduce` | Reduction size | Path | Why | |---|---|---|---| -| no peer access (the usual case for consumer cards), or `GGML_CUDA_P2P` not set | below 512 KiB (decode, small batches) | internal pipeline: a single kernel with a handshake through pinned memory | lower latency than NCCL | -| no peer access, or `GGML_CUDA_P2P` not set | from 512 KiB (prompt processing) | copy engines: each GPU copies its part to pinned memory in chunks, the other copies them back as they arrive | no SM is busy, both directions of a link are used at once | -| peer access over PCIe and `GGML_CUDA_P2P` set (recipe 6) | below 512 KiB | peer-to-peer kernel: each GPU writes its part and a flag into the memory of the other and polls only its own | lower latency than NCCL and the internal pipeline | -| peer access over PCIe and `GGML_CUDA_P2P` set | from 512 KiB | copy engines from GPU to GPU | as above, without the copy through host memory | -| NVLink (native peer atomics), or anything the paths above do not handle | | NCCL | | +| `host` | below 512 KiB (decode, small batches) | internal pipeline: a single kernel with a handshake through pinned memory | lower latency than NCCL | +| `host` | from 512 KiB (prompt processing) | copy engines: each GPU copies its part to pinned memory in chunks, the other copies them back as they arrive | no SM is busy, both directions of a link are used at once | +| `p2p` | below 512 KiB | peer-to-peer kernel: each GPU writes its part and a flag into the memory of the other and polls only its own | lower latency than NCCL and the internal pipeline | +| `p2p` | from 512 KiB | copy engines from GPU to GPU | as above, without the copy through host memory | +| `default` | all | NCCL | | -The paths matter most with `--split-overlap` (recipe 8): NCCL keeps an SM on each GPU busy, which slows down the compute of the other ubatch. Over PCIe, the bandwidth of the slower link limits a large reduction with or without peer access, so peer access mostly lowers the latency of the small ones. +`host` works without peer access, which consumer cards usually lack. `p2p` needs peer access between the GPUs (on consumer cards only with a patched driver) and falls back to `host` without it. Some platforms report peer access but corrupt or hang the transfers (IOMMU, see recipe 6). Over NVLink, keep `default`: NCCL uses the link natively. -Environment variables: +```bash +llama-server -m model.gguf -sm tensor --allreduce host --split-overlap -ub 256 +``` + +The paths matter most with `--split-overlap` (recipe 8): NCCL keeps an SM on each GPU busy, which slows down the compute of the other ubatch. Over PCIe, the bandwidth of the slower link limits a large reduction with or without peer access, so peer access mostly lowers the latency of the small ones. Measured with an RTX 4070 (PCIe 4.0 x16) and an RTX 3060 (PCIe 3.0 x4), Qwen3.8-27B, `--split-overlap -ub 256`, `default`/`host`/`p2p`: prompt processing 866/1178/1181 t/s, 3-token batches 88.8/97.4/99.5 t/s, generation 34.2/36.9/37.6 t/s (`default`: NCCL over peer access). + +For tuning, two environment variables move the thresholds: | Variable | Effect | |---|---| -| `GGML_CUDA_ALLREDUCE=nccl` | NCCL for every reduction | -| `GGML_CUDA_AR_NCCL_MIN_BYTES=N` | size below which the small-reduction path is used, and the default of `GGML_CUDA_AR_CE_MIN_BYTES` (0: NCCL for every reduction). The peer-to-peer kernel handles at most 512 KiB; the copy engines then start there. Setting it also enables the internal pipeline with peer access when the peer-to-peer kernel is off | -| `GGML_CUDA_AR_CE=0` | without peer access: NCCL instead of the copy engines for large reductions | -| `GGML_CUDA_AR_CE_P2P=0` or `1` | with peer access in use: copy engines from GPU to GPU off or on (default: on over PCIe) | -| `GGML_CUDA_AR_P2P_SMALL=0` or `1` | with peer access in use: peer-to-peer kernel off or on (default: on over PCIe) | -| `GGML_CUDA_AR_CE_MIN_BYTES=N` | size from which the copy engines are used | +| `GGML_CUDA_AR_NCCL_MIN_BYTES=N` | size below which the small path is used (default 512 KiB, the most the peer-to-peer kernel handles; 0: NCCL for every reduction) | +| `GGML_CUDA_AR_CE_MIN_BYTES=N` | size from which the copy engines are used (default: where the small path ends) | + +`GGML_CUDA_ALLREDUCE=nccl|internal|none` overrides `--allreduce`. The copy engines need 16 MiB of VRAM per GPU, plus 8 MiB of pinned memory without peer access, per context that has large reductions. They are allocated with the first large reduction. An MTP draft context processes the prompt batches as well, so it has its own buffers: 32 MiB of VRAM per GPU in total. The peer-to-peer kernel needs 512 KiB of VRAM per GPU and context. @@ -160,7 +164,7 @@ Notes: - Each scheduler reserves the compute buffers for a full ubatch. With half the `--ubatch-size`, the memory stays the same, and two ubatches are in flight at once. On the GPUs below, 256 per scheduler was faster than 128, 384 or 512. - Each scheduler captures its own CUDA graphs, which takes some more VRAM at run time (about 50 MiB here). -- The overlap works best with the copy-engine AllReduce of recipe 7: NCCL kernels hold an SM each, and the compute of the other ubatch then runs slower. +- The overlap works best with `--allreduce host` or `p2p` (recipe 7): NCCL kernels hold an SM each, and the compute of the other ubatch then runs slower. Qwen3.8-27B on an RTX 4070 (PCIe 4.0 x16) + RTX 3060 (PCIe 3.0 x4), `-ts 0.625,0.375 -as 1,3`, 262144 tokens of q5_1 cache in 2 slots, with MTP, prompt processing in t/s: diff --git a/ggml/include/ggml-backend.h b/ggml/include/ggml-backend.h index 273ee39c4ebf..e097ef5db60b 100644 --- a/ggml/include/ggml-backend.h +++ b/ggml/include/ggml-backend.h @@ -210,6 +210,9 @@ extern "C" { typedef void * (*ggml_backend_comm_init_t)(ggml_backend_t * backends, size_t n_backends); typedef void (*ggml_backend_comm_free_t)(void * comm_ctx); typedef bool (*ggml_backend_comm_allreduce_tensor_t)(void * comm_ctx, struct ggml_tensor ** tensors); + // the AllReduce paths of the comm contexts created after the call: "default" or a backend-specific mode (CUDA: "host", + // "p2p"), returns false for an unknown mode + typedef bool (*ggml_backend_comm_set_mode_t)(const char * mode); // Split buffer type for tensor parallelism (old) typedef ggml_backend_buffer_type_t (*ggml_backend_split_buffer_type_t)(int main_device, const float * tensor_split); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index cddeb0c0f8a6..10e4e851e721 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1247,76 +1247,86 @@ static void ggml_backend_cuda_comm_init_internal(ggml_backend_cuda_comm_context ggml_backend_cuda_comm_init_none(ret); } +// The AllReduce paths of the comm contexts created from now on (llama: --allreduce): "default" leaves every reduction +// to NCCL (or the internal pipeline without NCCL), "host" and "p2p" pick a path by size between two GPUs, see +// ggml_backend_cuda_comm_init_hybrid(). +enum ggml_cuda_comm_mode { + GGML_CUDA_COMM_MODE_DEFAULT, + GGML_CUDA_COMM_MODE_HOST, + GGML_CUDA_COMM_MODE_P2P, +}; + +static std::atomic ggml_cuda_comm_mode_cur { GGML_CUDA_COMM_MODE_DEFAULT }; + +static bool ggml_backend_cuda_comm_set_mode(const char * mode) { + if (mode == nullptr || strcmp(mode, "default") == 0) { + ggml_cuda_comm_mode_cur = GGML_CUDA_COMM_MODE_DEFAULT; + } else if (strcmp(mode, "host") == 0) { + ggml_cuda_comm_mode_cur = GGML_CUDA_COMM_MODE_HOST; + } else if (strcmp(mode, "p2p") == 0) { + ggml_cuda_comm_mode_cur = GGML_CUDA_COMM_MODE_P2P; + } else { + return false; + } + return true; +} + #ifdef GGML_USE_NCCL // Without peer access NCCL exchanges the data through host memory as well. A small reduction is then bound by the // latency of the handshake, where the in-kernel handshake of the internal pipeline is faster: with an RTX 4070 on // PCIe 4.0 x16 and an RTX 3060 on PCIe 3.0 x4, a Qwen3.8-27B decode step of 1-4 tokens runs 9-12% faster through it. -// From 32 tokens (640 KiB per reduction) on NCCL is faster. GGML_CUDA_AR_NCCL_MIN_BYTES sets the size from which -// NCCL is used, 0 uses NCCL for everything. +// From 32 tokens (640 KiB per reduction) on NCCL is faster. // // Larger reductions go through the copy engines instead of NCCL (see allreduce-ce.cuh): NCCL keeps an SM busy on each // GPU during the transfer, which pushes a kernel with one block per SM into a second wave, and does not use both -// directions of a link at the same time. On the same GPUs prompt processing gets 16% faster. GGML_CUDA_AR_CE=0 keeps -// NCCL for them, GGML_CUDA_AR_CE_MIN_BYTES sets the size from which the copy engines are used. +// directions of a link at the same time. On the same GPUs prompt processing gets 16% faster. +// +// With peer access ("p2p") the same holds: NCCL's small reductions have more latency than a kernel that writes into +// the other GPU and polls only its own memory, and a large one is bound by the link either way. The same GPUs with +// BAR1 P2P: tg 36.9 -> 37.6 t/s and pp3 97.5 -> 99.7 against the internal pipeline, with --split-overlap prompt +// processing 865 -> 1183 t/s against NCCL. Peer access is used only on request: some platforms report it but corrupt +// or hang the transfers (IOMMU). +// +// GGML_CUDA_AR_NCCL_MIN_BYTES sets the size from which the small path ends (default 512 KiB, the peer-to-peer kernel +// handles at most that, 0: NCCL for everything), GGML_CUDA_AR_CE_MIN_BYTES the size from which the copy engines are +// used (default: where the small path ends). static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * ret) { - if (ret->dev_ids.size() != 2) { + const int mode = ggml_cuda_comm_mode_cur; + if (mode == GGML_CUDA_COMM_MODE_DEFAULT || ret->dev_ids.size() != 2) { return; } - int can_access = 0; - CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access, ret->dev_ids[0], ret->dev_ids[1])); - - // With peer access over PCIe the same holds: NCCL's kernels hold an SM, its small reductions have more latency - // than a kernel that writes into the other GPU and polls only its own memory, and a large one is bound by the - // link either way. The same GPUs with BAR1 P2P: tg 36.9 -> 37.6 t/s and pp3 97.5 -> 99.7 against the internal - // pipeline, with --split-overlap prompt processing 865 -> 1183 t/s against NCCL. A link with native peer atomics - // (NVLink) keeps NCCL. GGML_CUDA_AR_P2P_SMALL and GGML_CUDA_AR_CE_P2P (0 or 1) override the choice. - // Peer access over PCIe is used only on request (GGML_CUDA_P2P, docs/multi-gpu.md), as in the rest of this - // backend: some platforms report it but corrupt or hang the transfers (IOMMU). NCCL_P2P_DISABLE also turns it - // off. Without it the GPUs are treated like GPUs without peer access. - int native_atomics = 0; -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) - if (can_access) { - CUDA_CHECK(cudaDeviceGetP2PAttribute(&native_atomics, cudaDevP2PAttrNativeAtomicSupported, ret->dev_ids[0], ret->dev_ids[1])); - } - const bool p2p_pcie = can_access && !native_atomics && ggml_cuda_p2p_enabled() && !ggml_cuda_env_flag("NCCL_P2P_DISABLE"); - const bool peer = p2p_pcie || (can_access && native_atomics); -#else - const bool p2p_pcie = false; - const bool peer = can_access != 0; -#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) - - auto env_flag = [](const char * name, bool def) { - const char * env = getenv(name); - return env != nullptr ? atoi(env) != 0 : def; - }; - - const char * env_min = getenv("GGML_CUDA_AR_NCCL_MIN_BYTES"); - const bool small_p2p = peer && env_flag("GGML_CUDA_AR_P2P_SMALL", p2p_pcie); - const bool small_internal = !small_p2p && (!peer || env_min != nullptr); - const bool large_ce = peer ? env_flag("GGML_CUDA_AR_CE_P2P", p2p_pcie) : env_flag("GGML_CUDA_AR_CE", true); - if (!small_p2p && !small_internal && !large_ce) { - return; + bool peer = false; + if (mode == GGML_CUDA_COMM_MODE_P2P) { + int can_access = 0; + CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access, ret->dev_ids[0], ret->dev_ids[1])); +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + can_access = 0; // the peer-to-peer kernel is CUDA only +#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + peer = can_access != 0; + if (!peer) { + GGML_LOG_WARN("%s: no peer access between CUDA%d and CUDA%d, AllReduce through host memory\n", + __func__, ret->dev_ids[0], ret->dev_ids[1]); + } } size_t min_bytes = 512*1024; - if (env_min != nullptr) { - min_bytes = std::strtoull(env_min, nullptr, 10); + if (const char * env = getenv("GGML_CUDA_AR_NCCL_MIN_BYTES")) { + min_bytes = std::strtoull(env, nullptr, 10); } - if (min_bytes == 0) { return; } - // the reductions it gets stay below the copy-engine threshold, its 32 MiB staging per GPU is not needed - // it needs Volta or newer, the copy engines do not - if (small_p2p) { + // the small path: the peer-to-peer kernel (it needs Volta or newer), otherwise the internal pipeline + if (peer) { ret->p2p_ar = ggml_cuda_p2p_ar_init(ret->dev_ids.data(), ret->dev_ids.size()); if (ret->p2p_ar != nullptr) { ret->nccl_min_bytes = std::min(min_bytes, ggml_cuda_p2p_ar_max_bytes()); } } - if (small_internal) { + if (ret->p2p_ar == nullptr) { + // the reductions it gets stay below the copy-engine threshold, its 32 MiB staging per GPU is not needed ret->ar_pipeline = ggml_cuda_ar_pipeline_init(ret->dev_ids.data(), ret->dev_ids.size(), /*small_only =*/ true); if (ret->ar_pipeline == nullptr) { (void) cudaGetLastError(); @@ -1325,13 +1335,11 @@ static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * } } - if (large_ce) { - ret->ce_ar = ggml_cuda_ce_ar_init(ret->dev_ids.data(), ret->dev_ids.size(), /*peer =*/ peer); - // the copy engines take over where the small path ends: the peer-to-peer kernel handles at most 512 KiB - ret->ce_min_bytes = small_p2p && ret->p2p_ar != nullptr ? ret->nccl_min_bytes : min_bytes; - if (const char * env = getenv("GGML_CUDA_AR_CE_MIN_BYTES")) { - ret->ce_min_bytes = std::strtoull(env, nullptr, 10); - } + // the large path: the copy engines take over where the small path ends + ret->ce_ar = ggml_cuda_ce_ar_init(ret->dev_ids.data(), ret->dev_ids.size(), /*peer =*/ peer); + ret->ce_min_bytes = ret->nccl_min_bytes > 0 ? ret->nccl_min_bytes : min_bytes; + if (const char * env = getenv("GGML_CUDA_AR_CE_MIN_BYTES")) { + ret->ce_min_bytes = std::strtoull(env, nullptr, 10); } if (ret->p2p_ar != nullptr) { @@ -1340,9 +1348,9 @@ static void ggml_backend_cuda_comm_init_hybrid(ggml_backend_cuda_comm_context * ret->ce_ar != nullptr ? "the copy engines" : "NCCL"); } else if (ret->ar_pipeline != nullptr && ret->ce_ar != nullptr) { GGML_LOG_INFO("%s: AllReduce below %zu bytes through the internal pipeline, from %zu bytes through the copy engines\n", - __func__, min_bytes, ret->ce_min_bytes); + __func__, ret->nccl_min_bytes, ret->ce_min_bytes); } else if (ret->ar_pipeline != nullptr) { - GGML_LOG_INFO("%s: AllReduce below %zu bytes through the internal pipeline, NCCL above\n", __func__, min_bytes); + GGML_LOG_INFO("%s: AllReduce below %zu bytes through the internal pipeline, NCCL above\n", __func__, ret->nccl_min_bytes); } else if (ret->ce_ar != nullptr) { GGML_LOG_INFO("%s: AllReduce from %zu bytes through the copy engines, NCCL below\n", __func__, ret->ce_min_bytes); } @@ -6193,6 +6201,9 @@ static void * ggml_backend_cuda_reg_get_proc_address(ggml_backend_reg_t reg, con if (strcmp(name, "ggml_backend_comm_allreduce_tensor") == 0) { return (void *)ggml_backend_cuda_comm_allreduce_tensor; } + if (strcmp(name, "ggml_backend_comm_set_mode") == 0) { + return (void *)ggml_backend_cuda_comm_set_mode; + } if (strcmp(name, "ggml_backend_register_host_buffer") == 0) { return (void *)ggml_backend_cuda_register_host_buffer; } diff --git a/include/llama.h b/include/llama.h index 7f325eb4b1f5..afda00d54b13 100644 --- a/include/llama.h +++ b/include/llama.h @@ -197,6 +197,18 @@ extern "C" { LLAMA_API const char * llama_flash_attn_type_name(enum llama_flash_attn_type flash_attn_type); + // how the GPUs of split mode tensor exchange their partial results (two CUDA GPUs, NCCL compiled in) + enum llama_allreduce_type { + LLAMA_ALLREDUCE_TYPE_DEFAULT = 0, // NCCL (the internal pipeline without NCCL) + LLAMA_ALLREDUCE_TYPE_HOST = 1, // by size, through host memory: small through a kernel with a handshake + // in pinned memory, large through the copy engines + LLAMA_ALLREDUCE_TYPE_P2P = 2, // by size, with peer access: small through a kernel that writes into the + // other GPU, large through the copy engines from GPU to GPU (HOST without + // peer access) + }; + + LLAMA_API const char * llama_allreduce_type_name(enum llama_allreduce_type allreduce_type); + enum llama_split_mode { LLAMA_SPLIT_MODE_NONE = 0, // single GPU LLAMA_SPLIT_MODE_LAYER = 1, // split layers and KV across GPUs @@ -383,6 +395,7 @@ extern "C" { enum llama_pooling_type pooling_type; // whether to pool (sum) embedding results by sequence id enum llama_attention_type attention_type; // attention type to use for embeddings enum llama_flash_attn_type flash_attn_type; // when to enable Flash Attention + enum llama_allreduce_type allreduce_type; // split mode tensor: the AllReduce paths between the GPUs // ref: https://github.com/ggml-org/llama.cpp/pull/2054 float rope_freq_base; // RoPE base frequency, 0 = from model diff --git a/src/llama-context.cpp b/src/llama-context.cpp index bf417af82589..bcedc64bd969 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -401,6 +401,9 @@ llama_context::llama_context( LLAMA_LOG_INFO("%s: n_ubatch = %u\n", __func__, cparams.n_ubatch); LLAMA_LOG_INFO("%s: causal_attn = %d\n", __func__, cparams.causal_attn); LLAMA_LOG_INFO("%s: flash_attn = %s\n", __func__, llama_flash_attn_type_name(params.flash_attn_type)); + if (params.allreduce_type != LLAMA_ALLREDUCE_TYPE_DEFAULT) { + LLAMA_LOG_INFO("%s: allreduce = %s\n", __func__, llama_allreduce_type_name(params.allreduce_type)); + } LLAMA_LOG_INFO("%s: kv_unified = %s\n", __func__, cparams.kv_unified ? "true" : "false"); LLAMA_LOG_INFO("%s: fa_native_quants = %s\n", __func__, cparams.flash_attn_native_quants ? "true" : "false"); LLAMA_LOG_INFO("%s: freq_base = %.1f\n", __func__, cparams.rope_freq_base); @@ -420,6 +423,15 @@ llama_context::llama_context( } if (!hparams.vocab_only) { + // the AllReduce paths of the meta backends created below, in the backends that offer a choice (process-wide) + for (size_t i = 0; i < ggml_backend_reg_count(); ++i) { + auto * set_mode = (ggml_backend_comm_set_mode_t) ggml_backend_reg_get_proc_address(ggml_backend_reg_get(i), "ggml_backend_comm_set_mode"); + if (set_mode != nullptr && !set_mode(llama_allreduce_type_name(params.allreduce_type))) { + LLAMA_LOG_WARN("%s: %s does not support allreduce '%s'\n", __func__, + ggml_backend_reg_name(ggml_backend_reg_get(i)), llama_allreduce_type_name(params.allreduce_type)); + } + } + // GPU backends for (const auto & dev : model.devices) { ggml_backend_t backend = ggml_backend_dev_init(dev.dev, nullptr); @@ -4130,6 +4142,7 @@ llama_context_params llama_context_default_params() { /*.pooling_type =*/ LLAMA_POOLING_TYPE_UNSPECIFIED, /*.attention_type =*/ LLAMA_ATTENTION_TYPE_UNSPECIFIED, /*.flash_attn_type =*/ LLAMA_FLASH_ATTN_TYPE_AUTO, + /*.allreduce_type =*/ LLAMA_ALLREDUCE_TYPE_DEFAULT, /*.rope_freq_base =*/ 0.0f, /*.rope_freq_scale =*/ 0.0f, /*.yarn_ext_factor =*/ -1.0f, diff --git a/src/llama.cpp b/src/llama.cpp index 9beb14841f98..741a08bb84d9 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -47,6 +47,18 @@ const char * llama_flash_attn_type_name(enum llama_flash_attn_type flash_attn_ty GGML_ABORT("fatal error"); } +const char * llama_allreduce_type_name(enum llama_allreduce_type allreduce_type) { + switch (allreduce_type) { + case LLAMA_ALLREDUCE_TYPE_DEFAULT: + return "default"; + case LLAMA_ALLREDUCE_TYPE_HOST: + return "host"; + case LLAMA_ALLREDUCE_TYPE_P2P: + return "p2p"; + } + GGML_ABORT("fatal error"); +} + const char * llama_load_mode_name(enum llama_load_mode load_mode) { switch (load_mode) { case LLAMA_LOAD_MODE_AUTO: diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp index f6193014896e..f770fb83b629 100644 --- a/tests/test-llama-archs.cpp +++ b/tests/test-llama-archs.cpp @@ -516,7 +516,8 @@ static std::pair get_model_and_ctx( const std::vector & devs, const llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER, bool encode = false, const llama_model_tensor_buft_override * tensor_buft_overrides = nullptr, const size_t moe_cache_size = 0, - const std::vector & attn_split = {}, const bool split_overlap = false) { + const std::vector & attn_split = {}, const bool split_overlap = false, + const llama_allreduce_type allreduce_type = LLAMA_ALLREDUCE_TYPE_DEFAULT) { GGML_ASSERT((gguf_ctx == nullptr) != (file == nullptr)); llama_model_params model_params = llama_model_default_params(); model_params.progress_callback = silent_model_load_progress; @@ -537,6 +538,7 @@ static std::pair get_model_and_ctx( ctx_params.n_threads_batch = 4; // the ubatches of 64 tokens alternate between the lanes of the split overlap ctx_params.split_overlap = split_overlap; + ctx_params.allreduce_type = allreduce_type; if (!encode) { ctx_params.n_ubatch = 64; } @@ -962,11 +964,12 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con size_t moe_cache_size; std::vector attn_split; // empty: follow the tensor split bool split_overlap; + llama_allreduce_type allreduce_type; device_config(std::vector devs, std::string name, llama_split_mode split_mode, bool host_experts = false, size_t moe_cache_size = 0, - std::vector attn_split = {}, bool split_overlap = false) + std::vector attn_split = {}, bool split_overlap = false, llama_allreduce_type allreduce_type = LLAMA_ALLREDUCE_TYPE_DEFAULT) : devs(std::move(devs)), label(std::move(name)), split_mode(split_mode), host_experts(host_experts), moe_cache_size(moe_cache_size), - attn_split(std::move(attn_split)), split_overlap(split_overlap) {} + attn_split(std::move(attn_split)), split_overlap(split_overlap), allreduce_type(allreduce_type) {} }; const llama_model_tensor_buft_override host_experts_overrides[] = { @@ -1011,6 +1014,13 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con // consecutive ubatches on two schedulers whose AllReduce overlaps dev_configs.emplace_back(devices_meta, "Meta -so", LLAMA_SPLIT_MODE_TENSOR, false, 0, std::vector{}, true); + + // the AllReduce paths by size (two CUDA GPUs with NCCL, otherwise the same as "Meta") + dev_configs.emplace_back(devices_meta, "Meta -ar host", LLAMA_SPLIT_MODE_TENSOR, false, 0, std::vector{}, false, + LLAMA_ALLREDUCE_TYPE_HOST); + dev_configs.emplace_back(devices_meta, "Meta -ar p2p -so", LLAMA_SPLIT_MODE_TENSOR, false, 0, std::vector{}, true, + LLAMA_ALLREDUCE_TYPE_P2P); + max_device_label_length = std::max(max_device_label_length, dev_configs.back().label.length()); } } @@ -1117,7 +1127,7 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con } if (dc.split_mode != LLAMA_SPLIT_MODE_TENSOR || llm_arch_supports_sm_tensor(arch)) { test_executed = true; - model_and_ctx_dev = get_model_and_ctx(gguf_ctx.get(), nullptr, seed, stdev, dc.devs, dc.split_mode, encode, overrides, dc.moe_cache_size, dc.attn_split, dc.split_overlap); + model_and_ctx_dev = get_model_and_ctx(gguf_ctx.get(), nullptr, seed, stdev, dc.devs, dc.split_mode, encode, overrides, dc.moe_cache_size, dc.attn_split, dc.split_overlap, dc.allreduce_type); logits_dev = get_logits(model_and_ctx_dev.first.get(), model_and_ctx_dev.second.get(), tokens, encode); const double nmse_val = nmse(logits_cpu, logits_dev); snprintf(nmse_str, sizeof(nmse_str), "(%.2e)", nmse_val); diff --git a/tools/llama-bench/llama-bench.cpp b/tools/llama-bench/llama-bench.cpp index 330671f16c79..1a24a806b435 100644 --- a/tools/llama-bench/llama-bench.cpp +++ b/tools/llama-bench/llama-bench.cpp @@ -360,6 +360,7 @@ struct cmd_params { std::vector flash_attn; std::vector flash_attn_native_quants; std::vector split_overlap; + std::vector allreduce; std::vector> devices; std::vector> tensor_split; std::vector> attn_split; @@ -409,6 +410,7 @@ static const cmd_params cmd_params_defaults = { /* flash_attn */ { LLAMA_FLASH_ATTN_TYPE_AUTO }, /* flash_attn_native_quants */ { false }, /* split_overlap */ { false }, + /* allreduce */ { LLAMA_ALLREDUCE_TYPE_DEFAULT }, /* devices */ { {} }, /* tensor_split */ { std::vector(llama_max_devices(), 0.0f) }, /* attn_split */ { std::vector(llama_max_devices(), 0.0f) }, @@ -484,6 +486,7 @@ static void print_usage(int /* argc */, char ** argv) { printf(" -fa, --flash-attn (default: %s)\n", join(transform_to_str(cmd_params_defaults.flash_attn, llama_flash_attn_type_name), ",").c_str()); printf(" -fanq, --flash-attn-native-quants <0|1> (default: %s)\n", join(cmd_params_defaults.flash_attn_native_quants, ",").c_str()); printf(" -so, --split-overlap <0|1> (default: %s)\n", join(cmd_params_defaults.split_overlap, ",").c_str()); + printf(" -ar, --allreduce (default: %s)\n", join(transform_to_str(cmd_params_defaults.allreduce, llama_allreduce_type_name), ",").c_str()); printf(" -dev, --device (default: auto)\n"); printf(" -lm, --load-mode (default: %s)\n", join(transform_to_str(cmd_params_defaults.load_mode, llama_load_mode_name), ",").c_str()); printf(" -lzm, --lazy-mode (default: %s)\n", join(transform_to_str(cmd_params_defaults.lazy_mode, lazy_mode_str), ",").c_str()); @@ -862,6 +865,23 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { } auto p = string_split(argv[i], split_delim); params.split_overlap.insert(params.split_overlap.end(), p.begin(), p.end()); + } else if (arg == "-ar" || arg == "--allreduce") { + if (++i >= argc) { + invalid_param = true; + break; + } + for (const auto & v : string_split(argv[i], split_delim)) { + if (v == "default") { + params.allreduce.push_back(LLAMA_ALLREDUCE_TYPE_DEFAULT); + } else if (v == "host") { + params.allreduce.push_back(LLAMA_ALLREDUCE_TYPE_HOST); + } else if (v == "p2p") { + params.allreduce.push_back(LLAMA_ALLREDUCE_TYPE_P2P); + } else { + invalid_param = true; + break; + } + } } else if (arg == "-nkvo" || arg == "--no-kv-offload") { if (++i >= argc) { invalid_param = true; @@ -1195,6 +1215,9 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { if (params.split_overlap.empty()) { params.split_overlap = cmd_params_defaults.split_overlap; } + if (params.allreduce.empty()) { + params.allreduce = cmd_params_defaults.allreduce; + } if (params.devices.empty()) { params.devices = cmd_params_defaults.devices; } @@ -1264,6 +1287,7 @@ struct cmd_params_instance { llama_flash_attn_type flash_attn; bool flash_attn_native_quants; bool split_overlap; + llama_allreduce_type allreduce; std::vector devices; std::vector tensor_split; std::vector attn_split; @@ -1352,6 +1376,7 @@ struct cmd_params_instance { cparams.flash_attn_type = flash_attn; cparams.flash_attn_native_quants = flash_attn_native_quants; cparams.split_overlap = split_overlap; + cparams.allreduce_type = allreduce; cparams.embeddings = embeddings; cparams.op_offload = !no_op_offload; cparams.swa_full = false; @@ -1390,6 +1415,7 @@ static std::vector get_cmd_params_instances(const cmd_param for (const auto & fa : params.flash_attn) for (const auto & fanq : params.flash_attn_native_quants) for (const auto & so : params.split_overlap) + for (const auto & ar : params.allreduce) for (const auto & nt : params.n_threads) for (const auto & cm : params.cpu_mask) for (const auto & cs : params.cpu_strict) @@ -1422,6 +1448,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .flash_attn = */ fa, /* .flash_attn_native_quants = */ fanq, /* .split_overlap = */ so, + /* .allreduce = */ ar, /* .devices = */ devs, /* .tensor_split = */ ts, /* .attn_split = */ as, @@ -1463,6 +1490,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .flash_attn = */ fa, /* .flash_attn_native_quants = */ fanq, /* .split_overlap = */ so, + /* .allreduce = */ ar, /* .devices = */ devs, /* .tensor_split = */ ts, /* .attn_split = */ as, @@ -1504,6 +1532,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .flash_attn = */ fa, /* .flash_attn_native_quants = */ fanq, /* .split_overlap = */ so, + /* .allreduce = */ ar, /* .devices = */ devs, /* .tensor_split = */ ts, /* .attn_split = */ as, @@ -1550,6 +1579,7 @@ struct test { llama_flash_attn_type flash_attn; bool flash_attn_native_quants; bool split_overlap; + llama_allreduce_type allreduce; std::vector devices; std::vector tensor_split; std::vector attn_split; @@ -1594,6 +1624,7 @@ struct test { flash_attn = inst.flash_attn; flash_attn_native_quants = inst.flash_attn_native_quants; split_overlap = inst.split_overlap; + allreduce = inst.allreduce; devices = inst.devices; tensor_split = inst.tensor_split; attn_split = inst.attn_split; @@ -1659,7 +1690,7 @@ struct test { "model_filename", "model_type", "model_size", "model_n_params", "n_batch", "n_ubatch", "n_threads", "cpu_mask", "cpu_strict", "poll", "type_k", "type_v", "n_gpu_layers", "n_cpu_moe", "split_mode", - "main_gpu", "no_kv_offload", "flash_attn", "flash_attn_native_quants", "split_overlap", "devices", "tensor_split", + "main_gpu", "no_kv_offload", "flash_attn", "flash_attn_native_quants", "split_overlap", "allreduce", "devices", "tensor_split", "attn_split", "tensor_buft_overrides", "load_mode", "lazy_mode", "embeddings", "no_op_offload", "no_host", "repack", "fit_target", "fit_min_ctx", @@ -1760,6 +1791,7 @@ struct test { std::to_string((int) flash_attn), std::to_string(flash_attn_native_quants), std::to_string(split_overlap), + llama_allreduce_type_name(allreduce), devices_to_string(devices), tensor_split_str, attn_split_str, @@ -1997,6 +2029,9 @@ struct markdown_printer : public printer { if (field == "split_overlap") { return "so"; } + if (field == "allreduce") { + return "ar"; + } if (field == "load_mode") { return "lm"; } @@ -2090,6 +2125,9 @@ struct markdown_printer : public printer { if (params.split_overlap.size() > 1 || params.split_overlap != cmd_params_defaults.split_overlap) { fields.emplace_back("split_overlap"); } + if (params.allreduce.size() > 1 || params.allreduce != cmd_params_defaults.allreduce) { + fields.emplace_back("allreduce"); + } if (params.devices.size() > 1 || params.devices != cmd_params_defaults.devices) { fields.emplace_back("devices"); } From 15b9b0a1ea5eb4974973e76c2aab1a33cb24be97 Mon Sep 17 00:00:00 2001 From: piggidragon Date: Fri, 16 Oct 2026 17:54:52 +0200 Subject: [PATCH 4/4] tests, readme : run the copy engines in the -ar configs, list the always-on server changes - test-llama-archs: the reductions of the test models stay below 512 KiB, so the -ar configs never used the copy engines; they now create their contexts with GGML_CUDA_AR_CE_MIN_BYTES=4096 - README: the RAM prompt cache changes of the server are on without a flag Co-Authored-By: Claude Opus 5.5 --- README.md | 9 +++++++-- tests/test-llama-archs.cpp | 27 +++++++++++++++++++++++++++ 2 files changed, 34 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 31017dbab6e4..42fecbb6470e 100644 --- a/README.md +++ b/README.md @@ -1,6 +1,6 @@ # What this fork adds -Branch `master` mirrors [ggml-org/llama.cpp](https://github.com/ggml-org/llama.cpp). Branch `main` is upstream plus the features below. Each one is off until you enable it; without the flags everything runs as upstream. The rest of this file is the upstream README. +Branch `master` mirrors [ggml-org/llama.cpp](https://github.com/ggml-org/llama.cpp). Branch `main` is upstream plus the features below. Each one is off until its flag enables it, except the changes to the server's RAM prompt cache, which are always on (see the prompt cache item). The rest of this file is the upstream README. ## Features @@ -9,7 +9,12 @@ Branch `master` mirrors [ggml-org/llama.cpp](https://github.com/ggml-org/llama.c - Enable: `--flash-attn-native-quants` (`-fanq 1` in llama-bench) with a quantized cache, the same type for K and V (`-ctk q5_1 -ctv q5_1`). `q4_1`, `q5_0` and `q5_1` need `-DGGML_CUDA_FA_QUANTS=` or `all` at build time. [Details](docs/quantized-native-flash-attention.md) - **Prompt cache on disk.** A disk level below the RAM prompt cache. Prompts are stored as KV blocks keyed by their prefix, so conversations with a shared prefix share blocks, and a growing conversation writes only its new blocks. - Benefit: conversations survive RAM eviction, restarts and model swaps, and resume without processing the prompt again. The next turn of a 25k-token conversation writes 0.16 GB instead of 2.3 GB. - - Enable: `--cache-disk PATH` (server), sized with `--cache-disk-size` (MiB) and `--cache-block-size` (tokens). [Details](tools/server/README.md#prompt-cache-on-disk) + - Enable: `--cache-disk PATH` (server), sized with `--cache-disk-size` (MiB). [Details](tools/server/README.md#prompt-cache-on-disk) + - Always on, with the RAM prompt cache (`--cache-ram`, on by default): + - The RAM cache stores prompts as KV blocks of `--cache-block-size` tokens (default 2048). + - A cached prompt keeps only its newest 3 context checkpoints. + - The server restores a cached prompt into a slot when that prompt shares at least 256 more tokens with the request than the slot does. + - SIGINT and SIGTERM are handled on their own thread. - **Attention split** (`--split-mode tensor`). The attention heads and their KV cache get their own split across the GPUs, separate from `--tensor-split`. - Benefit: with GPUs of different speed, the faster one takes more of the weights while the other one holds more of the KV cache. - Enable: `--attn-split 1,3` (`-as`), next to e.g. `--tensor-split 0.625,0.375`. [Details](docs/multi-gpu.md) diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp index f770fb83b629..3c10a2f43ccb 100644 --- a/tests/test-llama-archs.cpp +++ b/tests/test-llama-archs.cpp @@ -15,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -511,6 +512,22 @@ static bool silent_model_load_progress(float /*progress*/, void * /*user_data*/) return true; } +static void test_setenv(const char * name, const char * value) { +#ifdef _WIN32 + _putenv_s(name, value); +#else + setenv(name, value, 1); +#endif // _WIN32 +} + +static void test_unsetenv(const char * name) { +#ifdef _WIN32 + _putenv_s(name, ""); +#else + unsetenv(name); +#endif // _WIN32 +} + static std::pair get_model_and_ctx( struct gguf_context * gguf_ctx, FILE * file, const size_t seed, const float stdev, const std::vector & devs, @@ -555,7 +572,17 @@ static std::pair get_model_and_ctx( if (!model) { throw std::runtime_error("failed to create llama model"); } + // the reductions of the test models stay below the 512 KiB from which the AllReduce paths by size use the copy + // engines: with a threshold of 4 KiB a single token still takes the small path and the ubatches the copy engines + // (read when the context creates its comm contexts) + const bool ce_small = allreduce_type != LLAMA_ALLREDUCE_TYPE_DEFAULT && getenv("GGML_CUDA_AR_CE_MIN_BYTES") == nullptr; + if (ce_small) { + test_setenv("GGML_CUDA_AR_CE_MIN_BYTES", "4096"); + } llama_context_ptr lctx(llama_init_from_model(model.get(), ctx_params)); + if (ce_small) { + test_unsetenv("GGML_CUDA_AR_CE_MIN_BYTES"); + } if (!lctx) { throw std::runtime_error("failed to create llama context"); }