chore(train): engine-10: an f16 recurrent state, GLM-5.3, graph reuse, q8_0 attention read raw, and the PQ2_0 L2 prefetch - #70
Merged
Conversation
…hen empty, and is left alone while busy Since ggml-org#24755 a request's id_slot goes through get_available_slot, and the cache update ran on the named slot as if the similarity loop had picked it: - An idle slot saved and cleared (--cache-idle-slots, the default, with a unified KV cache) holds no tokens, so f_keep was 0 / 0. NaN compares false against 0.5, the prompt cache was never consulted, and a conversation returning to its slot was processed again from its first token: 74.1 s for a 178K-token conversation on an RTX 5090 (8 x 786,432 unified). An empty named slot now always goes through the prompt cache, also with --slot-prompt-similarity 0. A request that names no slot takes the similarity or LRU path, which already did. - The similarity loop skips a named slot that is processing, so f_keep was 0 and the cache update ran on it: the running request's state was saved, and the best cached match for the new prompt was restored into its sequence mid-generation. A busy named slot is now returned as is, and the caller defers the task as it did before ggml-org#24755. test_kv_keep_only_active: a named slot that was saved and cleared restores from cache-ram, and a stream on slot 0 while another request asks for slot 0 generates the same text as alone.
…ties With speculative decoding on, every token after a round's verification was sent with prob 1 and no top list (`result.prob = 1.0f; // set later`, `// TODO: set result.probs`, as upstream still has it): a client asking n_probs or logprobs got logprob 0 and an empty top_logprobs for all of them. With MTP that is most of the tokens. common_sampler_sample_and_accept_n takes an optional on_sample hook, called for each returned token right after it is sampled and accepted, while the sampler still holds that token's row. With n_probs set, the server's verification fills each token's probabilities there with populate_token_probs, as the plain path does after its sample: the logits row for pre-sampling probabilities, the sampler's candidates for post-sampling ones. A verification that restores a checkpoint returns before sending anything, and the replay's verification fills them again. test_speculative: the model drafting for itself, the same request with and without the drafter gives the same tokens, each with its probability within 0.01 and a top list as long as plain decoding's, pre- and post-sampling.
…ts own The checkpoint dedup (7eae848) skipped the copy a resumed task made where its first batch starts, since a checkpoint was there already. That checkpoint belonged to an earlier task, and the min-step thinning in create_checkpoint keeps only the current task's checkpoints and pinned ones: a later checkpoint of the same task dropped it, where the duplicate, the task's own, would have stayed. The next resumes then started at the checkpoint before it. Served at one slot on an RTX 5080 (rig's argv, 24 greedy questions after a 245,755-token prompt), from the fifth question on every question re-processed 5 more tokens than with the duplicate (prompt_n 507 against 512 at q4, cache_n 245,239 against 245,244), and the different re-processed span changed the text on 19 of 24 questions through batch-boundary numerics. The task now takes the checkpoint where its first batch starts as its own, as it would a new one there: the thinning keeps it, the resume points are the duplicate's, and no state is copied a second time. LLAMA_CHECKPOINT_DEDUP_LEGACY=1 still makes the copy.
…leaves out the mask scans that cannot skip anything At one slot every row of a decode's mask attends a prefix of the cells, so neither the range scan (kv_range) nor #63's live steps find a tile to skip: at the served decode shape (245,760 q4_0 cells, head 256, GQA 6, bit mask) they cost 4.5 us of a 377 us call at 1 row and 6.7 us of 397 us at 4 rows with the list against it off, and the range scan another 3.2 us. A one-slot round makes 19 such calls (16 of the target's verify, 3 of the MTP head). ggml_flash_attn_ext_set_mask_prefix marks a flash attention whose rows attend a prefix of the cells, a hint only (op param 4). llama sets it with one sequence (n_seq_max 1), causal attention and a layer without a window. The CUDA launch then leaves out the range scan and the live steps at a decode-sized batch (16 rows or fewer), and the kernel applies the mask over the whole range, as it does with both scans off; a prefill keeps the live steps. GGML_CUDA_FATTN_MASK_PREFIX_LEGACY=1 ignores the hint. test-backend-ops: the hint on random and banded masks, which are no prefix, at 1, 4, 16 and 17 rows (the result must not depend on it), and the served decode shape at 1 and 4 rows with and without it in the perf cases.
…uffer, not uploaded with every graph With a quantized KV cache the K and V Hadamard rotations were built as graph inputs: for a head size of 256 a 256x256 and a 64x64 F32 matrix (278,528 bytes) that the scheduler copied host to device before every graph, each copy a synchronous one. A speculative round runs five graphs. The CPU and CUDA backends transform these sizes with FWHT and never read the matrix, so the copies bought nothing. The matrices are now tensors in a buffer of the first layer's buffer type, set when the cache is made, and the graphs take them like the K and V tensors themselves. A cache sharing another's cells takes that cache's tensors. With no_alloc they sit in a dummy buffer, as K and V do, and the memory breakdown counts them. LLAMA_ATTN_ROT_INPUT_LEGACY=1 builds them as graph inputs again.
…on the backend, not every logit on the CPU A chain drawing on the CPU whose top-k is the first sampler to change a logit (sampler_top_k_first) reads only the k largest logits of a row, yet every decode copied each output row's n_vocab logits to the host and the CPU scanned them all for the k: with a 248,320-token vocabulary and MTP drafting 3, 3.97 MB copied and four scans a round, and top_k_candidates held 34 % of the server thread's own CPU time (perf on an RTX 5080, one slot at 245,760 tokens). A slot that samples on the CPU now attaches a chain of that top-k alone to the context: the graph takes each row's top k, a decode copies k logits and ids a row, and the slot's chain and RNG draw from them as from the k the CPU took, the candidates sorted by logit and then id as top_k_candidates sorts them. Not with a reader of every logit (pre-sampling probabilities, the pull detector, the lens); taken off as the chain is when a lazy grammar triggers or a reasoning budget forces, a verification ending at the token that switched it on as it does for backend-sampled rows. LLAMA_TOP_K_PREFILTER_LEGACY=1 keeps every logit on the CPU. test_top_k_prefilter: the same tokens and post-sampling probabilities with the top k taken on the backend and on the CPU, seeded at temperature 1 with every candidate kept, and across a lazy grammar's trigger and a reasoning budget's end.
…ts in 64 bits With flash attention off the V cache is transposed, one row of kv_size cells per V dimension, and the state save and restore address row j at j * kv_size. Both factors are uint32_t, so the product wrapped at 2^32: - the save (state_write_data) also took the element size as uint32_t, so the whole byte offset was 32-bit and wrapped once a layer's V reached 4 GiB: an F16 V of 1024 dimensions at 2,097,152 cells wrote its row 1024 over row 0; - the restore (state_read_data, both the contiguous and the scattered path) wrapped at 2^32 elements. Each read or write stayed inside the tensor, so nothing failed: the saved state, or the restored cache, was silently wrong. The served cache runs flash attention (V not transposed) and is not affected. The product is now taken in size_t. No test: reaching the wrap takes a V of 2^31 elements or more in one layer, several GiB, beyond what CI holds. Found by rig-lead's review of c1518d4.
…t runs in its compute's queue, and says when it stops adc853e's staging sets a split's small host inputs asynchronously and trusts the backend's stream to run each copy after the work still reading its destination. A CUDA stream (CUDA, ROCm, MUSA) does. Vulkan runs an async set on its transfer queue when it has one (the default on an AMD dGPU past GCN, or with GGML_VK_ASYNC_USE_TRANSFER_QUEUE), and its semaphore orders the transfer before the next compute only, not after the last one: a copy could overwrite an input copy the previous compute still reads (a prompt's earlier ubatch, queued without a wait), changing the attention mask or the KV indices under it with no error. Such backends keep the synchronous copy, as before adc853e; CUDA's path is the same code as before. A backend that stopped staging also did so silently for the scheduler's life: a missing event or staging buffer now logs a warning, and a backend that cannot stage a debug line. Found in the review of the engine-8 and engine-9 commits.
…ens until they part a8d2290's test asked the model drafting for itself to sample the same 16 tokens as plain decoding at temperature 0.2. On a 128-CPU box with the server's default threads the two runs round differently and part at the 15th token (the drafted run takes 982 where plain decoding takes 3186), on engine-9 as on engine-10, so the test failed before it read a single probability; with 8 threads they agree. The probabilities are now compared over the tokens both runs share, at least 4 of them: a verified token without its probabilities still fails it at the second token.
… when it cannot list the process's threads b07f082 leaves the placement to the scheduler when the GPU's node has fewer CPUs than the CPU backend's threads, but counted the target's threads only: with --spec-draft-threads 32 --threads 8 on a node of 16 CPUs the process was bound and the draft's 32 threads spun in ggml_barrier on 16 CPUs. The draft's threads (they default to the target's) now count too. set_all_threads_affinity ignored the error of listing /proc/self/task: with the directory unreadable no thread moved and the bind still logged "0 threads on CPUs ...". The listing's error now logs a warning, and an increment's error ends the walk instead of throwing. No test: both need a host with more than one NUMA node. Found in the review of the engine-9 commits.
state_read caught every exception of state_read_data (a state blob cut short: "unexpectedly reached end of buffer", a failed allocation) with catch (...) and rethrew a bare "failed to restore kv cache", so the reason reached no log. It now logs the exception's message first; the cache is cleared and the restore fails as before. Found in the review of 9123de6.
…d which layouts it skipped
On the CI model (heads 48 wide) the served q4_0 layout is always skipped, and the verdict "all restores match" read the
same with 8 checks run as with 12: only a line further up said a layout was skipped. The verdict now counts the restores
checked, the failures and the layouts skipped ("all restores match: 8 restores checked, 0 failed, 1 of 3 layouts
skipped"); the exit code is unchanged. The q4_0 layout is checked with a model whose heads are a multiple of 32 wide,
such as the served pack (12 of 12 on an RTX 5070 Ti with 7d932e9).
Found in the review of the engine-9 commits.
…0, false, no and off are off everywhere The 48 off switches each parsed their variable by hand. 46 took atoi(value) != 0, so GGML_CUDA_PQ2_MMA_LEGACY=true or =yes left the new path on with no word, and an A/B run whose control leg set the switch that way measured the change against itself. GGML_CUDA_FWHT_LEGACY and GGML_CUDA_GRAPH_KEY_LEGACY took any value, 0 included, so =0 turned them on while =0 turns every other switch off. ggml_env_switch (ggml.h) now reads them all: unset, "", 0, false, no and off are off, 1, true, yes, on and any other number are on, words in any case, and any other value is on with a warning naming it. =1, the value TORAD.md gives every switch, reads as before, and nothing changes with the variables unset. test-arg-parser checks the values. Found in the review of the engine-8 and engine-9 commits.
…v file with CRLF lines stays off 4384bbb's ggml_env_switch matched the value whole: "0\r" (every value of an env file saved with CRLF lines, as docker's --env-file passes it) or "0 " was no word and no number, so it was taken as on with a warning, where atoi read it as 0. The spaces around the value are dropped before it is matched, and the warning prints the value without them. test-arg-parser: " \t", "0\r" and " off\n" are off, "1\r" and " yes " on, "0 0" on with a warning.
…ect, so a free-form tool argument parses
json-schema-to-grammar gave every empty schema the "object" rule, before its own next branch (no type and no
structural keyword accepts any value, as JSON Schema has it) could take {}. Below the root that is wrong: a property,
an array's items or additionalProperties given as {} (zod's z.any(), z.record(z.any())) accepted only objects. Exa's
agent_run declares outputSchema as {"type": "object", "additionalProperties": {}}, so the served grammar forced each
of its values into an object: {"type": "object", ...} came out as {"type": {...}, "required": {...}} and ran to
max_tokens, and Exa rejected 43 of 43 calls with INVALID_OUTPUT_SCHEMA (rig-lead's report).
{} now takes the "value" rule, as in upstream's ggml-org#28736, without its rewrite. An empty schema still accepts any object
where it is the whole payload: json_schema_to_grammar's root and the Python twin's (response_format json_object, an
empty json_schema), and a chat parser's schema node for tool arguments or a response format. A tool argument's value
(tool_arg_value, tool_arg_json_value) marks its schema node, so an argument whose own schema is {} accepts any value too.
Tests: test-json-schema-to-grammar's "empty sub-schemas (any value)", C++ and the Python twin (a property, items and
additionalProperties of {} give value; "empty schema (object)" still gives object), and test-chat's Qwen3-Coder case,
outputSchema {"type": "object", "required": ["a"]} and 42 for an argument whose schema is {}.
… gives {} below the root
"array with empty items" and "array with empty items and prefixItems" (upstream's ggml-org#19968, written when every {}
took the object rule) still expected item ::= object. Since 737eba9 an array's items of {} accept any value, so
test-json-schema-to-grammar failed at the first of them, in C++ and in the Python twin alike. The rest of each
grammar is unchanged: the object rule already pulled in value and its rules. Found by feature-dev's code-reviewer
on 737eba9.
… too, as with GBNF and as -j documents
json_schema_to_grammar handed llguidance the schema as given, so with LLAMA_LLGUIDANCE a root of {} accepted any
value: -j '{}' and -jf (documented in common/arg.cpp as "{} for any JSON object"), json_schema {} on /completion
(server-schema.cpp) and the legacy chat path's schema (chat.cpp). GBNF takes such a root as any object, through
visit's empty-schema branch before 737eba9 and through json_schema_to_grammar's rewrite since, and that rewrite
ran after the llguidance return. It now runs first, and both take the same schema. No llguidance build here; the
rewrite is the one test-json-schema-to-grammar's "empty schema (object)" checks. Found by pr-review-toolkit's
silent-failure-hunter on 737eba9.
…l fails a probability never set A drafted token's probabilities come out of a verification batch and plain decoding's out of one-token decodes. On the CPU the two round apart: on seed 4242 they were equal to the last bit through the first verification round and then differed by up to 0.017 in logprob and 0.057 post-sampling, over the 0.01 the test allowed. Neither graph reuse (LLAMA_GRAPH_REUSE_DISABLE=1) nor the fork's KQ mask scans (both *_LEGACY) changed a digit of it, and a prompt batch of the same tokens without drafting gave the drafted values to 4 decimals on every token it listed, on seeds 4242, 1, 2 and 3 in both modes. Over seed 4242 and 0 to 63 on an AMD EPYC 7532, 17 of 65 runs part in each mode, by up to 0.096 in logprob and 0.112 post-sampling, and the drafted values equal the prompt batch's to 4 decimals wherever the batch lists the token. The tolerance is now 0.25, about twice that. The server sends a token whose probability it never set with exactly 1 (logprob 0), and 0.25 takes that wherever plain decoding gives 0.75 or more (found by pr-review-toolkit's silent-failure-hunter). A drafted value of exactly 1 (logprob 0) now passes only where plain decoding's is within 1e-3 of it: rounding cannot make exactly 1 of 0.999, which would take the other tokens' share from 1e-3 to under 6e-8. The comment put the runs' parting on the server's thread count; the run behind that had servers that never bound their port, and every answer came from a leftover server. It now says what was measured. Engine-9's bug, drafted tokens with no probabilities, stays outside the test: logprob 0 against -1.58 at the first drafted token, and an empty top list against 3.
…as CUDA and the CPU scale only f32 a941add gave the recurrent state narrower types (--cache-type-s f32|f16|bf16|q8_0) and zeroed a fresh state's row (rs_z) in the graph with ggml_scale_inplace for every type but the quantized ones, on the premise that GGML_OP_SCALE takes f32 and f16. It takes f32 only: ggml-cuda/scale.cu asserts an f32 source and destination and the CPU's ggml_compute_forward_scale aborts on any other type, while CUDA's supports_op claims SCALE for every type. So -cts f16 or bf16 aborted at the first decode that zeroes a row (rig-lead, on an RTX 5090: three legs, each at that assert). Only an f32 state is now scaled in the graph; zero_rs_z zeroes every other type's row on the host, as it did for q8_0, and all-zero bytes are 0.0 in f16 and bf16. f32 and q8_0 take the same path as before. rs_z is set only when a free state cell lies in the batch's range, so a steady decode zeroes nothing. rig-lead's tap at depth 131,072 (737eba9 with this change, 8,192 tokens decoded, full-vocab KL against f16 K/V and an f32 state): a q8_0 state 0.001839 and climbing, bf16 0.000292, f16 0.000108. An f16 state drifts 17 times less than q8_0, for 34 MiB more a state copy.
… is exactly 1 7f1a3be's comment argued a real exact 1 only from the other tokens' share rounding away. feature-dev's code-reviewer found the second way: llama_sampler_dist_apply sets exactly 1 when top-p or min-p leave one candidate (src/llama-sampler.cpp:1164). The check holds for it too. Top-p and min-p cut before the temperature, so where one run keeps one candidate and the other two, the second has at most 0.053 of the first's raw probability, 4e-7 of the mass at temperature 0.2, and plain decoding stays within 1e-3 of 1. Plain decoding under 0.999 needs a second candidate above 0.25 of the first's, 1.5 nats of logit from any cut, where the runs differ by 0.1. The comment now says so.
… it fails, as df4f8e5 made the KV cache's llama_memory_recurrent::state_read caught every exception of state_read_data with catch (...) and rethrew a bare "failed to restore kv cache": the pattern df4f8e5 fixed in llama_kv_cache::state_read only. A hybrid model (Gated DeltaNet layers and attention layers) restores through both, so a recurrent state cut short still failed with its reason in no log. It now logs the exception's message first; the state is cleared and the restore fails as before. These two were the only catch-alls in src/. Found by feature-dev's code-reviewer in its review of engine-10's release notes.
…ay to a real probability of exactly 1 9e64f11's comment named top-p or min-p leaving one candidate and the rest rounding away after the temperature, both post-sampling only. The same check runs on logprobs, which get_token_probabilities computes as a softmax of every logit with no sampler and no temperature (server-common.cpp:1498): there an exact 1 comes only from the other tokens' share rounding to nothing, and parting it from 0.999 takes about 9.7 nats of that share, above the 1.5 the comment gives as the least. The comment now says which way applies where. Found by pr-review-toolkit's silent-failure-hunter in its review of 9e64f11.
…th, as rows mode reads f32 only 84cba7d sent a quantized state cache to the gathered path because rows mode reads the state row in place and ggml_gated_delta_net_rows asserts an f32 state (ggml.c:6408). It tested !ggml_is_quantized for that, so f16 and bf16 (a941add's -cts f16|bf16) still took rows mode: on a CPU-only or all-Metal run with an MTP, EAGLE3, DFlash or DSpark draft (n_rs_seq > 0) the graph build aborted at that assert. Rows mode is now selected for an f32 cache only, as 84cba7d's message says; f16 and bf16 take the gathered path, whose get_rows converts the row to f32 and whose set_rows converts it back. CUDA never selects rows mode and is unchanged; f32 and q8_0 take the same path as before. Found while measuring 1fd214c's f16 state: the same premise, not quantized taken for f32, as its rs_z scale.
…h it Graph reuse compares the sampler map by address. Since 3ab698f and 383771d a sampler that leaves or attaches costs no scheduler reserve, and the reserve was what reset the previous graph: a chain freed after its request and the next request's chain allocated at the same address before any decode compared equal to the graph built with the old one. That graph was reused with the old chain's nodes (top-k's k and a temperature are baked into it) and the inputs the detach had released: a dist chain aborts on its empty uniforms (llama-sampler.cpp:1342), and the server's top-k prefilter chain, attached at every launch without -bs, draws from the previous request's k. set_sampler now resets the previous graph when it releases a sequence's sampler. A changed map rebuilt the graph anyway, so this costs a build only where the addresses matched, and no reserve. Tests: test-backend-sampler reattach_same_address (a chain that leaves and comes back before a decode, the deterministic form of the address reuse; the same map and shape reuse the graph first, so the check can fail): reused before the fix, rebuilt after. All 22 tests pass on the CI model (Qwen3-0.6B-Base q4_0, CPU); logit_bias and penalties fail on stories15M on the base sha too, from that model's logits.
…not dropped - --spec-draft-mtp-vocab, -swa, -decode-only and -window without --spec-type draft-mtp: no draft context is built, so the vocabulary file was never opened and the server served as if it had been. The server now refuses them at load. - -bs with --pull-layers: the pull detector reads, and may edit, the served logits before the sampler, so every slot sampled on the CPU with nothing said. The server now warns once at load. - The init warning that turns backend sampling off named the grammar whenever there was one. A forcing reasoning budget holds a lazy grammar off, so the warning now names the budget in that case. Tests: test_mtp_vocab_without_draft_mtp_is_refused, test_backend_sampling_with_pull_detector_is_named_at_load (with and without -bs) and test_backend_sampling_off_names_the_constraint (a budget of 0 whose start tag is the generation prompt, beside a lazy grammar awaiting its trigger). All three fail on 48ebd21 and pass here; test_speculative.py and test_completion.py: 58 passed, 1 skipped (not slow, CPU).
…rolls back in place An n-gram drafting ahead of a draft model copies runs longer than the draft model's n_max. The recurrent state keeps one snapshot per slot per token up to n_max, so every longer draft took the host checkpoint, the sampler clone and, when rejected past n_max, the replay. --spec-rollback N keeps N snapshots instead; 0 (the default) keeps the draft model's n_max, as before.
…to the reference state Covers the conv-state change that landed in 5531e53 (build_conv_state writes only the last min(n_seq_tokens, K) snapshot slots, as build_recurrent_attn does for the state): a 4-row verify under n_rs_seq 8, rolled back by 1, 2 and 3 rows, against a context that decoded only the accepted rows. The recurrent state is compared as the context serializes it (conv windows and state), then the logits over a correction token and two more. The synthetic qwen35 model's recurrent branch sits below the rounding of its residual stream: its logits stay bitwise equal with a wrong conv window, so the state comparison is the check. It fails on qwen35 with the window one token early (2.1e-5) and with the slots past the first left unwritten (1.5e-5), against 0 to 7e-9 with the change, on qwen35, nemotron-h and dsv4.
The table kept the token that followed each n-gram and nothing of the n-gram, so a lookup landing on a slot another n-gram filled returned that n-gram's token. The drafter's table is shared across requests and reset at 25% occupancy, so up to a quarter of the lookups at text it never saw drafted a stranger's continuation, and the n-gram drafts ahead of a draft model: each such draft took the draft model's round and verified to the one token the target samples. A slot now keeps the high half of its n-gram's 64-bit hash beside the token (the low half still picks the slot), and a lookup whose high half differs finds nothing. The table grows from 16 to 32 MiB. test-ngram-mod: a 4M-slot table filled to the reset threshold found 21,793 of 100,000 unseen 16-grams before, 0 after.
… batch wrote
A batch of n rows writes rollback snapshots 0..min(n, n_rs_seq + 1) - 1 (the state s rows before its end). The state
before the batch is kept in no slot, and neither is one from an earlier batch. seq_rm accepted any rollback up to
n_rs_seq and read the slot anyway: rolling back a whole batch, or past it into an earlier one, restored a stale state
with no error. The upstream multi-seq test rolled back a whole 3-row batch and its logits matched only because a small
model's recurrent branch sits below the rounding of its residual stream; its serialized state was 2.2e-5 (qwen35) and
5.9e-4 (nemotron-h) off the truth.
The contract, stated in llama.h next to n_rs_seq: a rollback reaches min(n_rs_seq, n - 1) rows into the seq's last
batch, and seq_rm refuses one reaching further. Both memories keep it as a per-seq rs_reach:
- llama_memory_recurrent sets it in the context's apply() from the ubatch's rows. It is 0 after clear, state_read and
seq_cp's destination, which rolls back only into a batch of its own.
- llama_kv_cache_dsv4 sets it for the seqs a ubatch snapshots (a unified stream, its first seq only). dsv4's own
plane d = rows is right, but the reach stops at rows - 1 so every memory keeps the one contract. dsv4's 0.042 after
single-token replays was this defect: dsv4_build_comp_plan writes plane d > rows with the ubatch's starting state.
Refusal, not a copy of the starting state into slot n, because no served path rolls back to a batch's first row:
- a verify rolls back draft + 1 - accepted <= draft = rows - 1 rows (server-context.cpp:4570), and a draft longer
than n_rs_seq takes the checkpoint path;
- the prompt path truncates at the cached end or after a checkpoint load (server-context.cpp:3941);
- draft contexts are built with n_rs_seq = 0 (speculative.cpp:2662).
The copy would cost one full recurrent state per layer on every verify (about 0.5 % of served throughput) for a path
nothing takes. A caller that does step out of the contract now aborts in common_context_seq_rm instead of reading a
stale state. That includes llama-context's cleanup after a failed ubatch, which used to roll back the whole ubatch
into a stale slot; it is now refused and the next decode fails on the positions.
The test holds every rollback to a context that decoded only the kept rows:
- main's replay is one verify-shaped batch; rolling it back whole, or one row past it, must be refused, then all but
its first row roll back and match the truth;
- a rollback right after a partial restore, and right after a full one over a dirty context, must be refused;
- the multi-seq test's tail is one row longer than its rollback; the whole tail must be refused.
It passes on every test model with n_rs_seq (bailingmoe3, deepseek4, lfm2, lfm2moe, nemotron_h, nemotron_h_moe,
qwen35, qwen35moe) and fails on the old code and on six one-line mutants: the recurrent memory's and dsv4's bound
back to n_rs_seq, their reach counting the batch's first row, and each memory's restore keeping the reach.
examples/rs-rollback{,-multi} roll back across single-token batches, the path the GDN op documents as unwritten
("trailing slots are left unwritten"). On the served Bonsai pack the old code let them read those slots: 109 of 128
logits rows differed from the reference (multi: 79 of 96 on the rolled-back seq, 0 on the other). They now report
ROLLBACK REFUSED.
… reads q4_0 The raw path (60feea0) reads q4_0 K and V straight from the cache into shared memory, once per GQA group, where the stock kernels either copy the whole cache to f16 first (MMA) or read it once per Q head (vec). A q8_0 cache still took the stock kernels. The raw path now takes its cache type as a template parameter (type_K, type_V; f16 is the stock kernel) and gains q8_0: - a row is 256 values in 8 blocks of 34 bytes (an f16 scale, 32 int8), 272 bytes, copied in 16-byte chunks; - V is dequantized in shared memory with a byte flip under the f16 exponent 0x64 (1152 + q, minus 1152, times the scale), bit-identical to dequantize_block_q8_0; - K*Q runs on int8 tensor cores (mma m16n8k32 s8.s8) against Q quantized per 32 values, the accumulator seeded with the float bias 2^23 + 2^22 so the int32 sum lands in the mantissa exactly (|dot| <= 32*127*128 < 2^22); q4_0 keeps s8.u8 with the bias less 8 * sum(q). Head size 256, GQA > 4, K and V of one type, NVIDIA Ampere and newer; everything else takes the stock kernels. GGML_CUDA_FATTN_Q8_0_LEGACY=1 restores them for q8_0, as GGML_CUDA_FATTN_Q4_0_LEGACY=1 does for q4_0. Decode takes the raw kernel from 4,096 cells of context on, as q4_0 does. RTX 5090, Ternary Bonsai 2 27B (GQA 6, 16 attention layers), llama-bench -r 3, one binary A/B through the switch: - q8_0 decode tg64 at depth 245,760: 44.95 -> 83.74 tok/s (1.86x); at 4,096 / 8,192 / 16,384: 151.32 -> 153.34, 145.54 -> 151.11, 133.88 -> 147.20; - q8_0 prefill pp512 at depth 65,536: 2,214.63 -> 2,800.09 tok/s (+26 %); - q4_0 unchanged by the rework: tg64 at 245,760 103.22 before, 103.29 and 103.36 after; pp512 at 65,536 3,174.14 before, 3,202.41 after. A layout that keeps only the K block scales where the f16 K tile was (15.9 KB less shared memory per block) measured 83.70 and 103.22 tok/s at 245,760 and 2,941 / 3,198 pp512, inside the spread of the kernel as committed; not taken. test-backend-ops FLASH_ATTN_EXT, the GQA-6 cases (q4_0 and q8_0; kv 2,048 / 4,096; 1, 3, 8, 13, 32, 75, 512 tokens, now including the 8 token x 8 head tile a drafted verify batch takes): 276/276 OK. Corrupting the V dequant or the K*Q operand turns the q8_0 cases FAIL, the 8 and 13 token ones included.
…uler llama_context reused only the previous graph (gf_res_prev). A speculative verify alternates shapes, e.g. an n-gram draft of 7 rows and then MTP's 3, and every change of shape rebuilt the graph. On the 27B hybrid on a 5080, one rebuild costs ~5 ms of host time to build, split and allocate the graph. The CUDA backend then resets that graph's warmup, so its ~1200 kernels launch uncaptured, which adds another ~5 ms. A verify round takes ~15 ms. In an nsys trace of the B7k config at 13-30K, rounds that change shape took 33-35 ms against 16-18 ms for rounds that keep it. The MTP draft context pays the same way every round: its catch-up and draft graphs differ in rows. Graphs of up to 32 rows now live in up to 3 slots. Each slot holds a graph result and a scheduler of its own, because ggml_backend_sched_split_graph rewrites a graph's node sources to that scheduler's input copies (ggml-backend.cpp:1431). A graph stays computable only while the scheduler that split it holds it, so a second graph cannot share the context's scheduler. process_ubatch tries gf_res_prev, then the slots. On a miss it builds into a new slot while there is room, else over the least recently used one. Larger graphs (prompt batches) stay in gf_res_prev on the context's scheduler. Outputs are read through the scheduler that computed them (sched_res). The slots are dropped on reserve, and their graphs reset wherever gf_res_prev's are. A slot's compute buffer is sized by its own graph on first allocation. LLAMA_GRAPH_REUSE_DISABLE keeps the old path. A backend sampler with inputs (the dist sampler's uniforms, penalties, a logit bias) holds those of the graph built last, where build_sampling bound them, so while one is set no slot is used and only gf_res_prev is reused, as before. The draft context's top-k chain has no inputs and keeps its slots. Measured on the 27B hybrid, RTX 5080, the served flags (-cts f16, -ot token_embd=CUDA0, --kv-unified, the draft vocabulary, n_max 3, --attn-mask-bits) at served depth: the six deep requests of 94-121K prompt tokens, greedy, 512 tokens, against the same tree without this change (6510894): - MTP alone (the served config): 184.19 -> 185.38 tok/s, same text 6/6, same-text geomean +0.65% (se 0.17%), all six faster. The draft context alternated its 4-row catch-up and 1-row draft graphs and rebuilt both every round. - MTP with n-gram drafts (ngram-mod n-max 7, rollback 7): 177.68 -> 180.36 tok/s, same text 6/6, +1.55% (se 0.35%), all six faster; target rebuilds 83 -> 46 (the rest are n_kv crossing a 256-cell pad). - Noise floor, the same slots leg repeated: +0.05% (se 0.09%) MTP, +0.14% (se 0.21%) hybrid, texts 6/6. At 13-30K on the round-4 set (-cts q8_0), MTP alone +0.33% and the hybrid +2.88% (same text 24/24 each). The kernels are unchanged: every pair above produced identical text. The slots cost ~82 MiB of compute buffers. Tests: test-recurrent-state-rollback on 8 archs; test-llama-archs (455 rows, 0 fails); test-thread-safety, test-save-load-state, test-state-restore-fragmented, test-kv-restore-runs on stories15M with CUDA. test-backend-sampler on CPU: 22/22 on Qwen3-0.6B q4_0 (the model upstream CI runs it with), with and without this change. On stories15M, multi_output_dist_transaction (n_ubatch 2: each decode alternates a 2-row and a 1-row graph) failed a first version that kept slots with a dist sampler set, its uniforms written into the other slot's graph; logit_bias and penalties fail there with and without this change (a +10 bias under dist is not certain; 806 tokens tie at the top-p cut).
…nt cache itself, as it does q8_0 Under -cts f16, the served state, every recurrent layer paid two copies around the kernel that -cts q8_0 fuses: a GET_ROWS widening the sequence's cache row into an f32 temp for the kernel to read, and a CPY narrowing the f32 snapshots the kernel wrote into its output's tail back into the cache. Both matchers took f32 and q8_0 only (ggml_cuda_try_gdn_gather_skip, ggml_cuda_try_gdn_cache_fusion). On the 27B hybrid at 94-121K on a 5080 (nsys, the served flags), per layer and verify: the CPY 39.3 us and the GET_ROWS 10.1 us, against 13.4 us for the kernel itself. The kernel now converts as it loads and stores, with __half2float and __float2half: the conversions ggml_cuda_cast gives get_rows and cpy, so the cache holds the same bytes and the kernel sees the same values. The fused-cache and gather records carry the cache's ggml_type in place of a q8_0 flag, and the kernel's state template parameter is that type. f16, like q8_0, is fused for the scalar gate on the recurrent kernel only (the chunked prefill pipeline reads and writes f32, so its GET_ROWS and CPY stay); unlike q8_0 it takes any head width. GGML_CUDA_GDN_F16_CACHE_LEGACY=1 keeps both copies. Measured on the 27B hybrid, RTX 5080, the served flags (-cts f16, MTP n_max 3, the draft vocabulary) at served depth: the six deep requests of 94-121K prompt tokens, greedy, 512 tokens, one binary on 9ddd463, A-B-A with the off switch in the middle leg: 185.77 -> 202.09 / 202.07 tok/s, same text 6/6 on both pairs with identical draft counters (913 rounds), same-text geomean +8.81 % (se 0.30 %) and +8.80 % (se 0.32 %), every request +7.3 % to +9.9 %; A against A -0.01 %. A round takes 16.65 ms against 18.11. Tests: test-backend-ops on the 5080, GATED_DELTA_NET_CACHE_FUSION 58/58 (the f16 cases: the served layer at decode and verify over 1-4 sequences, activated gates, 16/64/128-wide heads, the served cache view, the fused gather from two rows, and the chunked prefill ubatch that keeps its copies) and GATED_DELTA_NET 63/63. Two mutations fail them: the f16 store scaled by 1.001 fails the 20 fused f16 cases, the f16 load scaled by 1.001 the 4 f16 gather cases, and the 36 f32 and q8_0 cases pass under both.
…e buffers A graph slot's scheduler sizes its buffers by the graphs it holds, so what the slots add to a context's memory grew with traffic and was known only from a peak: +182 MiB at the 5080 tier (4 x 294,912, -cts f16) and +198 at the 4-slot desktop tier (4 x 262,144, -cts q8_0), both on an RTX 5090 under four concurrent requests, where a single request had shown ~82. sched_reserve now measures, without allocating (ggml_backend_sched_reserve_size), the largest graph a slot takes: graph_slot_max_tokens rows over every sequence, with as many outputs as such a batch can have, on the full memory. It logs that size per buffer type beside the compute buffer's, for the target context and a draft context alike: graph_slots_max of them bound what the slots add, so a head can charge it from its load log. Measured on an RTX 5080 at load, the 27B hybrid with its MTP draft context (9760e96 + this change), CUDA0, per slot, target / draft: 4 x 294,912, -cts f16, n_max 3 n_outputs_max 16 67.81 / 11.92 MiB 3 slots: <= 239.2 MiB (peak +182 on 9ddd463) 4 x 294,912, -cts f16, n_max 4 n_outputs_max 20 79.81 / 11.92 MiB 3 slots: <= 275.2 MiB 4 x 262,144, -cts q8_0, n_max 3 n_outputs_max 16 67.69 / 11.79 MiB 3 slots: <= 238.4 MiB (peak +198 on 9ddd463) 2 x 262,144, -cts q8_0, n_max 3 n_outputs_max 8 37.22 / 9.90 MiB 3 slots: <= 141.4 MiB 1 x 262,144, -cts q8_0, n_max 3 n_outputs_max 4 21.99 / 8.95 MiB 3 slots: <= 92.8 MiB The target's bound is linear in what the graph holds: 3.00 MiB an output row (n_max 3 -> 4 at four sequences), 3.23 a sequence, 4 B a KV cell (the bit mask at 32 rows: +0.12 MiB for 32,768 more cells) over 5.76; the draft's 0.95 a sequence over 7.00, whatever n_max. With the log, the same six greedy requests at 4 x 294,912 give the text and the draft counters of 9760e96 without it, 6/6.
…y used evicted A context kept one CUDA graph executable per graph shape, and per graph slot for the same shape, each holding device memory beside the buffers the scheduler sizes, and nothing bounded them but the 10 s eviction: under four concurrent streams a 27B hybrid's target context on a 5080 held 14 (the largest 10 MiB to instantiate) and its draft context 12, 84 MiB in all, a count that follows the traffic. GGML_CUDA_GRAPH_MAX (default 8, 0 for no cap) bounds the executables a context keeps: before one is instantiated, the least recently used others are evicted until it fits. The device memory each instantiation takes is measured (cudaMemGetInfo around cudaGraphInstantiate), and the most held and the largest are logged as they grow, the context named; the cap is logged when the first CUDA backend is made. test-cuda-graph-cap: ten row counts of y = w x under a cap of 4, each computed three times, then the first again; every shape captured, 4 held at most, the evicted first captured anew, every output matching the CPU. It passes on a 5080 and a 5090; without the eviction all ten are held and it fails. The 27B hybrid, n_max 3: on the 5080 one request at a time, the same text at 216.3 tok/s uncapped and at 16, 214.6 at 8 and 214.8 at 2, the five requests after the first within 0.3 % at 8 (3,572 against 3,561 ms). On a 5090 at 4 x 294,912 under four concurrent requests 15,400 MiB at peak with CUDA graphs off, 15,430 at 8, 15,446 uncapped; at 8 x 786,432 under eight 25,916, 25,938 and 25,938.
…, in the state's own type test-recurrent-state-rollback holds a rolled-back recurrent state to a context that decoded only the kept rows within 1e-6 absolute, reading the serialized state as f32 words: exact where a batch's shape does not change its arithmetic (the CPU), meaningless for an f16 or q8_0 state (-cts), and too tight on a GPU, whose kernels differ by batch size: the same prefix reached through a 9-row batch rolled back to 6 and through one 6-row batch differ by rounding amplified through the layers. test-recurrent-rollback-rounding decodes the recurrent rows in their own type (ggml's to_float) and compares them in normalized squared error. A rolled-back state must be within 4x the rounding spread of its prefix (decoded one and three rows at a time, against one batch) of the kept prefix, and further than that from the prefixes one token shorter and one token longer, so a rollback a token early or late fails, and so does a bound too loose to tell them apart. Four rollbacks: 3 rows off a 9-row batch, and 1, 2 and 3 rows off the speculative verify's 4 rows after a 5-token prompt. A state it cannot read fails it; DeepSeek V4, whose memory writes its own layout, is not registered. On the CI's tiny qwen35 and nemotron_h, on the CPU and on a 5080 as ctest runs them: nmse at most 1e-16 against the kept prefix, a token off 0.013-1.98. The 27B hybrid on the 5080 with f32, f16 and q8_0 states: at most 1.68e-4 against bounds of 6.7e-4-8.7e-4, a token off 0.41-0.66. A rollback that restores a token early or late (set_rs_idx at rollback + 1 and - 1) fails every case, CPU and GPU.
llm_graph_input_kpool had no can_reuse override, so the base returned false and poisoned the whole llama-graph reuse check on every decode step: fresh cgraph per step, fresh CUDA capture per step, zero replays. Serving evidence (rental 52836549, b9e1b78, GRAPH_MAX=8, -lv 5, LLAMA_GRAPH_INPUT_DEBUG/RESULT_DEBUG=2, 150-token fixed probe): - 479/479 reuse verdicts false; per-input 561 false / 1656 true - positional split 219x11101 + 82x11001 (5-input) + 178x1110 (4-input): index 3 false in all 479 checks = kpool in both build orders (glm5next.cpp:577-591, :710-754); index 2 intermittent = mem_hybrid_k - CUDA Graph id reused: 0; update failed: 0; disabling-on-node-type: 0 - 19 captures to cap 8 on both contexts, never replayed The override mirrors can_reuse_kq_mask: all kpool tensor shapes are deterministic functions of n_kv (padded to >=256, stable across steps), n_tokens/n_stream/n_tps/n_ps, and the kpool-dirty flag; scoring flips (n_ctx gate) refuse. Pure shape predicate is CPU-unit-tested (tests/test-kpool-can-reuse.cpp): identical/rebuild/non-scoring reuse, 8 mutation refusals; impl mutation probe fails 3 cases; neighbor test-kv-mask green (8560 masks, 0 failures).
… logged in Release current_dims counted the pools with its own copy of the builder's formula; it now calls llama_kpool_n_pools, so the check cannot drift from the shapes build_inp_kpool makes and turn graph reuse off again unnoticed. The two lines that say why a CUDA graph is not replayed (a MUL_MAT_ID that needs a stream sync, an executable update that failed) were compiled out under NDEBUG, so a Release server printed neither and their absence read as evidence. They are GGML_LOG_DEBUG, printed only under -lv, and the first now names the node, its op, src0's type and its batch.
Diagnostic only, no behavior change. The per-input verdict printed the literal string 'placeholder', so a 479-false serving leg could only be mapped to inputs by build order. It now prints the input index, and mem_hybrid_k prints k_idxs/kq_mask/rs plus stored-vs-current head/rs_z. Rides the held prebuilt; the verdict leg reads the names off the -lv5 log to decide the mem_hybrid structural fix.
LLAMA_GRAPH_INPUT_DEBUG=2 now prints, when can_reuse_rs refuses, which of its six checks did (s_copy, main, extra, the snapshot write rows, head, rs_z) with both sides of each value baked into the graph. Diagnostic only; every check is still evaluated. bonsai-2-27b drafted on a 5090 at bf76114 refused its hybrid memory input on 322 of 1,278 reuse checks, and a refused reuse costs ~0.46 ms a graph: graph reuse is worth 289.8 against 216.6 tok/s there.
rs_z enters a graph only as the view an f32 state's in-graph zeroing scales (build_rs, build_rs_cache_view). Any other state type is zeroed on the host before the graph runs (llama_memory_recurrent::zero_rs_z), and a context whose graphs hold no recurrent state (the GLM-5.3 NextN draft context, which filters every KDA layer out) never reads it. can_reuse_rs compared it regardless, so a draft context whose one cell seq_rm had emptied (src -1, rs_z 0 at the next slot search) refused every graph built at rs_z -1: 27 of 268 reuse checks on GLM-5.3 at bf76114, all rs_z -1 against 0 with head steady. The input now records whether build_rs baked rs_z and compares it only then; head and the copy sizes are compared as before. Also, LLAMA_GRAPH_INPUT_DEBUG=2 names which of llm_graph_input_attn_kv's checks refused (k_idxs, the KQ mask's shape against n_kv): bonsai-2-27b's MTP draft graph refuses its attention input on 322 of 1,278 checks.
…ght Systems
ggml_cuda_init logs each device's profile beside its name: SMs and clock, registers and shared memory per SM,
threads and blocks per SM, L2 and its persisting maximum, and the DRAM peak its memory clock and bus give
(the 5070 Ti: 70 SMs, 896 GB/s; the 5080: 84, 960). It is what a roofline or a launch shape is sized from.
GGML_CUDA_NVTX=1 puts an NVTX range around every graph evaluation ("graph") and every node that launches work,
named with the op, the node, the bytes it reads and writes and a matmul's weights; a launch that fused the
nodes after its own marks each of them inside its range (nvtx.cuh holds the format). An nsys capture with
-t cuda,nvtx then attributes each kernel to its node and its bytes. A Hadamard-hinted matmul is named FWHT and
reads its activation only, as ggml_cuda_op_fwht computes it; a CPY reads its source; a MUL_MAT_ID reads the
experts its tokens can select. Off, each node costs one test of a static flag.
Measured on an RTX 5070 Ti, Ternary Bonsai 2 27B, llama-bench tg64 with CUDA graphs off: 83.8 and 84.9 tok/s
with the ranges off, 84.3 and 79.6 +- 3.1 with them on and no profiler attached (the host's load ~17).
A parameter axis like -ctk and -ctv, the context's type_s (default f32, llama_context_default_params'), and a field in every printer. A bench of a Gated DeltaNet model now runs the state the server serves: Ternary Bonsai 2 27B serves an f16 state, and a bench at f32 moved twice its state bytes (48 layers x 48 x 128 x 128).
…zed from the card's profile Between two PQ2_0 matmuls a decode token runs a chain of small kernels (the norm and FWHT, the q8_1 quantize, the Gated DeltaNet layer's conv, recurrence and gated norm), 1-3 us each, and PDL lets the next matmul request its own weights only under the one kernel before it. On an RTX 5080 at depth 16,384 the chain is ~1.3 ms of a 10 ms token against floors of ~50 us (rig's roofline probe, NVTX node ranges of 1c83edd). The blocks of a PQ2_0 launch now take every gridDim.x-th tile from their index instead of a contiguous run, so what a launch reads first is its matrices' heads; each block of the launch before it, once all its boxes have landed, prefetches its share of those heads into L2 (cp.async.bulk.prefetch.L2), which DRAM streams while the chain runs. The graph evaluation plans it per node (ggml_cuda_pq2_prefetch_plan): a launch is the members that read one src1, a group or a gated pair its GLU reads (whose two heads it prefetches); the graph's last launch prefetches its first. The size is GGML_CUDA_PQ2_PREFETCH_US of the card's DRAM time (default 2, at most a quarter of its L2), from the device profile's DRAM peak and L2 size, now kept in the device info. A hint: which block computes a tile never changes its arithmetic, and a stale plan costs only bandwidth. Ternary Bonsai 2 27B, llama-bench tg64, served state (q4_0 K/V, f16 recurrent state), graphs on, per-rep medians of 20 reps a cell, the order rotated each round after a discarded warm-up, against the parent's libraries: RTX 5080: +2.4 % at depth 16,384 (100.54 -> 102.97 tok/s), +2.7 % at 0 (103.75 -> 106.55) RTX 5070 Ti: +1.6 % (78.98 -> 80.21), +1.1 % (79.83 -> 80.71) 1 us: +1.2 to +1.7 %; 4 us: +0.6 to +2.8 %; the interleaved tiles alone: 0.0 to +0.7 % Requested with the producer's last box instead, the prefetch took DRAM from the launch's own tail (+4 us a gate + up on the 5070 Ti, -2.8 % at 8 us on the 5080), so it waits for the block's last box. Greedy 128 tokens: byte-identical to the parent at 2 and 8 us on the 5080 and at the default on the 5070 Ti. test-backend-ops: the 296 PQ2_0 cases (MUL_MAT, MUL_MAT_VEC_FUSION among them) and the 24 MUL_MAT_GROUP cases pass with the prefetch off, at 8 us and at 64 us (L2-capped). Registers per instance: +0 to +7 over the parent's; the most any instance takes is still 168.
llama-bench's own type-name table had no f32, so `-cts f32` exited on its arguments though f32 is the state's default and a type the engine serves, and `-ctk f32 -ctv f32` did the same for a K/V pair the CUDA flash attention runs. A bench that passes the formats a tier serves, whatever they are, needs every one of them to parse. Measured on the RTX 5080: `-cts f32,f16` runs both rows (tg16 99.26 and 99.20), `-ctk f32 -ctv f32` runs, `-cts f64` still exits with the invalid-parameter error.
marcospaulo
changed the title
…d reuse, q8_0 read raw, the PQ2_0 L2 prefetch) Thirty rows for the 48 commits past main, each change with its measurement and its off switch as the commits and the rig pin records give them. The line's gates, each on a rented RTX 5090 against the pin before it, bit-exact plain and drafted and the swapped-out conversation token for token: e11 through e19. The last, 4104c47 against 02512a3: plain decode on the 245K question 90.9 -> 92.8 tok/s; llama-bench tg64 as served with the prefetch against 02512a3's CUDA library +4.84 % at 16,384 and +4.51 % at 0.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
The train engine-10, 48 commits past main (c1518d4). rig pins its tip, 4104c47.
build_rszeroed a state row withGGML_OP_SCALE, which CUDA and the CPU run on f32 only, so-cts f16aborted on its first decode. Over 8,192 tokens decoded three a call after 131,072 on Ternary Bonsai 2 27B, the served q4_0 K/V with a q8_0 state read 0.002488 mean KLD and with an f16 state 0.001724; an f16 state is f32's within noise.{}below a JSON schema's root accepts any value, so a free-form tool argument parses under the lazy tool grammar, GBNF and llguidance (737eba9, c206d74).GGML_CUDA_FATTN_MASK_PREFIX_LEGACY), the attention rotations set once (b0dd41c,LLAMA_ATTN_ROT_INPUT_LEGACY) and the backend top-k prefilter (6c0e372,LLAMA_TOP_K_PREFILTER_LEGACY). Every*_LEGACYswitch reads 0/false/no/off as off (4384bbb, 198e77e).--spec-type draft-mtpare refused, not dropped.Gate on 48ebd21's prebuilt (RTX 5090, the 245,752-token prompt's four questions greedy with top-5 log-probabilities, against engine-c1518d4):
Since 48ebd21 (21 commits). rig pinned each step. Every pin's prebuilt was gated on a rented RTX 5090 against the pin before it, on the same four 245K questions. Every gate passed: bit-exact plain and drafted, the same without
n_probs, and a 245K conversation swapped out and back answering token for token.rs_zonly when a graph baked it (02512a3). Refusals are named underLLAMA_GRAPH_INPUT_DEBUG=2(bf76114, cc9e415), and CUDA-graph refusals are logged in Release (7a17ff5). On 2 × RTX PRO 6000, GLM-5.3 IQ3_XXS decodes at 101.5 and 108.5 tok/s, against 87 before.--spec-rollbacksets how long a draft a recurrent target rolls back in place (17a4374).seq_rmrefuses a rollback past the snapshots the last batch wrote (6510894), tested at its batch shapes' rounding (b9e1b78).GGML_CUDA_GRAPH_MAX).-cts(b2eb433) and takesf32for-ctk/-ctv/-cts(4104c47).Gate on 4104c47's prebuilt (RTX 5090, against engine-02512a3, engine-4104c47): bit-exact plain and drafted (1,536 tokens, log-probabilities, top-5 and draft counts), the same without
n_probs, and a 245K conversation swapped out and back token for token (first token 1.01 s). Plain decode on the 245K question: 90.9 → 92.8 tok/s.e2e-driver-only.shon a 5070 Ti: PASS.