cache: conversations are stored against their learned prefix

cache: learning outlasts restarts, pins after two prompts, learned prefixes stay in RAM

cache: keep learned prefixes learned through the entry api

cache: learned prefixes are restore points for splices

cache: learn the prefix that distinct prompts share

cache: replay from a learned prefix inside reused turns

cache: learn and pin repeated system prefixes

scr: retain up to 64 matching spans for fragmented transcripts

scr: flush pending MTP rows before reused runs and aborts

docs: add README and standalone patches archive

bonsai: relax prefix cache and SCR floors for one-shot services (cache-tuning)

Relax prefix cache and SCR thresholds for short-prefix one-shot services (web-extract, openwebui): - SCR reuse floor 512 -> 128 tokens - keep warm slot state on failed cache load if common prefix > 0 instead of wiping slot - allow matching prefix cache entries if lcp >= 32 tokens even when f_keep < 0.25 - route to warm slot on system prompt match (lcp >= 32)

bonsai: drop rejected MTP rows on fair step to prevent draft desync (mtp-fair-deferred)

MTP drops the rejected rows of a verification from its deferred queue at accept. Without it a fair prompt step of the other slot flushed them into the draft context, and every later draft of the decoding slot failed (draft positions ahead of the target).

bonsai: q8_0 decode/verify flash attention raw K/V lead staging (fattn-verify-lead)

q8_0 decode/verify flash attention (the 4x8 tile) stages raw K/V with cp.async and starts the next tile's K and V a phase early. Bit-identical; decode +2.9% at 64K depth.

bonsai: fused CUDA kernel for DeltaNet conv path (gdn-conv)

One CUDA kernel for the DeltaNet conv path (concat -> ssm_conv -> silu -> conv-state writes): the concat is never materialized; same arithmetic. Bit-identical; MTP decode +1.5%. GGML_CUDA_SSM_CONV_CONCAT_FUSION=0 turns it off.

bonsai: in-place GDN recurrent state reading from cache rows (gdn-rows)

The DeltaNet layers read each seq's recurrent state straight from its cache row instead of gathering it first (only while every seq's row is its own cell; otherwise the old gather). Bit-identical; MTP decode +1.6-2.3%. GGML_GDN_STATE_GATHER=1 restores the gather.

bonsai: two-slot owed-time fairness scheduler (fairness)

Two-slot fairness: while a prompt is pending beside another active request, each server step is either a decode step or one prompt chunk (cut at the solo ubatch boundaries), picked by owed GPU time, so chat keeps ~22 t/s during a long prefill instead of ~1. No contention runs the old path unchanged. --fair-chunk (default n_ubatch, 0 = off), --fair-decode-weight (default 1).

bonsai: PTQ1_0 tensor-core verify mat-vec register reduction (decode-verify)

PTQ1_0 tensor-core verify mat-vec (MTP verify, 2-8 columns): the 32 residue slots fold in registers in warp_reduce_sum's pairing and only log2(warps) steps go through smem, with CTA geometry per shape. Bit-identical to the 1-column kernel; MTP decode +8.7% (two slots verifying +15%).

bonsai: per-request telemetry and latency metrics (telemetry)

Request telemetry: one record per request (cache/SCR/prefill split, TTFT, t/s, draft, images) plus minute samples, in a ring and capped JSONL under <cache-disk>/telemetry (on by default with a disk tier, --no-telemetry turns it off); GET /stats and /stats/requests, shown on the status page.

bonsai: image embedding cache and NVMe .kvc v3 format (imgcache)

Image conversations in the NVMe tier (.kvc v3: the media chunk map next to the tokens, keyed by image sha256), an image-embedding cache (256 MiB RAM + 4 GiB under <cache-disk>/embd), and a fix for SCR consuming RAM entries that only shared a few chat-template tokens with a new conversation.

bonsai: suffix cache reuse (SCR) for linear attention hybrids (scr)

Suffix cache reuse: after an edit in the middle of a cached prompt (compaction, a trimmed tool output, a system prompt change), the attention KV of the text after the edit is moved to its new positions (q8_0 keys re-rotated) and the DeltaNet layers continue from the pre-edit state, so only the new text is prefilled. Approximate (KL ~0.017 nats/token after a compaction); per request "scr", server default --scr.

bonsai: prompt family classification by system prompt block (cache-family)

A prompt family is the system block (the user turn carries each call's query or page), or without one the opening of the first user turn, so a service's one-shot calls share a family and learn "ephemeral".

bonsai: prompt cache retention policy, classes, and learned families (cache-policy)

Prompt-cache retention: per-entry class (skip, ephemeral, default, keep, pin) from the request's "cache" field, learned per prompt family from reuse, or an optional classifier; TTLs, eviction by class and reuse, pins.

bonsai: HTTP cache inspection and eviction API (cache-api)

GET /cache (slots and both prompt-cache tiers) and POST /cache/clear?scope=slots|ram|disk|all, run on the server loop like the slot actions; a slot that is generating is left alone.

bonsai: pair-tiled 6-head GQA q8_0 prefill flash attention (fattn-q8-prefill)

Prompt batches over a q8_0 cache tile the 6 query heads of a KV head in pairs instead of padding them to 8, stage raw q8_0 rows with cp.async and dequantize exactly in half2. Bit-identical; prefill +8% at 16K, +14% at 32K, +21% at 64K. Decode and verify keep the 8-wide tile.

bonsai: PTQ1_0 MMQ main loop staged per 128-value block (mmq-staged)

PTQ1_0 MMQ main loop staged one 128-value block at a time: one barrier per block, x and y copied with cp.async, trits decoded between the MMAs, B fragments via ldmatrix. Bit-identical; +10% prefill. GGML_CUDA_NO_PTQ1_0_STAGED=1 restores the old loop.

bonsai: Hadamard transform writes per-128 q8 tiles directly (fwht-q8)

The Hadamard transform in front of a PTQ1_0 matmul that takes the MMQ path also writes the per-128 q8 tiles MMQ reads, so the matmul skips its own quantize pass (and the f32 copy when nothing else reads it). Bit-identical; +3% prefill.

bonsai: PTQ1_0 MMQ 128-value activation scaling and int32 accumulation (mmq-y128)

PTQ1_0 MMQ quantizes activations with one scale per 128 values and accumulates a whole weight block in int32, so the float epilogue runs once per block instead of per 32: +15-18% on the matmuls, ~+11% prefill. Not bit-identical: PPL +0.12%, top token unchanged at 99.42% of positions. Decode and speculative verify (mat-vec) are untouched.

bonsai: PTQ1_0 MMQ cp.async double-buffering (mmq-pipeline)

PTQ1_0 MMQ: activations double-buffered with cp.async and the next weight tile fetched into registers during compute. Bit-identical, +3.3% on the matmuls.

bonsai: prefill kernel fusions (ADD+RMS_NORM+MUL and SWIGLU+FWHT) (prefill-fusions)

Prefill fusions, bit-identical: ADD+RMS_NORM+MUL on every NVIDIA GPU (PrismML-Eng/llama.cpp#209 by @cklxx, from professorpalmer's 0040), and the SWIGLU feeding a Hadamard-rotated matmul computed inside the transform's load. ~+2% prefill.

bonsai: NVMe disk tier for prompt cache with DeltaNet checkpoints (prompt-cache-disk)

--cache-disk DIR / --cache-disk-size MiB: prompts the RAM prompt cache evicts, and everything it and the slots hold at shutdown, go to files that a later prefix match reads back (an 8K-token conversation with its DeltaNet checkpoints is ~1 GB and restores in ~0.6 s instead of a ~14 s prefill).

bonsai: publish h for output rows to unblock MTP draft chain (mtp-chain-h)

The MTP graph publishes h only for output rows, as a masked draft context reads it back. With the catch-up rows riding in the draft batch, the second draft step was conditioned on the first catch-up row's h: n-max 2 acceptance 56% -> 73%.

bonsai: PTQ1_0 tensor-core mat-vec for 3-8 columns (ptq1-tc-verify)

PTQ1_0 mat-vec for 3 to 8 columns on the tensor cores (mma.m16n8k32 over decoded trits), so a speculative verify costs about one column instead of ~1.8x at 4. Same bits as the dp4a kernel.

bonsai: in-place q4_0/q8_0 flash attention on tensor cores (q8-kv)

Flash attention reads a q4_0/q8_0 cache in place on the tensor cores (professorpalmer's bonsai-ada-surgery 0009 + 0029), also under batch-invariant mode: a q8_0 cache decodes as fast as F16 at depth in about half the memory.

bonsai: MTP catch-up rows ride in next draft batch (mtp-decode)

Decode: MTP catch-up rows of a verify batch ride in the next draft batch (one MTP pass per round instead of two, rows computed identically), and the conv-state concat takes the transposing kernel. Greedy and sampled output identical; ~+6% decode on the 3060.

bonsai: GDN prefill warp reductions with ahead next-token loads (gdn-prefill)

Gated delta net prefill: next-token loads issued ahead, column-batched warp reductions with the same summation tree, 8 columns per warp. Bit-identical, ~1.9x faster per 1024-token ubatch on the 3060.

bonsai: optimize time-to-first-token on cached prompts (ttft)

Time to first token on cached prompts: no scheduler reserve per request with backend sampling, no rewrite of a just-restored checkpoint, recycled checkpoint buffers, no tiny batch split for a nearby last user message.

bonsai: batch-invariant mode over F16 cache (bi-mma-f16)

Batch-invariant mode sends 1 to 4 query tokens over an F16 cache through one tensor-core tile, instead of the vector kernel that slows down at long contexts.

Test: add head-size-256 decode cases for the FA vector kernel

Fix: keep the q4_0 V dequant of the FA vector kernel in registers

dequantize_V_q4_0 copied the quants into a local int and read them back through an int8_t pointer to it. nvcc 12.8 for sm_120 kept that local on the stack, with 128 generic stores per KV step in the one-query q4_0 kernel, and q4_0 decode fell to 12.4 tok/s at 64K depth on an RTX 5060 Ti against 30.5 tok/s with q8_0. Build the value in a register from two 2-byte loads and take each byte with a shift. Same arithmetic, byte-identical output; tg128 at 64K goes from 12.41 to 29.40 tok/s and the q4_0 kernels lose their stack at head sizes 64, 128 and 256.

Test: add a PTQ1_0 mat-vec case that fails when the kernel skips its PDL wait

Fix: call ggml_cuda_pdl_sync before the PTQ1_0 mat-vec reads its activations

Fix: cap the PTQ1_0 mat-vec at 4 columns and send 5 and above to the MMQ tile path

With the branch-free PTQ1_0 MMQ tile loader (#214) in the tree, the MMQ path is the faster one from 5 columns on. On an RTX 3060 with the K = 5120 projections of Bonsai 2 27B, llama-bench -p 8 runs a batch of 8 in 64.4 ms through the new MMQ tiles against 112.7 ms through the 8-column mat-vec (five-arm a/b, r=3, github.com/sudoingX/bonsai2-small-gpu kernel/ab/ab_215_214_rtx3060.md). The mat-vec keeps its lead at 2 to 4 columns (1.20x / 1.49x / 1.77x of a single pass against 1.46x / 2.21x / 2.75x for the tile path), which is the range speculative verification uses. ggml_cuda_should_use_mmvq now routes PTQ1_0 to the mat-vec for ne11 <= 4 (PTQ1_0_PT_MAX_COLS, was 8) and the 5 to 8 column instantiations of the dedicated kernel are gone. GGML_CUDA_BATCH_INVARIANT keeps its 1 to 4 column guarantee; the comment no longer mentions 5 to 8, which the flag never covered once those batches take the tile path.

Docs: scope the GGML_CUDA_BATCH_INVARIANT comment to the paths the flag changes

The comment in common.cuh read as a whole-model promise. The flag only changes the F16 and BF16 mat-vec paths it selects, the PTQ1_0 mat-vec and flash attention up to 8 queries, for batches of 1 to 4 columns; other weight types and attention shapes outside those kernels can still pick batch-dependent kernels. Say that, and keep the 1 to 4 wording.

Fix: budget the PTQ1_0 mat-vec shared memory the launch will request

The entry guard of mul_mat_vec_ptq1_0_pt_switch rejected a shape when 2 * (K / 128) * ncols exceeded PTQ1_0_PT_SMEM_FLOATS (4096 floats), that is, it budgeted two rows per CTA. The launcher uses 4 rows per work item up to 4 columns and ptq1_0_pt_rows_per_cta never returns fewer rows than one item, so for a long K the launch asked for ncols * 4 * (K / 128) * 4 bytes, twice that with gate fusion, without opting in to more than the 48 KiB default. With a gate and one column that is over 48 KiB from K = 196736 (1537 blocks, 49184 bytes) up to the K = 262144 the guard still let through (2048 blocks, 65536 bytes), and the launch failed with "CUDA error: invalid argument". The guard and the launcher now share ptq1_0_pt_smem_bytes(), which computes the request from the real rows_per_cta and the gate factor, and the guard rejects a shape whose request exceeds the device's default dynamic shared memory per block (cuda_device_info::smpb), so it falls back to the generic kernel. PTQ1_0_PT_SMEM_FLOATS stays the target used when choosing rows per CTA. test-backend-ops gains four eval cases at the boundary: MUL_MAT ptq1_0 m=64 n=1 at K = 393216 (49152 bytes, the last K that fits) and 393344 (falls back), and MUL_MAT_VEC_FUSION ptq1_0 m=1 n=64 swiglu with a gate at K = 196608 (49152 bytes) and 196736 (falls back). On the unfixed head the gate case at K = 196736 aborts with "CUDA error: invalid argument" (ggml-cuda.cu:107, exit 134); with this change all four pass against the CPU reference on an RTX 3060 (sm_86, CUDA 12.4). No Bonsai 2 27B shape was ever affected: its largest K is 17408 (136 blocks), the largest request any of its projections makes is 1 column x 16 rows x 136 blocks x 4 bytes = 8704 bytes, and the gate-fused ffn_up and ffn_gate at K = 5120 request 3840 bytes, all far under both the 16 KiB target and the 48 KiB limit.

Docs: state the batch-invariance guarantee as 1 to 4 columns

Whole-model bit-identical logits under GGML_CUDA_BATCH_INVARIANT=1 are established for batches of 1 to 4 columns; 5 to 8 agree with each other but can differ from 1 to 4. The comments promised 1 to 8.

Fix: keep GGML_CUDA_RESTRICT off the PTQ1_0 mat-vec kernel signature

The macro expands to __restrict__ in the host pass and to nothing in the Hopper-or-newer device pass when PDL is on, so the generated host stub for mul_mat_vec_ptq1_0_pt no longer matched its template declaration and the build failed for sm_90 and sm_120 (CUDA 12.8 with GCC 13.3, CUDA 13.0 with MSVC 14.44). Take plain pointers in the signature and alias them with GGML_CUDA_RESTRICT inside the body, as mul_mat_vec_q already does. No change in arithmetic.

Fix: use the mat-vec kernel for bf16 matrices under 64 rows at 2 to 8 columns

The qwen35 gated-delta-net gate projections are bf16 [5120 x 48]. At 2 to 8 columns on Ampere they took a 10.5 us path where mul_mat_vec_f takes 4.4 to 7.3 us (test-backend-ops perf, RTX 3060); at one column both already used mul_mat_vec_f (3.3 us). 96 calls per token on the 27B model.

Add: GGML_CUDA_BATCH_INVARIANT for batch-invariant small-batch kernels

Off by default, nothing changes. When set, kernels are picked so that the per-column arithmetic of a token does not depend on how many columns (1 to 8) share the launch: F16 and BF16 matrices with at most 512 rows take the mul_mat_vec_f kernel for 1 to 8 columns on Ampere and newer (its K partition depends on K only), flash attention with up to 8 queries takes the vector kernel on Turing and newer instead of switching to the tensor-core kernel at 2 queries, and the KV split (parallel_blocks) is sized as for one query tile so the partial-softmax combine order does not depend on the batch. Together with the batch-invariant PTQ1_0 mat-vec this makes the logits of a token bit-identical whether it is decoded alone or verified inside a speculative batch of 2, 3 or 4 tokens (measured on Ternary Bonsai 2 27B: greedy output with draft-mtp equals greedy output without it on three prompts at n-max 1, 2 and 3). Batches of 5 to 8 agree with each other but not with 1 to 4. Costs throughput at depth (the vector attention kernel at 3 queries over 42K keys).

Test: add Bonsai 2 projection shapes to the mul_mat perf cases

Add: dedicated PTQ1_0 mat-vec kernel with full lane utilization

With 128-element blocks the generic kernel's 128 threads cover 128 K blocks per iteration, so a K = 5120 projection (40 blocks) keeps 31% of the lanes busy and K = 17408 (136 blocks) 53%. This kernel handles plain 2D MUL_MAT (no batch dims, no expert ids, K a multiple of 128, up to 8 columns): it flattens (row group, K block) pairs into one index space with rows_per_cta chosen on the host to fill whole 128-thread iterations, gives each thread 4 rows per K block (2 at 5 to 8 columns) for latency hiding and activation reuse, stores one fp32 partial per (row, column, K block) in dynamic shared memory and has one warp per (row, column) sum them in a fixed order (lane-strided sequential, then a butterfly). The per-block partial is d * sum_k d8_k * sumi_k with exact integer sumi_k, written as __fmaf_rn and __fmul_rn so every instantiation rounds alike, and the order depends only on the weight shape: a column's result is bit-identical for every column count 1 to 8. Fusion (bias, gated GLU epilogue) is supported for one column as in the generic kernel, which remains the path for batched and MoE calls. Bonsai 2 shapes, RTX 3060 12GB, us per call, 1 column: attn_qkv 5120 x 10240 72.4 to 41.8, ffn_up 5120 x 17408 120.0 to 66.7, attn_gate 5120 x 6144 44.8 to 27.7. Whole model llama-bench tg32 26.1 to 39.8 tok/s, pp2 34.1 to 61.1, pp3 33.3 to 72.0.