Files
LocalAI/backend/cpp/llama-cpp/patches
Ettore Di Giacinto 2dd5d68e6d docs(paged): A.2 Phase 2 - locate the real decode lever (gated-DeltaNet SSM path)
Phase 1 ruled out CUDA graphs as the paged-decode lever (GPU 99.4% busy,
decode_agg flat graphs on-vs-off) and attributed the 2.6x gap to vLLM to the
per-step GPU kernel work (FP4 GEMM + attention at batch 128). Phase 2 decomposed
that kernel work directly on the Phase-1 nsys reps and corrects the attribution.

Findings (q36-27b-nvfp4 = gguf arch qwen35, a 48:16 hybrid gated-DeltaNet
linear-attention + full-attention model; DGX GB10 sm_121, fusion off):
- Graphs re-confirmed not the lever: fresh paged graphs-ON 146.03 vs OFF 144.90
  t/s (+0.78%, noise); the captured rep is 99.5% busy with the same ~3267ms
  memcpy (graphs capture memcpy nodes too).
- The 99.4% busy is real but ~19% of it is D2D memcpy, not compute: an
  overlap-correct interval-union sweep gives kernels-only 80.2% busy, the gap
  filled by 1584 D2D copies/run (~80/step, ~230MB each = the gated-DeltaNet
  recurrent state). Phase 1's cuda_gpu_trace lumped this into compute.
- Decode GPU-time decomposition (% of kernel+memcpy busy): gated_delta_net 23.4%,
  get_rows 21.9%, D2D state copy 18.9%, FP4 GEMV 15.5%, FP4 GEMM 10.4%,
  full attention 0.4%. Grouped: SSM/gated-DeltaNet machinery ~67%, FP4 matmul
  ~28%, full attention (all paged-attn optimizes) ~0.4%.

Verdict: not graphs, not the host loop, not primarily FP4 GEMM, not attention.
Paged attention touches ~0.4% of decode on this model, so no paged/graph/
block-table change can move decode_agg. The lever is the ggml qwen35
gated-DeltaNet decode: kill the per-layer recurrent-state D2D copy and fuse the
get_rows gather into the recurrence (vLLM's fused_recurrent_gated_delta_rule
keeps state in place). Ceiling: -copy ~146->180; -copy-and-gather ~146->247 t/s.

No code patch (the lever is an SSM-path rewrite, orthogonal to paged attention);
patches/paged/0018 stays free.

Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>
2026-06-24 21:44:22 +00:00
..

llama.cpp patch series — paged attention (vLLM-parity engine)

A stacking series: each patch is a small, self-contained, independently-buildable step toward an in-model paged-attention engine. They apply in numeric order on top of the pinned LLAMA_VERSION (backend/cpp/llama-cpp/Makefile). The build applies them automatically after checkout (see the llama.cpp: target). Keeping the work as ordered patches — rather than one big diff — is what lets us rebase cleanly across llama.cpp bumps and avoid drift: when a patch stops applying, only that small patch needs fixing, and the failure points at exactly which step the upstream change touched.

Base

  • LLAMA_VERSION pin in ../Makefile. All patches are generated against that exact commit. Bumping the pin = re-run the regen workflow below and fix only the patches that no longer apply.

The series (phases → patches)

# Patch What Verifies
0001 0001-vendor-paged-kv-manager.patch Add src/paged-kv-manager.{h,cpp} (vLLM-parity block manager, CPU foundation) + CMake; no behavior change builds; unit-tested separately under ../paged/
0002 0002-paged-kv-storage.patch Shared block-pool KV tensor + set_rows-by-slot writes, behind LLAMA_KV_PAGED builds; write/gather round-trip
0003 0003-paged-gather-read.patch build_attn_paged gather-read in llama-graph.cpp Gate 0: token-identical greedy gen, single + multi-seq
0004 0004-paged-ondemand-alloc.patch On-demand block allocation via PagedKVManager max concurrent seqs before OOM
0005 0005-paged-continuous-batching.patch Block-granular admit/evict in the server slot path tok/s vs concurrency, mixed-length
0006 0006-paged-prefix-caching.patch Block-hash cross-request prefix dedup TTFT + memory on shared prefixes

Each row is a separate git commit on the dev branch (below), exported 1:1 as a patch. Default off (LLAMA_KV_PAGED) until Gate 0 (0003) is green, so partial series never changes stock behavior.

Regen workflow (the anti-drift recipe)

# 1. check out the exact pin into a dev tree
git -C /tmp clone https://github.com/ggml-org/llama.cpp llama-dev && cd /tmp/llama-dev
git checkout <LLAMA_VERSION from ../Makefile>
git checkout -b paged

# 2. apply the current series (each becomes a commit), or develop the next patch
git am /path/to/backend/cpp/llama-cpp/patches/00*.patch     # or `git apply` + commit per patch

# 3. iterate a phase as ONE commit, then export the whole series 1:1
git format-patch <LLAMA_VERSION>..paged -o /path/to/backend/cpp/llama-cpp/patches/ --zero-commit -N

# 4. on a pin bump: rebase `paged` onto the new pin; only conflicting patches need edits; re-export.

Build integration

../Makefile's llama.cpp: target runs, after git checkout -b build $(LLAMA_VERSION):

for p in $(CURRENT_MAKEFILE_DIR)/patches/0*.patch; do git apply --verbose "$p"; done

All variants (avx/avx2/avx512/cuda/…) copy the patched llama.cpp/ tree, so the series ships everywhere.

Status

  • 0001 vendor manager — DONE. Applies clean to the pin; builds into libllama.
  • 0002 block placement — DONE + VERIFIED. Built llama-simple at the pin; greedy generation is token-identical stock vs LLAMA_KV_PAGED=1 (Qwen3-0.6B), paged branch confirmed firing.
  • 0003 gather-read — DONE + VERIFIED (Gate 0 green). Implemented in the additive form (ADDITIVE_DESIGN.md): all logic in new src/paged-attn.{h,cpp} (a llm_graph_input_i gather-index subclass + the K/V/mask gather), hooked by one line in build_attn + two thin accessors on llama_kv_cache_context + 1 CMake line (216 insertions; no edit to llm_graph_input_attn_kv or llama-graph.h). Greedy generation is token-identical stock vs LLAMA_KV_PAGED=1 (Qwen3-0.6B, 9/9 across 3 prompts × {32,96,128} tokens), with n_gather=71 < n_kv=256 confirming real compaction. Patch: 0003-paged-gather-read-env-LLAMA_KV_PAGED.patch.
    • Key correctness finding: get_gather_idxs must emit cells sorted by token position. The CPU flash-attn online softmax reduces cells in physical-array order and is FP-order-sensitive, so 0002's scattered placement alone (full-window read, no gather) diverges from stock once a sequence crosses the first 16-cell block. The position-sorted gather reproduces stock's exact reduction order -> bit- identical, not merely mathematically equivalent. So 0002 is the placement substrate; 0003 is what makes paged placement token-identical under flash-attn.
  • 00040006 follow.

Honest parity note (important)

This series delivers the paged-attention engine (capacity + scheduling + prefix sharing). It does not by itself reach vLLM throughput parity, because the measured prefill bottleneck is the FP4 MoE GEMM kernel (Lever 3: mul_mat_q<MXFP4> ~22 TFLOP/s, ~27× behind vLLM) — a per-token compute gap that paging does not touch. Paged attention closes the concurrency/memory gap (more sequences, prefix reuse); the prefill/throughput gap additionally needs the tcgen05/CUTLASS grouped-GEMM (deferred, upstream-grade, no shortcut — see ../paged/UPSTREAM_GGML_ISSUE.md and DGX_BLACKWELL_PLAN.md). So full vLLM parity = this series AND the kernel; neither alone suffices.