common_speculative_on_target_batch stored this batch's last hidden into target_hidden_by_seq before reading the map back for the shifted warmup conditioning, so row 0 was conditioned on this batch's last hidden (a future state) instead of the previous call's, and the position-0 zeros fallback was unreachable. Snapshot the previous value before the store; other readers are unaffected. Warmup-only; affects draft acceptance, not correctness.
read_kv_cache_data gated the restore on kv_self.v_trans != (v_state == 1).
v_state == 2 records that the writer had no V cache, while v_trans tracks
flash attention rather than V allocation, so with -fa 0 on a K-only or MLA
cache the two disagree and the restore is refused.
Compare V cache presence in both directions, and transposition only when a
V cache exists on both sides. The write path and serialized layout are
unchanged.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* Short aliases for context checkpoints
* common: fix typos and add/document obvious short aliases
Typos fixed in --help output:
- Swapped long names for -ctv-last/-ctk-last: --cache-type-k-last and
--cache-type-v-last now match their K/V semantics
- --scheduler-async (was --scheduler_async) and fix stray ')' in its help text
- Remove stray trailing commas in --merge-qkv, --merge-up-gate-experts,
--k-cache-hadamard, --v-cache-hadamard, --split-mode-f16, --split-mode-f32,
--split-mode-graph-scheduling
- 'top-n-sigma parmeter' -> 'parameter'
- 'embendings' -> 'embeddings' (2x)
Short aliases:
- New: -okv alias for --override-kv
- Document previously undocumented aliases in --help: -rtr, -cmoe, -ncmoe,
-thp (also adds the previously missing --transparent-huge-pages entry),
-ofreq, -to, -spf
* common: standardize arg parser to short-alias-first ordering
gpt_params_find_arg now consistently lists the short alias before the
full --long-name argument, matching the dominant convention (103 short-first
lines). Reordered 19 options that had the long name first: -gan, -gaw, -dt,
-mea, -ps, -mtprot, -mg, -sm, -ts, -ot, -gfm, -cmoe, -ncmoe, -dr, -op,
-no-ooae, -to, -sps, -wb. Pure style change, no functional impact.
* common: systematize cache-type help order, fix -cram-n-min help
- Reorder the mixed KV cache-type help entries to k-first, k-last,
v-first, v-last for logical grouping (short alias, then full name,
matching the parser convention)
- -cram-n-min now shows its N argument in --help (was missing)
llama_sampler_init_grammar_impl built the trigger_words pattern in a block-scoped std::string, stored a pointer into it (trigger_pattern_c), pointed trigger_patterns at that pointer, then let the block drop both locals before llama_grammar_init_impl dereferenced trigger_patterns -- a read of a dangling pointer into freed std::string storage (UB). Hoist trigger_pattern and trigger_pattern_c to function scope so both outlive the call.
* DSV4 checkpoints WIP
* DSV4 checkpoints: fix per-sequence state save/restore writing all streams
Critical bug: llama_state_seq_get_data() and llama_state_seq_set_data()
were serializing/deserializing ALL DSV4 compressed cache streams instead
of only the stream for the target sequence. This corrupted other active
sequences' compressed indexer state during checkpoint restore.
Fix:
- Add dsv4_stream_offset_size() helper to compute per-stream byte
offset and size for any DSV4 cache tensor (CSA K, LID K, HCA K,
and all state tensors)
- write_kv_cache_data: emit dsv4_single_stream flag + stream_idx
so per-sequence saves only write that stream's portion
- write_dsv4_cache: accept stream_idx parameter (-1 = full tensor,
>= 0 = single stream at computed offset)
- read_kv_cache_data: read the new format, validate consistency
(per-stream data needs dest seq_id, full data needs seq_id=-1),
and restore only the destination stream's bytes into the correct
tensor offset
Format change (WIP, backward compat not required):
[has_dsv4_cache] [n_layer] [single_stream] [stream_idx] [n_stream]
[per-layer: layer_type + stream tensor data]
* DSV4 checkpoints: fix checkpoint search using pos_max instead of pos_min
Bug: checkpoint search condition cur.pos_min < n_past || cur.pos_min == 0
always matched all DSV4 checkpoints because they all have pos_min=0 (the KV
cache starts at 0 and never evicts). The reverse-iterator always picked the
NEWEST checkpoint regardless of how far past the intended rewind/divergence
point it extended, causing n_past to be set to the checkpoint's pos_max
(e.g. 9500) instead of the rewind point (e.g. 5000). This made the system
skip reprocessing tokens between the rewind point and the checkpoint's
pos_max.
Fix in both batch_pending_prompt and apply_checkpoint:
- Condition changed to cur.pos_max < n_past — only checkpoints that end
BEFORE the divergence/rewind point are eligible
- Post-restore n_past uses it->pos_max directly instead of the incorrect
max(pos_min + 1, pos_max) which always returned pos_max for DSV4 anyway
Now when rewinding to position 5000 with checkpoints at pos_max=8000+:
no checkpoint matches, falls through to full reprocessing (correct).
When rewinding to position 8000 with a checkpoint at pos_max=7500: only
500 tokens need reprocessing (optimal).
* DSV4 checkpoints: throttle creation via interval gating on all paths
When ctx_checkpoints_interval > 0, DSV4 checkpoints (~145 MiB) were still
created at every transition point (PP done, TG start, release) because
direct create_checkpoint() calls bypassed the interval gate.
Fix:
- Modified create_checkpoint_at_interval() to handle interval <= 0 as
'always create' (preserving recurrent model behavior for small state)
- Replaced all 4 external direct create_checkpoint() calls with
create_checkpoint_at_interval() so the interval gate is respected
- Unified the prompt-loading branch that had split
create_checkpoint / at_interval calls
Now with --ctx-checkpoints-interval N, all checkpoint creation is
throttled to at most 1 per N positions regardless of the transition
phase.
* DSV4 checkpoints: clarify divergence log message for models with state checkpoints
The message 'does not support partial KV reuse' was misleading for DSV4,
which now supports checkpoint-based state restoration. Split the fallback
message: models with state checkpoints (DSV4, recurrent, hybrid) now print
'no checkpoint before divergence point' instead of 'does not support
partial KV reuse', explaining that the restore failed due to missing
checkpoints at the right position, not due to lack of support.
* DSV4 checkpoints: don't erase the just-restored checkpoint
The erasure condition pos_max > pos_min_thold was equivalent to
pos_max >= pos_next, erasing any checkpoint whose data touched or
went past the current write position. The checkpoint just restored
from (pos_max == pos_next) was immediately erased, wasting a ~145 MiB
checkpoint that was perfectly valid.
Fix:
- Changed erasure condition to pos_max > pos_next (strictly greater
than the next write position). Checkpoints at exactly the current
position (pos_max == pos_next, e.g. the one we just restored from)
are kept.
- Preserved cache-aligned pos_next through the restore block so the
erasure compares against cache positions, not prompt positions
(pre-existing bug where the prompt-tokens call at line 3677
overwrote pos_next with a prompt position).
* DSV4 checkpoints: update interval gate position after restore
After a checkpoint restore, slot.checkpoint_pos was still 0 (from
slot.release()), so the interval gate in create_checkpoint_at_interval
always passed (0 + 2048 <= pos + 1), creating a new ~145 MiB checkpoint
immediately after every restore — even just 5 tokens past the restored
checkpoint's position.
Fix: set slot.checkpoint_pos = it->pos_max in both restore paths
(apply_checkpoint generic restore and batch_pending_prompt DSV4
restore). This tells the gate that a checkpoint already exists at the
restored position, and no new one is needed until another interval
(2048 tokens) has elapsed.
* Missing info
* DSV4 checkpoints: document float-reduction-order reproducibility after restore
After a checkpoint restore, the PP batch loop processes remaining tokens
sequentially from n_past_prompt in chunks of n_batch. Because the loop
is stateless with no carry-over from earlier batches, the chunk boundaries
at and after the restore point are identical to a full-reprocess control
arm. This ensures float-reduction-order reproducibility between arms
when performing correctness validation.
Addresses joelfarthing's finding on openPangu, where mismatched chunk
boundaries between restore+reprocess and full-reprocess controls caused
bit-level differences that masked actual restore bugs.
* DSV4 checkpoints: verify position after restore
After both restore paths (apply_checkpoint and DSV4 batch_pending_prompt),
verify that llama_kv_cache_seq_pos_max() matches the checkpoint's pos_max.
A size-matched but misplaced restore can silently corrupt the KV cache;
on mismatch, force a full reset.
The DSV4 path pre-sets restored = true on byte-level success, then the
position check can revert it to false. Only if (restored) proceeds with
the restored state, matching the apply_checkpoint pattern.
* DSV4 checkpoints: correct misleading comment about chunk-boundary reproducibility
* DSV4 checkpoints: add FNV-1a checksum integrity check for checkpoint data
Sanity check (pos_max): catches misplaced restores (wrong stream offset,
partial overwrite) where the byte count matches but the cache position
doesn't.
Correctness check (FNV-1a hash of serialized data): catches in-memory
corruption of the checkpoint data vector between creation and restore.
Both checks are applied in the standard (apply_checkpoint) and DSV4
(batch_pending_prompt) restore paths. A mismatch in either causes the
restore to be treated as failed, falling back to full reprocess.
File format bumped to CKPT v2 (magic 0x434b5054, version 2) with a
data_hash field per checkpoint. Old-format files (LLAMA_STATE_SEQ_MAGIC)
are still loaded: the hash is computed on load so validation works
uniformly.
* Defer checkpoint hash computation to offload creation path
The FNV-1a hash (~5-15ms per 150 MiB checkpoint) is no longer computed
during checkpoint creation. Instead, data_hash is set to 0 and
hash_computed to false. The hash is computed lazily on first access via
ensure_checkpoint_hash(), called from:
- apply_checkpoint (before the integrity check during restore)
- save_checkpoints_to_file (before writing to disk)
This removes the hash computation from the time-critical checkpoint
creation path, reducing the pause between batches.
* Reuse pre-allocated scratch buffer for checkpoint serialization
Adds a reusable std::vector<uint8_t> scratch buffer to server_context,
eliminating the per-checkpoint zero-init allocation (~150 MiB memset)
from ckpt.data.resize(). The scratch is grown on demand and handed
off to the checkpoint via swap() — a zero-copy move.
Also removes default arguments from server_prompt_checkpoint_update()
since all callers already pass every parameter explicitly.
* Compute checkpoint data hash incrementally during serialization
Instead of a second pass over the serialized buffer (costly for 145 MiB DSV4
checkpoints) or deferring to save/restore time (breaks in-memory verification),
compute the FNV-1a hash as a streaming operation during llama_state_seq_get_data.
llama_data_write_buffer gains an optional fnv_hash pointer and updates it
during write() and write_tensor_data() — the hash is computed from bytes as
they land in the output buffer, with zero extra memory reads.
The server then obtains the hash at creation time by passing &ckpt.data_hash
(pre-initialized to the FNV-1a offset basis) to llama_state_seq_get_data.
This replaces the deferred hash approach (ensure_checkpoint_hash / hash_computed)
and restores in-memory round-trip verification.
* DSV4 checkpoints: add round-trip serialization verification
After restore, re-serialize the KV cache and compare byte-for-byte against
the original checkpoint data. This directly catches serialization bugs
that produce internally-consistent but wrong values (Joel's 59/64 case:
correct position, corrupted tensor data).
The check is added to both restore paths (standard apply_checkpoint and
DSV4 batch_pending_prompt) and runs after the pos_max sanity check and
FNV-1a hash integrity check. Cost: one extra llama_state_seq_get_size +
llama_state_seq_get_data + memcmp of the checkpoint data.
* Remove dead _ckpt_max_size member
_ckpt_max_size was set by server_prompt_checkpoint_update but never read.
Removed the member, the function parameter, and the call site.
* Remove no-op resize after swap in server_prompt_checkpoint_update
After swap, ckpt.data holds the scratch buffer which was already
resized to checkpoint_size. Since n == checkpoint_size (asserted),
the resize is a no-op.
* remove dead (void)has_hash cast
The variable is actually used later (for old-format file detection),
so the unused-variable suppression cast is misleading.
* add missing const qualifiers on to_json() methods
Both server_prompt_checkpoint::to_json() and server_prompt::to_json()
were missing const, preventing use on const references.
* remove duplicate n_kept_prompt assignment in server_prompt::from_json()
n_kept_prompt was assigned twice with the same value, clobbering the
slot where n_discarded_prompt should have been read.
* remove redundant params_base parameter from create_checkpoint_at_interval()
The parameter is already accessible as a member of server_context.
All callers were passing this->params_base, so the indirection was
unnecessary.
* factor duplicated restore verification into verify_restored_checkpoint() helper
The 3-step verification (pos_max sanity, FNV-1a hash, and round-trip
memcmp) was duplicated verbatim across apply_checkpoint() and
batch_pending_prompt(). Extract it into a shared static helper with
a label parameter for context-specific log messages.
Also eliminates the pos_next save/restore dance in apply_checkpoint
by using a local variable for the prompt-limit computation, and
removes a stray commented-out debug printf.
* fix two comments: fnv1a_hash comment was misleading, erasure comment imprecise
- fnv1a_hash() is used for all checkpoint verification (not just
backward-compat file loading) — broadened the description.
- 'may contain stale per-position state' → 'its per-position state
is stale' — the erasure is unconditional when pos_max > pos_next,
so the staleness is definite, not possible.
* remove FNV-1a hash and file-format bump (perf, Joelfarthing's review feedback)
The streaming FNV-1a hash added 65 ms to checkpoint creation and 107 ms
to restore (75 MiB checkpoints; roughly double at DSV4's 145 MiB). The
pos_max sanity check alone is sufficient for catching the real failure
modes (wrong stream offset, partial overwrite), and the initial byte-
count check from llama_state_seq_set_data catches outright corruption.
Removed:
- Streaming hash from llama_data_write_buffer (fnv_hash, fnv_update)
- hash_out parameter from llama_state_seq_get_data / llama.h API
- data_hash field from server_prompt_checkpoint struct
- FNV-1a computation during checkpoint creation and verification
- CKPT v2 file format (revert to LLAMA_STATE_SEQ_MAGIC/version)
- fnv1a_hash() helper function
Kept:
- pos_max sanity check in verify_restored_checkpoint (cheap, catches
misplaced restores)
- Scratch buffer reuse via swap() in server_prompt_checkpoint_update
(pure perf win, independent of hash)
* fix: restore off-by-one in n_past calculation after checkpoint restore
size_up_to_pos(pos_max) returns the number of cached tokens at positions
STRICTLY LESS THAN pos_max (non-mtmd: min(pos_max, size)). Since the
checkpoint encodes state for positions [pos_min, pos_max], the next
position to process is pos_max + 1, not pos_max.
This matters for DSV4 whose accumulator state is not position-indexed:
reprocessing the token at pos_max would double-count it in the compressed
indexer. For recurrent models the old pos_min+1 workaround happened to
give the right answer (since pos_min == pos_max there), but using
pos_max + 1 is correct for both.
Fixes both restore paths (apply_checkpoint and DSV4 in batch_pending_prompt).
* fix: only write/read DSV4 cache section for DSV4 models
The has_dsv4_cache uint32 was emitted unconditionally, changing the
serialized state layout for every model architecture without bumping
LLAMA_STATE_SEQ_VERSION. Old state/session files (which end before
this field) would fail with 'unexpectedly reached end of file' when
read by the new code.
Fix: guard the entire DSV4 section on both write and read sides with
ctx->model.arch == LLM_ARCH_DEEPSEEK4. Non-DSV4 models see the
identical layout they always had.
* fix: validate stream_idx < n_stream in dsv4_stream_offset_size
stream_idx was only checked >= 0 via GGML_ASSERT, but never checked
against n_stream. An invalid seq_id could compute an out-of-range
tensor offset or size, leading to memory corruption.
Now asserts 0 <= stream_idx < n_stream.
* fix: scratch buffer reuse — copy instead of swap
swap(scratch) moved the written data into ckpt.data but left scratch
empty. The next call's resize would then re-allocate from scratch,
defeating the purpose.
Now copies the data (ckpt.data = scratch) so scratch retains its size
and capacity across calls. resize becomes an in-place extension when
needed rather than a fresh allocation.
* fix: restore interval<=0 = disable semantics, split unconditional paths
The interval gate was inverted: interval <= 0 opened the gate, so every
call to create_checkpoint_at_interval created a checkpoint (PP, TG,
release, speculative). This changed the documented behavior ('<=0
disable' per --help) and created extra checkpoints on every decoded
token for recurrent models.
Fix:
- create_checkpoint_at_interval returns immediately when interval <= 0
(restoring the no-op semantics from the original code)
- Unconditional paths (release, PP end, PP start with slot.do_checkpoint)
call create_checkpoint(slot) directly, matching the original layout
- Interval-gated paths (TG tokens, PP start without slot.do_checkpoint,
speculative decoding) stay behind create_checkpoint_at_interval
* fix: gate ALL checkpoint creation by interval, not just TG paths
Three call sites bypassed the interval gate by calling create_checkpoint(slot)
directly instead of create_checkpoint_at_interval(slot):
- PP batch-boundary (was creating mid-PP checkpoints at unpredictable positions)
- PP end (created a checkpoint at every end-of-prompt, even if within the interval)
- release (created a checkpoint at every release, even if just 5 tokens later)
This caused checkpoints 5 and 6 in the log to be created only 5 tokens apart
(pos_max=8488 and pos_max=8493), and checkpoint 7 at release 455 tokens later,
all with interval=2048.
The original code had all checkpoint creation gated by interval (single
create_checkpoint_at_interval function called everywhere). The 'unconditional'
paths were introduced by our earlier fix that split create_checkpoint_at_interval
into a no-op for interval<=0 — but the split was too aggressive, making release,
PP-end, and PP-batch-boundary always fire.
Fix: route all checkpoint creation through create_checkpoint_at_interval, which
already handles do_checkpoint (early return) and interval <= 0 (no-op) correctly.
create_checkpoint is now an internal helper called only from
create_checkpoint_at_interval.
Result: with interval=2048 and an 8494-token prompt, checkpoints are created at
2048, 4096, 6144, 8192 only — the original semantics.
* fix: off-by-one in checkpoint gate condition, use -1 sentinel
The gate condition 'checkpoint_pos + interval <= 1 + pos' opened one
position early for non-first intervals. With checkpoint_pos=6143,
interval=2048: 6143+2048=8191, and pos=8190 gives 8191 <= 1+8190=8191
→ TRUE, creating a checkpoint at pos_max=8190 instead of 8191.
Root cause: checkpoint_pos=0 served dual duty ('no checkpoint yet' and
'checkpoint at position 0'). The '1 +' in the condition compensated
for this at startup but overcompensated later.
Fix:
- Change checkpoint_pos from size_t to llama_pos, initialized to -1
- Drop the '1 +' — condition is now checkpoint_pos + interval <= pos
With checkpoint_pos=-1: -1+2048=2047 <= 2047 → first checkpoint after
2048 tokens (correct).
With checkpoint_pos=6143: 6143+2048=8191 <= 8190 → FALSE (no early
open), 8191 <= 8191 → TRUE (opens at correct position).
Also fixes the mixed signed/unsigned comparison that existed with
size_t checkpoint_pos vs llama_pos pos.
* perf: serialize directly into ckpt.data, drop scratch buffer
The ckpt.data = scratch copy added ~10ms to checkpoint creation
(memcpy of 145 MiB). The scratch buffer was originally introduced to
avoid per-checkpoint resize allocation, but the lazy-zero paging of
modern OSes makes the resize essentially free.
Drop the _ckpt_scratch member entirely. Serialize directly into
ckpt.data after resize — same allocation cost, no extra copy.
* fix: replace GGML_ASSERT with runtime check in dsv4_stream_offset_size
GGML_ASSERT is compiled out in release builds (NDEBUG). An invalid
non-negative seq_id from the public state API would then compute
out-of-range tensor offsets and sizes, leading to memory corruption.
Replace with a runtime conditional that logs the error and sets
safe fallback values (offset=0, size=0). The caller that reads/writes
0 bytes will fail downstream in a defined way.
* fix: restore tolerance mechanism, slot.do_checkpoint bypasses interval gate
Samuel reviewed that we removed the slot.do_checkpoint branch from PP
batch-boundary, but batch_pending_prompt still sets slot.do_checkpoint
when the tolerance threshold is reached. Nowhere checks it, so the
tolerance checkpoint for short prompts (shorter than interval) is dead.
Fix: create_checkpoint_at_interval now checks slot.do_checkpoint — if
true, the interval gate is bypassed. After a successful creation the
flag is cleared so normal interval gating resumes for subsequent
checkpoints. Also handles interval <= 0 + slot.do_checkpoint correctly:
the early-return for disabled interval is itself gated by
!slot.do_checkpoint.
* revert: erasure condition back to cur.pos_max > pos_min_thold
The change from pos_min_thold to pos_next affected all models, not just
DSV4. Revert to the original condition (cur.pos_max >= pos_next after
integer simplification) which correctly erases checkpoints at or past
the write position.
* restore unconditional release checkpoint per firecoperana review
Release is a lifecycle boundary. The interval gate is for throttling
mid-processing checkpoints; the release should always capture the final
state (when do_checkpoint is enabled).
* restore original PP batch-boundary branching per firecoperana review
The explicit slot.do_checkpoint branch in the PP batch-boundary is
restored. The tolerance bypass is removed from create_checkpoint_at_interval
since it was only ever intended for the PP batch-boundary path (the
original code checked slot.do_checkpoint exclusively there). This keeps
the tolerance mechanism from leaking into TG, PP-end, and other paths.
* restore original PP-end checkpoint condition per firecoperana review
The original created an unconditional checkpoint at PP end when tolerance
is disabled (<=0). When tolerance > 0, the tolerance mechanism in the PP
loop handles the end-of-prompt capture at the tolerance point, so no
additional PP-end checkpoint is needed.
* consolidate DSV4 restore path into apply_checkpoint per firecoperana review
The DSV4-specific restore path in batch_pending_prompt duplicated the core
logic of apply_checkpoint (search, restore, verify) with different search
conditions and missing erasure. Consolidate by:
- Adding is_state_ckpt_model flag to apply_checkpoint
- Bypassing the pos_min >= pos_min_thold guard for state-checkpoint models
(DSV4 always has pos_min=0 from no eviction, so the guard blocked entry)
- Using pos_next instead of pos_min_thold for the search condition when
is_state_ckpt_model (allows finding checkpoints at pos_max == n_past - 1)
- Differentiating the reset log message per model type
- Recomputing n_past_offset and n_discarded_prompt after apply_checkpoint
(previously handled in the DSV4-specific path)
* conditional pos_next formula
State-checkpoint models (DSV4, recurrent) use pos_max + 1 — correct
for DSV4's multi-position checkpoints where pos_min=0, po neviction
max(pos_min+1, pos_max) = pos_max, which undercounts by 1.
For recurrent models pos_min==pos_max so both formulas agree.
Non-state-checkpoint models keep the original
max(pos_min + 1, pos_max) formula unchanged.
* remove redundant n_past_offset / n_discarded_prompt after apply_checkpoint
Both n_past and n_past_prompt are shifted by the same delta from the
restored checkpoint, so the difference (n_past_offset) is unchanged.
n_discarded_prompt is not used in the critical path.
* remove redundant speculative-decoding checkpoint per firecoperana review
speculative_decoding_accept is called from within the TG generation loop
which already creates interval-gated checkpoints at n_decoded > 1 (line
4779). The inner call would double-create.
* narrow DSV4-specific search and pos_next formula to DSV4 only per firecoperana review
is_state_ckpt_model includes recurrent models (e.g. Qwen 3.6) where
pos_max < pos_next search semantics may not be appropriate. Only
DSV4 needs pos_max+1 formula and pos_next-based search threshold.
* revert divergence-reset guard to original per firecoperana review
Unnecessary wrapping of the OpenPangu-only divergence path inside
!llama_model_supports_state_checkpoints. The condition is already
specific enough (!llama_model_supports_partial_kv_reuse is
OpenPangu-only), and OpenPangu does not support state checkpoints,
so the original code was functionally identical.
* narrow guard bypass to DSV4 only per firecoperana review
Recurrent state-checkpoint models don't need the pos_min >=
pos_min_thold guard bypass — only DSV4 (which always has pos_min=0
due to no KV cache eviction) requires it.
* narrow reset log message to DSV4 only per firecoperana review
Replace remaining is_state_ckpt_model with is_dsv4 in the
do_reset log branch; remove the now-unused variable.
* cleanup: revert unnecessary newlines, spacing, and comment changes
* fix: restore partial KV reuse for DSV4 in llama_model_supports_partial_kv_reuse
DSV4 has private per-position state but uses state checkpoints to
restore after a mid-sequence divergence. The function was returning
false, causing batch_pending_prompt to reset n_past=0 before
apply_checkpoint could restore from a checkpoint, which broke the
entire checkpoint mechanism.
* Remove bloat
* Reinstate deleted comment
* replace strcmp(arch_string) with llama_model_is_deepseek4()
SamuelOliveirads added the helper upstream — cleaner and avoids
the fragile string comparison.
* inline llama_model_supports_state_checkpoints into call site
Replaced with the inline expression
llama_model_has_recurrent(model) || llama_model_is_deepseek4(model)
and removed the now-unused function from llama.h and llama-model.cpp.
* fix: GCC 13.3 variadic macro trailing comma in SLT_WRN
SLT_WRN expands to LOG_WRN with __VA_ARGS__ at the end. When no extra
args follow the format string, the dangling comma causes GCC 13.3 to
error with 'expected primary-expression before')' token. Use '%s'
pattern consistent with all other zero-arg SLT_WRN callers.
* dsv4_stream_offset_size: bool return, GGML_ASSERT on write, graceful abort on read
dsv4_stream_offset_size silently returned offset=0, size=0 for invalid
stream indices. Now returns bool — writer hard-aborts via GGML_ASSERT
(prevents writing corrupt checkpoints), reader aborts the restore via
return false (handles corrupt checkpoints gracefully).
* feat: add initial speculative benchmark logic
* feat: enhance speculative benchmark with position tracking and JSONL prompt support
* feat: enhance speculative benchmark with batch processing and parameter limiting
* Refactor spec-bench to support Markdown output and new prompt files
* spec-bench: finalize interface inputs and reports
* spec-bench: finish report cleanup
* spec-bench: remove unused code
* spec-bench: improve docs and output details for metrics clarity
* spec-bench: align checkpoint mode with current main
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 04f9b42532.
* 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
