From 2db092842655bfff572758c82118a7e6fe10c863 Mon Sep 17 00:00:00 2001 From: Leonccaa <166551845+Leonccaa@users.noreply.github.com> Date: Tue, 27 Oct 2026 16:20:02 -0700 Subject: [PATCH 1/2] docs: 0.1.4 follow-up fixes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A consistency review of the docs after the 0.1.4 release found stale versions and a few statements that no longer match the release notes or each other. - Versions: getting started, BUILDING and the DFlash2 page now describe 0.1.4 (install commands, image tag, recommended DFlash2 configuration and its KV pool: 1.37 million tokens, 2.61× at 512K with probabilistic drafting). - Host memory: the default pinned PLE table needs at least 80 GiB (smallest tested limit that passed memory stress, KV offload off), 128 GiB recommended; the experimental whole-table PLE disk option ran under 64 GiB and 32 GiB limits. README, support scope, configuration and known limits now say the same. - Positioning: four or more V100 32 GB GPUs, validated on four, in the support scope as in the README. - New: `--watchdog-config` in launch options; the 0.1.4 known issue for truncated `tool_use` on `/v1/messages`; the 0.1.4 speed statements from the release notes in performance; the first start after upgrading from 0.1.3. - Numbers: whole-table PLE cold first token 1.6–2.3 times (from the table); DFlash2 needs about 2.4 times MTP4's 21 KB per token in the CPU pool and on disk, so a 48 GiB pool holds about 1M tokens with DFlash2 (estimate); the 2.2–2.4 GB/s figure is the disk read rate inside the restore. - Links: a support matrix anchor in known limits; CONTRIBUTING explains that the large-files list is optional. Co-Authored-By: Claude Opus 5.5 Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com> --- .github/ISSUE_TEMPLATE/bug_report.yml | 2 +- .github/MULTIMODAL_INPUTS.md | 13 +++++++---- .github/SUPPORT.md | 7 +++--- BUILDING.md | 4 ++-- CONTRIBUTING.md | 6 +++--- README.md | 10 ++++----- docs/configuration.md | 28 ++++++++++++++++++++---- docs/dflash2.md | 9 ++++++-- docs/getting-started.md | 31 +++++++++++++++------------ docs/known-limits.md | 21 +++++++++++++----- docs/kv-offload.md | 11 ++++++---- docs/performance.md | 22 +++++++++++++++---- 12 files changed, 113 insertions(+), 51 deletions(-) diff --git a/.github/ISSUE_TEMPLATE/bug_report.yml b/.github/ISSUE_TEMPLATE/bug_report.yml index 19d865b6e..63e0c9c2c 100644 --- a/.github/ISSUE_TEMPLATE/bug_report.yml +++ b/.github/ISSUE_TEMPLATE/bug_report.yml @@ -15,7 +15,7 @@ body: - **Unverified, best effort:** TP2/TP8, V100 16 GB, other GPU counts, other models and other topologies. - **Out of scope:** GPUs other than SM70, ROCm/CPU/XPU/TPU, GGUF. - INT8 KV with the SSD tier is not supported in 0.1.0. + INT8 KV with the SSD tier is not supported in 0.1.4 either. Remove credentials, tokens and private paths from logs. Report security vulnerabilities through GitHub private vulnerability diff --git a/.github/MULTIMODAL_INPUTS.md b/.github/MULTIMODAL_INPUTS.md index 832750ea5..735d0df30 100644 --- a/.github/MULTIMODAL_INPUTS.md +++ b/.github/MULTIMODAL_INPUTS.md @@ -34,10 +34,15 @@ defaults above: - With `--limit-mm-per-prompt` that leaves out `"video"`, video is set to 0 (disabled). Other modalities the value leaves out keep the defaults above. -This V100 behaviour is the same as in 0.1.0 and 0.1.1, and 0.1.2 keeps it. Set -`--limit-mm-per-prompt` explicitly and include `"video"`. To copy: - -- the generic 0.1.2 default: `--limit-mm-per-prompt '{"image": 8, "video": 1, "audio": 1}'`; +This V100 behaviour is the same as in 0.1.0 and 0.1.1, and 0.1.2 through 0.1.4 keep it. +Set `--limit-mm-per-prompt` explicitly and include `"video"`. To copy: + +- the generic default counts (8 images, 1 video, 1 audio item): + `--limit-mm-per-prompt '{"image": 8, "video": 1, "audio": 1}'`. Qwen3.8 Flash-Next has + no audio input, so the shorter `--limit-mm-per-prompt '{"image": 8, "video": 1}'` — the + form used in [Configuration](../docs/configuration.md#launch-options) and + [Support and known limits](../docs/known-limits.md) — gives the same image and video + counts; - the reference deployment, as in the README launch command: `--limit-mm-per-prompt '{"image": 16, "video": 1}'` (Qwen3.8 Flash-Next has no audio input). diff --git a/.github/SUPPORT.md b/.github/SUPPORT.md index dd8d75509..ce43a2077 100644 --- a/.github/SUPPORT.md +++ b/.github/SUPPORT.md @@ -1,15 +1,16 @@ # ReVolta support scope -The [support matrix in the README](../README.md#support-matrix) is the reference. +The [support matrix](../docs/known-limits.md#support-matrix) in +[Support and known limits](../docs/known-limits.md) is the reference. In short: | Area | Scope | | --- | --- | -| GPUs | 4 × NVIDIA V100 32 GB (PCIe or SXM2) | +| GPUs | Four or more NVIDIA V100 32 GB (PCIe or SXM2); validated on four | | PCIe switch with P2P | Recommended; validated | | PCIe without P2P | Validated by disabling P2P on P2P-capable hardware; about 30–55% of P2P speed | | NVLink | Recommended; tested with community results on 4 × V100-SXM2 32 GB ([#2](https://github.com/Leonccaa/ReVolta/issues/2), [#38](https://github.com/Leonccaa/ReVolta/issues/38), [#37](https://github.com/Leonccaa/ReVolta/issues/37)) | -| Host memory | 128 GiB recommended; 80 GiB measured minimum with KV offload off | +| Host memory | 128 GiB recommended; 80 GiB measured minimum for the default pinned PLE table with KV offload off; the experimental whole-table PLE disk option ran under 64 GiB and 32 GiB limits | | KV offload | Off, CPU, or CPU + SSD tier | | KV cache | E4M3 (default) or INT8; INT8 with the SSD tier is not supported | | Models | Qwen3.8 Flash-Next (QSA, multimodal) | diff --git a/BUILDING.md b/BUILDING.md index 2d86b85dc..3f3f66428 100644 --- a/BUILDING.md +++ b/BUILDING.md @@ -5,7 +5,7 @@ The release version is `RELEASE_VERSION` in `setup.py` (initially `0.1.0`). Set `VLLM_VERSION_OVERRIDE` for an exact wheel and runtime version, or pass `SETUPTOOLS_SCM_PRETEND_VERSION` to the Docker build. Inherited Git tags do not select the ReVolta version. Use a PEP 440 local version for private deployments, -for example `0.1.3+deployment.20261025.g0123456789`. +for example `0.1.4+deployment.20261025.g0123456789`. ## SM70 source build @@ -15,7 +15,7 @@ Use CUDA 12.8; CUDA 13 cannot compile SM70. Limit parallelism to available memor NVCC thread per job are a conservative starting point on a 32 GB host. ```bash -version=0.1.3 +version=0.1.4 docker build --target build -f docker/Dockerfile \ --build-arg torch_cuda_arch_list=7.0 \ --build-arg max_jobs=4 --build-arg nvcc_threads=1 \ diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index d8e8c6ea2..b00539284 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -26,9 +26,9 @@ GPU validation may require maintainer hardware. Every pull request runs these checks, and all must pass before it can merge: public text (no private addresses, local paths or credentials in the PR, its commit messages or added lines), sign-off and added files (every commit signed -off; no added file over 1 MiB unless listed in -`.github/large-files-allowed.txt`), pre-commit on the changed files, and mypy -for Python 3.12. +off; no added file over 1 MiB unless its path is listed one per line in +`.github/large-files-allowed.txt`, which the check reads only when it exists), +pre-commit on the changed files, and mypy for Python 3.12. Disclose AI assistance. The submitting human must review every changed line, understand the change, run relevant checks, and be able to explain the result. diff --git a/README.md b/README.md index 32fd8fb16..c18d607fd 100644 --- a/README.md +++ b/README.md @@ -6,7 +6,7 @@ ReVolta turns a server with V100 32 GB GPUs and plenty of system memory into a d It is built for four or more V100 32 GB GPUs and validated on four. We recommend, and test, the two interconnects that give the best speed: NVLink, or all GPUs under one PCIe switch with P2P. Every release passes the same gate on our own server before it ships: single-request greedy output identical to the previous release token for token, 4 and 8 concurrent requests, images and video, CPU cache readback, and 512K-token prompts. That server runs ReVolta every day. -**You need:** 4 × V100 32 GB (NVLink or PCIe; NVLink or one PCIe switch with P2P recommended), at least 64 GB of system memory (128 GB recommended), Linux x86-64 with glibc 2.34 or newer, and an NVIDIA driver for CUDA 12.8. Tested hardware and limits: [Support and known limits](docs/known-limits.md). +**You need:** four or more V100 32 GB GPUs (validated on four), on NVLink or with all GPUs under one PCIe switch with P2P; at least 80 GiB of system memory for the default pinned PLE table (128 GiB recommended; the experimental whole-table PLE disk option ran under 64 GiB and 32 GiB limits); Linux x86-64 with glibc 2.34 or newer; and an NVIDIA driver for CUDA 12.8. Tested hardware and limits: [Support and known limits](docs/known-limits.md). Current release: [0.1.4](https://github.com/Leonccaa/ReVolta/releases/tag/v0.1.4). @@ -23,7 +23,7 @@ Current release: [0.1.4](https://github.com/Leonccaa/ReVolta/releases/tag/v0.1.4 Prefixes evicted from the GPU stay in CPU memory or on disk: a cached 32K-token document comes back from CPU memory in **0.6 s**, and after an engine restart from a local NVMe disk tier in **1.5 s**, instead of **10.1 s** when recomputed (0.1.4, DFlash2). See [KV offload](docs/kv-offload.md). -**Speed** (total output tokens/s, 512-token responses; best of MTP4 and DFlash2) +**Speed** (total output tokens/s, 512-token responses; best of MTP4 and DFlash2; the PCIe row was measured on 0.1.2 and the NVLink row combines 0.1.2 and 0.1.1) | System | 1 request | 2 | 4 | 8 | Prefill, 8K / 131K prompt | | --- | ---: | ---: | ---: | ---: | ---: | @@ -37,7 +37,7 @@ Versions, power caps and the per-method breakdown are in [Performance](docs/perf | Model | Weights (bits per weight: routed experts / whole model) | Download | PLE table in host memory | GPU KV pool (MTP4) | Quality vs its BF16 parent: KL ↓ / top-1 ↑ | Notes | | --- | --- | ---: | ---: | ---: | ---: | --- | | [Uncensored AWQ g32](https://huggingface.co/leoncca/Qwen3.8-Flash-Next-Uncensored-AWQ-g32) **(recommended)** | 4-bit AWQ g32 experts, FP8 PLE, BF16 rest (**4.63 / 5.95**) | 142 GB, MTP draft included | 48 GiB | 1.15 M (1.37 M with DFlash2) | 0.0214 / 97.1% | Uncensored fine-tune; the reference setup | -| [Base AWQ g32](https://huggingface.co/leoncca/Qwen3.8-Flash-Next-AWQ-g32) | 4-bit AWQ g32 experts, FP8 PLE, BF16 rest (**4.63 / 5.95**) | 142 GB | 48 GiB | Same layout as Uncensored | 0.0207 / 97.1% | The original model | +| [Base AWQ g32](https://huggingface.co/leoncca/Qwen3.8-Flash-Next-AWQ-g32) | 4-bit AWQ g32 experts, FP8 PLE, BF16 rest (**4.63 / 5.95**) | 142 GB, MTP draft included | 48 GiB | Same layout as Uncensored | 0.0207 / 97.1% | The original model | | [Standard NVFP4](https://huggingface.co/RadixArk/Qwen3.8-Flash-Next-NVFP4) | NVFP4 experts, FP8 PLE, BF16 rest (**4.50 / 5.86**) | 135 GB + 5 GB draft pack | 48 GiB | 1.25 M | 0.0303 / 96.4% | Widely used community format | | [NVIDIA FP4-Local](https://huggingface.co/nvidia/Qwen3.8-Flash-Next-FP4-Local) | NVFP4 experts and PLE, MXFP8 attention and shared experts (**4.26 / 4.53**) | 102 GB, MTP layer built in | **27 GiB** | **1.47 M** | 0.0380 / 95.9% | NVIDIA's official checkpoint; least host memory | @@ -86,7 +86,7 @@ For another checkpoint, use its [preparation guide](docs/checkpoints.md) before - V100 16 GB, TP2/TP8, and other models have not been tested. - This release builds SM70 only; other GPU architectures, ROCm, CPU, XPU, TPU, and GGUF are outside the build. - Greedy outputs at C2–C8 can change between runs; fresh compiler autotuning can also change outputs. -- The AWQ PLE table needs about 48 GiB of host memory by default. An experimental option keeps the whole table on a local SSD instead: in 0.1.4 it served under 64 GiB and 32 GiB memory limits, also with a 32 GiB CPU KV pool under 64 GiB, and greedy output matched the default path. Single-request decoding is about 23–25% slower, and a cold first token takes 1.7–2.2 times as long. See [Support and known limits](docs/known-limits.md). +- The AWQ PLE table needs about 48 GiB of host memory by default. An experimental option keeps the whole table on a local SSD instead: in 0.1.4 it served under 64 GiB and 32 GiB memory limits, also with a 32 GiB CPU KV pool under 64 GiB, and greedy output matched the default path. Single-request decoding is about 23–25% slower, and a cold first token takes 1.6–2.3 times as long. See [Support and known limits](docs/known-limits.md). - PLE cascade (spilling part of the table to disk, `ple_disk_cascade=true`) is not supported yet and is off by default in the ReVolta profile; full support is planned for 0.2. - INT8 KV cannot be combined with DFlash2; the disk tier does not support INT8 yet. @@ -188,7 +188,7 @@ See [Support and known limits](docs/known-limits.md). ### Support matrix -See [Support matrix](docs/known-limits.md). +See [Support matrix](docs/known-limits.md#support-matrix). ### Reporting issues diff --git a/docs/configuration.md b/docs/configuration.md index 5610069bc..90bc0bfb3 100644 --- a/docs/configuration.md +++ b/docs/configuration.md @@ -30,6 +30,20 @@ or disable custom all-reduce with `--disable-custom-all-reduce`. to eight concurrent requests, at temperature 1 with thinking, including tool calls and structured output. `greedy` still works with DFlash2. This setting is about drafting and draft verification only; clients can still use any sampling parameters. +- **`--watchdog-config` (opt-in stack-dump watchdog, from 0.1.4, #33).** JSON with three + fields: `dump_dir` (directory for the dump files; the empty default disables the + watchdog), `timeout` (seconds without progress after which the process counts as stuck, + default 300, must be greater than 0) and `check_interval` (seconds between checks, + default 10). Example: + `--watchdog-config '{"timeout": 300, "check_interval": 10, "dump_dir": "/data/revolta/dumps"}'`. + When enabled it runs a background thread in the engine core process (`engine_`) + and in each worker (`worker_`); if that process makes no progress for `timeout` + seconds it appends a numbered dump of all its Python thread stacks to + `/VLLM_STACK_DUMP_for__.log`, at most once per stall, and logs + `[Watchdog]Dumped stack to ...`. The engine core also dumps its stacks on SIGTERM. + Waiting for incoming work does not count as a stall: the timeout window restarts when + the process resumes. This is a diagnostic aid, not a recovery mechanism: the process is + not restarted or killed. - **`--gpu-memory-utilization 0.95`** passed the full gate, including the multimodal stress test with 8 concurrent requests. Higher values ran out of memory in earlier multimodal stress tests. If you change concurrency or media limits, measure again. @@ -39,10 +53,12 @@ or disable custom all-reduce with `--disable-custom-all-reduce`. - `cpu_bytes_to_use` is the total for all four workers (48 GiB = 12 GiB per worker). - **Keep `--limit-mm-per-prompt` on V100.** With FlashAttention-V100, the engine allows 1 image and no video per request without it, and a value that leaves out `"video"` - disables video. 0.1.3 keeps this V100 default. `{"image": 16, "video": 1}` + disables video. 0.1.4 keeps this V100 default. `{"image": 16, "video": 1}` keeps the reference deployment's limits; Qwen3.8 Flash-Next takes images and video, not audio, so no audio limit is set. To use the generic image count instead, set - `--limit-mm-per-prompt '{"image": 8, "video": 1}'`; a lower image count leaves more + `--limit-mm-per-prompt '{"image": 8, "video": 1}'` (the inherited generic default is + 8 images, 1 video and 1 audio item, and the audio field is left out here because this + model has no audio input); a lower image count leaves more host memory for bursts of large images. Large images are resized either way (see [Support and known limits](known-limits.md)). - `--structured-outputs-config '{"backend": "xgrammar"}'` pins the structured-output @@ -103,8 +119,12 @@ embedding table uses about 47.7 GiB of pinned host memory even with KV offload d | 80 GiB | Off | Memory stress passed; full acceptance not run | | 64 GiB | Off | Out of memory during startup | -80 GiB is the smallest tested limit that passed memory stress, without full acceptance. With less than 128 GiB, -reduce or disable the CPU KV pool. +80 GiB is the smallest tested limit that passed memory stress for the default pinned PLE +table, with KV offload off and without full acceptance. With less than 128 GiB, +reduce or disable the CPU KV pool. The experimental whole-table PLE disk option +([Support and known limits](known-limits.md)) is a different path: on 0.1.4 it started and +served under a 64 GiB and a 32 GiB container memory limit, and under 64 GiB with a 32 GiB +CPU KV pool. Image and video preprocessing also uses host memory, in the API server process. In one test, 8 concurrent requests with 8 different 2048-pixel images each used about 33 GB at diff --git a/docs/dflash2.md b/docs/dflash2.md index 16b8618cd..f91c3b5c1 100644 --- a/docs/dflash2.md +++ b/docs/dflash2.md @@ -9,7 +9,10 @@ Use this guide to replace MTP4 with DFlash2 and choose its settings. First prepa > Volta (SM70) GPUs, with V100 attention kernels and CUDA graphs for the drafter and the > verifier, in the 512K-context setup with prefix caching and CPU KV offload. On one > request it writes English code at up to **234 tokens/s instead of 180** with MTP4 (1.30×) -> on community NVLink hardware. In 0.1.3, the GPU KV pool reaches **1.38 million tokens (2.63× at 512K)** with the recommended DFlash2 settings, and 1.54 million (2.94×) with state replay on. MTP4 +> on community NVLink hardware. With 0.1.4 and the recommended DFlash2 settings (probabilistic +> drafting, state replay off) the GPU KV pool reaches **1.37 million tokens (2.61× at 512K)**; +> with greedy drafting it is 1.38 million (2.63×), and the 0.1.3 measurement with state replay +> on (greedy drafting) reached 1.54 million (2.94×). MTP4 > stays the default: Chinese text and eight concurrent requests are still faster with MTP4. > [How to enable it](#enable-dflash2). @@ -19,7 +22,9 @@ instead (tested revision `a22a1744`): 5 sliding-window layers, an 8-token block 2048-token window. It is off by default. Without a `dflash` speculative config, the KV cache layout, block size and MoE routes are the same as in 0.1.1. -The configuration below describes 0.1.3. With 0.1.2, also set `VLLM_V2_CUDAGRAPH_MEM_MIB=1500`; its profile does not reserve enough memory for DFlash2's graphs. +The configuration below is the recommended 0.1.4 configuration. With 0.1.2, also set +`VLLM_V2_CUDAGRAPH_MEM_MIB=1500`; its profile does not reserve enough memory for DFlash2's +graphs. ## Enable DFlash2 diff --git a/docs/getting-started.md b/docs/getting-started.md index cc3dc4d47..e9be6e0a9 100644 --- a/docs/getting-started.md +++ b/docs/getting-started.md @@ -9,7 +9,7 @@ concurrent requests in testing, and a 48 GiB CPU KV pool. ## 1. Install -The distribution name is `revolta` (version `0.1.3`). The Python import name and the +The distribution name is `revolta` (version `0.1.4`). The Python import name and the CLI are still **`vllm`** (`import vllm`, `vllm serve`); this is intentional. Only compute capability 7.0 is built, with CUDA 12.8, Python 3.12 and PyTorch 2.10. @@ -18,8 +18,8 @@ From PyPI (Linux x86-64 with glibc 2.34 or newer, Python 3.12): ```bash python3.12 -m venv /data/revolta/venv . /data/revolta/venv/bin/activate -pip install revolta==0.1.3 -vllm --version # 0.1.3 +pip install revolta==0.1.4 +vllm --version # 0.1.4 ``` The wheel pins `torch==2.10.0`, `torchaudio==2.10.0` and `torchvision==0.25.0`, @@ -31,7 +31,7 @@ differs: it pins the same versions to the CUDA 12.8 builds on `download.pytorch. and has the generic `linux_x86_64` tag. ```bash -pip install "https://github.com/Leonccaa/ReVolta/releases/download/v0.1.3/revolta-0.1.3-cp312-cp312-linux_x86_64.whl" +pip install "https://github.com/Leonccaa/ReVolta/releases/download/v0.1.4/revolta-0.1.4-cp312-cp312-linux_x86_64.whl" ``` ReVolta does not send the anonymous usage statistics that upstream vLLM reports. @@ -39,7 +39,7 @@ ReVolta does not send the anonymous usage statistics that upstream vLLM reports. Or use the Docker image: ```bash -docker pull ghcr.io/leonccaa/revolta:0.1.3 +docker pull ghcr.io/leonccaa/revolta:0.1.4 ``` ## 2. Download the checkpoint and make 512K views @@ -164,11 +164,13 @@ vllm serve /data/revolta/views/target-512k \ Notes: - **Startup** takes roughly 10–15 minutes: weight and PLE loading, `torch.compile` - and CUDA graph capture. Later starts reuse the compile cache. Upgrading from 0.1.2 - to 0.1.3 invalidates that cache, even with MTP or no speculation, because the traced - GDN source changed; allow a cold compile after upgrading. On a first start with - an empty compile cache, the reported GPU KV capacity can be one allocation unit - (7,182 tokens) lower. + and CUDA graph capture. Later starts reuse the compile cache. After an upgrade that + cache can still be incomplete: keeping the 0.1.3 cache, MTP4 and DFlash2 each compiled + four graphs again on the first 0.1.4 start (about 5 s on the reference machine), and + later starts loaded them from the cache. The 0.1.2 to 0.1.3 upgrade invalidated the + whole cache, even with MTP or no speculation, because the traced GDN source changed; + allow a cold compile after upgrading. On a first start with an empty compile cache, + the reported GPU KV capacity can be one allocation unit (7,182 tokens) lower. See [Launch options](configuration.md#launch-options) before changing these settings. @@ -195,9 +197,10 @@ For request defaults and the model vendor's recommendations, see [Sampling parameters](configuration.md#sampling-parameters). The server also provides Anthropic-compatible `/v1/messages` (including streaming) -and `/v1/messages/count_tokens`. On current main, an overlong prompt returns HTTP 400 -with `invalid_request_error`; an unknown model returns HTTP 404 with `not_found_error`. -In 0.1.3, the overlong-prompt case returns HTTP 500 with `api_error` (#26). +and `/v1/messages/count_tokens`. From 0.1.4, an overlong prompt returns HTTP 400 with +`invalid_request_error` and an unknown model returns HTTP 404 with `not_found_error`; +internal errors stay HTTP 500 (#32). In 0.1.3 the overlong-prompt case returned HTTP 500 +with `api_error` (#26). ## Docker @@ -219,7 +222,7 @@ docker run --rm --gpus all \ -e VLLM_PLE_OFFLOAD_READY_TIMEOUT=1200 \ -e VLLM_QWEN4EXP_QSA_E4M3_STRICT_SCALES=1 \ -e TILELANG_CACHE_DIR=/data/revolta/cache/tilelang \ - ghcr.io/leonccaa/revolta:0.1.3 \ + ghcr.io/leonccaa/revolta:0.1.4 \ /data/revolta/views/target-512k --served-model-name Qwen3.8-Flash-Next ... ``` diff --git a/docs/known-limits.md b/docs/known-limits.md index 7f7ae82c2..d3471afb6 100644 --- a/docs/known-limits.md +++ b/docs/known-limits.md @@ -9,6 +9,8 @@ all GPUs under one PCIe switch with P2P, are the recommended and tested intercon The reference system uses four V100 32 GB PCIe cards behind one PCIe switch (PIX/PXB in `nvidia-smi topo -m`), with P2P enabled and a 128 GiB container memory limit. +## Support matrix + | Configuration | Test status | | --- | --- | | 4 × V100 32 GB PCIe, TP4 + DCP2, P2P | Full acceptance passed on the reference configuration | @@ -18,8 +20,10 @@ The reference system uses four V100 32 GB PCIe cards behind one PCIe switch **Full acceptance on the reference setup:** default Uncensored AWQ, Standard NVFP4, and FP4-Local with their documented settings. Base AWQ has quality measurements. -80 GiB RAM has memory-stress results. Disk caching has readback results; its full -release gate passed with an 8 GiB CPU pool and a 40 GiB disk budget. +80 GiB RAM has memory-stress results for the default pinned PLE table with KV offload +off; the experimental whole-table disk option below ran under lower limits. Disk caching +has readback results; its full release gate passed with an 8 GiB CPU pool and a 40 GiB +disk budget. V100 16 GB, TP2/TP8, and other models have not been tested. This release builds SM70 only; other GPU architectures, ROCm, CPU, XPU, TPU, and @@ -39,8 +43,10 @@ GGUF are outside the build. build retrieved 20/20 needles at 489K. At 512K, the test covers capacity and the input/output boundary. - **Memory.** The default GPU utilization of 0.95 passed multimodal stress. Retest - memory use when changing concurrency or media limits. By default the AWQ PLE - table needs about 48 GiB of host memory; see the experimental disk option below. + memory use when changing concurrency or media limits. By default the AWQ PLE table + needs about 48 GiB of pinned host memory, and 80 GiB is the smallest tested limit that + passed memory stress on that path with KV offload off; see the experimental disk option + below. - **PLE cascade.** Cascade (`ple_disk_cascade=true`) spills the part of the PLE table that does not fit in host memory to the checkpoint files. It is not supported yet and is off by default in the ReVolta profile; full support is planned for 0.2. @@ -72,7 +78,7 @@ GGUF are outside the build. Single-request decoding was 101.5–104.9 tokens/s against 135.5 with the pinned table, about 23–25% slower. A cold first token took 5.4–5.9 / 16.8–17.9 / 65.6–69.8 s at 8K / 32K / 131K prompt tokens, against 2.61 / 9.66 / 40.40 s, so - 1.7–2.2 times as long; once the table pages for a prompt had been read, the first + 1.6–2.3 times as long; once the table pages for a prompt had been read, the first token was within 2–3% of the pinned path (5–8% with the 32 GiB CPU KV pool). NVMe and SATA, and the 32 and 64 GiB limits, differed little. Peak memory of the serving container was about 57 GiB under the 64 GiB limit (62 GiB with the CPU KV @@ -100,6 +106,11 @@ GGUF are outside the build. `required` parse can still report `finish_reason="tool_calls"`. Stop tokens inside JSON strings and very large parallel-call deltas remain problematic; strict native XML structural tags remain deferred. +- **Anthropic endpoint.** On `/v1/messages`, a response that `max_tokens` cuts inside a + tool call streams the partial `tool_use` input, but a non-streaming response returns no + `tool_use` block (0.1.4 known issue). `/v1/messages` and `/v1/messages/count_tokens` + return HTTP 400 for a prompt longer than `--max-model-len` and HTTP 404 for a missing + model from 0.1.4 (#32); internal errors stay HTTP 500. - **Image placeholder text.** Text that contains the literal image placeholder `<|image_pad|>` can take the place of a real image in the same request. Upstream vLLM and HF behave the same (vllm-project/vllm#57740), and ReVolta follows them. Remove such tokens from diff --git a/docs/kv-offload.md b/docs/kv-offload.md index 3186b6333..731d4075e 100644 --- a/docs/kv-offload.md +++ b/docs/kv-offload.md @@ -34,7 +34,7 @@ computed again. holds about 2.4M tokens. With DFlash2, the drafter uses a different cache block size, so every cache group takes pages of the full block size and a prefix needs about 50 KB per token in the CPU pool and on disk (measured in 0.1.4: 1.84 GB for a 32K-token prefix, -6.36 GB for 131K), about 2.5 times as much as with MTP4. `mamba_state_slots_reference_tokens` sets how many +6.36 GB for 131K), about 2.4 times MTP4's about 21 KB per token. `mamba_state_slots_reference_tokens` sets how many recurrent-state slots each layer group gets. A smaller reference gives more slots, which helps when many short sessions share the pool. @@ -83,7 +83,8 @@ use a larger CPU pool and set `max_size_bytes` to your disk budget. be larger than the largest prefix you expect to restore from disk (about 21 KB per token with MTP4 and about 50 KB with DFlash2; a 100K-token prefix needs more than 2 GiB with MTP4 and about 5 GB with DFlash2). Otherwise the request is computed - again. The quick-start 48 GiB pool holds about 2.4M tokens. + again. The quick-start 48 GiB pool holds about 2.4M tokens with MTP4; with DFlash2 the + same pool holds about 1M tokens, an estimate from the same about 50 KB per token. - The SSD tier is **not supported with INT8 KV** yet. That combination is rejected at configuration time. @@ -92,8 +93,10 @@ system (8 GiB CPU pool, 40 GiB disk budget), including a 99K-token prefix restor disk (96,288 cached tokens, output identical). In 0.1.4 the tier was also measured on a local NVMe SSD with DFlash2 (7 GiB CPU pool, 64 GiB disk budget, one read and one write thread). After an engine restart, 11 of 11 documents were read back from disk with the -same 64-token greedy output as a cold run, at about 2.2–2.4 GB/s per request; a -32K / 99K / 131K-token prefix came back in 1.47 / 2.93 / 3.77 s instead of +same 64-token greedy output as a cold run, reading the pages from disk at about +2.2–2.4 GB/s inside the restore. That rate covers the disk read on its own: it is not the +stored bytes divided by the time to the first token, and that quotient is the lower of +the two. A 32K / 99K / 131K-token prefix came back in 1.47 / 2.93 / 3.77 s instead of 10.12 / 31.88 / 42.43 s when recomputed. A root holding 27.8 GB in 1,088 files did not slow the restart. See [Coming back to a document](performance.md#coming-back-to-a-document). diff --git a/docs/performance.md b/docs/performance.md index 225814448..813df8159 100644 --- a/docs/performance.md +++ b/docs/performance.md @@ -9,7 +9,8 @@ Read the [hardware scope](known-limits.md) first; each group keeps its own versi Version and test conditions: the PCIe values use 0.1.2; the NVLink overview combines 0.1.2 and 0.1.1. See [DFlash2 against MTP4](#dflash2-against-mtp4), [Cold prefill](#cold-prefill) and -[NVLink](#nvlink-community-result) for the per-group conditions. +[NVLink](#nvlink-community-result) for the per-group conditions. The 0.1.4 release-gate +speed comparison against 0.1.3 is in [Reference system, MTP4](#reference-system-mtp4). Best measured output speed in tokens/s, 512-token responses. The breakdown by speculative method (MTP4 and DFlash2) and prompt type is in [Performance](#performance). @@ -39,7 +40,9 @@ Both readbacks hit the whole cacheable prefix, and the disk readbacks gave the s | 99,000 | 31.88 s | 0.77 s | 2.93 s | | 131,072 | 42.43 s | 0.99 s | 3.77 s | -Times are to the first token, one request at a time. Eight more 32K-token documents +Times are to the first token, one request at a time. The about 2.2–2.4 GB/s quoted in +[KV offload modes](kv-offload.md) is the disk read rate inside the restore, not these +stored bytes divided by these times. Eight more 32K-token documents gave 0.62–0.66 s from CPU memory and 1.41–1.46 s from NVMe for the first four of them. Four 32K-token documents read from NVMe at the same time after the restart took 1.48–3.84 s each (median 2.67 s). @@ -99,14 +102,14 @@ Measured on our 4x V100 PCIe deployment with TP4 + DCP2, `--gpu-memory-utilizati | Standard NVFP4 + E4M3 KV | 0.1.2 / 0.1.3 | MTP4 | 1,250,448 | 2.39× | | Standard NVFP4 + E4M3 KV | 0.1.2 | DFlash2 | 1,337,745 | 2.55× | -The 0.1.3 AWQ DFlash2 measurements include sharded selector codebooks (#9) and the 1200 MiB graph reserve with tail graphs off (#23), with state replay (#21) off (recommended) or on. The MTP4 rows marked 0.1.2 / 0.1.3 were remeasured unchanged on the 0.1.3 release image; the FP4-Local E4M3-draft row is a separate smoke with the `v1.1-base-b9ce18a6` scale pack. The FP16 KV, FP4-Local and NVFP4 DFlash2 values were measured on 0.1.2; they have not been remeasured with the 0.1.3 DFlash2 changes. The 0.1.2 AWQ DFlash2 measurements use a 1500 MiB graph reserve. INT8 target KV is not supported with DFlash2; current main rejects it during configuration even with a separate FP16 draft cache (#25). +The 0.1.3 AWQ DFlash2 measurements include sharded selector codebooks (#9) and the 1200 MiB graph reserve with tail graphs off (#23), with state replay (#21) off (recommended) or on. The MTP4 rows marked 0.1.2 / 0.1.3 were remeasured unchanged on the 0.1.3 release image; the FP4-Local E4M3-draft row is a separate smoke with the `v1.1-base-b9ce18a6` scale pack. The FP16 KV, FP4-Local and NVFP4 DFlash2 values were measured on 0.1.2; they have not been remeasured with the 0.1.3 DFlash2 changes. The 0.1.2 AWQ DFlash2 measurements use a 1500 MiB graph reserve. INT8 target KV is not supported with DFlash2; 0.1.4 rejects it during configuration even with a separate FP16 draft cache (#33, fixes #25). ## Performance ### Reference system, MTP4 Release configuration: Uncensored AWQ g32, E4M3 KV, TP4 + DCP2, MTP4, on the reference system. -The first table records 0.1.0; the later release comparison follows below. +The first table records 0.1.0; the release-to-release comparisons follow below. | Metric | 0.1.0 | | --- | ---: | @@ -149,6 +152,17 @@ per concurrency level: | C4 | 263.4 | 268.3 | | C8 | 389.3 | 390.4 | +The 0.1.4 release gate measured MTP4 output against 0.1.3 on the reference system: two +concurrent requests produced about 3% more output tokens per second, and the other MTP4 +levels were unchanged within 1% (#34). The gain comes from the MoE all-reduce no longer +using the slow generic sum2 kernel at two and three requests behind a PCIe switch. MTP4 +output at two and three concurrent requests is therefore not bitwise identical to 0.1.3, +because the summation order changed, while single-request output and DFlash2 output are +unchanged. A repeated A/B measurement at three to five requests found DFlash2 within about +2% per step of 0.1.3, with the same acceptance length. The best-of entries in +[Speed](#speed) come from DFlash2 on the reference system and do not change: the MTP4 +values in these windows stay below them at every concurrency level shown. + ### Cold prefill The PCIe reference measured 0.1.2 with MTP4 at 250 W per GPU: **3,382 tok/s at 8K** and **3,016 tok/s at 131K**. With DFlash2 on 0.1.2, the same window measured 3,120 and 3,093 tok/s. These use 512K views, one request at a time and a cold prefix cache; throughput is computed prompt tokens divided by time to first token. The Measurements table uses the MTP4 values so both systems use the same speculative method. See [NVLink](#nvlink-community-result) for the community method and longer prompts. From f9168e6efeb0a57202e19bae1e0ab4e07366e238 Mon Sep 17 00:00:00 2001 From: Leonccaa <166551845+Leonccaa@users.noreply.github.com> Date: Tue, 27 Oct 2026 16:30:23 -0700 Subject: [PATCH 2/2] docs: 32 GiB is the tested host memory minimum The tested minimum is 32 GiB, with the experimental whole-table PLE disk option (0.1.4 served under a 32 GiB limit and greedy output matched the pinned path). 80 GiB is the minimum for the default pinned PLE table; 128 GiB stays recommended. README, support scope, configuration (two rows added to the memory table) and known limits now lead with the 32 GiB figure and state the speed cost of the disk option. Co-Authored-By: Claude Opus 5.5 Signed-off-by: Leonccaa <166551845+Leonccaa@users.noreply.github.com> --- .github/SUPPORT.md | 2 +- README.md | 2 +- docs/configuration.md | 12 +++++++----- docs/known-limits.md | 8 ++++---- 4 files changed, 13 insertions(+), 11 deletions(-) diff --git a/.github/SUPPORT.md b/.github/SUPPORT.md index ce43a2077..5751b6b0c 100644 --- a/.github/SUPPORT.md +++ b/.github/SUPPORT.md @@ -10,7 +10,7 @@ In short: | PCIe switch with P2P | Recommended; validated | | PCIe without P2P | Validated by disabling P2P on P2P-capable hardware; about 30–55% of P2P speed | | NVLink | Recommended; tested with community results on 4 × V100-SXM2 32 GB ([#2](https://github.com/Leonccaa/ReVolta/issues/2), [#38](https://github.com/Leonccaa/ReVolta/issues/38), [#37](https://github.com/Leonccaa/ReVolta/issues/37)) | -| Host memory | 128 GiB recommended; 80 GiB measured minimum for the default pinned PLE table with KV offload off; the experimental whole-table PLE disk option ran under 64 GiB and 32 GiB limits | +| Host memory | 32 GiB tested minimum with the experimental whole-table PLE disk option; 80 GiB with the default pinned PLE table (KV offload off); 128 GiB recommended | | KV offload | Off, CPU, or CPU + SSD tier | | KV cache | E4M3 (default) or INT8; INT8 with the SSD tier is not supported | | Models | Qwen3.8 Flash-Next (QSA, multimodal) | diff --git a/README.md b/README.md index c18d607fd..3645fbe2c 100644 --- a/README.md +++ b/README.md @@ -6,7 +6,7 @@ ReVolta turns a server with V100 32 GB GPUs and plenty of system memory into a d It is built for four or more V100 32 GB GPUs and validated on four. We recommend, and test, the two interconnects that give the best speed: NVLink, or all GPUs under one PCIe switch with P2P. Every release passes the same gate on our own server before it ships: single-request greedy output identical to the previous release token for token, 4 and 8 concurrent requests, images and video, CPU cache readback, and 512K-token prompts. That server runs ReVolta every day. -**You need:** four or more V100 32 GB GPUs (validated on four), on NVLink or with all GPUs under one PCIe switch with P2P; at least 80 GiB of system memory for the default pinned PLE table (128 GiB recommended; the experimental whole-table PLE disk option ran under 64 GiB and 32 GiB limits); Linux x86-64 with glibc 2.34 or newer; and an NVIDIA driver for CUDA 12.8. Tested hardware and limits: [Support and known limits](docs/known-limits.md). +**You need:** four or more V100 32 GB GPUs (validated on four), on NVLink or with all GPUs under one PCIe switch with P2P; at least 32 GiB of system memory with the experimental whole-table PLE disk option (about 23–25% slower single-request decoding), 80 GiB with the default pinned PLE table, 128 GiB recommended; Linux x86-64 with glibc 2.34 or newer; and an NVIDIA driver for CUDA 12.8. Tested hardware and limits: [Support and known limits](docs/known-limits.md). Current release: [0.1.4](https://github.com/Leonccaa/ReVolta/releases/tag/v0.1.4). diff --git a/docs/configuration.md b/docs/configuration.md index 90bc0bfb3..4efda5fb5 100644 --- a/docs/configuration.md +++ b/docs/configuration.md @@ -118,13 +118,15 @@ embedding table uses about 47.7 GiB of pinned host memory even with KV offload d | 128 GiB | Off, or 48 GiB CPU pool | Full acceptance passed | | 80 GiB | Off | Memory stress passed; full acceptance not run | | 64 GiB | Off | Out of memory during startup | +| 64 GiB, whole PLE table on disk (experimental) | Off, or 32 GiB CPU pool | Served; greedy output matched the pinned path (0.1.4) | +| 32 GiB, whole PLE table on disk (experimental) | Off | Served; greedy output matched the pinned path (0.1.4) | -80 GiB is the smallest tested limit that passed memory stress for the default pinned PLE -table, with KV offload off and without full acceptance. With less than 128 GiB, +With the default pinned PLE table, 80 GiB is the smallest tested limit that passed memory +stress, with KV offload off and without full acceptance. With less than 128 GiB, reduce or disable the CPU KV pool. The experimental whole-table PLE disk option -([Support and known limits](known-limits.md)) is a different path: on 0.1.4 it started and -served under a 64 GiB and a 32 GiB container memory limit, and under 64 GiB with a 32 GiB -CPU KV pool. +([Support and known limits](known-limits.md)) keeps the table on a local SSD instead and +brings the tested minimum down to 32 GiB; single-request decoding is about 23–25% slower +and a cold first token takes 1.6–2.3 times as long. Image and video preprocessing also uses host memory, in the API server process. In one test, 8 concurrent requests with 8 different 2048-pixel images each used about 33 GB at diff --git a/docs/known-limits.md b/docs/known-limits.md index d3471afb6..db0bf6f0b 100644 --- a/docs/known-limits.md +++ b/docs/known-limits.md @@ -20,8 +20,8 @@ The reference system uses four V100 32 GB PCIe cards behind one PCIe switch **Full acceptance on the reference setup:** default Uncensored AWQ, Standard NVFP4, and FP4-Local with their documented settings. Base AWQ has quality measurements. -80 GiB RAM has memory-stress results for the default pinned PLE table with KV offload -off; the experimental whole-table disk option below ran under lower limits. Disk caching +32 GiB RAM served with the experimental whole-table PLE disk option below; 80 GiB RAM +has memory-stress results for the default pinned PLE table with KV offload off. Disk caching has readback results; its full release gate passed with an 8 GiB CPU pool and a 40 GiB disk budget. @@ -45,8 +45,8 @@ GGUF are outside the build. - **Memory.** The default GPU utilization of 0.95 passed multimodal stress. Retest memory use when changing concurrency or media limits. By default the AWQ PLE table needs about 48 GiB of pinned host memory, and 80 GiB is the smallest tested limit that - passed memory stress on that path with KV offload off; see the experimental disk option - below. + passed memory stress on that path with KV offload off. The experimental disk option + below served under a 32 GiB limit. - **PLE cascade.** Cascade (`ple_disk_cascade=true`) spills the part of the PLE table that does not fit in host memory to the checkpoint files. It is not supported yet and is off by default in the ReVolta profile; full support is planned for 0.2.