While using the server with the Mistral Vibe agent, I ran into an issue where every single prompt triggered an immediate context auto-compaction, showing that the model's context window was fully maxed out.
The root cause is a mismatch with the OpenAI API spec. In streaming mode (stream: true), every intermediate SSE chunk incorrectly includes a usage object containing prompt_tokens and completion_tokens.
According to the OpenAI specification, usage should only appear once in the final chunk (alongside an empty choices array). Upstream llama.cpp handles this correctly — if you look at examples/server/server-task.cpp, server_task_result_cmpl_partial::to_json_oaicompat_chat() builds delta chunks without any usage field.
However, in this fork, both to_json_oaicompat_partial() and to_json_oaicompat_chat_partial() embed a usage object in every single intermediate chunk. Since standard agents aggregate prompt_tokens from streaming chunks, they end up multiplying the real token count by the number of chunks.
For example:
Prompt: 9,000 tokens x 80 chunks = 720,000 "tokens" reported
This triggers fake "100% context" errors and forces an auto-compaction on every request.
What changed:
I removed the usage block from the partial chunk methods. The _final methods (to_json_oaicompat_chat_stream(), to_json_oaicompat_final()) already send usage correctly in the last chunk with empty choices, so I left those untouched.
(Debugged with some assistance from Mistral Vibe Qwen3.6-27B)
* DS4 optimizations (part 2)
* This is slightly better
* Another minor tweak
* Increase max. number of graph splitinputs to 64
Else with DS4 we can trun into an assert for specific offload
situations with more than one GPU.
llama_token_nl() can return LLAMA_TOKEN_NULL (-1). Falcon3's BPE tokenizer maps
"\n" to zero tokens, so the loader takes its fallback (linefeed_id =
special_pad_id), and the two variants tested reach null by different routes. On
Falcon3-7B-Instruct that copy runs before LLM_KV_TOKENIZER_PAD_ID is read from
the GGUF, so it copies the BPE default, which is itself LLAMA_TOKEN_NULL, even
though the model has a pad token. Falcon3-7B-Base carries no pad id at all and
lands on null whatever the ordering, so a load-order fix alone would not close
this.
llama_sampling_prepare_impl then evaluated logits[-1], an out-of-bounds read one
float before the current position's logit row. Whether that address is mapped
depends on allocation layout, so the crash is configuration-dependent rather
than universal.
This is a crash risk only and cannot change output: the value read is written
back only to a candidate whose id equals nl_token, and no real candidate id is
-1, so it never reaches the sampler.
The fix caches the token once, skips the read when it is null, and skips the
penalize-newline restore, since there is nothing to restore. For a vocab with a
real newline token the block is unchanged.
Repro on Falcon3-7B-Instruct-Q4_K_M, four P100s with the layers split across all
four, -ngl 99 -fa 1 -c 8192, one chat request per trial with a fresh server each
trial: main segfaults 4/4, this change returns HTTP 200 4/4. Across eight models
and both --penalize-nl polarities, 26 greedy comparisons of generated text show
no difference between main and this change on any vocab that has a real newline
token.
Mainline carried this block verbatim until ggml-org/llama.cpp#9294 moved the
penalty stage into the sampler chain, which dropped the raw read and handles the
null id at sampler init instead. ggml-org/llama.cpp#10803 later removed
penalize_nl entirely, so there is no upstream counterpart to port this to.
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* Adding ds4_comp op with CPU implementation
* ds4_comp on CUDA
* ds4_comp: ratio = 4 specialization
Surprisingly small performance gain
* Also handle HCA via ds4_comp
But much smaller gain, if any.
* Delete commented out stuff
* Remove the [(size_t) il] noise
* Minor
* Fix quantized cache
Adds ggml_latent_attn_prefix_ext / ggml_latent_attn_indexed_ext: MLA
latent-cache attention with an always-visible learned K/V prefix
(openPangu's 128 param_sink rows), joint softmax over [prefix | cache],
reading the raw F32/F16/Q8_0 latent cache directly. CUDA implementation
plus a scalar CPU reference that pins the op's semantics; the CPU
backend reports support truthfully, and openPangu adopts the op only on
a non-CPU backend as builder policy.
openPangu routes its dense/SWA/MTP full-span attention and the gathered
DSA path through the op, capability-gated per layer on the attention
output projection's scheduled backend, with the latent cache required
resident on that same backend (--no-kv-offload keeps the unfused
chain); any layer whose backend cannot run the candidate keeps the
exact unfused chain.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* hotswap: keep load-time-derived and transformed tensors coherent after reloads
- Re-derive the MLA combined attn_kv_b (computed_wkv_b) in place when a layer's
attn_k_b/attn_v_b are hot-swapped: the mla>1 prompt-processing path consumes
the derived tensor, so swaps of the source tensors previously had no effect
(KLD stayed exactly 0 in per-tensor benchmarks, e.g. GLM-5.2 attn_k_b/v_b).
- Refuse (loudly) hot-swaps that cannot be correct: views into -mqkv/-muge
merged tensors, khad-folded MLA weights, in-place-scaled ffn_gate_inp_s,
BitNet fused scales, OpenPangu parameter-sink sources, and same-dtype swaps
of mmap-backed tensors. A refused reload produces no 'reloaded tensor' line,
so benchmark drivers quarantine the round instead of recording wrong data.
- Propagate reloaded data to same-name duplicate instances (tied lm head copy
of token_embd, per-layer rope_freqs/rope_factors copies, expert-bias dups),
warning when a duplicate cannot be refreshed.
- Warn that derived state stays stale where a refresh is not possible:
pre-transposed wk_b_pp under -sm graph/attn, requantized MTP head
(output_extra.weight), and k_b/v_b derived from a reloaded attn_kv_b.
- Warn at registration time when -rtr is enabled (restores cannot reproduce
the run-time-repacked state; F16 -> BF16_R16 is lossy).
- server: only attempt the /health hot-swap reload when no slot is processing,
and clear the KV cache + cached prompts after a successful reload (they were
computed with the previous weights).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* Deduplicate llm_compute_wkv_b
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* indexer_topk: fix quantized q8_1 scratch sizing on CUDA
The quantized-K path under-sized its q8_1 scratch buffer: it allocated q->ne[1]*max_rows blocks using the unpadded head dim, but quantize_mmq_q8_1_cuda writes q_padded/QK8_1 blocks per row and must process all q->ne[1]*nrows rows. Size the buffer by (q_padded/QK8_1) blocks x (q->ne[1]*max_rows) rows and pass the full q->ne[1]*nrows row count so the scratch cannot be overrun and every query row is quantized.
CUDA graph identity: INDEXER_TOPK dispatches a source-type-specialized kernel (dense F16 vs quantized cache, F32 vs F16 mask). Source addresses alone do not identify the captured kernel, so snapshot each source type in ggml_graph_node_properties and force re-capture when an INDEXER_TOPK source type changes, preventing a reused graph from replaying the wrong kernel variant when sources are reallocated at the same address.
CPU backend: report INDEXER_TOPK support via iqk_indexer_topk_supported (guarded by GGML_USE_IQK_MULMAT) so the scheduler places the node on a backend that can run it. Harden the scheduler's pass-5 node-assignment check from assert to GGML_ASSERT so a node no backend supports fails as a defined abort under NDEBUG instead of indexing sched->backends[-1].
* indexer_topk: drop CPU capability predicate
* indexer_topk: fall back when unsupported
* cuda: drop speculative INDEXER_TOPK graph type matching
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* initial map to load deepseek 4 arch
* wip
* wip: match graph build and attn logic for dpv4
* wip: Enhance DeepSeek-V4 architecture with new tensor types and sqrtsoftplus gating function
* Update DeepSeek-V4 to support raw key indexing with read/write indices
* fix mismatch in attn_raw
* Enable FA with CSA/HCA
* Fix logit mismatch with FA path
* Clean traces and logs for debug
* Refactor DSV4 tensor handling for MTP execution and improve raw context management
* Refactor DeepSeek4 tensor operations: replace manual weighted sum and post-processing with new helper functions
* Share mHC pre-projection and fix packed DSV4 writes
* DSV4: add shared top-k selection and improve mask handling
* Fix DSV4 c2048 view stride and duplicate loader instantiation
* Reuse shared RMS normalization in DSV4 graph
* Replace DSV4 indexer rotation with shared Hadamard
* Share CSA visibility mask with DSV4 LID
* dsv4: document dependency ordering and reset state
* Remove DSV4 zero-dependency graph shim
* Fix DSV4 packed stream execution
* Remove DSV4 l_out backend override
* Enable DSV4 quantized K-only cache
* Revert "Enable DSV4 quantized K-only cache"
This reverts commit 04f9b425321f62ba60e16d1bea2f8de714cfe855.
* Fix DSV4 quantized cache accounting
* Fail closed on unsupported DSV4 cache lifecycle operations
* Various optimizations
* llama: fix GGML_METAL=ON build - missing ggml-metal.h include in llama-dflash.cpp (#2134)
llama-dflash.cpp calls ggml_backend_is_metal() and
ggml_backend_metal_set_n_cb() inside an #ifdef GGML_USE_METAL block but
never includes ggml-metal.h, so any Metal-enabled build fails to
compile. Add the same guarded include llama.cpp already uses.
* New op: ggml_sum_rows_ext (#2132)
* Add ggml_sum_rows_ext
* openPangu: use ggml_sum_rows_ext also in mhc_post
* openPangu: use ggml_sum_rows_ext also in mhc_tail
* Minor
* Reuse shared inverse RoPE operation for DSV4
* Reuse maintainer CUDA concat implementation
* WIP
* hc_pre
* hc_post
* Remove unnecessary mask manipulations
* WIP
* Take into account swiglu limits
* Turn on fused indexer by default
* Give names to mat mul results
* More named ops
* dsv4: do not uselessly copy the KV cache
+20% TG at 32k tokens
* mask_to_index and make CPU FA work with that
* Much better CPU-only, CUDA still not functional
* Better CPU TG
I'm now at 9.7 t/s for zero context and 6.5 t/s for context of 32k.
PP is 120 t/s for short context and 101 t/s at 32k.
* Even better CPU TG
I'm now at 8.1 t/s for context of 32k tokens.
* Turn off DSA on CUDA for now
* Fix CUDA DSA
* Remove again the unnecessary softmax result buffer
* Experiments
* Various
* More named ops
* Forgot to uncomment
---------
Co-authored-by: samuel <samueloliveira32df@gmail.com>
Co-authored-by: hchengit <95317477+hchengit@users.noreply.github.com>
Several buffer sizes and pointer offsets in examples/perplexity compute counts as
n_ctx / n_token / n_chunk / index times nv (or n_vocab) using int operands. At 16k
context with a >=131k-vocab model, n_ctx*nv exceeds INT_MAX (16384*131076 = 2^31+),
overflowing to a negative int that becomes a huge size_t in resize/ctor -> std::length_error
crash. Affects:
- --kl-divergence-base / --kl-divergence: log_probs.resize, the token/logits writes,
the compare-path base pointer (segfault at 32k) and inner offsets, and the read reserve
- hellaswag_score / winogrande_score / multiple_choice_score: batch_logits(n_vocab*n_ctx)
Cast the counts/indices to size_t at each site. Byte-identical below the overflow threshold;
the offset accesses that were already size_t (i_logits, eval_pairs.first) are unchanged.
Repro: llama-perplexity -c 16384 --kl-divergence-base out.dat on any >=131k-vocab model
(the hellaswag/winogrande/multiple_choice allocs overflow the same way at large -c).
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
For IMROPE models (Qwen3.5/QWEN35MOE) the RoPE op reads 4 position sections per
token, but the legacy null-pos decode fallback (llama_batch_get_one path, used by
llama-perplexity/llama-cli/llama-eval-callback) sized the position vector to n_tokens.
llama_set_inputs then copies n_tokens*4 out of that n_tokens-sized vector, an
out-of-bounds read that feeds garbage into 3 of the 4 rope sections, giving wrong and
run-to-run non-deterministic RoPE. Build the sections like the explicit-pos path (text
t,t,t,0). Standard-rope (NEOX) unchanged; VL image input uses the explicit-pos builder.
Qwen3.5-4B self-vs-self same-top: 91.3% -> 100.0% (CUDA), single-thread CPU likewise.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The Metal backend lacked GGML_ROPE_TYPE_MROPE/IMROPE support, so models
whose GGUF carries rope_sections (e.g. Qwen 3.5 hybrids) could not run
fully offloaded on Apple Silicon.
Add rope_multi_f32/f16 kernels with section-based position handling
(t/h/w/e blocks, 4 position ids per token), imrope's interleaved section
selection, and the corresponding dispatch in ggml-metal.m. Vision-mode
mrope is not implemented and is asserted out explicitly.
Validated on M2: kernel output matches the CPU backend, and Qwen 3.5-9B
(Q4_K_M, -ngl 99) WikiText-2 perplexity over 145 chunks lands within
0.006 of the same model's CPU baseline.
Faster fp32-class replacement for the tile_f32 flash-attention inner path on P100
(sm_60). Bit-identical fp32 arithmetic to the un-retiled kernel (same-top 99.28% =
the float-reorder floor; QK differs only in accumulation order, ~1 ulp; P.V
bit-identical), restructured via half2 K/V shared-memory staging that both cuts
shared memory and leans the inner loop.
Measured +4-9% vs the un-retiled fp32 tile kernel, back-to-back. The speedup is a
COMPOUND of two effects the staging produces together (shares not isolated):
(1) occupancy: smem 36992 -> 28800 B/block admits 2 blocks/SM where the un-retiled
kernel fits only 1 (2*28800=57600<=65536; 80 regs would allow 3, smem is the
ceiling); a genuine 1->2 gain, corroborated by a cross-family perf panel;
(2) a leaner inner loop: ~2x fewer QK-loop smem loads, the dropped score round-trip,
one fewer barrier.
__launch_bounds__(...,2) only PINS the target the smem reduction already reaches (a
(...,1) bound compiles bit-identically), so the directive is not itself a lever, and a
3rd block does not help (cost-curve + an implemented regs-80->124 attempt confirm smem
caps occupancy at 2).
Validated on P100 as the drop-in for the carve-out path (pr-p100-fp16). Reviewed by a
6-model Claude council + a cross-family cloud perf panel; the panel corrected an earlier
'not occupancy' framing to this compound one.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
On P100 (GP100, sm_60) the fp16 vec kernel used for decode (batch<=8)
accumulates the online-softmax denominator and the P*V product in fp16,
flipping ~3-4% of decode top-1 tokens vs an all-fp32 reference
(llama.cpp#25593). Decode is memory-bandwidth-bound on P100, so routing
sm_60 decode to the in-tree vec_f32 kernel is free (tg128 ~identical).
Gated on cc == CC_PASCAL && Q->ne[1] <= 8 (decode only) inside the
!fp16_mma_available block, so the prefill tile_f16 path, the D=256 prefill
vec path, and fast_fp16_available() are untouched, and the
is_pascal_mla_absorbed_decode early-return (MLA) is unaffected.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
ggml_graph_node_has_matching_properties exempted every GGML_OP_CPY node from the source-address check, for both operands. The exemption is needed for KV-cache write copies, whose destination advances each step through the indirect-destination path. It also skips a CPY whose read source (src[0]) moves between graph replays while the surrounding subgraph stays shape-stable, so the captured kernel replays against a stale source pointer and reads the previous step's bytes.
Keep the exemption for the destination operand (src[1]) only, and compare a CPY's read source (src[0]) like any other node. Stable-source copies are unaffected and force no additional re-captures; a CPY whose read source genuinely moves now re-captures instead of reading stale memory.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
llama-dflash.cpp calls ggml_backend_is_metal() and
ggml_backend_metal_set_n_cb() inside an #ifdef GGML_USE_METAL block but
never includes ggml-metal.h, so any Metal-enabled build fails to
compile. Add the same guarded include llama.cpp already uses.
* perplexity: add signal-driven hot-swap mode for persistent KLD/PPL benchmarking
llama-perplexity can now stay resident and be driven through control/status files (reload/compute/exit): it reloads only the tensors that changed on disk and recomputes PPL/KLD without ever reloading the full model. File-based signalling works on Windows, macOS and Linux. The reload returning-to-original path now refreshes tensor data from disk instead of
reattaching stale weights.
* Not Cygwin specific
Loading a deepseek2 model with -sm graph across 2 or more GPUs segfaults
during model load, in the per-head split of the MLA attn_k_b/attn_v_b tensors.
In the ik-converted GGUFs these ship 2D with the heads folded into ne[1]
(DSV2-Lite attn_k_b is [128, 8192]). distribute_mla_tensors_for_split_mode_graph
splits them with prepare_split_tensors(2, ...), which takes the head count
from ne[2]. With the heads still in ne[1] each split comes out n_head times
too large and the upload reads past the end of the host buffer. It is
host-side (compute-sanitizer is clean) and happens at any device count >= 2.
Fix: reshape the 2D layout to 3D [head_dim, inner, n_head] before the split,
the same head-outermost interpretation ggml_reshape_3d already uses on the
non-split path. Metadata-only, guarded on ne[2]==1 so it is a no-op for GGUFs
that already ship these tensors 3D.
Tested on DeepSeek-V2-Lite-Chat Q4_K_M (only what fits across the GPUs here):
-sm graph PPL matches -sm layer and CPU within stderr (7.26 / 7.24 / 7.24,
wikitext-2 n_ctx 2048), coherent generation at 2, 3 and 4 GPU splits, and a
REAP-pruned DSV2-Lite in parity too. GLM-DSA, Mistral4 and larger deepseek
GGUFs share the split path so I expect the same fix to cover them, untested here.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* WIP: indexer_topk on CUDA
* Forgot these
* WIP
* WIP
* This seems to work
* Minor
* Fix bug. Fix suggested by @sayap using GLM-5.2
* GLM-DSA: much better PP long context performance (CUDA)
* DSA: Better way to build the attention mask
* llama-swap latest version
- Sample config still works
* llama-swap latest version ik_llama-cuda.Containerfile
- Sample config still works
* Update README.md
- The prebuilt docker images are fully working, move their instructions to the front
This, to avoid an intermediary loading of the split tensors in RAM.
Commit c32c3819f7 (PR 2102) introduced mmap for split-mode graph tensor
loading to reduce memory usage via MADV_DONTNEED. However, on Windows
dontneed_fragment() is a no-op (no madvise equivalent), so the entire
mmap'd file stays resident in RAM until load_all_data() returns.
This causes excessive RAM usage with larger models than the available RAM on Windows.
This PR fixes it by guarding the mmap path with !defined(_WIN32) and using the original read-buffer approach on Windows, where per-tensor memory is bounded to the largest single tensor size.
* WIP: indexer_topk on CUDA
* Forgot these
* WIP
* WIP
* This seems to work
* Minor
* Fix bug. Fix suggested by @sayap using GLM-5.2
* GLM-DSA: much better PP long context performance (CUDA)
* ggml: add fused sinkhorn op (eps + output-layout params); use it for openPangu mHC
* openpangu: call ggml_sinkhorn directly from mhc_post
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
Route DSA indexer selection through ggml_indexer_topk when -fidx is set: one op computes the weighted-relu head sum plus causal mask and returns top-k rows without materializing the [n_kv, n_ihead, T] score tensor. Off by default; the unfused chunked/full paths are unchanged. Composes with gathered DSA, deferred attention-chunk masks, and the set_rows mask path.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
Kawrakow
Samuel Oliveira Alves
Joel Farthing
pbrejtfus
dmaivel
mb8565
Nexesenex
Thireus ☠
usrlocalben
hchengit
firecoperana
Marian M.
cora4
Jun Yamog