diff --git a/README.md b/README.md index a8f38ffa9ca..42fecbb6470 100644 --- a/README.md +++ b/README.md @@ -1,55 +1,45 @@ # 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 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 + +- **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). [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) +- **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 1eead7c308b..606cc136ffc 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 8d58f1d40d8..193d1904d64 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 0b98923331a..063f8a205fe 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 c6514373315..cb9fe78b97d 100644 --- a/docs/multi-gpu.md +++ b/docs/multi-gpu.md @@ -113,28 +113,38 @@ 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 +`--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. -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: +### 7. Two GPUs over PCIe (`--allreduce`) -| 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 | | +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: + +| `--allreduce` | Reduction size | Path | Why | +|---|---|---|---| +| `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 | | -With peer access, NCCL handles every reduction. +`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 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_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 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`) @@ -154,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 273ee39c4eb..e097ef5db60 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/allreduce-ce.cu b/ggml/src/ggml-cuda/allreduce-ce.cu index 0106b2c16e2..4f02b993ae6 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 38302dcd25e..baaa88724b2 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 00000000000..3dec00ce86d --- /dev/null +++ b/ggml/src/ggml-cuda/allreduce-p2p.cu @@ -0,0 +1,169 @@ +#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. 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; +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 + 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, 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; + + 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 unsigned int *) &peer_flags[blockIdx.x] = token; + + while (*(const volatile unsigned int *) &own_flags[blockIdx.x] - token > 0x7fffffffu) { + } + __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(unsigned int)) == cudaSuccess; + ok = ok && cudaMemset(ar->flags[i], 0, GGML_CUDA_P2P_AR_BLOCKS*sizeof(unsigned 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 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) { + 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 00000000000..0b8a7047b47 --- /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 e9d25ec5fe1..10e4e851e72 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; @@ -1238,61 +1247,110 @@ 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])); + 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 (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 - return; } - 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; + // 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()); + } } - - // 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()); - 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->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(); + } else { + ret->nccl_min_bytes = min_bytes; } } - if (ret->ar_pipeline != nullptr && ret->ce_ar != nullptr) { + // 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) { + 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); + __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); } @@ -6143,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 7f325eb4b1f..afda00d54b1 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 bf417af8258..bcedc64bd96 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 9beb14841f9..741a08bb84d 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 f6193014896..3c10a2f43cc 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,12 +512,29 @@ 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, 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 +555,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; } @@ -553,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"); } @@ -962,11 +991,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 +1041,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 +1154,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 330671f16c7..1a24a806b43 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"); }