From 0e7ec0a418f06767df3d19da7d38743c2721b9d5 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Tue, 27 Oct 2026 10:23:19 +0100 Subject: [PATCH] [Docs] Record BeNice! KV sharing and AsymLoop explorations as not qualified Both were measured as serving proxies on Ouro-1.4B with one L40S and pre-registered criteria. BeNice! shared KV gains overload throughput but regresses TPOT 1.7-3.4x, fails low-load parity and prefix-hit criteria, and collapses GSM8K-87 (60/87 -> 0/87). AsymLoop at the same loop count holds 1.33x the tokens but gives only 1.08-1.10x throughput at 16K and 0.99x on a Mooncake trace; the paper's depth-matched setting gives 0.48-0.63x. The page records the data, the post-hoc Triton recompilation confound and how to reproduce from the research branch on a fork; no code is proposed. Co-Authored-By: Claude Opus 5.5 Signed-off-by: 0z5a --- docs/README.md | 2 + docs/kv-sharing-asymloop-exploration.md | 249 ++++++++++++++++++++++++ 2 files changed, 251 insertions(+) create mode 100644 docs/kv-sharing-asymloop-exploration.md diff --git a/docs/README.md b/docs/README.md index d83666cb..a21b3e79 100644 --- a/docs/README.md +++ b/docs/README.md @@ -27,6 +27,8 @@ an environment to receiving the first generated response. The Earlier measurements and known limitations are retained in the [historical validation record](runtime-history.md); they do not qualify the current revision. +Explorations that did not qualify for merge are recorded with their data and +reproduction steps: [BeNice! KV sharing and AsymLoop on Ouro-1.4B](kv-sharing-asymloop-exploration.md). ## Support and Runtime Notes diff --git a/docs/kv-sharing-asymloop-exploration.md b/docs/kv-sharing-asymloop-exploration.md new file mode 100644 index 00000000..13612c57 --- /dev/null +++ b/docs/kv-sharing-asymloop-exploration.md @@ -0,0 +1,249 @@ +# Not qualified: BeNice! KV sharing and AsymLoop on Ouro-1.4B + +**Status: not qualified (both explorations).** Neither is proposed for merge. This +page records what was measured, so that the two routes are not repeated without +a new reason, and how to reproduce the measurements. + +| Exploration | Paper | What would have to hold | Outcome | +| --- | --- | --- | --- | +| BeNice! KV sharing | [arXiv 2610.31633](https://arxiv.org/abs/2610.31633) | Shared final-recurrence KV raises high-load throughput and lowers TTFT without hurting low load or quality | Not qualified: overload throughput +28–34%, but TPOT 1.7–3.4× worse, low-load parity and prefix-hit criteria failed, GSM8K-87 accuracy 60/87 → 0/87 | +| AsymLoop | [arXiv 2610.28489](https://arxiv.org/abs/2610.28489) | A one-pass encoder plus a looped decoder serves more load than Ouro's full loop | Not qualified: same loop count gives 1.33× KV capacity but only 1.08–1.10× throughput at 16K and 0.99× on a Mooncake trace; the paper's depth-matched setting gives 0.48–0.63× | + +Both are **serving proxies on Ouro-1.4B weights**. Neither paper releases weights +for these architectures, so quality numbers say nothing about trained BeNice! or +AsymLoop models. BeNice!'s own Table 1 reports that Ouro-1.4B is not shareable +without post-training. + +## Setup shared by both + +- One NVIDIA L40S (48 GB) per run; Ouro-1.4B at revision + `574fa66cb8bf5abdc979642d01cf2b79b16bfab1`, bf16; torch 2.13.0+cu130, Triton 3.7.1. +- KV pool pinned to 30 GiB in every arm, so capacity differences come only from + the layout. Triton attention, block size 16, refill scheduling, + `max_num_seqs` 64, 2048 batched tokens, 2048-token prefill chunks. +- `ouro_delayed` exit with threshold 1.0, so every token runs the full depth. +- Workloads: + - **W1**: bursts of C fresh random prompts of 8K or 16K tokens, C ∈ {1, 8, 32}. + - **W2**: open-loop replay of the Mooncake FAST'25 `toolagent_trace.jsonl`, with + 512-token hash blocks mapped to deterministic token blocks, so the trace's + prefix reuse is real. Prompts are capped at 16K and arrivals rescaled to a + fixed offered rate. +- Hypotheses and thresholds were written before any GPU result. Every cell ran + once unless a replicate is listed. +- Code: branch + [`research/benice-asymloop-20261026`](https://github.com/0z5a/vllm-rlt/tree/research/benice-asymloop-20261026) + on a fork, three commits on `8fbf4cd`: + - `6ae0277`: loop-order shared prefill; + - `8cfd342`: AsymLoop; + - `4f23cd6`: benchmark drivers. + +## BeNice! KV sharing + +BeNice! lets every recurrence read the final recurrence's KV, so the cache keeps +one plane per token instead of one per loop. vllm-rlt's existing `shared` layout +has this semantics, but its prompt runs position by position, and the layout +rejected prefix caching. The research branch adds two things: +- `CacheConfig.shared_prefill="loop"`: a shared-KV prompt chunk advances one loop + at a time through packed prefill. + - Inside a chunk, loop d reads the chunk's own loop-d KV; earlier chunks are + read at their final KV. + - The result therefore depends on chunk boundaries. A single chunk is bitwise + equal to `last_exited`. +- Prefix caching for `shared` KV. + +**Arms.** Both use async scheduling, CUDA graphs, static buffers and packed +prefill (`prefill_uva`), with prefix caching on. +- **A0**: `last_exited`, 40,960 tokens in 30 GiB. +- **B1**: `shared` with `shared_prefill="loop"`, 163,840 tokens (4.00×). + +| Criterion (pre-registered) | Cell | Result | Verdict | +| --- | --- | --- | --- | +| Capacity 4× | 30 GiB | 40,960 → 163,840 tokens | pass | +| Low-load parity: TTFT and TPOT ratio in [0.9, 1.1] | W1 8K C=1 / W1 16K C=1 | 1.03 and 1.02 / 0.98 and 0.99 | pass / pass | +| Low-load parity | W2 0.1 req/s | TTFT p50 0.89, TPOT p50 1.74 | fail | +| High load: throughput ≥ 1.15× and p90 TTFT ≤ 0.67× | W1 8K C=32 | 1.31×, 0.58× | pass | +| High load | W1 16K C=32 | 1.32×, 0.72× | fail | +| High load | W2 0.6 req/s | 1.28×, 0.42× (replicate on a second host: 1.30×, 0.39×) | pass | +| Prefix hit +5 points | W2 0.6 req/s | 22.8% → 25.0% (replicate 24.0% → 24.4%) | fail | +| Quality (expected to fail on Ouro) | GSM8K-87, greedy | 60/87 → 0/87, −69.0 pp (paired SE 5.0) | fail | + +TPOT at high load went from 0.078 to 0.244 s (W1 8K C=32), from 0.071 to 0.233 s +(W1 16K C=32) and from 0.095 to 0.319 s (W2 0.6). Raising the CUDA-graph cache cap +from 16 to 64 removed B1's graph fallbacks but changed TPOT by less than 3%. + +In the GSM8K run, 86 of 87 first tokens matched A0. Decoding then degenerates into +repetition within a few tokens: the prefill change is not the cause; reading +final-loop KV during decode is. + +**Why it does not qualify.** Sharing saves KV memory but not KV reads. +- ncu on the decode attention kernel (W1 16K, C=8): + - A0, 2 rows: 422 µs per launch; + - B1, 8 rows: 1.70 ms per launch; + - both at ~85% of DRAM peak, SM throughput ~7.5%, L2 hit rate ~0.3%. +- The extra capacity therefore becomes larger decode batches whose step time + grows with rows × context. +- Throughput gains appear only past saturation, and they cost a 1.7–3.4× TPOT + regression. +- On Ouro the quality loss is total. + +## AsymLoop + +AsymLoop runs the first E layers once per token as an encoder. Each decoder layer +projects the encoder output once to a fixed global KV. The D decoder layers loop +K times, and each one attends to the global KV and the current loop's KV under +one causal softmax. The proxy uses E = D = 12 on Ouro's 24 layers. + +**Implementation** (`8cfd342`): +- **KV planes.** `KVCacheManager(aux_planes=2)` adds an encoder plane and a + global plane next to the per-depth planes. + - The pool is cut into page groups: plane p of group g is block p·G + g. + - Any depth's block table therefore maps to the same page's encoder and global + blocks by arithmetic, and mixed-depth batches need no new metadata. +- **Attention.** Paged attention (the Torch reference, Triton decode and packed + prefill) gains `sources=2`: two tables with the same causal length under one + online softmax. `sources=1` is unchanged. +- **Model.** `AsymLoopOuroForCausalLM` uses the KV-bearing prelude path, which + this branch opens to packed prefill: + - prelude = embedding, encoder, and global KV writes; + - recurrent = the decoder; + - coda = the LM head. +- **Global projections.** Ouro has none, so each starts as a copy of its decoder + layer's K/V projections, applied after that layer's input norm. +- **Not wired.** Asynchronous scheduling and CUDA graphs are not supported for + this model. +- **Correctness.** CPU and GPU tests compare the engine (chunked, mixed-depth, + with and without `prefill_uva`) token for token with an unpaged AsymLoop + reference. + +**Arms.** All three use synchronous scheduling, eager execution, packed prefill +and no prefix cache. The baseline was re-measured in this configuration, so +these numbers are not comparable with the BeNice! table. + +| Arm | Structure | KV slots per token | Layer passes per token | Decode KV read per token (× context) | Tokens in 30 GiB | +| --- | --- | ---: | ---: | ---: | ---: | +| A0s | Ouro full loop, K = 4 | 96 | 96 | 96 | 40,960 | +| AL4 | AsymLoop, K = 4 (same loop count) | 72 | 60 | 108 | 54,608 | +| AL7 | AsymLoop, K = 7 (same effective depth 96, the paper's matching rule) | 108 | 96 | 180 | 36,400 | + +The workloads were W1 16K (C = 1, 8 and 32, 128 output tokens) and W2 at 0.6 req/s +(60 requests, at most 128 output tokens). Ratios are against A0s. + +| Criterion (pre-registered) | Arm | Cell | Result | Verdict | +| --- | --- | --- | --- | --- | +| Capacity matches the table | all | 30 GiB | exact | pass | +| Single-request TPOT ratio in [0.80, 1.00] / [1.30, 1.70] | AL4 / AL7 | W1 16K C=1 | 0.78 / 0.87 | fail / fail (not interpretable, see below) | +| Single-request TTFT ratio in [0.75, 0.95] / [1.20, 1.50] | AL4 / AL7 | W1 16K C=1 | 0.87 / 1.49 | pass / pass | +| Throughput ≥ 1.15× | AL4 | W1 16K C=32 | 1.10× (replicate 1.08×) | fail | +| Throughput ≥ 1.15× | AL4 | W2 0.6 | 0.75× (warm replicate 0.99×) | fail | +| Throughput ≤ 0.90× | AL7 | W1 16K C=32 / W2 0.6 | 0.63× / 0.61× (0.48× against the warm A0s replicate) | pass / pass | + +AL7 passing means the predicted loss was confirmed. In the setting where the +paper claims its quality advantage, serving is clearly more expensive. + +**Where the time goes.** These are post-hoc diagnostics and are not scored. +- **ncu.** On W1 16K C=4, every decode attention launch costs ~210 µs per row and + per KV source: A0s 2 rows 423 µs, AL4 decoder 3 rows × 2 sources 1.25 ms, AL4 + encoder 1 row 213 µs. SM throughput is 7–8%. + - Per token at 16K this is 20.3 ms for A0s and 22.8 ms for AL4, exactly 108/96. + - The two-source kernel adds nothing beyond the extra bytes. +- **Per-step timing, W1 16K C=32 replicate.** + - Prefill: AL4 spends 55 s, A0s 69 s. + - Decode loops: about equal, 121 s with 2.86 rows per step against 153 s with + 1.98 rows. + - AL4 also runs the encoder as a separate prelude step for every new token: + 28 s in total, 8.8 ms per step, 1.27 rows per step. In A0s the prelude is an + embedding lookup (0.7 s). + - The capacity gain is spent on two-source reads and on this step. +- **Single request.** AL4 is faster at C=1 (warm TPOT 0.048 s against 0.087 s) + only because the eager path pays a fixed host cost per layer, and AL4 runs 36 + fewer layers per token. +- **cProfile.** On W2 the synchronous eager engine is host-bound: only 29 of + 169 s wait on the GPU; the rest goes to per-row, per-layer written-position + and prefix bookkeeping and eager RMSNorm. + +**Confound found after scoring.** `_paged_attention_kernel` declares +`table_stride` (the block-table width) as `tl.constexpr`. +- Every new table width compiles a new kernel variant, and `sources=2` adds its + own variants. +- The per-cell warm-up uses 512-token prompts and does not reach 16K widths, so + the first cell at those lengths compiled inside the timed window. This was + seen in the cold run (`cc1plus` busy) and as 300–490 ms steps in a W2 replicate. +- The single-request TPOT ratios and the first W2 run are cold-versus-warm + confounded across arms. The pre-registered verdicts above are kept as scored. +- TTFT (the prefill kernel's stride is not a constant) and W1 C=32 (its widths + were compiled by the earlier cells; the replicate agreed within 3%) are not + affected. + +## Notes for related engine work + +- KV sharing or extra KV sources move work from capacity to KV bandwidth. On + L40S, decode attention is DRAM-bound at ~210 µs per 16K row and per source. + Any proposal should show decode-step cost at the target batch size, not only + capacity. +- A model whose prelude carries KV adds a per-token engine step in decode that + batches poorly. AsymLoop-style designs need that step fused into the + recurrent step, or graph-captured, before they can pay off. +- In eager runs, warm up at the measured context lengths, or make the Triton + decode `table_stride` a runtime argument; otherwise each new width recompiles + during measurement. + +## Reproducing + +Use the fork branch above. Install the repository as usual, download Ouro-1.4B at +the pinned revision, and download the trace from +`https://raw.githubusercontent.com/kvcache-ai/Mooncake/main/FAST25-release/traces/toolagent_trace.jsonl`. +Run on a single GPU, with `CUDA_VISIBLE_DEVICES` set to one device. Every command +writes one JSON record with a `CAPACITY` line, per-cell latency and throughput, +and optional per-step timing (`--step-trace`). + +Engine flags for the BeNice! arms (async, CUDA graphs): + +```bash +COMMON="--model /path/to/Ouro-1.4B --attention triton --dtype bfloat16 --block-size 16 \ + --kv-bytes 30 --max-num-seqs 64 --max-batched-tokens 2048 --chunk 2048 --mode refill \ + --async --cuda-graphs --graph-max 64 --static-buffers --pad-pow2 --prefill-uva \ + --exit-mode ouro_delayed --exit-threshold 1.0 --prefix-cache" +A0="--kv-layout last_exited --shared-prefill position" +B1="--kv-layout shared --shared-prefill loop" +``` + +Engine flags for the AsymLoop arms (synchronous, eager, no prefix cache): + +```bash +COMMON="--model /path/to/Ouro-1.4B --attention triton --dtype bfloat16 --block-size 16 \ + --kv-bytes 30 --max-num-seqs 64 --max-batched-tokens 2048 --chunk 2048 --mode refill \ + --prefill-uva --exit-mode ouro_delayed --exit-threshold 1.0 --kv-layout last_exited" +A0s="" +AL4="--asymloop-encoder 12 --asymloop-loops 4" +AL7="--asymloop-encoder 12 --asymloop-loops 7" +``` + +Workloads (the BeNice! run used decode 256, 30/60/90 requests for W2 at +0.1/0.3/0.6 req/s, and also W1 8K; the AsymLoop run used decode 128 and 60 +requests): + +```bash +python -m benchmarks.kv_sharing_asymloop $COMMON $AL4 --label AL4-w1 --out AL4-w1.json \ + --workload burst --prompt-len 16384 --decode 128 --concurrency 1,8,32 --warmup 1 +python -m benchmarks.kv_sharing_asymloop $COMMON $AL4 --label AL4-w2 --out AL4-w2.json \ + --workload mooncake --traces toolagent_trace.jsonl --tokens-per-block 512 \ + --max-prompt 16384 --max-decode 128 --requests 60 --concurrency 0.6 --warmup 1 +``` + +Run every arm on the same GPU. Before an eager measurement, run one cell at the +target lengths to keep Triton compiles out of the timed window. + +For the BeNice! quality gate, prepare the GSM8K-87 protocol as described in +[accuracy evaluation](accuracy.md), then run both layouts and compare them: + +```bash +python -m benchmarks.kv_sharing_asymloop_gsm8k --protocol protocol.json --out A0.json \ + --label A0 --max-num-seqs 32 --kv-layout last_exited --shared-prefill position +python -m benchmarks.kv_sharing_asymloop_gsm8k --protocol protocol.json --out B1.json \ + --label B1 --max-num-seqs 32 --kv-layout shared --shared-prefill loop +python -m benchmarks.kv_sharing_asymloop_gsm8k --compare A0.json B1.json +``` + +For the decode attention profile, run ncu on one cell with +`--kernel-name regex:_paged_attention_kernel` and a small `--launch-count`, about +24. Kernel replay backs up the 30 GiB KV pool for every profiled launch.