[Piggidragon]

This branch contains no code change. It is a measurement ground and a plan,
rebuilt from beellama/dev at 946c1e5b6 with a clean tree and a clean build.

It started as a search for new decode optimizations for a host-resident KV cache.

Current status update

• R1 / D2: the divergent CUDA FlashAttention barrier is independently reproduced on unmodified current upstream and reported as ggml-org/llama.cpp#27678. An upstream fix must be merged, incorporated here, and revalidated before R1 is closed.
• D3: the MTP + host-resident-KV abort remains confirmed, but D2 is only a correlated upstream defect. It is not proven to cause D3, so the owner and mechanism remain open.
• R9 / D7: PR #34 implements the two-line KVarN build repair. It is pending merge and a clean enabled-KVarN build on beellama/dev; D7 remains historical evidence for this PR's base commit, not a current-dev claim after fix: restore CUDA KVarN build #34 lands.
• R7: deliberately not planned. Making default-build validation cover this KVarN path would add unacceptable CUDA build time; run it only in an explicitly enabled KVarN build when needed.

The optimizations were built, measured, and all but one rejected — and the
measurements falsified two conclusions this lane had already recorded. What is
worth keeping is the evidence and the plan, so that is what this PR is.

docs/kv-offload-roadmap.md	Twelve items: owner, size, blast radius, blockers, acceptance gates
docs/kv-offload-defects.md	Ten defects, each with the exact command and the exact failing output
docs/kv-offload-measurements.md	Every number, and what each one rules out
docs/repro/	One script per defect and per measurement, with how each was verified
docs/probes/	Six patches, verified to apply to pristine 946c1e5b6

Defects, re-confirmed on this base

#	Defect	Status	Owner
D1	Head dim 320 routed to a vector kernel that has no D320 instance	confirmed, exit 134 at fattn.cu:412	BeeLlama
D2	Divergent __syncthreads() in the MMA FA kernel	confirmed, localized, and reported upstream as #27678	upstream
D3	MTP + host-resident KV aborts mid-generation	confirmed, 1 in 3 runs, twice independently; cause not established	investigation; D2 correlation only
D4	test-kvarn aborts on any default build with a CUDA device	confirmed; not planned	BeeLlama
D5	Host-resident KVarN: CPU fallback, 108x the host workspace	confirmed, both arms re-measured	BeeLlama + us
D6	Quantized K/V + logit_softcap	not reproduced — coverage is one case wide	BeeLlama
D7	GGML_CUDA_KVARN=ON does not compile	historical base failure; fixed by #34 pending merge/validation	BeeLlama
D8	q2_0s FlashAttention returns wrong values at hsk=64	confirmed, one case fails at three seeds	BeeLlama
D9	KVarN is refused for a device-resident cache on a default build	confirmed	BeeLlama
D10	test-upstream-merge-keepers-static fails on a clean checkout	confirmed	BeeLlama

D6 is listed precisely because it could not be confirmed. PR #4 records it; the
suite enumerates logit_softcap != 0 only at hsk == 128 and has exactly one
quantized softcap case, which passes.

D1 is broader than recorded. PR #4 gave the confirmed reproducer as GQA 32;
on 946c1e5b6 the suite aborts at nr23=[4,1], kv=512, nb=75.

D7: historical KVarN-build failure; PR #34 supplies the repair

On this PR's base, GGML_CUDA_KVARN=ON did not compile at all — 82 errors, all no instance of function template flash_attn_ext_f16_process_tile at fattn-mma-kvarn-case.cuh(178).

flash_attn_ext_f16_process_tile gained a parameter,
causal_prefix_first_bound_smem, in 0a85856ee. The non-KVarN caller passes it;
both KVarN call sites, at :178 and :245, still pass the old argument list.
git blame attributes the signature change and both stale call sites to the
same author.

The fix is two lines, one nullptr each, and it is verified rather than
asserted: with probes/06 the build
completes with exit 0 and zero errors, without it 82. PR #34
carries that repair. It remains pending merge and validation on beellama/dev.

D2, localized

compute-sanitizer --tool synccheck --show-backtrace device on a pristine
-lineinfo build gives 100 barrier errors, all at one line:

Barrier error detected. Divergent thread(s) in block.
  at flash_attn_ext_f16_process_tile<256,256,4,8,4,0,0,0,1,F16,F16>
     in fattn-mma-f16.cuh:1732

if (np > 1 && threadIdx.y % np == 0) {   // 1687
    __syncthreads();                     // 1732  <-- reported
} else if (np > 1) {                     // 1764
    __syncthreads();                     // 1768  balancing barrier
}

np == nwarps == 4, so only threadIdx.y == 0 takes the if. Balancing the
count is necessary but not sufficient — __syncthreads() is defined only where
its guard evaluates identically across the block, and threadIdx.y % np == 0
does not. git blame puts both lines on upstream's CUDA maintainer
(0208355f42, 95e18884fc); the fork owns neither.

This corrects an earlier note in our own record, which said the balancer "is
present and correct" and concluded the divergence was elsewhere. The first half
was right and the conclusion was wrong.

The measurement that changes the plan

The decode token is copy + compute in series. Two independent routes agree
to within 4%: the telemetry's blocking H2D, and subtracting a device-resident arm.

Depth	host KV	ms/tok	device KV	ms/tok	implied copy	measured H2D	agreement
4,096	31.72	31.53	39.5148	25.31	6.22 ms	6.31 ms	1.4%
16,384	19.69	50.79	36.2507	27.59	23.20 ms	24.05 ms	3.5%
32,768	13.03	76.75	32.7283	30.55	46.19 ms	47.10 ms	1.9%

Remove the ordering constraint — deliver on a copy stream nothing waits on,
incorrect output, upper bound only — and the copy engine hides 96–99% of it:

Depth	ordered	unordered	gain	share of max(copy,compute)
4,096	31.7287, 31.7165	38.4860, 38.5025	+21.3%	98.6%
16,384	19.7025, 19.6820	34.8800, 34.8439	+77.0%	96.5%
32,768	13.0142, 13.0137	20.8258, 20.8300	+60.0%	95.9%

The penalty is the ordering, not the bandwidth. Transfers already run at the
practical PCIe 4.0 x16 rate.

Two things this lane had wrong

• The store side is not the serializer. Removing all 32 blocking D2H copies,
  all 32 CPU set_rows nodes and 32 of 33 splits per token is worth +0.18%
  at 32,768 and nothing at 4,096. The store moves 34.8 KB per token, not the
  2.2 MB an earlier revision recorded — wrong by ~60x, and that error is what
  made the store look like a plausible culprit.
• The transfer can be hidden behind compute. Two earlier attempts concluded
  it could not. Both were argued from structure; both fell to measurement.

Rejected with evidence, each in measurements §5:
async H2D on the compute stream (+0.13%), cross-layer prefetch (+1.38% and not
exact), write-combined pinned memory (−0.85%), huge pages (null result, IOMMU
premise absent on this host), NUMA (untestable, one socket).

Roadmap

#	Item	Owner	Size	Blast radius	Blocked by
R1	Fix the divergent __syncthreads()	upstream	M–L	very high	—
R2	Fix D320 vector routing	BeeLlama	S–M	medium	—
R3	Interim guard for MTP + -nkvo	this fork	S–M	medium	—
R4	Correct copy/compute overlap	this fork	XL	high	design decision
R5	Re-measure the vector/MMA crossover	this fork	S	low	R1
R6	KVarN in the KV-offload line	BeeLlama + us	XL	high	R9, ownership decision
R7	test-kvarn on default builds	BeeLlama	—	—	not planned: unacceptable default CUDA build time
R8	Widen quantized × softcap coverage	BeeLlama	S	none	—
R9	Fix the GGML_CUDA_KVARN=ON build	BeeLlama	S	none	implemented in #34; pending merge and validation
R10	Fix q2_0s FlashAttention at hsk=64	BeeLlama	S-M	medium	—
R11	Make KVarN reachable on a device-resident cache	BeeLlama	M	high	PR #34 merge and validation
R12	Repair test-upstream-merge-keepers-static	BeeLlama	S	none	—

Order. R1 first: it gates the two largest numbers on record — MTP on a
host-resident cache is +75.9% on a 24,000-token prompt where it completes, and
PR #4's vector/MMA crossover is +11.59% on decode at 32,768 and cannot be taken
while lowering the threshold routes more work onto the diverging shapes. It is a
defect fix, not a feature, and it belongs upstream.

Then merge and validate R9 / PR #34. R7 is deliberately not planned because of default CUDA build time. Then the other small BeeLlama items. Then R4.

Two decisions are needed before anyone writes code. Where the staging ring
for R4 lives — fork-local or a generalization of upstream's n_copies
machinery. And who implements KVarN host-residency, which is a conversation
between us and BeeLlama rather than a technical question.

Supersedes PR #32

That PR added a fail-closed guard for KVarN + --no-kv-offload. The guard works
— verified at runtime, it refuses with a named error instead of degrading — but
it converts a slow configuration into an error message, which is not a fix and
gets in the way of the real one. The analysis is folded into D5 and the
integration is roadmap item R6. #32 is closed.

Validation

RTX 4070 (12 GB, sm_89, CUDA 13.3), i5-13400F, one socket, no IOMMU. All
processes under taskset -c 0,2,4, one GPU job at a time behind a lock —
concurrent jobs produced a CUDA OOM and an arm with σ = 1.72 against a normal
σ ≈ 0.06; both were discarded and re-run. Model Qwen3.8-27B-UD-IQ2_M.gguf,
q8_0/q8_0, -b 512 -ub 512 --parallel 1. llama-bench arms are --no-warmup
with reverse-order A/B/A/B pairs; server arms are greedy, seed 1234.

Every d* script was executed end to end and the output in the defect register
is what it printed. The m* scripts are tidied-up forms of the ad-hoc commands
that produced the measurements; m4 has since been run against build-clean,
m1–m3 are verified for argument handling but have not regenerated their
numbers. docs/repro/README.md states this per script.