feat(llama-cpp-localai-paged): paged KV cache llama.cpp backend + cross-request prefix sharing + GB10 decode optimization [WIP]

.agents/llama-cpp-localai-paged-backend.md

@@ -0,0 +1,143 @@
+# llama-cpp-localai-paged Backend (paged attention + Blackwell NVFP4 decode)
+
+`llama-cpp-localai-paged` is LocalAI's **CUDA-only** paged-attention variant of the
+llama.cpp backend. It targets high-concurrency decode for the Qwen3.6 hybrid
+gated-DeltaNet (SSM) models on Blackwell (GB10 / DGX Spark). It reuses the stock
+`llama-cpp` backend's sources and applies a vendored patch series on top at build
+time. It is **not** a fork: a source-only `*.patch` stack plus one canonical doc.
+
+**Canonical reference:** `backend/cpp/llama-cpp-localai-paged/README.md`
+(architecture, the patch series 0001-0030, benchmarks, dev notes, generality,
+pin/canary policy). Read it for any technical detail; this guide is the maintenance
+how-to.
+
+## Where things live
+
+- `backend/cpp/llama-cpp-localai-paged/Makefile` - the thin wrapper. It copies the
+  stock `backend/cpp/llama-cpp/` build infra into a build dir, clones llama.cpp at
+  this backend's **own** pin (`LLAMA_VERSION`), applies the paged series via the
+  `apply-paged-patches` define (strict `git apply`), then builds `grpc-server`.
+- `backend/cpp/llama-cpp-localai-paged/patches/paged/` - the source-only `.patch`
+  series (0001-0030), nothing else.
+- `backend/cpp/llama-cpp-localai-paged/README.md` - the canonical doc. The
+  operational docs (`PAGED_BITEXACT_NOTE.md`, `UPSTREAM_LAYER2_SCOPE.md`) and
+  dev artifacts live in
+  `backend/cpp/llama-cpp-localai-paged/docs/`.
+- `backend/Dockerfile.llama-cpp-localai-paged`, `.docker/llama-cpp-localai-paged-compile.sh`
+  - the CUDA build entry points.
+- `backend/cpp/llama-cpp/` - the **stock** backend, pure upstream. It carries no
+  paged patches.
+
+## Invariants (do not break these)
+
+- **Stock stays pure.** The paged patches live ONLY in this backend. Never add a
+  `patches/paged/` dir or `LLAMA_PAGED` logic to `backend/cpp/llama-cpp/`.
+- **CUDA-only.** Ship cublas/cuda targets only. Off-CUDA the fusions are gated off
+  (patch 0030) and NVFP4 falls back to dequant, so the backend is neutral-to-
+  slightly-negative there - non-CUDA users use the stock `llama-cpp`. Do not add
+  cpu/vulkan/sycl/metal rows for this backend in `.github/backend-matrix.yml`.
+  (Those builds also fail to link `grpc-server` on darwin/arm64 against upstream
+  `stream_*` server symbols - another reason it is CUDA-only.)
+- **Source-only patches.** A `.patch` may touch only llama.cpp source - never a
+  dev doc or `*.md`. Strict `git apply` on a clean checkout must reach exit 0. (A
+  stray `SSM_DECODE_FIX_RESULTS.md` hunk in patch 0019 once broke the CI build.)
+- **Bit-exact by default.** Every shipped patch is byte-identical to the f32
+  baseline. (The one opt-in precision trade, `ssm_bf16_tau` / patch 0026, was
+  DROPPED: it went flat once the decode fusions landed - forcing all gated-DeltaNet
+  heads to bf16 gave 780.6 vs 780.0 t/s, zero benefit - so the series is now
+  bit-exact end to end. Do not reintroduce a per-head SSM-precision lever; see the
+  rejected-levers note in the backend README section 5.)
+
+## Fork-first workflow (MANDATORY)
+
+The fork **`mudler/llama.cpp` branch `localai-paged`** is the CANONICAL source
+of truth for ALL paged-backend kernel and patch work. The vendored
+`patches/paged/*.patch` series is a **derivative**: the fork is the source, the
+series is a generated mirror of it.
+
+**Always update the fork FIRST, in this exact order:**
+
+1. **Commit the change on the `localai-paged` branch and push it.** Every
+   kernel or patch change lands as a fork commit first.
+2. **Then regenerate the LocalAI series from the fork** via `git format-patch`
+   (one patch per fork commit, source-only) into
+   `backend/cpp/llama-cpp-localai-paged/patches/paged/`, so the series stays a
+   **1:1, drift-free mirror** of the branch.
+
+Hard rules, no exceptions:
+
+- **NEVER edit the `patches/paged/*.patch` files directly.** They are generated
+  output, not source.
+- **NEVER add a patch to the series that has no corresponding fork-branch
+  commit.** Every `.patch` must be the `git format-patch` of a real commit on
+  `localai-paged`.
+- The fork branch is **where the build and the per-path bit-exact md5 gate
+  actually run**, so it is the **only** place a change is truly validated. A
+  patch living only in the LocalAI series has never been built or gated.
+
+Verify the mirror by tree hash: applying the full on-disk series on the pin
+must reproduce the fork branch tree byte-for-byte. (The patch maintenance
+detail is in `backend/cpp/llama-cpp-localai-paged/docs/PATCH_MAINTENANCE.md`;
+the hard-gate is section 2.5 of `docs/PARITY_HANDOFF.md`.)
+
+## Maintaining the pin against new llama.cpp
+
+The pin (`LLAMA_VERSION` in the wrapper Makefile) is advanced ONLY by the manual
+pin-sync. It is deliberately **excluded from the nightly auto-bumper**
+(`bump_deps.yaml`): a naive bump would shift the tree out from under the patches
+and break `git apply` at build time.
+
+1. **The canary tells you when to sync.** `.github/workflows/llama-cpp-paged-canary.yml`
+   runs weekly: it applies + builds the series against the latest upstream tip and
+   goes **red** when upstream drifts past the patches. Canary red -> run a pin-sync.
+2. **The pin-sync** (recorded in the README section 7 and git history): rebase the series onto the new
+   tip (resolve conflicts; re-export **source-only** with a pathspec like
+   `-- src/ ggml/ common/ include/ tools/ tests/ cmake/`), rebuild on a CUDA box,
+   pass the bit-exact gate on **every** path + `test-backend-ops`, **and confirm
+   the full grpc-server build/link is green on CI**, then bump `LLAMA_VERSION`.
+
+**Hard constraint: keep the pin == the stock `llama-cpp` pin.** `grpc-server.cpp`
+is shared with the stock backend and tracks the stock pin. A paged pin that
+diverges PAST an upstream server-API refactor breaks the grpc-server LINK even
+when the patches are byte-for-byte bit-exact - the bit-exact gate alone does NOT
+catch it. The `c299a92c` bump did exactly this (patches applied + greedy-md5
+bit-exact, but `grpc-server.cpp` failed to link with undefined `stream_*` server
+helpers the refactor pulled into its headers), so it was reverted to `9d5d882d`.
+A pin bump is shippable only once the full CI grpc-server build is green, which in
+practice means moving in lockstep with the stock pin (or vendoring a
+pin-matched grpc-server.cpp, which we deliberately do not, to keep stock pure).
+
+## The bit-exact gate (run for every change)
+
+- greedy md5: `llama-completion -m MODEL -ngl 99 -fa on -p "The capital of France is" -n 48 --temp 0 --seed 1 </dev/null | md5sum`,
+  paged paths prefixed `LLAMA_KV_PAGED=1` (+ `LLAMA_MOE_FORCE_GRAPHS=1` for paged
+  MoE). Must match the recorded baseline. Redirect stdin from `/dev/null` or
+  `llama-completion` hangs in conversation mode.
+- `test-backend-ops` (CUDA0 vs CPU oracle) for every touched op (`SSM_CONV*`,
+  `GATED_DELTA_NET`, `MUL_MAT`, `MUL_MAT_ID`).
+- **The gate is per-path.** The paged-MoE md5 differs from the non-paged md5 - a
+  benign, KL-validated FP-accumulation-order difference (see `docs/PAGED_BITEXACT_NOTE.md`).
+  Compare a paged-MoE change to the **paged** reference, not the non-paged one.
+
+## Encapsulating your work
+
+- When you change a kernel, follow the **Fork-first workflow** above: commit and
+  push on the `localai-paged` branch first, then regenerate the `.patch`
+  (source-only) from the fork so this worktree mirrors the branch byte-for-byte.
+  Commit with sign-off.
+- New optimization -> next patch number (gaps 0005/0027 are intentional). Update
+  the README's patch table and dev notes - keep the README the single doc; do not
+  scatter `*_RESULTS.md` files.
+- Record rejected/flat levers in the README too (they stop the next person from
+  re-running dead ends).
+
+## Follow-ups (Metal / SYCL / Vulkan)
+
+The decode fusions are implemented for **CUDA + CPU only**. The base
+gated-DeltaNet + SSM_CONV ops already exist upstream on Metal, SYCL, and Vulkan,
+so the models **run** there via the non-fused path - what is missing is the
+fusion speedup. Porting it (strictly mirroring the CUDA kernels, since we have no
+Metal/SYCL/Vulkan hardware to test on here) is scoped in `docs/UPSTREAM_LAYER2_SCOPE.md`
+(recommended order: Metal, then SYCL, then Vulkan; ops-first upstream PR, then one
+PR per backend, each gated by `test-backend-ops` on the target hardware). The
+methodology for that work is in [.agents/vllm-parity-methodology.md](vllm-parity-methodology.md).

.agents/vllm-parity-methodology.md

@@ -0,0 +1,101 @@
+# Methodology: Closing the vLLM Decode-Throughput Gap in llama.cpp
+
+This is the playbook that took the paged backend
+([.agents/llama-cpp-localai-paged-backend.md](llama-cpp-localai-paged-backend.md))
+from ~38% of vLLM decode to **parity-to-ahead on dense** (and a proven, honest
+ceiling on MoE) on GB10. Use it for any "make llama.cpp match or beat engine X on
+accelerator Y" effort. The *levers* are model- and hardware-specific; the
+*discipline* below is not. The worked example, with all numbers, is the paged
+backend README.
+
+## The core loop
+
+1. **Establish a bit-exact baseline and gate FIRST.** Record the greedy md5 (per
+   path) and an f32 reference. Every optimization must stay byte-identical to it -
+   or ship as an explicit, default-off precision opt-in. This is what lets you
+   optimize aggressively without silently regressing quality. Gate two ways:
+   greedy md5, and `test-backend-ops` against the CPU oracle.
+
+2. **Profile - do not assume.** nsys the steady-state decode step, broken down per
+   *kernel* AND per *memcpy*. Find the dominant cost. "It's the GEMM" was wrong
+   here: on hybrid gated-DeltaNet models the bottleneck was the recurrent-state
+   **plumbing** (state memcpy + gathers, ~67% of the step), not the weight GEMM.
+   Also sanity-check GPU-busy %: an early "low utilization" reading was a profiling
+   window artifact (decode was 96-99% GPU-busy), not real idle.
+
+3. **Ground-truth BOTH engines.** Decompose *your* decode step AND the
+   competitor's, side by side, per bucket, and compute the per-bucket delta. This
+   tells you WHERE the gap actually is - not where you would guess. It overturned
+   premises here: e.g. vLLM does NOT run the GDN/attn projections as NVFP4 (it
+   keeps them bf16, same as us); the MoE expert GEMM was a llama *win*, not the gap.
+
+4. **Per-lever discipline.** For each candidate: implement -> bit-exact gate ->
+   same-harness A/B bench. Use a runtime env-toggle (flag off vs on) ONLY for
+   levers that are actually runtime-gated; a lever **compiled into** the binary
+   (e.g. the SSM decode fusions here) is NOT isolated by a runtime flag, so measure
+   it build-vs-build. The full-patchset "stock" baseline likewise needs a
+   **separately-built unpatched binary at the same pin** - toggling the runtime
+   flag on the patched binary does not reproduce stock (it measures only the gated
+   part; here that was ~neutral, which is exactly how this gotcha hides). Bank only
+   what lifts AND gates. **Record every rejected or flat lever with the reason** -
+   over time this is the most valuable part: it stops the next person re-running
+   dead ends.
+
+5. **Name the structural floor.** Prove the bit-exact ceiling exhaustively (every
+   lever measured, not assumed). What remains is physical - the memory-bandwidth
+   floor, the irreducible serial-SSM host loop (sampling can't start until logits
+   land). Name it; do not claim more than you measured.
+
+## Hard rules learned
+
+- **Apples-to-apples, or label it.** Stock-vs-patched on the SAME harness
+  (`llama-batched-bench`) is exact - lead with it. But "stock" must be a
+  separately-built unpatched binary at the SAME pin, NOT the patched binary with
+  the runtime flag off (compiled-in wins survive the toggle). Cross-engine "% of vLLM"
+  (batched-bench vs vLLM server+client) is *indicative*; always caveat the harness
+  and config (context length alone shifted the MoE figure 76% <-> 86%).
+- **Re-measure a "win" after later levers land - it may evaporate.** bf16 SSM
+  state (the `ssm_bf16_tau` lever) benched +12% early and failed the f32 KL gate
+  (vLLM keeps f32 too), so it was kept default-off opt-in. Once the decode fusions
+  (recurrent-state gather-fusion + block-table cache) landed, a clean re-measure
+  forcing ALL gated-DeltaNet heads to bf16 (`tau=100000`) went **flat** - 780.6 vs
+  780.0 t/s. The "+12%" was subsumed by the fusions: the lever bought nothing, so
+  it was **dropped** (precision trade + bug surface + extra CUDA template-instantiation
+  compile cost, zero benefit). A win measured before the rest of the series is not a
+  win after it.
+- **Reject the obvious-but-wrong, with evidence.** A faster kernel that is off the
+  critical path benches FLAT (the freed time becomes idle). Quantizing the bf16
+  projections to NVFP4 cost ~6% PPL - and vLLM keeps them bf16 for the same reason.
+  Always measure before believing; a plausible mechanism is not a result.
+- **The gate can be per-path.** Paged vs non-paged attention legitimately produces
+  different (equivalent) FP-reduction orders; validate the difference is benign
+  (KLD to f32) and then gate each path against its own reference.
+
+## Orchestration (multi-agent)
+
+- **One GPU profiler/bencher at a time** (the GPU-contention rule). Parallel
+  design/analysis/read agents are fine; concurrent GPU benches pollute each other's
+  numbers.
+- **Adversarial verify.** Before banking a finding, spawn skeptics prompted to
+  *refute* it; majority-refute kills it. Prevents plausible-but-wrong results.
+- **Anti-punt.** Use foreground, blocking ssh loops with short benches and a
+  progress-file checkpoint. Agents that background work and "wait for the monitor
+  event" stall - forbid that pattern.
+- **GPU coexistence.** On a shared host, stop the user's deployments for a clean
+  benchmark window (with their OK) and ALWAYS restore them (wrap the bench so a
+  failure cannot strand them).
+
+## What generalizes (and what doesn't)
+
+The *speedups* may be hardware-specific (here: CUDA/Blackwell - the SSM fusions,
+NVFP4 FP4-MMA, the occupancy tune), which is why other accelerators did not
+benefit. But the *findings* often generalize and are worth upstreaming: the
+"decode is plumbing-bound, not GEMM-bound" insight and the bit-exact, CPU-mirrored
+fusion ops help any backend running these models. Separate "ship our tuned backend"
+from "upstream the portable op" - they are different deliverables.
+
+## The closing record
+
+Write up the result HONESTLY: the shipped wins, the rejected levers (with reasons),
+the structural ceiling, and the cross-backend / cross-quant generality. Negative
+results are as valuable as wins. The paged backend README is the template.

.docker/llama-cpp-localai-paged-compile.sh

@@ -0,0 +1,39 @@
+#!/usr/bin/env bash
+# Shared compile logic for backend/Dockerfile.llama-cpp-localai-paged.
+# Sourced (via bind mount) from both builder-fromsource and builder-prebuilt stages.
+
+set -euxo pipefail
+
+export CCACHE_DIR=/root/.ccache
+ccache --max-size=5G || true
+ccache -z || true
+
+export CMAKE_ARGS="${CMAKE_ARGS:-} -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DCMAKE_CUDA_COMPILER_LAUNCHER=ccache"
+
+if [[ -n "${CUDA_DOCKER_ARCH:-}" ]]; then
+  CUDA_ARCH_ESC="${CUDA_DOCKER_ARCH//;/\\;}"
+  export CMAKE_ARGS="${CMAKE_ARGS} -DCMAKE_CUDA_ARCHITECTURES=${CUDA_ARCH_ESC}"
+  echo "CMAKE_ARGS(env) = ${CMAKE_ARGS}"
+  rm -rf /LocalAI/backend/cpp/llama-cpp-localai-paged-*-build
+fi
+
+cd /LocalAI/backend/cpp/llama-cpp-localai-paged
+
+if [ -z "${BUILD_TYPE:-}" ]; then
+  # Pure CPU image: one ggml CPU_ALL_VARIANTS build replaces the per-microarch binaries.
+  # arm64: the armv9.2 SME variants need gcc-14 (gcc-13 rejects +sme).
+  if [ "${TARGETARCH}" = "arm64" ]; then
+    apt-get update -qq && apt-get install -y -qq gcc-14 g++-14
+    export CC=gcc-14 CXX=g++-14
+  fi
+  make llama-cpp-localai-paged-cpu-all
+else
+  # GPU build (cublas/hipblas/sycl/vulkan/...): single fallback CPU build, the accelerator
+  # does the compute. Keeps the GPU compile from also building the CPU variant matrix and
+  # avoids the gcc-14 apt step on GPU base images such as nvidia l4t.
+  make llama-cpp-localai-paged-fallback
+fi
+make llama-cpp-localai-paged-grpc
+make llama-cpp-localai-paged-rpc-server
+
+ccache -s || true

.github/backend-matrix.yml

@@ -5177,6 +5177,39 @@ include:
     dockerfile: "./backend/Dockerfile.golang"
     context: "./"
     ubuntu-version: '2404'
+  # llama-cpp-localai-paged: the LocalAI paged-attention llama.cpp variant. Each
+  # row mirrors the corresponding llama-cpp row with backend/dockerfile/tag-suffix
+  # swapped; builder-base-image is left UNCHANGED so these reuse the same
+  # base-grpc-* prebuilt bases (same gRPC + same toolchain), needing no new
+  # base-images.yml variant.
+  - build-type: 'cublas'
+    cuda-major-version: "13"
+    cuda-minor-version: "0"
+    platforms: 'linux/amd64'
+    tag-latest: 'auto'
+    tag-suffix: '-gpu-nvidia-cuda-13-llama-cpp-localai-paged'
+    builder-base-image: 'quay.io/go-skynet/ci-cache:base-grpc-cuda-13-amd64'
+    runs-on: 'bigger-runner'
+    base-image: "ubuntu:24.04"
+    skip-drivers: 'false'
+    backend: "llama-cpp-localai-paged"
+    dockerfile: "./backend/Dockerfile.llama-cpp-localai-paged"
+    context: "./"
+    ubuntu-version: '2404'
+  - build-type: 'cublas'
+    cuda-major-version: "13"
+    cuda-minor-version: "0"
+    platforms: 'linux/arm64'
+    skip-drivers: 'false'
+    tag-latest: 'auto'
+    tag-suffix: '-nvidia-l4t-cuda-13-arm64-llama-cpp-localai-paged'
+    builder-base-image: 'quay.io/go-skynet/ci-cache:base-grpc-cuda-13-arm64'
+    base-image: "ubuntu:24.04"
+    runs-on: 'ubuntu-24.04-arm'
+    ubuntu-version: '2404'
+    backend: "llama-cpp-localai-paged"
+    dockerfile: "./backend/Dockerfile.llama-cpp-localai-paged"
+    context: "./"
 
 # Darwin matrix (consumed by backend-jobs-darwin).
 includeDarwin: