From 48c2f751b4cf643d2bc1ca12a7f4a3b155e7d120 Mon Sep 17 00:00:00 2001 From: Georg Bauer Date: Fri, 11 Sep 2026 17:18:45 +0200 Subject: [PATCH] Align DeepSeek and GLM execution with DS4 --- build.rs | 1 + docs/deepseek-reference-followup-20260911.md | 519 +++++++++ docs/glm-scheduling-followup-20260911.md | 585 ++++++++++ ...ference-reference-parity-audit-20260905.md | 123 ++ native/metal/ds4_canary.m | 66 ++ native/metal/ds4_gpu.h | 3 + native/metal/ds4_metal.m | 42 - src/engine.rs | 154 ++- src/engine/ds4_sampling.rs | 345 ++++++ src/engine/gguf.rs | 9 +- src/engine/metal.rs | 1030 ++++++++++++----- src/engine/metal/checkpoint.rs | 4 +- src/engine/metal/glm.rs | 478 +++++++- src/engine/metal/gpu.rs | 3 + src/engine/metal/markov.rs | 327 ++++++ src/engine/metal/qwen_mtplx/request.rs | 6 +- src/engine/metal/qwen_mtplx/session_cache.rs | 4 +- src/engine/metal/qwen_mtplx_tests.rs | 3 +- src/engine/tokenizer.rs | 102 ++ src/main.rs | 8 + src/metrics.rs | 25 +- src/model_eval.rs | 170 ++- tests/fixtures/ds4-sampling-ec7642c.json | 1 + tools/build-ds4-session-reference.sh | 31 + tools/ds4-session-reference.rs | 775 +++++++++++++ tools/sampler-replay-benchmark.rs | 39 + tools/test-supervisor.rs | 64 +- 27 files changed, 4517 insertions(+), 400 deletions(-) create mode 100644 docs/deepseek-reference-followup-20260911.md create mode 100644 docs/glm-scheduling-followup-20260911.md create mode 100644 native/metal/ds4_canary.m create mode 100644 src/engine/ds4_sampling.rs create mode 100644 src/engine/metal/markov.rs create mode 100644 tests/fixtures/ds4-sampling-ec7642c.json create mode 100644 tools/build-ds4-session-reference.sh create mode 100644 tools/ds4-session-reference.rs create mode 100644 tools/sampler-replay-benchmark.rs diff --git a/build.rs b/build.rs index 17c7949..4bc679b 100644 --- a/build.rs +++ b/build.rs @@ -10,6 +10,7 @@ fn main() { cc::Build::new() .include("native/metal") .file(metal) + .file("native/metal/ds4_canary.m") .flag("-fobjc-arc") .flag("-ffast-math") .flag("-mcpu=native") diff --git a/docs/deepseek-reference-followup-20260911.md b/docs/deepseek-reference-followup-20260911.md new file mode 100644 index 0000000..ce382a4 --- /dev/null +++ b/docs/deepseek-reference-followup-20260911.md @@ -0,0 +1,519 @@ +# DeepSeek standalone-reference follow-up — 2026-09-11 + +The original antirez/ds4 process is used only under the user's explicit +reference-benchmark authorization. No original C engine is linked into DS4Server. +Evidence: `local-eval-results/deepseek-paired-20260911.nyUCsL/`. + +**Measurement qualification (user clarification, 2026-09-11):** The user was +watching videos in parallel, using part of the GPU. The exact overlap with each +earlier run is not recorded. The current sequential comparisons are therefore +not controlled idle-device performance acceptance or causal before/after proof. +Their outputs and functional checks remain useful; timing receipts are retained, +not discarded. Do not attribute every gap to video playback or claim thermal +throttling from this information. Further throughput acceptance runs are deferred +until competing GPU activity can be controlled; code work and functional tests +can continue. This does not waive the2% requirement. + +## Initial pair is not performance-parity evidence + +Both AR processes finished the full Summary README → lighthouse Story → Python +`is_prime` conversation after a separate32-token OK warmup. Installed0731 GGUF +only, power100, Low, temperature0.6/top-p0.95/min-p0/top-k0/seed42, context32768, +quality/SSD/canary/DSpark off. Startup30s/progress45s,112GiB process-memory guards, +no total timeout. No builds or other model tests overlapped the GPU runs. + +Two mismatches invalidate this initial pair (`ar-{reference,rust}.*`): + +- The standalone driver requested2048 prefill rows, whereas the UI/harness + requested automatic (0), which the DeepSeek executor resolves to4096 at this + context. Original DS4's automatic setting also resolves to4096. The reference + driver now requests0 and records the public session prefill-cap value instead + of silently assuming the same chunk for all models. GLM's automatic request + remains unchanged. No product chunk reduction or special harness path. +- The same first prompt contains2741 tokens in the original but2742 in the + production runtime. The direct Rust tokenizer and all three original + continuation fixtures agree exactly; the mismatch occurs later, in the shared + cached-prompt renderer. A valid system-bootstrap tag with empty chat history + took the continuation branch and prepended an EOS before the first user turn. + Both UI and headless generation use this function. + +## Root-cause correction and regression + +`render_text_prompt` now requires nonempty history for the cached continuation +branch. A first user turn uses canonical full rendering, while `align_prompt` +still reuses its prepared system prefix. Real ongoing histories retain the +existing append-only behavior. No artificial wait, power change, kernel change +or KV-persistence rewrite was introduced. + +The CPU-only `ds4_chat_matches_original_session_tokens` regression loads only +GGUF metadata/tokenizer and original recorded token IDs. It first checks the +three-turn tokenizer contract, then invokes the same shared renderer with the +actual system-bootstrap frontier and tag. Before the fix it fails with: + +```text +shared bootstrap renderer differs: 2742/2741 tokens; +starts [0, 1, 128803, 45585]/[0, 128803, 45585, 260] +``` + +`token-before.*` is the passing standalone-tokenizer diagnostic; +`bootstrap-before.*` is the deliberately failing shared-runtime regression. +Neither file is overwritten or treated as a successful full parity run. +The test also accepts the saved GLM fixture to cover the other shared consumer. + +Reproduction after building release tests: + +```sh +DS4SERVER_CHAT_REFERENCE=local-eval-results/deepseek-paired-20260911.nyUCsL/chat-token-fixture.json \ + target/release/test-supervisor 2048 30 30 \ + target/release/deps/ds4_server-6141856e5c0fa6df \ + ds4_chat_matches_original_session_tokens +``` + +The fixed renderer passes both saved DeepSeek and GLM three-turn fixtures +(`{chat,glm-chat}-token-fixture-fixed.*`,0.04/0.05s). Release all-target/all-feature +build, warnings-denied Clippy, rustfmt/diff checks,17 enabled sampling tests and +two enabled tokenizer tests pass. The model-dependent tests remain explicitly +ignored by default and were invoked with the recorded local fixtures above. +CLI SHA256: `9fb13c7f9049c3b5059b4fe79b68bf997dbe4e1bfdc4c631c130595cab8c359a`. +Standalone driver SHA256: +`ad7a4b94a5918505c58130d336d026e2c5b9640aab717ef185808bebd7536a25`. + +## Corrected complete AR pair + +`fixed-ar-{reference,rust}.*` and `fixed-ar-comparison.json` contain the new +original-then-Rust pair. The original session confirms prefill-cap4096 and +engine power100. Every text/thinking/token/cache comparison passes; all turns +end naturally. Counts are826/1444/199, prompts2741/3587/5057 and cached1/3567/5031. + +| Turn | Rust / original engine-prefill ms | Rust / original decode t/s | +| --- | ---: | ---: | +| Summary | 6561.316 / 5425.137 | 36.916 / 37.813 | +| Story | 286.826 / 290.380 | 32.712 / 34.246 | +| Python | 334.486 / 369.018 | 32.020 / 32.734 | + +This is functional agreement for the complete AR workload, **not** performance +acceptance: all three decode ratios are below98% in this single pair, and the +Summary prefill is slower. Different natural outputs in the invalid initial +pair must not be used as before/after throughput evidence. No GUI/canary run was +performed in this clean series. + +## DSpark pair: functional comparison fails + +Both full processes finished naturally and both watchdogs exited successfully. +The final comparison exits1 deliberately (`dspark-comparison.json`); this is +not an inference crash. Rust confirms opportunistic sampling, confidence0.6, +strict/exact off; the original loads the installed three-stage/block5 support +model and enables direct verifier commits. + +The first prompt/cache counts agree (2741/1), but its output already diverges: +647 Rust tokens versus1208 original tokens. Thus later speed ratios compare +different histories and **cannot establish matched-work performance parity**. +Counts in Story/Python are1675/195 versus1455/347. The raw rates are retained +in the comparison file, not promoted to valid performance results. + +An independent frontier defect is visible even within Rust's own history: +after2741 prompt +647 output tokens, the next cache is3389 instead of3388; +after3409 prompt +1675 output tokens, it is5085 instead of5084. A speculatively +returned EOS remains committed. The shared consumer currently rewinds only +GLM, while original `ds4_session_rewind` also truncates the DeepSeek logical +frontier and invalidates DSpark capture. This is separate from the now-fixed +first-prompt bootstrap EOS and from the first-answer decode divergence. + +### Next cohesive DSpark work unit + +1. Locate the first divergent cycle with identical recorded target history: + compare proposal IDs/confidence, GPU row argmax, accepted prefixes and + compression/DSpark capture state against original DS4. Cover full acceptance, + partial acceptance and rejection before declaring the full implementation + equivalent. Existing target-owned self-tests alone are not an oracle proof. +2. Match the consumer's EOS frontier and capture invalidation using DS4's + DeepSeek contract, including an accepted-EOS transition into the next user + turn. Do not transplant GLM's KDA rollback or merely hide the extra count. +3. Remove verified extra verifier work as part of that same comparison: + Rust `eval_batch_inner(collect_tops)` reads/clones every row's full vocabulary + and computes argmax on the CPU. Original opportunistic verification passes + `row_logits=NULL`, obtains GPU row tops, then reads only the committed row + (`ds4.c:69112–69124,69153–69159,69210–69216`). Preserve the separate exact- + sampling requirements and verify row/stop decisions before timing the change. + Also verify per-turn counter scope across warmup/resident switching before + using cycle/acceptance ratios as exclusive stage measurements. + +The four reviewed AR/DSpark Python answers (both implementations) each pass +their five generated assertions and5011 independent cases in isolated Python +with restricted builtins (`python-quality.jsonl`). No model-generated tools +were invoked. This quality check does not waive DSpark's token/state mismatch. + +The six-cell performance goal, repeated timing acceptance and actual UI +responsiveness acceptance remain open. The user's working bundle is unchanged. +All model/reference/canary processes from this follow-up have terminated. + +## DSpark cycle/state audit (second work unit) + +Evidence: `local-eval-results/ds4-cycle-audit-20260911.vfsUc8/`. The original +source and installed artifacts remain the same. `DS4_SPEC_CYCLE_TRACE=1` is an +opt-in diagnostic in the shared production consumer and standalone reference +driver; it records the sampled first token, returned IDs, generated count and +committed position. It is off by default. Original `DS4_DSPARK_SPEC_LOG=1` +adds draft lengths, full/partial acceptance and scheduler decisions without +changing the graph. These traced runs are **not clean throughput acceptance**. + +The initial complete traced pair reproduces the failure. In the first Summary +cycle, both targets sample671; Rust proposes12275 and emits only671, while DS4 +proposes3967 and emits `[671,3967]`. The original warmup first differs in the +fifth returned cycle, after earlier partial accepts. This is not evidence of +a context-size or power-setting cause. + +### Consolidated findings and implementation tasks + +This table scopes the inspected path to the installed0731 DSpark support, +resident M5 Max execution and default opportunistic sampling. It does not claim +completion for arbitrary SSD/strict/exact modes or the six-cell performance goal. +Original line references below are for pinned `ds4.c` at `ec7642c`. + +| Area | Original behavior / Rust discrepancy | Current action | +| --- | --- | --- | +| Initial support KV | `32627–32821`: target HC expansion, per-stage HC mixing and attention normalization precede KV projection. Rust projected stage-0 normalized hidden directly to KV. | Implemented the missing operations using existing Metal calls and batch scratch. First Summary proposal now matches. | +| Single accepted draft | `36205` and `69112`: even one draft uses the target batch verifier. Rust used scalar decode, with different arithmetic and no captured suffix for the next support cache. | Removed the resident single-draft scalar branch. First two Summary cycles then match. | +| Verifier capture | `27951–28053`: capture contains the preceding seed row plus all verified target rows. Rust retained only the suffix. | Added seed-row capture with the original slot stride and sufficient workspace. | +| Cache timing and gaps | `66665–66732`: seed from the captured batch only when its end equals the next proposal position; otherwise crop/reset the absolute cache window. Ring maintenance does not create or bridge a window (`33329`, `27712`). Rust eagerly seeded each batch, extended windows after failed proposals and tracked only modulo positions. | Added deferred seed ownership and absolute window start; no false append after rejected proposals. Replaced batch workspaces are released before allocation of the next batch. | +| Partial acceptance | `69221–69240`: ordinary partial commits invalidate row and batch captures, preserve compressor-prefix state, and do not claim support KV rows. Rust retained the captures and extended support counters; its old seed scratch also overwrote captured hidden rows. | Invalidate the deferred capture, preserve the original cache frontier, and stop overwriting hidden capture storage. Three and then32 recorded Summary cycles pass. | +| Scheduler | `53500`, `65215`: reset per request. `53666`: a window pause replaces, rather than maximizes with, the cold no-draft pause. | Reset in shared prompt alignment; regression covers7→4 and clearing prior-request counters. | +| EOS consumer | `73690`: logical DeepSeek rewind plus capture invalidation, not GLM KDA rollback. | Shared consumer now invokes the model-specific DeepSeek rewind; invalid forward rewind is rejected before mutation. Full next-turn verification still required. | +| Q8 confidence arithmetic | `34350` calls `matvec_any`, which dispatches type8 to activation-quantized Q8 matvec (`8144–8170`). The installed confidence tensor is Q8_0 `[4352,1]`; Rust instead used the dequantized-weight × F32-input reference calculation. | Quantize activations for the confidence dot too, and match the two alternating four-lane ARM SDOT/FMA accumulators (`7510`) in both confidence and Markov scoring. CPU regression distinguishes the quantized result from the old F32-input dot. | +| Confidence/output head | `66780–66835`: check first confidence before doing the vocabulary projection; read later rows only as required. Rust calculated all base logits and read all rows before the check. | Implemented the early gate and per-needed-row readback. Confidence-disabled execution keeps hidden/head submission fused. Full1755-cycle comparison remains exact; clean timing acceptance is separate. | +| Verifier output | `36325–36431`, `69153`, `69210`: fuse head into the layer command sequence, obtain GPU row tops, read only the committed logits row. Rust drained before the head and cloned all vocabulary rows for CPU argmax. | Implemented fused submission, existing GPU argmax/top-k reductions and committed-row readback for the ordinary verifier. Exact sampling retains all target distributions. Full1755-cycle default-mode comparison remains exact. | +| CPU Markov workers | `33889–33963`: persistent helper pool and fused Q8 argmax; `1912` defaults to min(online CPUs,12), including the caller. Rust created scoped OS threads on every dense argmax, using all18 logical CPUs here. | Implemented persistent Rust workers with the existing Q8 arithmetic, identical contiguous row partition and ordered first-tie reduction. The caller executes slot0; default12 total threads and positive integer `DS4_THREADS` overrides up to32. All1755 original cycles remain exact; timing acceptance remains separate. | +| Verifier allocations | Original graph retains verifier scratch and GPU frontier/prefix buffers; Rust constructed `BatchScratch` and snapshot buffers for every verifier. | Reuse baseline and high-water prefix buffers, plus one verifier batch per existing padded row shape. Scratch is recycled only after delayed seeding or capture invalidation. Extended32-cycle oracle regression proves native-buffer identity reuse and byte-exact rollback; full-chat and timing evidence below. | +| Acceptance telemetry | Rust's executor-wide cycle counts and resident support counters have different ownership scopes. | **Open:** reconcile request/session counters before deriving acceptance ratios or exclusive stage percentages. | + +Excluded after checking the guards: seed-batch fusion defaults to the ROCm +gfx1151 path, not this Mac (`53412`); the Markov GPU branch is under +`#ifndef __APPLE__` (`34266`); adaptive extra decode splits are pre-M5-only. +The output-head padding to8 rows is present in DS4's helper (`26086`) and is +not an unmatched Rust optimization. None of these were blindly enabled/removed. + +### Runnable regression evidence + +`dspark_matches_original_summary_cycles` loads only the installed GGUFs and +the recorded original JSONL, prepares the actual bootstrap boundary, samples +with the shared DS4 sampler and verifies each cycle's IDs and position. It is +ignored by default and explicitly supervised when run: + +```sh +DS4SERVER_DSPARK_REFERENCE=local-eval-results/ds4-cycle-audit-20260911.vfsUc8/before-reference.jsonl \ + target/release/test-supervisor 114688 30 45 \ + target/release/deps/ds4_server-6141856e5c0fa6df \ + dspark_matches_original_summary_cycles +``` + +- `cycle-test-before.*`: fails at generated2, Rust `[10059]` versus original + `[10059,260,13672,294,270]`. +- `cycle-test-batch.*`: after single-row batch verification, fails at generated7, + Rust `[4496,3051,943,30941,22]` versus original `[4496,3051,943,30941]`. +- `cycle-test-capture.*`: all three unchanged expected cycles pass (7.28s). +- `cycle-test-32.*`:32 original cycles, prior-request scheduler reset and + safe logical rewind pass (9.32s). These are functional diagnostics, not + stories truncated for a throughput comparison. + +The original full pair, intermediate cache-only full run and failing receipts +are retained. No original C implementation was added to the application; no +model download, bundle replacement, commit or push was performed. + +### Confidence follow-up from the full chat + +The subsequent `fixed-rust.*` full run matches451 complete returned cycles +(including warmup), then differs at Summary generated1018. The first divergent +decision is one cycle earlier: at position3759, original DS4 proposes `[588,699]` +and rejects the first draft, while Rust's confidence logit0.40281284 falls below +the0.6 sigmoid threshold and suppresses drafting entirely. The different +no-draft scheduler decision then changes the returned cycles. All turns still +end naturally, with counts1194/1424/296 and **no extra cached EOS**; this is +progress, not a matched-output performance pair (`fixed-cycle-comparison.json`). + +Read-only inspection of the installed support GGUF confirms the confidence +head type8/Q8_0,4352 inputs. Correcting the activation quantization yields +confidence0.4063788 at that exact position and the original two draft IDs. +The stable Rust SDOT intrinsic is unavailable in this toolchain; the two SDOT +instructions are isolated in guarded Rust inline assembly, with stdlib NEON +FMA/reduction and the existing scalar fallback. No C host code or dependency +was added. Six focused DSpark CPU/layout tests pass, including the new +quantized-confidence test. Release all-target/all-feature build, Clippy with +warnings denied, rustfmt and diff checks pass after the correction. + +Checkpoint-load invalidation now clears the new deferred capture and absolute +cache metadata through the same `reset_cache` helper as session reset. This +does not change the checkpoint format or disk KV policy. + +### Completed functional comparison, performance still open + +`quantized-rust.*` completes all three turns to natural EOS. All1755 returned +cycles match the original recording exactly, including warmup, IDs and committed +positions (`quantized-cycle-comparison.json`). All three text, thinking, +completion, prompt and cached-token comparisons pass +(`quantized-result-comparison.json`): + +| Turn | Completion tokens | Prompt / cached | Rust / original prefill ms | Rust / original decode t/s | Decode gap | +| --- | ---: | ---: | ---: | ---: | ---: | +| Summary | 1208 | 2741 /1 | 5562.979 /5519.685 | 36.128 /36.580 | −1.24% | +| Story | 1455 | 3969 /3949 | 261.651 /281.080 | 30.366 /31.321 | −3.05% | +| Python | 347 | 5450 /5424 | 323.474 /344.868 | 33.936 /34.777 | −2.42% | + +These are **diagnostic** timings with cycle logging, not repeat-median clean +performance acceptance. Story and Python still miss even the single-pair2% +threshold; the missing early-confidence gate, verifier submission/readback, +allocation and worker-lifetime tasks above remain explicit work. No UI canary +or real GUI event-loop acceptance is claimed for this series. + +The Story has a coherent title, narrative and ending, not a clarification or +broken output. The Python answer is byte-identical to the previous original +DSpark answer (`python-reference-unchanged.json`), whose five generated asserts +and5011 independent cases passed in `deepseek-paired-20260911.nyUCsL/python-quality.jsonl`. +Six DSpark CPU/layout tests,17 shared sampling tests, the separate heap-fallback +test and two tokenizer tests pass; model-dependent tests are not silently +counted as run. The explicit32-cycle live regression passed before the final +confidence correction, and the final full1755-cycle harness/oracle comparison +validates the production path after it. + +Final CLI SHA256: +`2d15458d12a14457519a22232061ce76182123b24f23d1ae595df424c36963e5`. +Traced standalone reference SHA256: +`395cc58f8ea59afa0a6e22f8d46b3fb3f129bc367b8ee3ee1a0dd4bc0f01bd4d`. +The working app bundle remains +`ea4d555c2faf0940d9cbcf76d8638ca614a9cb2c6b034e3b2f80aeef86b0b339`. + +### Early confidence and lazy verifier output + +The shared Rust executor now follows the original early confidence gate before +the draft vocabulary projection. Hidden and logits rows are read only when the +proposal loop needs them; disabling confidence keeps hidden/head encoding fused. +The ordinary verifier keeps its layer commands alive through the output head +and the existing GPU argmax/top-k reduction. It reads only the committed +distribution; exact stochastic sampling retains all row distributions. +No prefill-cap, power, artificial pause, kernel math or KV persistence change. + +`lazy-head-rust.*`, `lazy-head-cycle-comparison.json` and +`lazy-head-result-comparison.json` establish the unchanged complete1755-cycle +recording and all three answers/thinking/token/cache/stop boundaries. Seven +focused CPU/layout tests pass (two installed-model tests remain ignored by +default), including the committed-row boundary regression. Release all-target, +all-feature build and warnings-denied Clippy pass. This full default-mode +comparison is not a new exact-stochastic oracle claim. + +CLI SHA256: `de10b537c2cf2bfbb680ebc350e810305def3ca2ba59f0bea89f1d8c980699cb`. +The original driver and working app bundle hashes above are unchanged. +`run-clean-pairs.sh` records serial three-turn timing pairs without cycle logging, +original proposal logging or canary, with alternating reference-first/Rust-first +order. Both workers retain startup/progress/memory supervision, no total timeout. +The first postprocessing command used the wrong reference event name (`result` +instead of `reference_result`) and failed after both workers had finished; the +comparison was corrected against the unchanged receipts, with no inference rerun. + +All six processes completed successfully, with natural EOS for all nine Rust +answers and matching reference content, thinking, prompt, cache and completion +counts (1208/1455/347 in every pair). `clean-{1,2,3}-comparison.json` and +`clean-pairs-summary.json` retain all results; none were discarded: + +| Pair / order | Turn | Rust / DS4 prefill ms | Rust / DS4 decode t/s | Decode delta | +| --- | --- | ---: | ---: | ---: | +| 1 DS4→Rust | Summary | 6396.150 /5435.701 | 34.878 /37.298 | −6.49% | +| 1 | Story | 262.284 /289.008 | 28.623 /31.716 | −9.75% | +| 1 | Python | 350.909 /332.242 | 30.372 /35.511 | −14.47% | +| 2 Rust→DS4 | Summary | 7905.861 /7910.498 | 29.783 /28.784 | +3.47% | +| 2 | Story | 307.278 /376.104 | 25.918 /24.574 | +5.47% | +| 2 | Python | 363.522 /488.866 | 29.186 /27.959 | +4.39% | +| 3 DS4→Rust | Summary | 8389.686 /8292.276 | 27.523 /27.801 | −1.00% | +| 3 | Story | 301.627 /383.962 | 24.399 /24.151 | +1.02% | +| 3 | Python | 379.914 /458.934 | 27.995 /27.801 | +0.70% | + +Pair3 is within2% in decode, with Summary prefill1.17% slower and the two short +continuation prefills faster. **It is not sufficient acceptance for this series:** +DS4 itself slows from37.30 to27.80 Summary t/s (about25.5%), and pair order changes +the sign of the relative gap. Median aggregation cannot establish a causal code +speedup in this nonstationary series. Do not compare the clean first pair against +the earlier logged pair as a before/after regression claim. Remaining worker and +allocation discrepancies still need their reference-aligned implementation and +an appropriately controlled follow-up. + +Read-only system checks during pair2 reported no recorded thermal/performance +warning and25% system memory free; our benchmark processes ran strictly serially. +These observations do not prove stable GPU clocks or exclude other system GPU +activity. No application was stopped or configuration changed in response. +Rust supervisor totals were107.691/121.168/129.178s, maximum observed progress +gaps6.396/7.905/8.389s in prefill, supervisor lag60/59/59ms and exit0 throughout. +Canary was off: these are not GPU-canary or actual GUI event-loop measurements. +Original supervisor totals were96.236/124.665/127.935s, all with `error:null`. + +The existing ignored `flash_0731_runs_exact_sampled_dspark` initially failed its +draft-count assertion: its four-token allowance takes the scheduler's existing +`max_tokens <10` no-draft branch. The test now allows16 tokens and explicitly +disables the confidence gate to exercise drafting independently of this short +fixture. It retains the original assertions and additionally invokes two-row +verification, checking both complete finite distributions, the GPU top ID +against CPU argmax, and the retained last-row logits. The supervised rerun +passes in1.60s (`lazy-head-exact-fixed.*`); the initial failure is retained in +`lazy-head-exact-test.*`. This is a regression check, not a new exact-sampling +reference/performance claim. No production setting was changed for this test. +Final warnings-denied Clippy, rustfmt and diff checks pass after this test-only +edit. No model or benchmark process remains from this work unit. + +### Persistent CPU Markov workers + +`src/engine/metal/markov.rs` replaces per-draft scoped thread creation with a +model-owned Rust worker pool. Original source contracts are `ds4.c:1912–1947` +(default min(online,12), caller plus helpers), `1971–2005` (contiguous partitions, +serial execution below512 rows), and `33859–33963` (ordered first-tie Q8 argmax). +Positive integer `DS4_THREADS` settings are supported up to the original32-thread +limit. No CUDA/non-Apple Markov branch is enabled. + +The GGUF mapping is shared through `Arc` without remapping or copying the +weights. Worker inputs own their temporary data: the full logits `Vec` is moved +into shared read-only ownership and returned after all dispatched jobs finish. +There are no borrowed raw pointers between worker lifetimes. Workers release +their input before signalling completion; results are drained on failure too, +and pool destruction closes and joins every helper. Existing dot/quantization +functions are reused without arithmetic changes. No dependency or native host +code was added. + +`markov-rust.*` completes the full warmup and ongoing three-turn chat. All1755 +original cycles and all answer/thinking/token/cache/EOS checks remain exact +(`markov-cycle-comparison.json`, `markov-result-comparison.json`). Two focused +Markov tests and seven GGUF tests pass, including thread and logits-buffer +reuse, ordered ties, shape validation and mapping bounds. Release all-target/ +all-feature build and warnings-denied Clippy pass. The full diagnostic run is +not throughput acceptance. CLI SHA256: +`0e24c6e8102c2035cc89e93f77eede4235f5aa4611c36ce5fb225770641485f2`. + +The remaining allocation task is concrete: `BatchScratch::allocate` reserves +workspace based on both row count and context position; `snapshot_spec_frontier` +allocates compressor/indexer/target snapshots every verifier cycle. Reuse must +preserve delayed support seeding, partial acceptance, error rollback and padded +output-head behavior. Merely retaining an arbitrary previous batch is not a +safe implementation of the original persistent scratch contract. + +Two clean full pairs (`clean-markov-{1,2}-*`) ran DS4→Rust→Rust→DS4, power100, +same installed files/settings/warmup/ongoing chat, trace and canary off. All +content/thinking/prompt/cache/token checks pass; both workers in both pairs exit +successfully, without downloads or overlapping model processes: + +| Pair | Turn | Rust / DS4 prefill ms | Rust / DS4 decode t/s | Decode delta | +| --- | --- | ---: | ---: | ---: | +| 1 | Summary | 6686.486 /5445.567 | 32.046 /36.968 | −13.32% | +| 1 | Story | 296.457 /280.371 | 26.201 /31.162 | −15.92% | +| 1 | Python | 376.741 /362.034 | 28.251 /34.122 | −17.21% | +| 2 | Summary | 8407.889 /8535.688 | 26.901 /27.488 | −2.14% | +| 2 | Story | 332.174 /383.535 | 23.301 /24.322 | −4.20% | +| 2 | Python | 398.755 /466.153 | 26.975 /28.215 | −4.40% | + +The second pair still misses decode parity. DS4 Summary throughput again falls +by about25.6% over the series; no systemwide throttling cause is asserted. A +single `sudo -n powermetrics` query failed immediately because a password was +required (`markov-power-sample.txt`); it started no sampler. Rust supervisor +totals117.095/132.871s versus DS4 totals98.054/127.837s include different +frontend/finalization work and are not pure decode. Rust maximum progress gaps +were6.686/8.407s and supervisor lag59/58ms. No canary/UI-responsiveness claim. + +To separate dispatch cost from GPU drift, the ignored CPU-only +`installed_markov_worker_dispatch` test uses installed Markov W1 row671 and the +actual W2 mapping with a fixed logits row. It compares per-call18-thread spawning +against persistent12 and18 workers, holding row arithmetic/input ownership +constant. Four alternating orders,128 calls per mode per round, all1536 argmax +results identical. Median times per128 calls: + +| Dispatch | Median ms | ms/call | +| --- | ---: | ---: | +| Scoped18 | 40.356 | 0.3153 | +| Persistent12 (production default) | 36.077 | 0.2819 | +| Persistent18 (diagnostic only) | 30.496 | 0.2382 | + +The default pool reduces this isolated dispatch/calculation time by10.6%, about +0.034ms per call; it does not explain seconds of whole-chat difference. This +diagnostic is neither a C Markov microbenchmark nor end-to-end parity evidence. +Production retains the reference's default12-thread policy. The same runnable +worker regression also closes one helper and checks that all other jobs drain, +an error is returned and the caller's original logits buffer is preserved. +Receipts: `markov-dispatch-test.*`, `markov-worker-failure-tests.txt`. + +Final verification after the additional test-only coverage: seven enabled +DSpark tests pass; the explicitly supervised exact-sampling test passes in1.35s; +warnings-denied Clippy, rustfmt and diff checks pass. All processes from this +work unit have finished. No app bundle replacement, commit or push was made. +The full six-cell performance goal remains open. + +### Persistent verifier buffers: functional checks pass, timing remains unaccepted + +The next implementation reuses baseline compressor/indexer snapshots, high-water +prefix snapshots and one verifier workspace per existing padded row shape. +Deferred seed ownership determines when a batch can be recycled; full-prefill +workspaces are not retained. Position-sensitive verifier storage reserves the +session context. Existing output-head row padding, sampling, power and command +submission boundaries are unchanged. Error rollback keeps its previous behavior; +an error may drop scratch and require allocation on the next attempt. + +Evidence remains in `local-eval-results/ds4-cycle-audit-20260911.vfsUc8/`. +CLI SHA256: `4f3643bd854cad1d317fdd6bb93649327902b19cb3cd7b9cfff24a760d57e880`. +The prior CLI is retained as `before-verifier-reuse-ds4-server` for a later +controlled comparison; merely comparing older sequential runs is insufficient. + +- `reuse-cycle-test.txt`: the32-cycle original oracle passes, now also checking + native buffer identities for both padded verifier shapes and byte-exact + compressor/indexer rollback after another target step. +- `reuse-cycle-comparison.json` and `reuse-result-comparison.json`: the complete + traced chat matches all1755 original cycles and all text/thinking/token/cache + checks. Traced timing is diagnostic only. +- `reuse-exact-test.txt`: supervised exact-sampling regression passes after the + reuse changes, including complete target-distribution readback. +- `reuse-build.txt`, `reuse-clippy.txt`, `reuse-dspark-tests.txt`: release build, + warnings-denied Clippy and seven enabled DSpark tests pass. +- GPU allocation at each measured turn end is95,039,750,144bytes, approximately + 73MiB more retained than the prior Summary endpoint. This is retained scratch, + not proof of a throughput improvement. + +Two complete pairs ran in DS4→Rust→Rust→DS4 order, with trace/canary off and no +overlapping model processes or builds. The user subsequently confirmed concurrent +video playback; "clean" in these artifact filenames means instrumentation off, +**not** an idle GPU. Both pairs preserve all outputs and exit successfully. + +| Pair | Turn | Rust / DS4 prefill ms | Rust / DS4 decode t/s | Decode delta | +| --- | --- | ---: | ---: | ---: | +| 1 | Summary | 6679.941 /5718.646 | 34.204 /38.273 | −10.63% | +| 1 | Story | 259.070 /273.257 | 28.306 /33.200 | −14.74% | +| 1 | Python | 355.567 /333.715 | 29.343 /36.744 | −20.14% | +| 2 | Summary | 8968.416 /10283.116 | 26.678 /24.830 | +7.44% | +| 2 | Story | 324.111 /445.256 | 21.926 /22.225 | −1.35% | +| 2 | Python | 365.097 /560.790 | 23.092 /26.814 | −13.88% | + +Receipts: `clean-reuse-{1,2}-{comparison.json,rust.jsonl,reference.jsonl}` and +matching stderr. Rust supervisor totals110.153/140.256s versus original +93.671/141.037s include frontend/finalization differences, not just decode. +Rust progress gaps6.679/8.968s and supervisor lag45/45ms are not GPU-canary or +GUI-eventloop delays. No new responsiveness acceptance, app bundle replacement, +commit or push. Remaining work includes telemetry ownership reconciliation, +reference code/cost audit and the controlled full six-cell comparison. + +### Interim checkpoint requested by the user + +The current work unit is being closed for commit/push, not declared full parity. +The commit-gate test run exposed an already-committed stale Qwen source-inventory +assertion:22 runtime units were expected although both HEAD's Metal export and +the pinned generator contain26. The read-only command below verifies the entire +export byte for byte against pinned sources (17 custom bodies,26 runtime units, +dynamic QSA sources). Only the stale expected count is corrected; no kernel, +hash, fixture or per-body assertion is changed or removed. + +```sh +python3 tools/mtplx-kernel-source.py local-eval-results/mtplx-reference-e652d55 --gated-delta-source local-eval-results/mtplx-reference-env-0.32.2/lib/python3.12/site-packages/mlx_lm/models/gated_delta.py --check +``` + +The app bundle is rebuilt as required by the commit gates. This supersedes the +earlier statements that this follow-up had not yet replaced the bundle; no GUI +or new throughput series is launched. Unrelated `tools/__pycache__/` files are +left untouched and excluded from the commit. + +Commit verification: `cargo fmt --all -- --check`, warnings-denied all-target/ +all-feature Clippy, `make bundle`, and `cargo test --all-features` pass. The full +suite reports303 main tests plus4 supervisor and4 integration tests passed, +zero failures,202 explicitly ignored main tests. The separately supervised +32-cycle and exact-sampling results above remain distinct from these normal +gates. `codesign --verify --deep --strict` and the bundle's `model-eval --help` +also pass without opening the GUI. Final suite receipt: +`local-eval-results/ds4-cycle-audit-20260911.vfsUc8/interim-commit-tests.txt`. diff --git a/docs/glm-scheduling-followup-20260911.md b/docs/glm-scheduling-followup-20260911.md new file mode 100644 index 0000000..d8e96c9 --- /dev/null +++ b/docs/glm-scheduling-followup-20260911.md @@ -0,0 +1,585 @@ +# GLM execution and responsiveness follow-up — 2026-09-11 + +## User acceptance and scope + +- Qwen is confirmed good in normal interactive use. +- GLM decode is now also confirmed good interactively. GLM prefill remains + usable, but feels less smooth than Qwen and affects other applications. + This is not a claim of a complete freeze or a new confirmed beachball. +- The user authorized resuming the outstanding work and explicitly authorized + compiling/running antirez/ds4 as a standalone, supervised reference benchmark. + No DS4 C objects are linked into DS4Server or its application bundle. +- Qwen's golden master remains MTPLX; DeepSeek/GLM remain antirez/ds4. + The accepted 2.6% Qwen Summary AR exception is not a general tolerance. + +## Implemented execution changes + +Both scalar GLM loops now flush periodically every four completed layers, +excluding the final layer and SSD expert streaming. Previously they flushed +only once at layer four. This follows the active indexed DS4 graph, including +scalar MTP fallback/rejection replay. The reference's dynamic per-layer mapping +fallback must not be confused with Rust's static non-expert decode mapping: +`glm_streaming_model_spans` retains non-expert tensors, the configured resident +expert prefix, and incompatible expert layouts; selected experts are loaded +through the existing native cache. No new per-layer SSD waits were introduced. +Low-memory dynamic mapping fallback parity is not established by this patch. + +GLM 5.3 prefill progress now advances at existing completed GPU drains and after +the final output evaluation, not after every submitted layer. Chunk selection, +prefill flush/drain placement, Metal kernels, sampling and power policy are +unchanged. This corrects progress accounting; it does not by itself fix the +remaining prefill smoothness issue. + +Targeted checks passed: periodic/final/SSD decode boundary test; existing +prefill boundary test; live two-row verifier acceptance, rejection, rewind, +scalar fallback, recurrent state, unused HC workspace guards and lifetime +counters. The live test additionally verifies completed-prefill progress points. + +## Standalone reference and instrumentation + +`tools/ds4-session-reference.rs` is a separate Rust benchmark driver for the +unchanged public DS4 engine/session interface. It is deliberately not a Cargo +target and is never included in the app. Its build script verifies reference +commit `ec7642cdd9ec81d01ad4b1fd8f8a3d1511533748`, the pinned header hash, +unchanged tracked engine sources and current reference objects. The arm64 ABI +layout is checked against Clang's record layout (engine options 280 bytes, +distributed offset 152, TP offset 216). + +The public DS4 CLI/API maps `low` to `high`. The driver instead constructs the +Low system prefix through the public chat API and passes tokens to the original +session implementation. It also reproduces the UI's separate system-prefix +prefill (9 tokens for GLM), and retains generated token history in one session. +`glm53_reference_prompt_tokens_match_shared_runtime` checks every token of all +four prompt streams, including continued turns, against the production Rust +tokenizer. The first complete matched-bootstrap AR reference passed this check. + +Build from the DS4Server checkout, with already-built reference objects: + +```sh +bash tools/build-ds4-session-reference.sh /Users/gb/Projects/ds4 /absolute/path/reference +``` + +Run the resulting binary from the reference checkout under `test-supervisor`: + +```text +test-supervisor 114688 30 45 --command /absolute/path/reference MODEL_GGUF glm on README_PATH +``` + +The driver uses power100, Low, context32768, temperature0.6, top-p0.95, +top-k0, min-p0, seed42, SSD streaming off and graph-selected GLM chunks. +Warmup is a separate session (up to32 tokens), followed by README Summary, +lighthouse Story, and Python `is_prime` in one ongoing chat, to natural EOS. +There is no total-runtime watchdog. Missing models fail; no downloads occur. +DeepSeek is supported by the driver's `deepseek` family argument, with the +installed DSpark support GGUF required for acceleration-on; it has not yet been +validated by this follow-up's GLM runs. + +Optional `DS4_REFERENCE_CANARY=/absolute/path/ds4-server` starts the **same native +probe and monitor implementation** as the UI/harness in a separate process. +The `gpu-canary` CLI accepts phase labels on stdin and ends on EOF. A readiness +handshake waits for the first successful probe before model loading. Its +readiness/stall/sample reports are not model-progress watchdog heartbeats. +Clean throughput runs leave this variable unset. An external observer and an +in-process observer must not be treated as identical OS scheduling conditions. + +`DS4_REFERENCE_IN_PROCESS_CANARY=1` instead enables a native probe thread inside +the reference process. It links the same `native/metal/ds4_canary.m` used by +DS4Server, with the same4096-byte blit, separate queue and100ms cadence, rather +than duplicating a Metal implementation. This bridge contains no model code. +Its Rust monitor logs per-sample phases/timing but is not the UI event loop. +The optional external observer can also be enabled simultaneously. Neither is +enabled for clean reference throughput. `reference --canary-self-test` exercises +readiness, two phases and clean shutdown without loading a model. + +The initial external-probe integration test exposed a startup race: a phase +could end before the executable initialized. It was not worked around by +loosening the assertion; the adapter now requires a readiness handshake. The +initial diagnostic without that handshake is retained, not a full-startup proof. + +## Measurements and limitations + +All raw receipts, outputs and failed attempts are retained under +`local-eval-results/glm-scheduling-20260911.7VSaeW/` (ignored, local evidence). +The before binary is the secured `02db096` implementation, SHA256 +`ab289f067a7a01c22113eec76aa896638d83e392242192e1440d14ed11524d5c`. + +Fresh complete AR runs (after, then before), same output/reasoning/tokens/EOS: + +| Turn | Tokens | Before decode t/s | After decode t/s | Before prefill ms | After prefill ms | +| --- | ---: | ---: | ---: | ---: | ---: | +| Summary |640|24.467|25.403|7561|6034| +| Story |1102|23.703|24.135|262|255| +| Python |200|24.726|25.401|305|302| + +These are single sequential pairs, not drift-controlled medians. The large +summary-prefill difference cannot be attributed to a decode-only flush change. +No hard decode regression was observed; full performance parity is not proven. + +The matched-bootstrap original DS4 AR reference completed naturally at +26.618/25.123/25.825 decode t/s, with779/1047/198 output tokens. Its generated +text differs from Rust despite matching initial prompt tokens/settings; later +contexts therefore also differ. This is not an exact-output performance pair. +Reference prefill timers measure session sync; Rust's current GLM `prefill_ms` +still includes the observed UI phase. Do not silently equate those intervals. + +The new Rust MTP run retains the previous625/1026/196 completion tokens and +identical output/reasoning. Its draft acceptance fractions are241/385 (62.6%), +342/685 (49.9%) and95/102 (93.1%). Python therefore has the expected higher +acceptance; MTP's benefit is workload-dependent, not uniformly absent. + +The fresh MTP before/after pair also preserves every output/reasoning token +and natural completion: + +| Turn | Tokens | Before decode t/s | After decode t/s | Before prefill ms | After prefill ms | +| --- | ---: | ---: | ---: | ---: | ---: | +| Summary |625|19.988|23.154|9303|6190| +| Story |1026|17.728|19.295|306|272| +| Python |196|27.843|29.506|347|322| + +The same sequential-run/drift limitation applies. This establishes no observed +hard regression, not a controlled causal speedup or reference-parity acceptance. + +### Canary placement and timestamp attribution + +The full `after-mtp-dual-canary` run had simultaneous internal/external probes. +The internal probe recorded867 successful samples, with a prefill maximum +of490.740ms and decode maximum3.180ms. The external probe was ready before model +launch and continued until after termination:3004 successful samples, overall +maximum3.331ms (startup), and1.572ms while labelled `preparing` across the model +lifetime. That external label is deliberately not turn/phase attribution. +Both probes stopped cleanly; no sample failed or reached2s. This is diagnostic +evidence, not a clean throughput run or a compositor-frame test. + +Thus the previous `completed_ms` cannot be interpreted as a measured systemwide +GPU blockade. It includes host-side waiting and completion delivery. The optional +shared native probe now also records commit-to-GPU-start (`gpu_wait_ms`), +GPU-start-to-end (`gpu_interval_ms`), and GPU-end-to-host-return (`host_return_ms`). +Metal's GPU timestamps use system mach time; the probe uses `mach_absolute_time` +and the native timebase for those differences, not `CLOCK_MONOTONIC`. Missing or +inconsistent timestamps remain null, not zero. The GPU interval includes possible +GPU scheduling/preemption, not exclusively active blit execution. See Apple's +[GPUStartTime documentation](https://developer.apple.com/documentation/metal/mtlcommandbuffer/gpustarttime). + +The model-free Metal integration check verifies phase coverage, valid nonnegative +intervals and their bounds against wall completion; the synthetic unit check +retains null timing when unavailable. The existing UI uses the same enhanced +native probe, but its stats panel still displays the existing wall latency fields. +Neither model work nor disabled-canary execution invokes the new timestamp work. + +The first full Rust timestamp run (`after-mtp-timeline`) preserved every MTP +output/reasoning token and EOS. All902 samples had valid Metal timestamps and +none failed. Its worst prefill sample was300.003ms:299.892ms before GPU start, +0.001917ms GPU interval, and0.105958ms after GPU end. The decode maximum was +3.733ms. This directly rules out delayed host return as the dominant cause of +that prefill sample; the queued probe waits for GPU execution. It does not show +that a different application's rendering queue is delayed by the same amount. + +The subsequent extraction into the shared native object changes no probe work: +direct `[cb commit]` replaces the wrapper whose model-queue-only hook never +applied to this separate canary queue. Both native bindings pass their model-free +checks after extraction. Final current product binary SHA256: +`50b4b4abbbdc45ff600c1f46d0bec611879249ac8e4d8291d22d656b9c6e9a5d`; +standalone reference binary: +`ccc7a8a774cb1c202add6dba60b04dffe3597822b15a34e22c7e4a5574b50adf`; +shared probe source: +`dd3abc34088ee27ba0759f01a291b9b714114420295252d63e85fd6f326fddab`. + +Answer correctness is checked separately from natural termination. Rust AR, +Rust MTP and reference AR passed their generated assertions plus5011 `is_prime` +cases (-10 through5000). The preliminary reference MTP output passed its own +five assertions but failed347 additional cases, first at49: it omits the +`i + 2` divisor test. This is a failed generated Python answer, not by itself +evidence of an engine defect. It must not be reported as a successful code +benchmark merely because EOS was reached. Details are in `python-check.json`. + +### Reference clean MTP and record integrity + +The clean `reference-mtp-clean` run (both canaries disabled) completed all turns +with the same593/872/167 tokens, text, stop tokens and failed Python answer as +the diagnostic reference run. Its prefill times were5462.306/351.803/434.414ms; +decode21.483/16.754/26.143t/s. The preceding in-process diagnostic measured +23.623/19.287/30.440t/s. This spread must not be disguised as a port speedup or +accepted2% parity: it is one sequential comparison with different probe state, +not controlled repeated clean medians. Canary-on throughput is not the baseline. + +The first internal reference run reported788 successful probes, but only787 +were independently parseable: a watchdog resource record interrupted one +canary JSON record at a pipe-read boundary. That failed record is preserved in +`reference-mtp-inline.stderr.log`, not silently counted as missing/zero latency. +The supervisor now forwards complete lines in one locked stream write (with a +64KiB cap for newline-free output), while watchdog progress still consumes every +incoming chunk immediately. EOF flushes partial output. A split-record regression +test and all existing memory/start/continuation/long-run watchdog tests pass. +The reference diagnostic is repeated as `reference-mtp-inline-records` for a +fully parseable receipt; the earlier run is retained as the failure evidence. + +That repeated reference run completed with **771/771 parseable, successful, +fully timestamped samples** and identical593/872/167 generated tokens/text/EOS. +Prefill p95/max was233.106/264.292ms (48 samples); decode p95/max was +0.226/19.338ms (713 samples). The worst prefill probe waited264.195ms before +GPU start, ran over0.001750ms, and returned to the host0.092083ms after GPU end. +No probe reached2s. The reference's prefill samples also include its short +warmup; the worst sample occurred during the measured summary prefill. +Its diagnostic throughput was23.893/19.177/32.090t/s, not the clean baseline. + +| Matched native in-process probe | Prefill p95 ms | Prefill max ms | Decode max ms | +| --- | ---: | ---: | ---: | +| DS4Server, `after-mtp-timeline` |289.978|300.003|3.733| +| Original DS4, `reference-mtp-inline-records` |233.106|264.292|19.338| + +These sequential diagnostics reproduce the same GPU-start-wait phenomenon in +the golden master. They do not excuse the remaining Rust prefill cost, establish +statistical latency equivalence, or measure another application's compositor. +Moving inference to another thread cannot by itself reproduce the independent +process's scheduling conditions; process isolation is a distinct architectural +option, not implemented or declared proven as a UI fix here. + +Verification at this checkpoint: release all-target/all-feature build; release +all-target/all-feature Clippy with warnings denied; rustfmt and diff checks; +11 model-eval unit tests; both model-free native probe bindings; all4 supervisor +tests; earlier live GLM verifier/progress/HC guards and the full AR/MTP chats. +The updated supervisor fixes measurement transport, not inference scheduling. + +## Sampling versus model execution — continued investigation + +The prior follow-up made concrete progress (execution fixes plus a fair native +in-process latency reference), but did not establish the full three-model, +AR/speculative2% goal. This continuation addresses the different GLM outputs +before treating their different ongoing histories as matched performance work. + +`reference --sampler-fixture` runs the original public `ds4_sample_logits` without +loading a model or using Metal. The checked-in +`tests/fixtures/ds4-sampling-ec7642c.json` contains64 cases: four vocabulary sizes, +eight temperature/top-k/top-p/min-p settings, seeds0/42,32 consecutive tokens +per case and the final RNG state. The original Rust test failed40 of64 cases. +The shared DS4/GLM sampler now preserves the first argmax tie and original +negative sentinel, skips RNG consumption for greedy/all-invalid and the DS4 +full-vocabulary min-p fallback, and preserves seed0 until the original RNG's +zero-state substitution. Qwen's independent MTPLX sampler is untouched. +All64 oracle cases and the16 enabled sampling tests pass. Crucially, the positive +temperature/top-p benchmark cases at seed42 already passed before the fix: +these edge corrections are not the explanation for the observed GLM chat gap. + +Optional `DS4_REFERENCE_LOGITS_TRACE` records the first32 summary logit rows +through the public original session API. It requires AR mode, creates a new +file rather than overwriting one, and does not change generated tokens or RNG. +The full `reference-ar-logits` chat retained exactly the779/1047/198 tokens, +text and stop tokens of `reference-ar-bootstrap`. Its timings are diagnostic, +not a clean performance baseline. The binary trace contains19,824,640 bytes +(32 rows of154,880 little-endian floats). Its path is serialized as an OsString +and decoded losslessly by the replay test. + +`glm53_reference_logits_replay_separates_sampling_from_execution` first samples +those original C-produced rows through the production Rust sampler: **all32 +tokens match**. It then opens the installed GLM at Power100/context32768, +prefills the same9-token bootstrap and exact summary suffix, and advances only +with reference-selected tokens. Thus histories never diverge during comparison. +On the Rust-generated rows the test **fails at step17**, choosing906 instead of +the reference320. Already the first post-prefill row has max absolute difference +5.722162 and RMS difference0.851167. All32 per-step row errors are retained in +`logits-replay.stderr.log`; the watched test terminates normally with failure +status in9s. This is a new, deliberately retained red parity test, not a passed +live validation or a speed result. No DS4/GLM/Metal/CPU diagnostic override was +present in the parent environment. + +The next localization belongs in the model execution path: compare existing +original DS4 per-layer tensor dumps with the corresponding Rust HC/KDA/DSA/FFN +stages, starting at the first bootstrap/prefill block. Do not explain this away +as stochastic output variation or hide it with a lower chunk/power setting. +No speculative numerical tolerance or new scheduling workaround was applied. + +## Root cause: GLM 5.2 chunk boundary applied to GLM 5.3 + +The active original indexed GLM 5.3 path deliberately keeps full2048-token +chunks across both the old2048 indexer threshold and the4096/8192 dense-attention +threshold. Rust was still applying the GLM 5.2 top-k boundary: after the9-token +bootstrap it evaluated2039 tokens, whereas DS4 evaluated2048. This changes the +recurrent prefill computation, despite identical total prompt tokens. + +The original layer0 bootstrap `attn_out` and `ffn_out` dumps matched Rust +bit-for-bit. The original position9 dumps contain2048*4096 floats, establishing +the actual chunk geometry rather than inferring it from configuration. +Detailed HC dump hooks elsewhere in DS4 belong to an inactive dense path and +were not used as evidence for the active indexed execution. + +Rust now retains complete GLM 5.3 chunks and splits only the attention slices +at the dense/sparse boundary, as DS4 does. This also removes the incorrect +whole-pair sparse override for a two-row verifier crossing that boundary. +GLM 5.2 retains its old top-k splitting. Unit checks cover both families and +the4096/8192 attention transitions. No smaller chunk, delay, or power reduction +was introduced. + +After this correction, the same fixed-history replay is green: **all32 full +154880-value logit rows are bit-identical** to the original trace (max absolute +and RMS error both0), and all sampled tokens agree. This run had no stage +instrumentation enabled. Evidence is retained under +`local-eval-results/glm-stage-20260911.rwQBaJ/mixed-replay.*.log`. +The earlier red replay remains historical evidence, not the current result. +The optional Rust stage reader exists only under `cfg(test)` and validates +tensor geometry before comparing values; it adds no production GPU drains. + +The live verifier at frontier4095/context32768 passed across the4096 boundary, +including acceptance, rejection, rewind to either retained frontier, scalar +fallback and recurrent-state restoration (`mixed-boundary.*.log`,70.64s). +The32-row replay alone is not a complete performance or output-parity claim. +The initial source-only note about one-token suffixes was incomplete: the +shared UI/headless consumer already routes one-token extensions through scalar +execution. The actual remaining crossover was two/three-token extensions; +see the subsequent common-consumer correction below. + +### Complete chats after the chunk correction + +Fresh clean runs used the same ongoing workload, Power100/Low, native EOS, +separate warmup and no active canary. Rust executable SHA256: +`4b23c04325c931854b98c23bd2c98df8a5c2362927aa9b1faed65019d07fd40d`. +The original reference retained its prior tokens/text/stops exactly. + +| Mode / turn | Rust tokens | DS4 tokens | Rust decode t/s | DS4 decode t/s | Output + thinking identical | +| --- | ---: | ---: | ---: | ---: | --- | +| AR Summary |779|779|24.512|20.955|yes| +| AR Story |1047|1047|22.956|19.744|yes| +| AR Python |198|198|23.436|20.650|yes| +| MTP Summary |593|593|20.235|21.427|yes| +| MTP Story |905|872|18.039|18.007|no| +| MTP Python |169|167|29.443|29.511|no| + +All six Rust turns and six reference turns ended naturally. AR prompt/cached +counts also match exactly. Receipts: `clean-comparison.json`, +`mixed-ar-output-check.json`, `mixed-mtp-output-check.json` in the stage evidence +directory. The AR reference was materially slower than earlier clean runs; +these sequential pairs are not a controlled speedup or2% acceptance claim. +MTP Summary is about5.6% slower in Rust in this pair; the later MTP throughput +numbers do not compare identical histories. Prefill UI-phase and original +session-sync timers still have different boundaries (raw values in the receipt). +Control-loop maxima of40–59ms are not GPU canary or compositor measurements. + +### Second root cause: MTP stop token retained in the ongoing frontier + +Although MTP Summary text/thinking and593 emitted tokens match, Rust starts +Story with3229 cached tokens and3249 prompt tokens; DS4 uses3228/3248. +The shared Rust generation consumer returned on an MTP stop token without +rewinding the already evaluated block. Both normal and raw original DS4 agent +consumers call `ds4_session_rewind(block_start + ti)` at that point. The standalone +reference's stop handling therefore agrees with its real agent, not just an +arbitrary benchmark convention. + +The shared UI/headless consumer now calls `rewind_speculative_output`, a thin +GLM adapter over the existing two-row rollback, to keep exactly +`prompt_tokens + emitted_tokens` before retaining the chat. +This restores the saved two-row KDA state and replays the retained row; it does +not merely truncate IDs or re-render generated text. Invalid frontiers fail +explicitly. Both sampled and greedy generation use this consumer. Qwen's own +whole-turn controller is unchanged. Other model-specific speculative stop +contracts are not claimed validated by this GLM change. +The live verifier regression now exercises that same consumer rollback path. +`align_prompt` is intentionally not used: it retains one fewer token to force +logit recomputation during prompt synchronization, which is a different contract. +The full post-frontier-fix MTP measurement (`frontier-mtp.*.log`) now matches +the original for **all three turns**: text, thinking, emitted token count, +prompt count, cached frontier and natural stop. Emitted counts are593/872/167; +Story starts at3228 cached/3248 prompt tokens, Python at4120/4145. The executable +SHA256 is `b964336d64fbb90b3a9ca595a4705eda02e7afe9c39aedb4ea775e0d52fcf20e`. +`frontier-mtp-output-check.json` has three entries with every equality true; +the checked `jq -e` assertion requires all three entries and all five properties. +Decode rates are23.880/19.357/31.643t/s, versus21.427/18.007/29.511 in the directly +preceding clean original MTP run. This is one sequential pair, not repeated2% +acceptance. The Python answer is now exactly the reference's previously checked +incorrect answer (first counterexample49); matching the oracle does not waive +the independent generated-code quality failure. + +### Final regression and responsiveness diagnostics + +The final strict replay passes with bit-equal logits at all32 steps. The two +original layer0/position9 stage tensors each contain8388608 floats and also +match bit-for-bit (`final-replay.*.log`,10.21s). The updated live verifier at4095 +passes through the same rollback entrypoint used by the consumer, including +invalid/unchanged-frontier checks, rejection and both retained rows +(`final-boundary.*.log`,82.27s). The five enabled GLM unit tests pass. + +`final-canary` retained identical full MTP output/frontiers. Its in-memory +summary reports840 samples, no failures, prefill p95/max395.611/483.166ms, +decode max3.820ms and no sample crossing the configured2s threshold. However, +strict raw-log parsing found an interleaved canary/resource JSON record: the +model-eval parent inherited the child's stderr, and both processes serialized +JSON fragments to that descriptor. This raw file is retained as a **failed +record-integrity diagnostic**, not silently filtered into a complete sample set. + +The model-eval supervisor now pipes child stderr and forwards complete lines +under the parent's shared stderr lock, the same lock used by resource samples. +Diagnostics do not refresh inference progress deadlines. Reader failures are +reported on join. This fixes the app harness counterpart of the earlier +standalone watchdog forwarding issue; it changes measurement transport, not +GPU scheduling or the UI inference graph. + +The directly following original DS4 in-process probe run +(`final-reference-canary`) has841/841 parseable samples, no failures, unchanged +reference tokens/text/stops, prefill p95/max373.220/388.090ms and decode max4.370ms. +The worst prefill sample spent387.964ms before GPU start,0.002875ms over its +GPU interval and0.122ms returning to the host. Thus substantial prefill queue +waiting still occurs in the original oracle; the larger Rust spike is not +declared equivalent or explained away. + +The repeated Rust run after the forwarding correction (`final-canary-records`) +completed the entire chat in83.744s and preserved all output/frontier fields. +Every JSON record beginning with `{` in stderr was parsed with `fromjson` +(no error suppression): **776/776 canary records and82/82 resource records** +match the independently reported totals. The checked receipt is +`final-canary-records-check.json`. There are no probe failures or observed2s +threshold crossings. Prefill p95/max is119.507/247.181ms; decode max1.854ms. +The worst sample waits247.062ms before GPU start, spans0.001750ms on the GPU, +and returns after0.115458ms. This lower maximum is not attributed to the +transport-only fix: the prior483ms Rust and388ms original spikes remain recorded, +and scheduling/throughput variability still requires repeated paired testing. +The optional probe remains off by default; no negligible-overhead claim is made. + +Final source verification: release all-target/all-feature build and Clippy +with warnings denied; rustfmt/diff checks; five GLM unit tests; eleven +model-eval unit tests; sixteen sampling tests including the64-case original +sampler fixture; strict live logits/stage and consumer-rollback boundary tests. +The final CLI SHA256 is +`cbe04f8ce8f8d2fcb6c82b97c3d85b7bed561418893621a6a653d344d1aa6d85`. +The previously good bundle remains unchanged at SHA256 +`ea4d555c2faf0940d9cbcf76d8638ca614a9cb2c6b034e3b2f80aeef86b0b339`. + +## Common prompt timing and DS4 CPU sampling follow-up + +Evidence for this continuation is under +`local-eval-results/glm-paired-20260911.eClfCS/`. The preceding goal turn made +verified progress (chunk scheduling and stop-token frontier fixes); it did not +establish the full six-cell performance goal. + +The shared consumer now uses DS4's GLM5.3 resumed-prefill crossover of2 tokens, +not the generic4-token threshold. DS4 explicitly documents this choice as +measured on M5 Max/GB10 (`ds4.c:36784`). One-token continuations were already +scalar; cold/vision paths and the separate MTPLX whole-turn controller are +unchanged. The enabled crossover test covers GLM5.3 versus GLM5.2/DeepSeek. + +DeepSeek/GLM now publish the existing `PromptTiming` at the shared prompt- +evaluation boundary: after restoration/bootstrap, around actual suffix execution +including its progress callbacks, before decode/checkpoint storage. Exact cache +hits report zero evaluated work. Separately unmeasured restore/history components +are `null`, not fabricated zeros; Qwen continues reporting the same measured +numeric values through `Some`. The new metric test and existing Qwen progress/ +decode-timer test pass. The ordinary UI-prefill timer remains separately visible. + +A fresh clean AR pair kept all three outputs/thinking/token counts/frontiers +identical. Rust's engine-prefill times were5235.183/269.769/320.767ms, original +DS4 session-sync8160.750/495.857/479.216ms; Rust decode24.811/23.114/23.410t/s +versus16.633/17.333/19.054. These large sequential-run differences are not a +controlled speedup or a completed repeat matrix (`baseline-ar-comparison.json`). +The reference driver now additionally queries and checks actual engine power100 +after load, rather than only recording its requested options. + +The CPU sampler still differed algorithmically: Rust sorted the full vocabulary +and drew from renormalized probabilities, while DS4 first tries a512-candidate +heap and draws from raw retained weights. A CPU-only replay uses the existing32 +full logit rows, one32-draw warmup and16 measured batches (512 draws). The same +small runner serves the independent original public `ds4_sample_logits` and the +production Rust sampler. It loads no model and performs no Metal work; both are +supervised with1GiB memory/start30s/idle30s limits. The original public function +allocates a scratch buffer per call, unlike its session API, so its microbenchmark +is not an exact measure of session-sampler overhead. + +Before alignment Rust took2.645ms/draw versus original0.834ms, with all512 tokens +equal. The aligned Rust path initially measured0.401ms/draw with the same512 +tokens (`sampler-{before,after}-rust.json`, `sampler-reference.stdout.log`). +It uses stdlib `BinaryHeap`, DS4's logit/index tie order, bounded-nucleus fallback +without advancing RNG, original raw cumulative sampling, full-vocabulary/min-p +fallback and the original expf-verified log-space rejection boundary. Top-k +retains the original1024 cap. Separate distribution materialization for +speculative correction and Qwen's MTPLX sampler are untouched. + +All64 original sampler fixture cases and17 enabled sampling tests pass, as does +the added missing-mass/near-one fallback, RNG and signed-zero tie check. Release +all-target/all-feature build and warnings-denied Clippy pass. The new executable +SHA256 is `ece6aed3601fb402e6dba6ac2e289d6e0c2dc86663600c3d4b1a4cc07e8fb42c`. +The first full post-sampler AR and MTP pairs both preserve all three outputs, +thinking, completion/prompt/cached counts and natural stops. The independently +queried reference engine reports power100. Receipts are +`sampler-{ar,mtp}-comparison.json`; these are single pairs, not the repeat matrix. + +| Mode / turn | Rust / original engine-prefill ms | Rust / original decode t/s | +| --- | ---: | ---: | +| AR Summary | 5208.845 / 5548.775 | 26.129 / 24.711 | +| AR Story | 267.263 / 287.268 | 24.553 / 23.728 | +| AR Python | 319.127 / 341.355 | 24.964 / 24.520 | +| MTP Summary | 6574.400 / 5427.866 | 23.935 / 23.628 | +| MTP Story | 273.240 / 278.956 | 19.419 / 19.089 | +| MTP Python | 315.279 / 360.621 | 33.140 / 31.919 | + +The Summary MTP prefill regression in this pair remains visible despite the +slightly faster Rust decode. Reversed-order repetitions are needed to distinguish +run variability from a repeatable graph cost. AR before/after the sampler keeps +the entire chat output identical and improves decode by5.310/6.225/6.638% in this +one sequential comparison (`sampler-ar-before-after.json`); no controlled causal +end-to-end percentage is inferred from that pair alone. + +MTP is not universally beneficial in the original either: its Story decode is +19.089t/s versus23.728 AR, while Python is31.919 versus24.520. Rust's full MTP +cycle receipts show228/366,289/584 and82/86 accepted drafts respectively +(62.3%,49.5%,95.3%). The corresponding complete decode-loop time per cycle is +67.69/76.89/58.60ms. At1.62/1.49/1.94 emitted tokens per cycle, the Python case +amortizes the extra draft/verification work much better. These are whole-cycle +averages, not isolated kernel timings: the existing `verifier_ms` includes other +cycle work and must not be presented as an exclusive verification stage. +AR and MTP have different natural histories, so their t/s comparison is not a +matched-token microbenchmark. The previously recorded Python correctness failure +also remains open even though both implementations produce the same code. + +### Reversed-order pairs: acceptance still fails + +Both modes were repeated in original-then-Rust order, serially without builds +or canary probes. All twelve measured answers in these four processes again +match text/thinking/counts/cache frontiers and end naturally; all watchdogs +exit successfully. No slow run was discarded (`repeat2-*-comparison.json`). + +| Mode / turn | Rust / original engine-prefill ms | Rust / original decode t/s | +| --- | ---: | ---: | +| AR Summary | 6738.280 / 5283.041 | 23.569 / 25.535 | +| AR Story | 309.647 / 281.591 | 21.743 / 24.186 | +| AR Python | 387.381 / 329.667 | 21.107 / 24.898 | +| MTP Summary | 9240.178 / 8835.102 | 17.706 / 17.678 | +| MTP Story | 358.905 / 402.447 | 15.146 / 14.581 | +| MTP Python | 403.605 / 491.547 | 25.843 / 23.655 | + +AR decode now misses by7.70/10.10/15.23%; MTP Summary prefill misses by4.38%. +The subsequent original MTP run is itself much slower than its first run. +This excludes neither a Rust scheduling difference nor changing device clocks; +it does preclude a pass based on the favorable first pair or a selected median. +The required third pair and full six-cell acceptance remain outstanding. + +Rust AR emits exactly9445/12612/2424 command buffers in both repetitions, with +the same outputs, but its GPU timestamp-interval sums increase from +34536/42061/8084ms to39285/47601/9590ms (`ar-drift-comparison.json`). Those sums +are `GPUEndTime - GPUStartTime` and may include preemption; they are not exclusive +kernel or hardware-clock measurements. The slowdown is not explained by changed +token counts or extra command buffers, and is not declared thermal throttling. +During the sequence, a read-only process snapshot showed only the intended +reference model process. macOS reported no recorded thermal/performance warning +and normal VM pressure (1), which does not exclude frequency changes. The +AGX PerformanceStatistics snapshot exposes utilization but no frequency field. +Hardware was freshly checked: Apple M5 Max,128GiB,18 logical CPUs. + +## Remaining acceptance + +- Compare repeated clean throughput pairs; GLM ongoing histories now match in + both modes, but sequential run variability does not establish2% performance parity. +- Localize the remaining GLM prefill cost against original DS4's active indexed + path, now that in-process GPU-start waiting is observable on both sides. + The engine-prefill timer is now exposed separately from UI-phase timing; + use that aligned boundary in the paired comparisons. + No chunk reduction or extra waits are justified by these measurements alone. +- Verify the corrected short-extension crossover live where needed, and other + model-specific speculative stop contracts; the recorded GLM workload does not + cover every possible interaction. Full chats pass after CPU-sampler alignment; + repeated timing acceptance remains separate. +- Validate the remaining SSD expert-streaming cases separately from resident + scheduling. This is unrelated to replacing DS4 KV checkpoint persistence. +- Complete the DeepSeek AR/DSpark reference cells and Qwen residual performance + analysis. Interactive confirmations are not a substitute for the six-cell + numerical acceptance matrix. + +No bundle replacement, commit or push has been performed by this follow-up so far. +All processes have terminated. The subsequent DeepSeek comparison and its +separate bootstrap/DSpark findings are recorded in +[DeepSeek follow-up](deepseek-reference-followup-20260911.md). diff --git a/docs/inference-reference-parity-audit-20260905.md b/docs/inference-reference-parity-audit-20260905.md index 0753d08..07443b7 100644 --- a/docs/inference-reference-parity-audit-20260905.md +++ b/docs/inference-reference-parity-audit-20260905.md @@ -7931,6 +7931,19 @@ Verifikation 2026-09-05, ohne GPU-Modelllauf oder breite Gates: - `git diff --check`: bestanden. Keine Durchsatzverbesserung aus diesen Tests abgeleitet. Vor der neuen DSpark-Messung bleiben P02/P03/P04 maßgeblich. +**Gemeinsamer DS4/GLM-Sampler angeglichen (11.09.):** Der CPU-only-Replay mit +32 echten GLM-Logit-Zeilen und512 gemessenen Auswahlen zeigte2,645ms pro Schritt +in Rust gegenüber0,834ms über die originale öffentliche DS4-Funktion. Die Rust- +Implementierung benutzt nun den begrenzten512-Kandidaten-Heap, originale Logit- +Tie-Reihenfolge, rohe CDF-Summen und DS4-Fallbacks statt vollständiger Sortierung +mit erneuter Normalisierung. Initial0,401ms/Schritt bei identischen512 Tokens; +64 Original-Fixturefälle und17 Sampling-Tests bestanden. Qwens MTPLX-Sampler und +die separate Verteilungsberechnung für spekulative Korrektur bleiben unverändert. +Ein kompletter GLM-AR-Chat bewahrt Ausgabe/Thinking/Frontiers und zeigt im einzelnen +Vorher-/Nachher-Paar5,3–6,6% höheren Decode-Durchsatz. Das ersetzt weder die +DSpark-Modusabnahme noch die wiederholte Gesamtmatrix. Details im +[GLM follow-up](glm-scheduling-followup-20260911.md). + ### P14 — GLM-MTP tatsächlich batchen wie DS4 **Produktpfad umgesetzt, Referenz-Performanceabnahme offen (10.09.):** @@ -7996,6 +8009,61 @@ Priorität P1; nach P04, P13; Quellcodedifferenz belegt. ### P15 — DS/GLM-Scheduling, Kaltstart und gemeinsame Runtime abgleichen +**Numerische GLM-Abweichung reproduziert (Fortsetzung11.09.):** +64 modellfreie Original-DS4-Sampler-Fälle sichern Tokens und RNG-Zustand ab; +Greedy-Tie/RNG- und Seed0-Abweichungen sind korrigiert, ohne Qwen zu ändern. +Die verwendeten positiven Seed42-Benchmarkfälle waren davon nicht betroffen. +Der neue feste-Historie-Test auf32 echten Original-DS4-Logit-Zeilen trennt die +Ursachen: Rust-Sampling derselben Zeilen liefert32/32 identische Tokens; Rust- +GPU-Ausführung derselben Eingaben liefert bereits nach Prefill max5,722/RMS0,851 +Logit-Abweichung und bei Schritt17 ein anderes Sample. Der Test bleibt bewusst +rot und ignoriert für normale modellfreie Läufe. Die Referenz selbst bewahrt +den vollständigen AR-Chat exakt trotz optionalem Trace. Nächste Priorität ist +die erste divergierende HC/KDA/DSA/FFN-Stufe, nicht das Durchsatzetikett auf +auseinanderlaufenden Chathistorien. Belege/Kommandos und Grenzen stehen im +[GLM follow-up](glm-scheduling-followup-20260911.md). + +**Fortsetzung 11.09.:** Der Benutzer bestätigt jetzt normalen GLM-Decode und +benutzbaren, aber gegenüber Qwen weniger geschmeidigen Prefill. Kein vollständiger +Freeze behauptet. Die unten am10.09. identifizierte residente skalare Flush- +Abweichung ist in beiden GLM-Schleifen korrigiert (periodisch alle4, ohne letzten +Layer/SSD-Streaming); Prefill-Fortschritt meldet nur abgeschlossene bestehende +Drains und den Batchabschluss. Frische vollständige AR/MTP-Vorher-/Nachher-Chats +bewahren Tokens/Text/Thinking/EOS ohne beobachtete harte Regression; einzelne +Laufpaare sind kein kontrollierter Speedup-Beweis. Der ausdrücklich genehmigte +Standalone-DS4-Benchmark benutzt originale DS4-Objekte ausschließlich außerhalb +des Produkts und reproduziert Low, separaten Bootstrap und laufende Chathistorie. +Die separate Canary-Prozessplatzierung war kein fairer Vergleich zum internen +Canary: gleichzeitig intern490,740ms versus extern höchstens3,331ms. Neue native +Metal-Zeitstempel lokalisieren eine300ms-Prefill-Probe fast vollständig vor dem +GPU-Start, nicht in der Host-Rückmeldung. Dieselbe Probe innerhalb des originalen +DS4-Prozesses reproduziert264,292ms Prefill-Spitze, ebenfalls vor dem GPU-Start; +771/771 Datensätze sind nach Korrektur der Watchdog-Ausgabeweiterleitung sauber +auswertbar. Die restliche Kosten-/Latenzdifferenz und weitere Abnahme bleiben in +[GLM follow-up](glm-scheduling-followup-20260911.md). P15 bleibt offen. + +**GLM-Ausführungsursache lokalisiert (11.09.):** Rust übernahm die GLM-5.2- +Top-k-Chunkgrenze für GLM5.3: nach9 Bootstrap-Tokens2039 statt der originalen2048 +Tokens. Nun bleiben GLM5.3-Chunks wie im aktiven DS4-Pfad vollständig; nur die +Attention-Slices werden an der Dense/Sparse-Grenze geteilt, auch für kreuzende +Verifier-Paare. Im unverändert vorgegebenen32-Schritt-Replay sind danach alle +Logit-Vektoren bitidentisch zur Originalreferenz, vorher lagen deutliche Fehler +bereits im ersten Vektor vor. Vollständige AR/MTP-Chats, Grenzfall-Livetest und +Performance-Neumessung bleiben erforderlich; dies allein schließt P15 nicht. + +**Nachprüfung desselben Tages:** Der vollständige AR-Chat stimmt nun in allen +drei Antworten samt Thinking, Tokenzahlen und Prompt-/Cache-Grenzen überein. +Bei MTP wurde zusätzlich ein im gemeinsamen UI-/Harness-Consumer behaltenes +Stop-Token gefunden: der zweite Prompt hatte3249 statt3248 Tokens. Der Consumer +nutzt nun GLMs vorhandenen Zwei-Zeilen-Rollback, analog zur Stop-Behandlung des +originalen DS4-Agenten. Danach stimmt auch der vollständige MTP-Chat in allen +drei Antworten und Grenzen überein (593/872/167 ausgegebene Tokens). Der +Attention-Grenzfall4095+2 wurde live geprüft. Der strikte32-Logit-Replay und die +großen Layer0-Tensorvergleiche sind bitidentisch. Einzelne neue Geschwindigkeits- +paare und diese konkrete Workload schließen nicht die gesamte2%-Matrix oder +alle Interaktionsfälle ab; P15 bleibt offen. Ergebnisse und Einschränkungen im +[GLM follow-up](glm-scheduling-followup-20260911.md). + **Aktiver GLM-Pfad und durchgehender Canary geprüft (10.09.):** `local-eval-results/glm-scheduling-canary-20260910.T1ABR1/{manifest.md,comparison.json}` enthält zwei vollständige AR/MTP-Chats bis EOS bei Power100/Low und Canary an. @@ -8061,6 +8129,61 @@ Priorität P1; nach P01, P04, P13; Source-Inventar vorhanden, Detailbeweis offen Priorität P1; nach den relevanten P01–P15; offen. +**Original-DS4-Vergleich erweitert (11.09.):** GLM-AR/MTP bewahren nach dem +CPU-Sampler-Abgleich in je zwei vollständigen Paaren Ausgabe/Thinking/Frontiers. +Die umgekehrte Reihenfolge zeigt dennoch AR-Decode-Rückstände bis15,23% und +MTP-Summary-Prefill bis4,38%; auch die Originalreferenz schwankt stark. Kein +Median-Pass, keine thermische Erklärung ohne Taktbeleg. Details und unveränderte +Command-Buffer-Zahlen im [GLM follow-up](glm-scheduling-followup-20260911.md). + +Der erste DeepSeek-Vergleich ist ausdrücklich ungültig: Der Referenztreiber +forderte2048 statt des UI-/DS4-Automatikwerts4096 an. Zusätzlich fügte der +gemeinsame Produkt-Renderer bei einem reinen Bootstrap-Cache fälschlich EOS vor +dem ersten User-Prompt ein (2742 statt2741 Tokens). Der Tokenizer selbst stimmt. +Ein gezielter CPU-Test reproduziert diesen UI-/Headless-Fehler vor der Korrektur; +nach dem Nichtleer-History-Guard bestehen DeepSeek- und GLM-Originalfixtures. +Referenztreiber auf Automatik angeglichen, keine Produkt-Chunkverkleinerung. +Das korrigierte vollständige DeepSeek-AR-Paar stimmt nun in allen Antworten, +Thinking-, Token- und Cache-Feldern überein, verfehlt aber im Einzelpaar noch +das Durchsatzziel. Das neue DSpark-Paar scheitert funktional: erste Antwort +bereits647 statt1208 Tokens bei identischem Prompt; anschließend jeweils ein +zusätzlich behaltenes EOS im Rust-Cache. Die normale Beendigung und bestandene +Python-Tests ändern diesen roten Vergleich nicht. Nächste zusammenhängende +DSpark-Einheit: erster divergenter Proposal-/Verify-Zyklus, Stop-/Capture-Frontier +und zusätzliche Vollvokabular-Readbacks gemeinsam gegen Original DS4 prüfen +und angleichen, dann neu messen. Konkrete Quellen, Rohdaten und Tests im +[DeepSeek follow-up](deepseek-reference-followup-20260911.md). + +**DSpark-Zyklus-/Zustandsabgleich fortgesetzt:** Der neue vollständige +Diagnoselauf stimmt jetzt in allen1755 Originalzyklen einschließlich Warmup, +Antworten, Thinking und Cachepositionen überein. Ursachen waren fehlende +HC-/Attention-Norm beim Support-KV-Aufbau, der abweichende Einzeldraft-Verifier, +Seed-/Teilannahme-Capture, Cachefenster-/Scheduler-/EOS-Übergänge und fehlende +Aktivierungsquantisierung des tatsächlich installierten Q8-Confidence-Kopfs. +Zusammenhängende Quellen-/Taskliste und Rot→Grün-Belege stehen im verlinkten +Follow-up. Early-Confidence-Gate und Verifier-Readbacks/Submission sind inzwischen +ebenfalls angeglichen; alle1755 Originalzyklen bleiben exakt gleich. Noch keine +Leistungsabnahme: saubere Wiederholungspaare zeigen starke zeitliche Drift in +beiden Implementierungen. Die CPU-Worker-Policy ist ebenfalls angeglichen: +dauerhafte12 Threads inklusive Aufrufer, weiterhin1755 identische Zyklen. +Ein CPU-only-Dispatchvergleich zeigt einen kleinen Gewinn, keine Erklärung für +den verbleibenden Gesamtgap. Verifier-/Frontier-Puffer werden inzwischen über +Zyklen wiederverwendet; der32-Zyklen-Test prüft Pufferidentität und bytegenaue +Rücknahme, der vollständige Chat weiterhin alle1755 Originalzyklen. Auch der +Exact-Sampling-Test besteht. Die jüngsten beiden Paare verfehlen jedoch weiter +Teile des Durchsatzziels; kein belegter Gesamtgewinn durch diese Änderung. + +**Messbedingung nach Nutzerklarstellung (11.09.):** Parallel laufen Videos und +beanspruchen einen Teil der GPU. Die genaue zeitliche Überlappung früherer +Messungen ist unbekannt. Die jüngsten seriellen Paare bleiben als Rohdaten +erhalten, gelten aber nicht als kontrollierte2%-Abnahme oder kausaler +Vorher/Nachher-Beleg. Weder alle Differenzen auf Videos schieben noch thermische +Drosselung behaupten. Keine weiteren Durchsatzserien während dieser Nutzung; +Funktionsprüfungen und Code-Abgleich bleiben möglich. Die vollständige +Sechsfeldmatrix, Responsivitätsabnahme und Zähler-Scope-Bereinigung bleiben offen. +Der Nutzer hat den Abschluss und Commit/Push dieses Zwischenstands freigegeben; +die vorgeschriebenen Commit-Gates umfassen auch den erneuten Bundle-Bau. + **Aktuelles interaktives Test-Bundle bereitgestellt (10.09.,22:32):** `make bundle` und strikte Codesign-Prüfung bestanden. Das Bundle enthält den zuletzt vermessenen Release-Code und identische Metal-Ressourcen; der genaue diff --git a/native/metal/ds4_canary.m b/native/metal/ds4_canary.m new file mode 100644 index 0000000..30fa8b8 --- /dev/null +++ b/native/metal/ds4_canary.m @@ -0,0 +1,66 @@ +#import +#import +#include +#include +#include "ds4_gpu.h" + +/* Shared optional UI/headless probe. Process placement matters: a separate + * queue in the model process is not a separate-process scheduling test. */ +static id g_canary_device; +static id g_canary_queue; +static id g_canary_buffer; + +static double ds4_monotonic_seconds(void) { + struct timespec time; + if (clock_gettime(CLOCK_MONOTONIC, &time) != 0) return 0.0; + return (double)time.tv_sec + (double)time.tv_nsec / 1000000000.0; +} + +int ds4_gpu_canary_probe(ds4_gpu_canary_sample *sample) { + if (!sample) return 0; + sample->scheduled_seconds = 0.0; + sample->completed_seconds = 0.0; + sample->gpu_wait_seconds = -1.0; + sample->gpu_interval_seconds = -1.0; + sample->host_return_seconds = -1.0; + @autoreleasepool { + if (!g_canary_device) g_canary_device = MTLCreateSystemDefaultDevice(); + if (!g_canary_queue && g_canary_device) { + g_canary_queue = [g_canary_device newCommandQueue]; + } + if (!g_canary_buffer && g_canary_device) { + g_canary_buffer = [g_canary_device newBufferWithLength:4096 + options:MTLResourceStorageModeShared]; + } + if (!g_canary_queue || !g_canary_buffer) return 0; + + id cb = [g_canary_queue commandBuffer]; + id blit = [cb blitCommandEncoder]; + if (!cb || !blit) return 0; + [blit fillBuffer:g_canary_buffer range:NSMakeRange(0, 4096) value:0]; + [blit endEncoding]; + mach_timebase_info_data_t timebase; + if (mach_timebase_info(&timebase) != KERN_SUCCESS || !timebase.denom) return 0; + const double scale = (double)timebase.numer / timebase.denom / 1e9; + const double started = ds4_monotonic_seconds(); + const double mach_started = (double)mach_absolute_time() * scale; + [cb commit]; + [cb waitUntilScheduled]; + sample->scheduled_seconds = ds4_monotonic_seconds() - started; + [cb waitUntilCompleted]; + const double mach_returned = (double)mach_absolute_time() * scale; + sample->completed_seconds = ds4_monotonic_seconds() - started; + // Metal GPU times use system mach time, unlike CLOCK_MONOTONIC on + // macOS. Read only after completion; unavailable timestamps stay -1. + const double gpu_start = cb.GPUStartTime; + const double gpu_end = cb.GPUEndTime; + if (gpu_start > 0.0 && gpu_start >= mach_started && + gpu_end >= gpu_start && mach_returned >= gpu_end) { + sample->gpu_wait_seconds = gpu_start - mach_started; + // Includes GPU scheduling/preemption, not exclusive busy time. + sample->gpu_interval_seconds = gpu_end - gpu_start; + sample->host_return_seconds = mach_returned - gpu_end; + } + return cb.status == MTLCommandBufferStatusCompleted; + } +} diff --git a/native/metal/ds4_gpu.h b/native/metal/ds4_gpu.h index d29834f..d412a3b 100644 --- a/native/metal/ds4_gpu.h +++ b/native/metal/ds4_gpu.h @@ -116,6 +116,9 @@ void ds4_gpu_busy_stats_get(ds4_gpu_busy_stats *stats); typedef struct { double scheduled_seconds; double completed_seconds; + double gpu_wait_seconds; + double gpu_interval_seconds; + double host_return_seconds; } ds4_gpu_canary_sample; int ds4_gpu_canary_probe(ds4_gpu_canary_sample *sample); #ifdef __APPLE__ diff --git a/native/metal/ds4_metal.m b/native/metal/ds4_metal.m index c8f9a9b..a421ec2 100644 --- a/native/metal/ds4_metal.m +++ b/native/metal/ds4_metal.m @@ -1201,48 +1201,6 @@ static void ds4_gpu_busy_stats_record(id cb, uint64_t epoch) { } } -/* A separate queue used by the headless evaluator's supervisor process. The - * tiny blit measures whether another process can still schedule UI-sized GPU - * work while the model owns its own Metal queue. */ -static id g_canary_device; -static id g_canary_queue; -static id g_canary_buffer; - -static double ds4_monotonic_seconds(void) { - struct timespec time; - if (clock_gettime(CLOCK_MONOTONIC, &time) != 0) return 0.0; - return (double)time.tv_sec + (double)time.tv_nsec / 1000000000.0; -} - -int ds4_gpu_canary_probe(ds4_gpu_canary_sample *sample) { - if (!sample) return 0; - sample->scheduled_seconds = 0.0; - sample->completed_seconds = 0.0; - @autoreleasepool { - if (!g_canary_device) g_canary_device = MTLCreateSystemDefaultDevice(); - if (!g_canary_queue && g_canary_device) { - g_canary_queue = [g_canary_device newCommandQueue]; - } - if (!g_canary_buffer && g_canary_device) { - g_canary_buffer = [g_canary_device newBufferWithLength:4096 - options:MTLResourceStorageModeShared]; - } - if (!g_canary_queue || !g_canary_buffer) return 0; - - id cb = [g_canary_queue commandBuffer]; - id blit = [cb blitCommandEncoder]; - if (!cb || !blit) return 0; - [blit fillBuffer:g_canary_buffer range:NSMakeRange(0, 4096) value:0]; - [blit endEncoding]; - const double started = ds4_monotonic_seconds(); - ds4_gpu_commit_command_buffer(cb); - [cb waitUntilScheduled]; - sample->scheduled_seconds = ds4_monotonic_seconds() - started; - [cb waitUntilCompleted]; - sample->completed_seconds = ds4_monotonic_seconds() - started; - return cb.status == MTLCommandBufferStatusCompleted; - } -} /* A failed command buffer can leave a cross-threadgroup arrival counter at an * arbitrary partial value. Drop cached ownership instead of CPU-resetting diff --git a/src/engine.rs b/src/engine.rs index 5b81cdf..0d13620 100644 --- a/src/engine.rs +++ b/src/engine.rs @@ -1,14 +1,19 @@ +#[cfg(any(target_os = "macos", test))] +mod ds4_sampling; mod gguf; #[cfg(any(target_os = "macos", test))] mod kvstore; #[cfg(target_os = "macos")] mod metal; mod qwen; +#[cfg(test)] +#[path = "../tools/sampler-replay-benchmark.rs"] +mod sampler_replay_benchmark; mod tokenizer; mod validation; #[cfg(target_os = "macos")] -use crate::metrics::{KvLookup, Metrics, SsdStats}; +use crate::metrics::{KvLookup, Metrics, PromptTiming, SsdStats}; use crate::model::{ModelChoice, validate_engine_artifacts}; #[cfg(target_os = "macos")] use crate::settings::TurnSettings; @@ -430,6 +435,8 @@ fn render_text_prompt( match messages.split_last() { Some((latest, history)) if latest.user + // A system-prefix cache is not a previous assistant turn. + && !history.is_empty() && checkpoint_tag == conversation_tag( &settings.system_prompt, @@ -476,6 +483,16 @@ pub(crate) struct ModelSummary { pub(crate) vision_loaded: bool, } +#[cfg(any(target_os = "macos", test))] +fn resume_prefill_min_tokens(model: ModelChoice) -> usize { + // DS4's GLM 5.3 crossover is measured on M5 Max; other DS4 graphs use four. + if model == ModelChoice::Glm53Flash { + 2 + } else { + 4 + } +} + impl Model { #[allow(dead_code)] pub(crate) fn open(settings: &EngineSettings) -> Result { @@ -1352,7 +1369,7 @@ impl Generator { if matches!(self.executor, metal::Executor::QwenMtplx(_)) || has_vision || (reused == 0 && tokens.len() > 1) - || suffix.len() >= 4 + || suffix.len() >= resume_prefill_min_tokens(self.executor.model().summary().model) { let context = self.executor.context(); self.executor.prefill(suffix, |used| { @@ -1531,7 +1548,18 @@ impl Generator { let mut generated = GeneratedText::new(settings.reasoning_mode); let prompt_tokens = tokens.len(); let suffix = &tokens[reused..]; + let prefill_started = Instant::now(); let completed = self.prefill_suffix(&tokens, reused, has_vision, cancelled, progress)?; + self.metrics.set_prompt_timing(Some(PromptTiming { + evaluated_tokens: completed, + eval_seconds: if suffix.is_empty() { + 0.0 + } else { + prefill_started.elapsed().as_secs_f64() + }, + mtp_history_seconds: None, + restore_seconds: None, + })); self.publish_execution_stats(); if completed != suffix.len() { return Ok(( @@ -1552,6 +1580,7 @@ impl Generator { let decode_metrics = Arc::clone(&self.metrics); let _decode_timer = decode_metrics.measure_decode(generation_started); let mut generated_tokens = 0_u32; + let trace_cycles = std::env::var_os("DS4_SPEC_CYCLE_TRACE").is_some(); let generation_limit = settings .max_generated_tokens .max(0) @@ -1638,6 +1667,14 @@ impl Generator { cancelled, )? }; + if trace_cycles { + eprintln!( + "{}", + serde_json::json!({"event":"spec_cycle", "prompt_tokens":prompt_tokens, + "generated":generated_tokens, "first":token, "accepted":cycle, + "position":self.executor.position()}) + ); + } self.publish_execution_stats(); for token in cycle { if generated_tokens >= generation_limit @@ -1646,6 +1683,18 @@ impl Generator { .model() .is_stop_token_for_reasoning(token, settings.reasoning_mode) { + // The consumer owns stop handling; each DS4 model restores + // its own state rather than retaining an extra EOS. + let keep = prompt_tokens + generated_tokens as usize; + match &mut self.executor { + metal::Executor::DeepSeek(executor) => { + executor.rewind_speculative_output(keep)? + } + metal::Executor::Glm(executor) => { + executor.rewind_speculative_output(keep)? + } + _ => {} + } generated.finish(&settings.stops, emit); return Ok(( GenerationOutput { @@ -1996,8 +2045,23 @@ fn sample( top_k: i32, rng: &mut Rng, ) -> i32 { - let probabilities = sampling_probabilities(logits, temperature, top_p, min_p, top_k); - sample_probabilities(&probabilities, rng, None) + ds4_sampling::sample(logits, temperature, top_p, min_p, top_k, rng) +} + +#[cfg(any(target_os = "macos", test))] +fn ds4_sample_argmax(logits: &[f32]) -> usize { + // Strict comparison preserves the first tie, ignores NaN and matches DS4's + // finite negative sentinel, including the all-invalid fallback to token0. + logits + .iter() + .enumerate() + .fold( + (0, -1.0e30_f32), + |best, (i, &v)| { + if v > best.1 { (i, v) } else { best } + }, + ) + .0 } #[cfg(any(target_os = "macos", test))] @@ -2125,16 +2189,7 @@ fn sampling_probabilities( min_p: f32, top_k: i32, ) -> Vec<(usize, f32)> { - let greedy = || { - vec![( - logits - .iter() - .enumerate() - .max_by(|a, b| a.1.total_cmp(b.1)) - .map_or(0, |(index, _)| index), - 1.0, - )] - }; + let greedy = || vec![(ds4_sample_argmax(logits), 1.0)]; if temperature <= 0.0 { return greedy(); } @@ -2363,7 +2418,7 @@ enum Rng { #[cfg(any(target_os = "macos", test))] impl Rng { fn new(seed: u64) -> Self { - Self::Ds4(seed.max(1)) + Self::Ds4(seed) } fn new_qwen(seed: u64) -> Self { @@ -2467,6 +2522,75 @@ impl Rng { mod sampling_tests { use super::*; + #[test] + #[ignore = "CPU-only; requires DS4_SAMPLER_RECORDING with the recorded GLM logits"] + fn ds4_sampler_real_logits_cpu_benchmark() { + let path = std::env::var_os("DS4_SAMPLER_RECORDING").unwrap(); + let mut rng = Rng::new(42); + let result = sampler_replay_benchmark::run(std::path::Path::new(&path), 154_880, |row| { + sample(row, 0.6, 0.95, 0.0, 0, &mut rng) + }) + .unwrap(); + assert_eq!(result["draws"], 512); + eprintln!("{result}"); + } + + #[test] + fn resumed_prefill_crossover_matches_each_ds4_graph() { + for (model, minimum) in [ + (ModelChoice::DeepSeekV4Flash0731, 4), + (ModelChoice::Glm52, 4), + (ModelChoice::Glm53Flash, 2), + ] { + let batched = (1..=4) + .filter(|&rows| rows >= resume_prefill_min_tokens(model)) + .collect::>(); + assert_eq!(batched, (minimum..=4).collect::>()); + } + } + + #[test] + fn ds4_sampler_matches_original_reference_tokens_and_rng() { + let fixture: serde_json::Value = + serde_json::from_str(include_str!("../tests/fixtures/ds4-sampling-ec7642c.json")) + .unwrap(); + let cases = fixture["cases"].as_array().unwrap(); + assert_eq!(cases.len(), 64); + let mut failures = Vec::new(); + for (index, case) in cases.iter().enumerate() { + let n = case["n"].as_u64().unwrap() as usize; + let logits: Vec = if n == 4 { + vec![2.0, 2.0, 0.0, -2.0] + } else { + (0..n) + .map(|i| ((i * 37 % 101) as f32 - 50.0) / 8.0) + .collect() + }; + let mut rng = Rng::new(case["seed"].as_u64().unwrap()); + let mut mismatches = 0; + for expected in case["tokens"].as_array().unwrap() { + let token = sample( + &logits, + case["temperature"].as_f64().unwrap() as f32, + case["top_p"].as_f64().unwrap() as f32, + case["min_p"].as_f64().unwrap() as f32, + case["top_k"].as_i64().unwrap() as i32, + &mut rng, + ); + mismatches += usize::from(i64::from(token) != expected.as_i64().unwrap()); + } + let Rng::Ds4(state) = rng else { unreachable!() }; + let rng_matches = state == case["rng_after"].as_u64().unwrap(); + if mismatches != 0 || !rng_matches { + failures.push((index, mismatches, rng_matches)); + } + } + assert!( + failures.is_empty(), + "(case, token differences, RNG matches): {failures:?}" + ); + } + #[test] fn qwen_rng_matches_numpy_pcg64() { let mut rng = Rng::new_qwen(12_345); diff --git a/src/engine/ds4_sampling.rs b/src/engine/ds4_sampling.rs new file mode 100644 index 0000000..754e327 --- /dev/null +++ b/src/engine/ds4_sampling.rs @@ -0,0 +1,345 @@ +//! DS4's sampled-token path. Distribution materialization for speculative +//! correction and MTPLX sampling are separate algorithms, not substitutes. +use super::{Rng, ds4_sample_argmax}; +use std::cmp::{Ordering, Reverse}; +use std::collections::BinaryHeap; + +#[derive(Clone, Copy)] +struct Candidate { + id: usize, + logit: f32, + probability: f32, +} + +impl PartialEq for Candidate { + fn eq(&self, other: &Self) -> bool { + self.logit == other.logit && self.id == other.id + } +} +impl Eq for Candidate {} +impl PartialOrd for Candidate { + fn partial_cmp(&self, other: &Self) -> Option { + Some(self.cmp(other)) + } +} +impl Ord for Candidate { + fn cmp(&self, other: &Self) -> Ordering { + // Candidates contain finite logits only. C treats signed zeros as + // equal and breaks logit ties by the earlier vocabulary index. + self.logit + .partial_cmp(&other.logit) + .unwrap() + .then_with(|| other.id.cmp(&self.id)) + } +} + +fn ranked(heap: BinaryHeap>) -> Vec { + let mut candidates = heap.into_vec().into_iter().map(|c| c.0).collect::>(); + candidates.sort_unstable_by(|a, b| b.cmp(a)); + candidates +} + +fn draw(candidates: &[Candidate], sum: f32, rng: &mut Rng) -> i32 { + let mut remaining = rng.unit() * sum; + for candidate in candidates { + remaining -= candidate.probability; + if remaining <= 0.0 { + return candidate.id as i32; + } + } + candidates.last().unwrap().id as i32 +} + +fn filtered(candidates: &[Candidate], total: f32, top_p: f32, min_p: f32) -> (usize, f32, bool) { + let minimum = (candidates[0].probability / total) * min_p; + let mut sum = 0.0; + let mut count = 0; + let mut stopped_by_min_p = false; + for (i, candidate) in candidates.iter().enumerate() { + if i > 0 && candidate.probability / total < minimum { + stopped_by_min_p = true; + break; + } + sum += candidate.probability; + count += 1; + if sum / total >= top_p { + break; + } + } + (count, sum, stopped_by_min_p) +} + +#[allow(clippy::too_many_arguments)] +fn fast_top_p( + logits: &[f32], + finite: usize, + maximum: f32, + best: i32, + temperature: f32, + top_p: f32, + min_p: f32, + rng: &mut Rng, +) -> Option { + const CAP: usize = 512; + if finite > CAP && top_p >= 0.999 { + return None; + } + let capacity = finite.min(CAP); + let mut heap = BinaryHeap::>::with_capacity(capacity); + let mut total = 0.0; + let mut heap_sum = 0.0; + for (id, &logit) in logits.iter().enumerate() { + if !logit.is_finite() { + continue; + } + let probability = ((logit - maximum) / temperature).exp(); + total += probability; + let candidate = Candidate { + id, + logit, + probability, + }; + if heap.len() < capacity { + heap_sum += probability; + heap.push(Reverse(candidate)); + } else if candidate > heap.peek().unwrap().0 { + let mut worst = heap.peek_mut().unwrap(); + heap_sum -= worst.0.probability; + worst.0 = candidate; + heap_sum += probability; + } + } + if total <= 0.0 || !total.is_finite() { + return Some(best); + } + if heap.len() < finite && heap_sum < top_p * total { + return None; + } + let candidates = ranked(heap); + let min_p = if min_p > 0.0 { min_p } else { 0.0 }; + let (count, sum, stopped_by_min_p) = filtered(&candidates, total, top_p, min_p); + // DS4 falls back if min-p stopped inside the heap but the unseen tail + // might still pass its raw threshold. Do not consume RNG on fallback. + if candidates.len() < finite + && stopped_by_min_p + && min_p > 0.0 + && candidates.last().unwrap().probability >= candidates[0].probability * min_p + { + return None; + } + Some(if count == 0 { + best + } else { + draw(&candidates[..count], sum, rng) + }) +} + +pub(super) fn sample( + logits: &[f32], + temperature: f32, + top_p: f32, + min_p: f32, + top_k: i32, + rng: &mut Rng, +) -> i32 { + if temperature <= 0.0 || logits.is_empty() { + return ds4_sample_argmax(logits) as i32; + } + let top_p = if top_p <= 0.0 || top_p > 1.0 { + 1.0 + } else { + top_p + }; + let min_p = if min_p < 0.0 { 0.0 } else { min_p }; + if top_k > 0 { + let capacity = (top_k as usize).min(1024).min(logits.len()); + let mut heap = BinaryHeap::>::with_capacity(capacity); + for (id, &logit) in logits.iter().enumerate() { + if !logit.is_finite() { + continue; + } + let candidate = Candidate { + id, + logit, + probability: 0.0, + }; + if heap.len() < capacity { + heap.push(Reverse(candidate)); + } else if candidate > heap.peek().unwrap().0 { + heap.peek_mut().unwrap().0 = candidate; + } + } + if heap.is_empty() { + return ds4_sample_argmax(logits) as i32; + } + let mut candidates = ranked(heap); + let maximum = candidates[0].logit; + let mut total = 0.0; + for candidate in &mut candidates { + candidate.probability = ((candidate.logit - maximum) / temperature).exp(); + total += candidate.probability; + } + if total <= 0.0 || !total.is_finite() { + return candidates[0].id as i32; + } + let (count, sum, _) = filtered(&candidates, total, top_p, min_p); + return if count == 0 { + candidates[0].id as i32 + } else { + draw(&candidates[..count], sum, rng) + }; + } + let mut maximum = -1.0e30_f32; + let mut best = 0; + let mut finite = 0; + for (id, &logit) in logits.iter().enumerate() { + if !logit.is_finite() { + continue; + } + finite += 1; + if logit > maximum { + maximum = logit; + best = id as i32; + } + } + if finite == 0 { + return ds4_sample_argmax(logits) as i32; + } + if top_p < 1.0 + && let Some(token) = fast_top_p( + logits, + finite, + maximum, + best, + temperature, + top_p, + min_p, + rng, + ) + { + return token; + } + let mut candidates = Vec::with_capacity(finite); + let mut total = 0.0; + if top_p >= 1.0 { + let min_relative = if min_p > 0.0 { min_p } else { 0.0 }; + if min_relative > 1.0 { + return best; + } + // Same conservative expf-verified rejection boundary as DS4: skip + // exponentials only when the value is guaranteed to fail min-p. + let mut reject_scaled = None; + if min_relative > 0.0 && min_relative.is_finite() { + let mut cutoff = min_relative.ln(); + for _ in 0..8 { + if !cutoff.is_finite() { + break; + } + cutoff = cutoff.next_down(); + if cutoff.exp() < min_relative { + reject_scaled = Some(cutoff); + break; + } + } + } + for (id, &logit) in logits.iter().enumerate() { + if !logit.is_finite() { + continue; + } + let scaled = (logit - maximum) / temperature; + if reject_scaled.is_some_and(|cutoff| scaled <= cutoff) { + continue; + } + let probability = scaled.exp(); + if probability < min_relative { + continue; + } + total += probability; + candidates.push(Candidate { + id, + logit, + probability, + }); + } + if total <= 0.0 || !total.is_finite() { + return best; + } + let mut remaining = rng.unit() * total; + for candidate in candidates { + remaining -= candidate.probability; + if remaining <= 0.0 { + return candidate.id as i32; + } + } + return best; + } + for (id, &logit) in logits.iter().enumerate() { + if !logit.is_finite() { + continue; + } + let probability = ((logit - maximum) / temperature).exp(); + total += probability; + candidates.push(Candidate { + id, + logit, + probability, + }); + } + if total <= 0.0 || !total.is_finite() { + return best; + } + if min_p > 0.0 && min_p <= 1.0 { + let minimum = (1.0 / total) * min_p; + candidates.retain(|c| c.probability / total >= minimum); + } + if candidates.is_empty() { + return best; + } + candidates.sort_unstable_by(|a, b| b.cmp(a)); + let min_p = if min_p > 0.0 { min_p } else { 0.0 }; + let (count, sum, _) = filtered(&candidates, total, top_p, min_p); + if count == 0 { + best + } else { + draw(&candidates[..count], sum, rng) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn heap_fallback_preserves_rng_and_signed_zero_ties() { + let flat = vec![0.0; 1024]; + let mut rng = Rng::new(42); + for top_p in [0.95, 0.999] { + assert_eq!( + fast_top_p(&flat, 1024, 0.0, 0, 0.6, top_p, 0.0, &mut rng), + None + ); + } + assert_eq!(rng.unit(), Rng::new(42).unit()); + let mut concentrated = vec![-100.0; 1024]; + concentrated[100] = 0.0; + assert_eq!( + fast_top_p(&concentrated, 1024, 0.0, 100, 0.6, 0.95, 0.0, &mut rng), + Some(100) + ); + for seed in 0..16 { + for top_k in [0, 1] { + assert_eq!( + sample( + &[-0.0, 0.0, -0.0, 0.0], + 1.0, + 0.95, + 1.1, + top_k, + &mut Rng::new(seed) + ), + 0 + ); + } + } + } +} diff --git a/src/engine/gguf.rs b/src/engine/gguf.rs index 3eadaeb..a13e13f 100644 --- a/src/engine/gguf.rs +++ b/src/engine/gguf.rs @@ -3,6 +3,7 @@ use sha2::{Digest, Sha256}; use std::collections::HashMap; use std::fs::File; use std::path::{Path, PathBuf}; +use std::sync::Arc; const MAGIC: u32 = 0x4655_4747; const MAX_COUNT: u64 = 1_000_000; @@ -48,7 +49,7 @@ pub(super) struct Tensor { pub(super) struct Gguf { path: PathBuf, - map: Mmap, + map: Arc, data_offset: u64, max_tensor_bytes: u64, pub(super) metadata: HashMap, @@ -169,7 +170,7 @@ impl Gguf { Ok(Self { path: path.to_owned(), - map, + map: Arc::new(map), data_offset: data_start, max_tensor_bytes, metadata, @@ -215,6 +216,10 @@ impl Gguf { self.map.as_ptr() } + pub(super) fn shared_map(&self) -> Arc { + Arc::clone(&self.map) + } + pub(super) fn data_offset(&self) -> u64 { self.data_offset } diff --git a/src/engine/metal.rs b/src/engine/metal.rs index 71cadf6..7865624 100644 --- a/src/engine/metal.rs +++ b/src/engine/metal.rs @@ -2,6 +2,7 @@ mod checkpoint; mod glm; mod gpu; mod hotlist; +mod markov; mod profile; mod qwen; #[cfg(test)] @@ -57,6 +58,9 @@ pub(crate) struct GpuBusySnapshot { pub(crate) struct GpuCanarySnapshot { pub(crate) scheduled_seconds: f64, pub(crate) completed_seconds: f64, + pub(crate) gpu_wait_seconds: Option, + pub(crate) gpu_interval_seconds: Option, + pub(crate) host_return_seconds: Option, } pub(crate) fn reset_gpu_busy_stats() { @@ -86,6 +90,11 @@ pub(crate) fn gpu_canary_probe() -> Result { Ok(GpuCanarySnapshot { scheduled_seconds: sample.scheduled_seconds, completed_seconds: sample.completed_seconds, + gpu_wait_seconds: (sample.gpu_wait_seconds >= 0.0).then_some(sample.gpu_wait_seconds), + gpu_interval_seconds: (sample.gpu_interval_seconds >= 0.0) + .then_some(sample.gpu_interval_seconds), + host_return_seconds: (sample.host_return_seconds >= 0.0) + .then_some(sample.host_return_seconds), }) } @@ -366,6 +375,7 @@ impl DsparkStageWeights { struct Dspark { config: DsparkConfig, + prefill_cap: u32, weights: Vec, mean_weights: Buffer, mean_rows: Buffer, @@ -383,6 +393,9 @@ struct Dspark { capture_mask: u32, cache_start: u32, cache_len: u32, + cache_token_start: u32, + pending_seed: Option<(u32, u32, BatchScratch)>, + spare_batches: Vec, strict: bool, drafted: u64, accepted: u64, @@ -393,11 +406,12 @@ struct Dspark { scheduler_lifetime_accepted: u32, scheduler_long_accept_seen: bool, last_confidence: Option, + markov: Option, } -fn dspark_scheduler_pause(cycles: u32, accepted: u32, no_draft: u32) -> u32 { +fn dspark_scheduler_pause(cycles: u32, accepted: u32, no_draft: u32, current_skip: u32) -> u32 { if cycles == 0 { - return 0; + return current_skip; } let low_acceptance = u64::from(accepted) * 1_000 < u64::from(cycles) * 1_500; let many_no_draft = no_draft * 2 >= cycles; @@ -406,7 +420,7 @@ fn dspark_scheduler_pause(cycles: u32, accepted: u32, no_draft: u32) -> u32 { } else if low_acceptance { 2 } else { - 0 + current_skip } } @@ -457,12 +471,16 @@ impl Dspark { scratch: BatchScratch::allocate(model, session.context, rows, false)?, logits: Buffer::floats(u64::from(config.block_size) * shape.vocab)?, config, + prefill_cap: session.prefill_cap, weights, mean_weights, mean_rows, capture_mask: 0, cache_start: 0, cache_len: 0, + cache_token_start: 0, + pending_seed: None, + spare_batches: Vec::new(), strict: settings.dspark_strict || quality, drafted: 0, accepted: 0, @@ -473,6 +491,7 @@ impl Dspark { scheduler_lifetime_accepted: 0, scheduler_long_accept_seen: false, last_confidence: None, + markov: None, }) } @@ -484,6 +503,16 @@ impl Dspark { true } + fn begin_request(&mut self) { + self.scheduler_cycles = 0; + self.scheduler_accepted = 0; + self.scheduler_no_draft = 0; + self.scheduler_skip = 0; + self.scheduler_lifetime_accepted = 0; + self.scheduler_long_accept_seen = false; + self.last_confidence = None; + } + fn scheduler_note(&mut self, accepted: u32, no_draft: bool) { self.scheduler_cycles += 1; self.scheduler_accepted = self.scheduler_accepted.saturating_add(accepted); @@ -506,11 +535,12 @@ impl Dspark { self.scheduler_skip = self.scheduler_skip.max(skip); } if self.scheduler_cycles >= 4 { - self.scheduler_skip = self.scheduler_skip.max(dspark_scheduler_pause( + self.scheduler_skip = dspark_scheduler_pause( self.scheduler_cycles, self.scheduler_accepted, self.scheduler_no_draft, - )); + self.scheduler_skip, + ); self.scheduler_cycles = 0; self.scheduler_accepted = 0; self.scheduler_no_draft = 0; @@ -529,6 +559,38 @@ impl Dspark { self.capture_mask = 0; } + fn reset_cache(&mut self) { + self.capture_mask = 0; + self.release_pending_seed(); + self.cache_start = 0; + self.cache_token_start = 0; + self.cache_len = 0; + } + + fn recycle_batch(&mut self, batch: BatchScratch) { + if batch.output_logits.is_some() { + // One workspace per verifier shape; never retain a full prefill + // batch merely because it was used to seed the support cache. + self.spare_batches.retain(|saved| saved.rows != batch.rows); + self.spare_batches.push(batch); + } + } + + fn release_pending_seed(&mut self) { + if let Some((_, _, batch)) = self.pending_seed.take() { + self.recycle_batch(batch); + } + } + + fn take_verifier_batch(&mut self, rows: u32, pos: u32) -> Option { + let rows = batch_workspace_rows(rows, true); + let index = self + .spare_batches + .iter() + .position(|batch| batch.rows == rows && batch.position_capacity >= pos)?; + Some(self.spare_batches.swap_remove(index)) + } + fn commit_proposed_prefix(&mut self, rows: u32, raw_cap: u32) { let added = rows.min(raw_cap); let total = self.cache_len.saturating_add(added); @@ -571,14 +633,25 @@ impl Dspark { layer: u32, hc: &Buffer, rows: u32, + prepend_seed: bool, prefill_cap: u32, shape: super::Shape, ) -> Result<(), String> { let Some(slot) = self.target_slot(layer) else { return Ok(()); }; + let slot_offset = u64::from(slot) * u64::from(prefill_cap) * shape.embd * 4; + if prepend_seed { + self.target_hidden_batch.copy_from( + slot_offset, + &self.target_hidden, + u64::from(slot) * shape.embd * 4, + shape.embd * 4, + "capturing the DSpark verifier seed row", + )?; + } let batch = self.target_hidden_batch.view( - u64::from(slot) * u64::from(prefill_cap) * shape.embd * 4, + slot_offset + u64::from(prepend_seed) * shape.embd * 4, u64::from(rows) * shape.embd * 4, )?; call( @@ -611,9 +684,9 @@ impl Dspark { fn seed_batch_cache( &mut self, support: &Gguf, + scratch: &BatchScratch, pos: u32, rows: u32, - prefill_cap: u32, raw_cap: u32, shape: super::Shape, ) -> Result<(), String> { @@ -628,15 +701,23 @@ impl Dspark { let map = support.map_ptr().cast(); let size = support.len(); let input = self.config.target_layers.len() as u64 * shape.embd; - let projected = self - .target_hidden_batch - .view(0, u64::from(rows) * shape.embd * 4)?; - let norm = self - .packed_target_hidden - .view(0, u64::from(rows) * shape.embd * 4)?; - let kv_bytes = u64::from(rows) * shape.head_dim * 4; - let kv_raw = self.target_hidden_batch.view(0, kv_bytes)?; - let kv = self.target_hidden_batch.view(kv_bytes, kv_bytes)?; + // DS4 seeds target rows through each support stage's HC/attention norm. + // Reuse the completed target batch's scratch, preserving captured hidden + // rows for partial verifier commits and current_hc for its output head. + let hc_dim = shape.hc * shape.embd; + let mix_hc = 2 * shape.hc + shape.hc * shape.hc; + let projected = scratch.current.view(0, u64::from(rows) * shape.embd * 4)?; + let norm = scratch.norm.view(0, u64::from(rows) * shape.embd * 4)?; + let target_hc = scratch + .after_attention_hc + .view(0, u64::from(rows) * hc_dim * 4)?; + let flat_hc = scratch.flat_hc.view(0, u64::from(rows) * hc_dim * 4)?; + let hc_mix = scratch.hc_mix.view(0, u64::from(rows) * mix_hc * 4)?; + let hc_split = scratch.hc_split.view(0, u64::from(rows) * mix_hc * 4)?; + let kv_raw = scratch + .kv_raw + .view(0, u64::from(rows) * shape.head_dim * 4)?; + let kv = scratch.kv.view(0, u64::from(rows) * shape.head_dim * 4)?; let commands = Commands::begin()?; call( unsafe { @@ -646,7 +727,7 @@ impl Dspark { rows, shape.embd as u32, self.config.target_layers.len() as u32, - prefill_cap, + self.prefill_cap, ) }, "packing DSpark target hidden states", @@ -676,7 +757,64 @@ impl Dspark { }, "normalizing DSpark target hidden states", )?; + call( + unsafe { + ds4_gpu_repeat_hc_rows_tensor( + target_hc.raw(), + norm.raw(), + rows, + shape.embd as u32, + shape.hc as u32, + ) + }, + "expanding DSpark target HC rows", + )?; for (stage, cache) in self.weights.iter().zip(&self.raw_caches) { + let block = &stage.block; + call( + unsafe { + ds4_gpu_rms_norm_plain_rows_tensor( + flat_hc.raw(), + target_hc.raw(), + hc_dim as u32, + rows, + shape.rms_epsilon, + ) + }, + "normalizing DSpark target attention HC rows", + )?; + matmul_rows( + &hc_mix, + block.hc_attn_fn, + hc_dim, + mix_hc, + &flat_hc, + rows, + map, + size, + )?; + call( + unsafe { + ds4_gpu_hc_split_weighted_sum_norm_tensor( + projected.raw(), + norm.raw(), + hc_split.raw(), + hc_mix.raw(), + target_hc.raw(), + map, + size, + block.hc_attn_scale.offset, + block.hc_attn_base.offset, + block.attn_norm.offset, + shape.embd as u32, + shape.hc as u32, + shape.hc_sinkhorn as u32, + shape.hc_epsilon, + shape.rms_epsilon, + ) + }, + "mixing and normalizing DSpark target cache rows", + )?; matmul_rows( &kv_raw, stage.block.attn_kv, @@ -751,6 +889,7 @@ impl Dspark { commands.finish()?; self.cache_start = pos % raw_cap; self.cache_len = rows; + self.cache_token_start = pos; Ok(()) } @@ -761,6 +900,10 @@ impl Dspark { raw_cap: u32, shape: super::Shape, ) -> Result<(), String> { + // DS4 ring maintenance never creates a new window or bridges a gap. + if self.cache_len == 0 || self.cache_token_start + self.cache_len != pos { + return Ok(()); + } if !self.capture_complete() { return Err("DSpark target-layer capture is incomplete".into()); } @@ -868,13 +1011,10 @@ impl Dspark { )?; } commands.finish()?; - let append = (self.cache_start + self.cache_len) % raw_cap; - if self.cache_len == 0 || append != pos % raw_cap { - self.cache_start = pos % raw_cap; - self.cache_len = 1; - } else if self.cache_len < raw_cap { + if self.cache_len < raw_cap { self.cache_len += 1; } else { + self.cache_token_start += 1; self.cache_start = (self.cache_start + 1) % raw_cap; } Ok(()) @@ -902,10 +1042,25 @@ impl Dspark { return Err("DSpark block exceeds its raw-cache capacity".into()); } let max_support = raw_cap - rows; + let mut seeded = false; + if let Some((start, count, scratch)) = self.pending_seed.take() { + if start + count == pos { + self.seed_batch_cache(support, &scratch, start, count, raw_cap, shape)?; + seeded = true; + } + self.recycle_batch(scratch); + } + if !seeded && self.cache_len != 0 { + if pos <= self.cache_token_start || pos > self.cache_token_start + self.cache_len { + self.cache_start = 0; + self.cache_token_start = 0; + self.cache_len = 0; + } else { + self.cache_len = pos - self.cache_token_start; + } + } if self.cache_len > max_support { - let discard = self.cache_len - max_support; - self.cache_start = (self.cache_start + discard) % raw_cap; - self.cache_len = max_support; + return Ok(Vec::new()); } let map = support.map_ptr().cast(); let size = support.len(); @@ -1212,7 +1367,12 @@ impl Dspark { }, "quantizing DSpark KV rows", )?; - let append = (self.cache_start + self.cache_len) % raw_cap; + let raw_start = if self.cache_len == 0 { + pos % raw_cap + } else { + self.cache_start + }; + let append = (raw_start + self.cache_len) % raw_cap; call( unsafe { ds4_gpu_store_raw_kv_batch_tensor( @@ -1238,7 +1398,7 @@ impl Dspark { draft, self.cache_len + rows, raw_cap, - self.cache_start, + raw_start, shape.heads as u32, shape.head_dim as u32, ) @@ -1605,50 +1765,68 @@ impl Dspark { }, "normalizing DSpark output rows", )?; - matmul_rows( - &self.logits, - base_weights.output, - shape.embd, - shape.vocab, - &output_norm, - draft, - base.main.map_ptr().cast(), - base.main.len(), - )?; - commands.finish()?; - let confidence = final_stage .confidence .ok_or("DSpark confidence head is missing")?; - let mut logits = vec![0.0; (u64::from(draft) * shape.vocab) as usize]; - let mut hidden = vec![0.0; (u64::from(draft) * shape.embd) as usize]; - self.logits.read_f32(&mut logits)?; - output_norm.read_f32(&mut hidden)?; + let mut commands = Some(commands); + if confidence_threshold > 0.0 { + // DS4 checks confidence before spending GPU work on the vocab head. + commands + .take() + .expect("DSpark hidden commands are active") + .finish()?; + } + let mut logits = vec![0.0; shape.vocab as usize]; + let mut features = vec![0.0; (shape.embd + markov_w1.dims[0]) as usize]; let mut proposals = Vec::with_capacity(draft as usize); let mut previous = first_token as u32; self.last_confidence = None; for row in 0..draft as usize { let state = dense_row(support, markov_w1, previous)?; - let mut features = Vec::with_capacity(shape.embd as usize + state.len()); - features.extend_from_slice( - &hidden[row * shape.embd as usize..(row + 1) * shape.embd as usize], - ); - features.extend_from_slice(&state); - let confidence_logit = dense_dot(support, confidence, 0, &features)?; - let confidence_value = if confidence_logit >= 0.0 { - 1.0 / (1.0 + (-confidence_logit).exp()) - } else { - let value = confidence_logit.exp(); - value / (1.0 + value) - }; + if confidence_threshold > 0.0 { + output_norm + .view(row as u64 * shape.embd * 4, shape.embd * 4)? + .read_f32(&mut features[..shape.embd as usize])?; + features[shape.embd as usize..].copy_from_slice(&state); + let confidence_logit = dense_dot(support, confidence, 0, &features)?; + let confidence_value = if confidence_logit >= 0.0 { + 1.0 / (1.0 + (-confidence_logit).exp()) + } else { + let value = confidence_logit.exp(); + value / (1.0 + value) + }; + if row == 0 { + self.last_confidence = Some(confidence_logit); + } + if confidence_value < confidence_threshold { + break; + } + } if row == 0 { - self.last_confidence = Some(confidence_logit); + let commands = commands.take().map_or_else(Commands::begin, Ok)?; + matmul_rows( + &self.logits, + base_weights.output, + shape.embd, + shape.vocab, + &output_norm, + draft, + base.main.map_ptr().cast(), + base.main.len(), + )?; + commands.finish()?; } - if confidence_threshold > 0.0 && confidence_value < confidence_threshold { - break; + self.logits + .view(row as u64 * shape.vocab * 4, shape.vocab * 4)? + .read_f32(&mut logits)?; + if self.markov.is_none() { + self.markov = Some(markov::MarkovPool::new(support, markov_w2)?); } - let row_logits = &logits[row * shape.vocab as usize..(row + 1) * shape.vocab as usize]; - let best = dense_argmax(support, markov_w2, &state, row_logits)?; + let best = self + .markov + .as_mut() + .expect("Markov pool initialized") + .argmax(&state, &mut logits)?; proposals.push(best); previous = best as u32; } @@ -2673,6 +2851,8 @@ struct Scratch { } struct BatchScratch { + rows: u32, + position_capacity: u32, tokens: Buffer, current_hc: Buffer, next_hc: Buffer, @@ -2716,14 +2896,19 @@ struct BatchScratch { output_logits: Option, } +fn batch_workspace_rows(rows: u32, output_logits: bool) -> u32 { + if output_logits && rows < 8 { + // Preserve both the existing head padding and preceding seed-row room. + if rows == 1 { 2 } else { 8 } + } else { + rows + } +} + impl BatchScratch { fn allocate(model: &Model, pos: u32, rows: u32, output_logits: bool) -> Result { let shape = model.shape; - let output_rows = if output_logits && rows > 1 && rows < 8 { - 8 - } else { - rows - }; + let output_rows = batch_workspace_rows(rows, output_logits); let rows = u64::from(output_rows); let hc_dim = shape.hc * shape.embd; let mix_hc = 2 * shape.hc + shape.hc * shape.hc; @@ -2732,6 +2917,8 @@ impl BatchScratch { let low_dim = shape.out_groups * shape.lora_o; let routed = shape.experts_used * shape.ff_expert; Ok(Self { + rows: output_rows, + position_capacity: pos, tokens: Buffer::bytes(rows * 4)?, current_hc: Buffer::floats(rows * hc_dim)?, next_hc: Buffer::floats(rows * hc_dim)?, @@ -2882,6 +3069,7 @@ struct SpecFrontier { dspark_capture_mask: u32, dspark_cache_start: u32, dspark_cache_len: u32, + dspark_cache_token_start: u32, } struct SpecPrefixFrontier { @@ -2898,17 +3086,25 @@ fn capture_compression_frontier( state: &CompressionState, bytes: u64, purpose: &str, + reusable: Option, ) -> Result { - let state_kv = Buffer::bytes(bytes)?; - let state_score = Buffer::bytes(bytes)?; - state_kv.copy_from(0, &state.state_kv, 0, bytes, purpose)?; - state_score.copy_from(0, &state.state_score, 0, bytes, purpose)?; - Ok(CompressionFrontier { - state_kv, - state_score, - bytes, - rows: state.rows, - }) + let mut saved = match reusable.filter(|saved| saved.bytes == bytes) { + Some(saved) => saved, + None => CompressionFrontier { + state_kv: Buffer::bytes(bytes)?, + state_score: Buffer::bytes(bytes)?, + bytes, + rows: 0, + }, + }; + saved + .state_kv + .copy_from(0, &state.state_kv, 0, bytes, purpose)?; + saved + .state_score + .copy_from(0, &state.state_score, 0, bytes, purpose)?; + saved.rows = state.rows; + Ok(saved) } fn restore_compression_frontier( @@ -3047,6 +3243,8 @@ pub(super) struct DeepSeekExecutor { speculative_cycles: u64, verifier_passes: u64, verifier_ns: u64, + spare_frontier: Option, + spare_prefixes: Vec, checkpoint_tag: [u8; 32], model_modified: (u64, u32), model_identity: [u8; 32], @@ -3190,6 +3388,8 @@ impl DeepSeekExecutor { speculative_cycles: 0, verifier_passes: 0, verifier_ns: 0, + spare_frontier: None, + spare_prefixes: Vec::new(), checkpoint_tag: [0; 32], model_modified, model_identity, @@ -3268,70 +3468,73 @@ impl DeepSeekExecutor { Ok(()) } - fn snapshot_spec_frontier(&self) -> Result { + fn snapshot_spec_frontier(&mut self) -> Result { let shape = self.model.shape; - let commands = Commands::begin()?; - let layers = self - .session - .layers - .iter() - .map(|layer| { - let compression = layer - .compression - .as_ref() - .map(|state| { - let coefficient = if state.ratio == 4 { 2 } else { 1 }; - capture_compression_frontier( - state, - coefficient * coefficient * state.ratio as u64 * shape.head_dim * 4, - "saving speculative compressor state", - ) - }) - .transpose()?; - let indexer = layer - .indexer - .as_ref() - .map(|state| { - capture_compression_frontier( - state, - 4 * state.ratio as u64 * shape.indexer_head_dim * 4, - "saving speculative indexer state", - ) - }) - .transpose()?; - Ok::<_, String>(LayerFrontier { - compression, - indexer, + let mut saved = self.spare_frontier.take().unwrap_or_else(|| SpecFrontier { + layers: (0..shape.layers) + .map(|_| LayerFrontier { + compression: None, + indexer: None, }) - }) - .collect::, _>>()?; - let dspark_target_hidden = self - .dspark - .as_ref() - .map(|dspark| { - let bytes = dspark.config.target_layers.len() as u64 * shape.embd * 4; - let saved = Buffer::bytes(bytes)?; - saved.copy_from( + .collect(), + position: 0, + token_len: 0, + logits: Vec::new(), + dspark_target_hidden: None, + dspark_capture_mask: 0, + dspark_cache_start: 0, + dspark_cache_len: 0, + dspark_cache_token_start: 0, + }); + let commands = Commands::begin()?; + for (layer, saved) in self.session.layers.iter().zip(&mut saved.layers) { + if let Some(state) = &layer.compression { + let coefficient = if state.ratio == 4 { 2 } else { 1 }; + saved.compression = Some(capture_compression_frontier( + state, + coefficient * coefficient * state.ratio as u64 * shape.head_dim * 4, + "saving speculative compressor state", + saved.compression.take(), + )?); + } + if let Some(state) = &layer.indexer { + saved.indexer = Some(capture_compression_frontier( + state, + 4 * state.ratio as u64 * shape.indexer_head_dim * 4, + "saving speculative indexer state", + saved.indexer.take(), + )?); + } + } + if let Some(dspark) = &self.dspark { + let bytes = dspark.config.target_layers.len() as u64 * shape.embd * 4; + if saved.dspark_target_hidden.is_none() { + saved.dspark_target_hidden = Some(Buffer::bytes(bytes)?); + } + saved + .dspark_target_hidden + .as_ref() + .expect("snapshot allocated") + .copy_from( 0, &dspark.target_hidden, 0, bytes, "saving speculative DSpark target state", )?; - Ok::<_, String>(saved) - }) - .transpose()?; + } commands.finish()?; - Ok(SpecFrontier { - layers, - position: self.session.position, - token_len: self.tokens.len(), - logits: self.logits.clone(), - dspark_target_hidden, - dspark_capture_mask: self.dspark.as_ref().map_or(0, |value| value.capture_mask), - dspark_cache_start: self.dspark.as_ref().map_or(0, |value| value.cache_start), - dspark_cache_len: self.dspark.as_ref().map_or(0, |value| value.cache_len), - }) + saved.position = self.session.position; + saved.token_len = self.tokens.len(); + saved.logits.clone_from(&self.logits); + saved.dspark_capture_mask = self.dspark.as_ref().map_or(0, |value| value.capture_mask); + saved.dspark_cache_start = self.dspark.as_ref().map_or(0, |value| value.cache_start); + saved.dspark_cache_len = self.dspark.as_ref().map_or(0, |value| value.cache_len); + saved.dspark_cache_token_start = self + .dspark + .as_ref() + .map_or(0, |value| value.cache_token_start); + Ok(saved) } fn restore_spec_frontier(&mut self, frontier: &SpecFrontier) -> Result<(), String> { @@ -3371,6 +3574,8 @@ impl DeepSeekExecutor { dspark.capture_mask = frontier.dspark_capture_mask; dspark.cache_start = frontier.dspark_cache_start; dspark.cache_len = frontier.dspark_cache_len; + dspark.cache_token_start = frontier.dspark_cache_token_start; + dspark.release_pending_seed(); } commands.finish()?; self.session.position = frontier.position; @@ -3413,20 +3618,12 @@ impl DeepSeekExecutor { } } if let Some(dspark) = &mut self.dspark { - let row = u64::from(count - 1); - for slot in 0..dspark.config.target_layers.len() as u64 { - dspark.target_hidden.copy_from( - slot * self.model.shape.embd * 4, - &dspark.target_hidden_batch, - (slot * u64::from(self.session.prefill_cap) + row) * self.model.shape.embd * 4, - self.model.shape.embd * 4, - "committing speculative DSpark target prefix", - )?; - } - dspark.capture_mask = (1_u32 << dspark.config.target_layers.len()) - 1; + // DS4 invalidates both capture forms on an ordinary partial commit. + dspark.capture_mask = 0; + dspark.release_pending_seed(); dspark.cache_start = baseline.dspark_cache_start; dspark.cache_len = baseline.dspark_cache_len; - dspark.commit_proposed_prefix(count, self.session.raw_cap); + dspark.cache_token_start = baseline.dspark_cache_token_start; } commands.finish()?; self.session.position = baseline.position + count; @@ -3449,7 +3646,6 @@ impl DeepSeekExecutor { } let started = Instant::now(); if self.quality - || proposals.len() == 1 || self .ssd .as_ref() @@ -3486,11 +3682,10 @@ impl DeepSeekExecutor { return Err(error); } }; - let mut commit = 1_usize; - while commit < proposals.len() && verification.tops[commit - 1] == proposals[commit] { - commit += 1; - } + let commit = greedy_verified_prefix_len(proposals, &verification.tops); if commit == proposals.len() { + self.spare_frontier = Some(frontier); + self.spare_prefixes = verification.prefixes; self.verifier_ns = self .verifier_ns .saturating_add(u64::try_from(started.elapsed().as_nanos()).unwrap_or(u64::MAX)); @@ -3502,6 +3697,8 @@ impl DeepSeekExecutor { verification.logits.get(commit - 1), ) { self.commit_spec_prefix(&frontier, prefix, &proposals[..commit], logits)?; + self.spare_frontier = Some(frontier); + self.spare_prefixes = verification.prefixes; self.verifier_ns = self .verifier_ns .saturating_add(u64::try_from(started.elapsed().as_nanos()).unwrap_or(u64::MAX)); @@ -3520,6 +3717,8 @@ impl DeepSeekExecutor { } self.verifier_passes += 1; } + self.spare_frontier = Some(frontier); + self.spare_prefixes = verification.prefixes; self.verifier_ns = self .verifier_ns .saturating_add(u64::try_from(started.elapsed().as_nanos()).unwrap_or(u64::MAX)); @@ -3627,6 +3826,8 @@ impl DeepSeekExecutor { accepted += 1; } if replacement.is_none() { + self.spare_frontier = Some(frontier); + self.spare_prefixes = verification.prefixes; self.verifier_ns = self .verifier_ns .saturating_add(u64::try_from(started.elapsed().as_nanos()).unwrap_or(u64::MAX)); @@ -3642,6 +3843,8 @@ impl DeepSeekExecutor { .get(accepted - 1) .ok_or("missing stochastic verifier logits")?; self.commit_spec_prefix(&frontier, prefix, &proposals[..accepted], logits)?; + self.spare_frontier = Some(frontier); + self.spare_prefixes = verification.prefixes; let replacement = replacement.expect("replacement disappeared"); self.eval_target(replacement)?; self.verifier_passes += 1; @@ -3705,11 +3908,13 @@ impl DeepSeekExecutor { } let no_draft = proposals.is_empty(); let verified = self.verify_target_suffix(&proposals, cancelled)?; - if verified.is_empty() { - self.dspark - .as_mut() - .expect("DSpark disappeared") - .commit_proposed_prefix(1, self.session.raw_cap); + if std::env::var_os("DS4_SPEC_CYCLE_TRACE").is_some() { + let dspark = self.dspark.as_ref().expect("DSpark disappeared"); + eprintln!( + "ds4: DSpark proposal pos={} first={first_token} confidence={:?} drafts={proposals:?} verified={verified:?}", + self.session.position - verified.len() as u32, + dspark.last_confidence + ); } accepted.extend_from_slice(&verified); let dspark = self.dspark.as_mut().expect("DSpark disappeared"); @@ -3908,8 +4113,35 @@ impl DeepSeekExecutor { { return Err("prefill contains a token outside the vocabulary".into()); } - let mut batch = - BatchScratch::allocate(&self.model, self.session.position, rows, collect_tops)?; + let prepend_seed = collect_tops + && self.session.position > 0 + && rows < self.session.prefill_cap + && self.dspark.as_ref().is_some_and(Dspark::capture_complete); + if let Some(dspark) = &mut self.dspark { + // The new capture replaces the old one; do not retain two large + // prefill workspaces while allocating the replacement batch. + dspark.release_pending_seed(); + } + let reusable = collect_tops + .then(|| { + self.dspark + .as_mut() + .and_then(|dspark| dspark.take_verifier_batch(rows, self.session.position)) + }) + .flatten(); + let mut batch = match reusable { + Some(batch) => batch, + None => BatchScratch::allocate( + &self.model, + if collect_tops { + self.session.context + } else { + self.session.position + }, + rows, + collect_tops, + )?, + }; if let Some(dspark) = &mut self.dspark { dspark.begin_capture(); } @@ -3918,16 +4150,21 @@ impl DeepSeekExecutor { let size = self.model.main.len(); let shape = self.model.shape; let pos = self.session.position; - let mut prefixes = (0..if collect_tops { rows } else { 0 }) - .map(|_| SpecPrefixFrontier { + let mut prefixes = if collect_tops { + std::mem::take(&mut self.spare_prefixes) + } else { + Vec::new() + }; + if collect_tops && prefixes.len() < rows as usize { + prefixes.resize_with(rows as usize, || SpecPrefixFrontier { layers: (0..shape.layers) .map(|_| LayerFrontier { compression: None, indexer: None, }) .collect(), - }) - .collect::>(); + }); + } if let Some(ssd) = &self.ssd { ssd.static_decode_map_current .store(false, std::sync::atomic::Ordering::Release); @@ -4040,6 +4277,7 @@ impl DeepSeekExecutor { index as u32, &batch.next_hc, rows, + prepend_seed, self.session.prefill_cap, shape, )?; @@ -4080,26 +4318,6 @@ impl DeepSeekExecutor { ); std::mem::swap(&mut batch.current_hc, &mut batch.next_hc); } - if pipelined_verifier { - commands - .take() - .expect("batch commands are active") - .finish()?; - } - - if self.dspark.is_some() && self.ssd.is_some() { - install_speculative_model_maps(&self.model, "DSpark prefill support mapping")?; - } - if let (Some(dspark), Some(support)) = (&mut self.dspark, self.model.support.as_ref()) { - dspark.seed_batch_cache( - support, - pos, - rows, - self.session.prefill_cap, - self.session.raw_cap, - shape, - )?; - } if self.ssd.is_some() { install_deepseek_model_spans( &self.model, @@ -4109,26 +4327,63 @@ impl DeepSeekExecutor { } let (tops, output_logits) = if collect_tops { - commands = Some(Commands::begin()?); + let commands = commands.take().map_or_else(Commands::begin, Ok)?; encode_batch_output(&batch, &self.weights, shape, map, size, rows)?; - commands - .take() - .expect("batch commands are active") - .finish()?; let logits = batch .output_logits .as_ref() .expect("batch output logits are allocated"); - let mut all_logits = vec![0.0; (u64::from(rows) * shape.vocab) as usize]; - logits.read_f32(&mut all_logits)?; - let output_logits = all_logits - .chunks_exact(shape.vocab as usize) - .map(<[f32]>::to_vec) - .collect::>(); - let tops = output_logits.iter().map(|logits| argmax(logits)).collect(); - self.logits - .clone_from(output_logits.last().expect("batch has an output row")); - (tops, output_logits) + let top_rows = rows - 1; + if top_rows > 0 { + call( + unsafe { + if top_rows == 1 { + ds4_gpu_argmax_tensor( + batch.indexer_selected.raw(), + logits.raw(), + shape.vocab as u32, + ) + } else { + ds4_gpu_indexer_topk_tensor( + batch.indexer_selected.raw(), + logits.raw(), + shape.vocab as u32, + top_rows, + 1, + ) + } + }, + "reducing verifier row tops", + )?; + } + commands.finish()?; + let mut tops = vec![0; top_rows as usize]; + if top_rows > 0 { + batch.indexer_selected.read_i32(&mut tops)?; + } + if self.speculative.dspark_exact_sampling { + // Exact sampling needs each target distribution, unlike DS4's + // ordinary verifier, which reads only the committed row. + let mut all_logits = vec![0.0; (u64::from(rows) * shape.vocab) as usize]; + logits.read_f32(&mut all_logits)?; + let output_logits = all_logits + .chunks_exact(shape.vocab as usize) + .map(<[f32]>::to_vec) + .collect::>(); + self.logits + .clone_from(output_logits.last().expect("batch has an output row")); + (tops, output_logits) + } else { + let commit = greedy_verified_prefix_len(tokens, &tops); + logits + .view((commit - 1) as u64 * shape.vocab * 4, shape.vocab * 4)? + .read_f32(&mut self.logits)?; + let mut output_logits = vec![Vec::new(); rows as usize]; + if commit < tokens.len() { + output_logits[commit - 1] = self.logits.clone(); + } + (tops, output_logits) + } } else { let row = rows - 1; let commands = Commands::begin()?; @@ -4146,6 +4401,14 @@ impl DeepSeekExecutor { }; self.session.position += rows; self.tokens.extend_from_slice(tokens); + if let Some(dspark) = &mut self.dspark { + // Keep the captured target rows until DS4's next proposal needs them. + dspark.pending_seed = Some(( + pos - u32::from(prepend_seed), + rows + u32::from(prepend_seed), + batch, + )); + } if let Some(profile) = &self.profile { profile.write()?; } @@ -4215,12 +4478,25 @@ impl DeepSeekExecutor { self.session.position } + pub(super) fn rewind_speculative_output(&mut self, keep: usize) -> Result<(), String> { + if keep > self.tokens.len() { + return Err("DeepSeek speculative rewind exceeds the committed frontier".into()); + } + // Match ds4_session_rewind: logical rewind and capture invalidation. + // Unlike GLM 5.3, DeepSeek does not restore a recurrent KDA snapshot. + self.tokens.truncate(keep); + self.session.position = keep as u32; + if let Some(dspark) = &mut self.dspark { + dspark.capture_mask = 0; + dspark.release_pending_seed(); + } + Ok(()) + } + pub(super) fn reset(&mut self) -> Result<(), String> { self.session = Session::new(&self.model, self.session.context, self.session.prefill_cap)?; if let Some(dspark) = &mut self.dspark { - dspark.capture_mask = 0; - dspark.cache_start = 0; - dspark.cache_len = 0; + dspark.reset_cache(); } self.tokens.clear(); self.checkpoint_tag = [0; 32]; @@ -4269,6 +4545,9 @@ impl DeepSeekExecutor { } pub(super) fn align_prompt(&mut self, tokens: &[i32]) -> Result { + if let Some(dspark) = &mut self.dspark { + dspark.begin_request(); + } if !tokens.starts_with(&self.tokens) { self.reset()?; } @@ -5086,6 +5365,7 @@ fn compress_attention_batch( state, coefficient * coefficient * u64::from(ratio) * shape.head_dim * 4, "capturing speculative compressor prefix", + prefixes[row as usize].layers[layer].compression.take(), )?); } } @@ -5272,6 +5552,7 @@ fn compress_index_batch( state, 4 * u64::from(ratio) * shape.indexer_head_dim * 4, "capturing speculative indexer prefix", + prefixes[row as usize].layers[layer].indexer.take(), )?); } } @@ -8178,89 +8459,13 @@ fn dense_dot_bytes(kind: u32, width: usize, bytes: &[u8], values: &[f32]) -> f32 }) .sum(), Q8_0 => { - let mut sum = 0.0; - for (block, bytes) in bytes.chunks_exact(34).enumerate() { - let scale = half_to_f32(u16::from_le_bytes([bytes[0], bytes[1]])); - for (index, quantized) in bytes[2..].iter().enumerate() { - let input = block * 32 + index; - if input == width { - break; - } - sum += scale * f32::from(*quantized as i8) * values[input]; - } - } - sum + let (quantized, scales) = quantize_q8_activation(values); + dense_dot_q8(bytes, &quantized, &scales, width) } _ => unreachable!(), } } -fn dense_argmax( - model: &Gguf, - weight: Weight, - values: &[f32], - logits: &[f32], -) -> Result { - if weight.dims[0] as usize != values.len() || weight.dims[1] as usize != logits.len() { - return Err("DSpark dense argmax has mismatched dimensions".into()); - } - let width = values.len(); - let row_bytes = match weight.kind { - F32 => width.checked_mul(4), - F16 => width.checked_mul(2), - Q8_0 => width.div_ceil(32).checked_mul(34), - _ => None, - } - .ok_or("unsupported DSpark dense tensor layout")?; - let bytes_len = row_bytes - .checked_mul(logits.len()) - .ok_or("DSpark dense argmax size overflow")?; - if weight.offset > model.len() || bytes_len as u64 > model.len() - weight.offset { - return Err("DSpark dense argmax is outside the GGUF mapping".into()); - } - let bytes = unsafe { - std::slice::from_raw_parts(model.map_ptr().add(weight.offset as usize), bytes_len) - }; - let quantized = (weight.kind == Q8_0).then(|| quantize_q8_activation(values)); - let workers = std::thread::available_parallelism() - .map_or(1, std::num::NonZero::get) - .min(logits.len()); - let chunk = logits.len().div_ceil(workers); - let best = std::thread::scope(|scope| { - let mut handles = Vec::with_capacity(workers); - let quantized = quantized.as_ref(); - for start in (0..logits.len()).step_by(chunk) { - let end = (start + chunk).min(logits.len()); - handles.push(scope.spawn(move || { - let mut best = (start, f32::NEG_INFINITY); - for token in start..end { - let row = &bytes[token * row_bytes..(token + 1) * row_bytes]; - let dot = quantized.as_ref().map_or_else( - || dense_dot_bytes(weight.kind, width, row, values), - |(values, scales)| dense_dot_q8(row, values, scales, width), - ); - let score = logits[token] + dot; - if score > best.1 { - best = (token, score); - } - } - best - })); - } - handles - .into_iter() - .map(|handle| handle.join().expect("DSpark argmax worker panicked")) - .fold((0, f32::NEG_INFINITY), |best, candidate| { - if candidate.1 > best.1 { - candidate - } else { - best - } - }) - }); - Ok(best.0 as i32) -} - fn quantize_q8_activation(values: &[f32]) -> (Vec, Vec) { let blocks = values.len().div_ceil(32); let mut quantized = vec![0; blocks * 32]; @@ -8280,6 +8485,15 @@ fn quantize_q8_activation(values: &[f32]) -> (Vec, Vec) { } fn dense_dot_q8(bytes: &[u8], values: &[i8], scales: &[f32], width: usize) -> f32 { + #[cfg(target_arch = "aarch64")] + if width.is_multiple_of(32) && std::arch::is_aarch64_feature_detected!("dotprod") { + assert!( + bytes.len() >= width / 32 * 34 && values.len() >= width && scales.len() >= width / 32 + ); + // DS4 accumulates alternating Q8 blocks in two four-lane FMA vectors. + // Preserve that reduction order for both confidence and Markov scores. + return unsafe { dense_dot_q8_neon(bytes, values, scales, width) }; + } bytes .chunks_exact(34) .zip(values.chunks_exact(32)) @@ -8299,6 +8513,37 @@ fn dense_dot_q8(bytes: &[u8], values: &[i8], scales: &[f32], width: usize) -> f3 .sum() } +#[cfg(target_arch = "aarch64")] +#[target_feature(enable = "dotprod")] +unsafe fn dense_dot_q8_neon(bytes: &[u8], values: &[i8], scales: &[f32], width: usize) -> f32 { + use std::arch::aarch64::*; + let mut sums = [vdupq_n_f32(0.0); 2]; + for block in 0..width / 32 { + let offset = block * 34; + let scale = + half_to_f32(u16::from_le_bytes([bytes[offset], bytes[offset + 1]])) * scales[block]; + let weight = unsafe { bytes.as_ptr().add(offset + 2).cast::() }; + let input = unsafe { values.as_ptr().add(block * 32) }; + let mut dot = vdupq_n_s32(0); + // Stable Rust does not yet expose the SDOT intrinsic. Keep the same + // two integer dot instructions here without adding native host code. + unsafe { + std::arch::asm!( + "sdot {dot:v}.4s, {w0:v}.16b, {x0:v}.16b", + "sdot {dot:v}.4s, {w1:v}.16b, {x1:v}.16b", + dot = inout(vreg) dot, + w0 = in(vreg) vld1q_s8(weight), + x0 = in(vreg) vld1q_s8(input), + w1 = in(vreg) vld1q_s8(weight.add(16)), + x1 = in(vreg) vld1q_s8(input.add(16)), + options(pure, nomem, nostack), + ); + } + sums[block % 2] = vfmaq_n_f32(sums[block % 2], vcvtq_f32_s32(dot), scale); + } + vaddvq_f32(vaddq_f32(sums[0], sums[1])) +} + fn half_to_f32(value: u16) -> f32 { let sign = u32::from(value & 0x8000) << 16; let exponent = u32::from((value >> 10) & 0x1f); @@ -8522,6 +8767,17 @@ fn power_throttle_keeps_ds4_ewma_and_full_power_bypass() { assert_eq!(average, 0.015); } +// The first proposal was checked against the preceding target distribution. +fn greedy_verified_prefix_len(proposals: &[i32], tops: &[i32]) -> usize { + usize::from(!proposals.is_empty()) + + proposals + .iter() + .skip(1) + .zip(tops) + .take_while(|(token, top)| token == top) + .count() +} + fn call(result: i32, operation: &str) -> Result<(), String> { if result == 0 { Err(format!("Metal failed while {operation}")) @@ -8560,9 +8816,21 @@ mod tests { #[test] fn dspark_scheduler_matches_the_ds4_default_window() { - assert_eq!(dspark_scheduler_pause(4, 8, 0), 0); - assert_eq!(dspark_scheduler_pause(4, 5, 0), 2); - assert_eq!(dspark_scheduler_pause(4, 8, 2), 4); + assert_eq!(dspark_scheduler_pause(4, 8, 0, 0), 0); + assert_eq!(dspark_scheduler_pause(4, 5, 0, 0), 2); + assert_eq!(dspark_scheduler_pause(4, 8, 2, 0), 4); + // The window decision replaces the cold no-draft pause, not max(7, 4). + assert_eq!(dspark_scheduler_pause(4, 0, 4, 7), 4); + } + + #[test] + fn dspark_verifier_selects_only_the_committed_row() { + use super::greedy_verified_prefix_len; + assert_eq!(greedy_verified_prefix_len(&[], &[]), 0); + assert_eq!(greedy_verified_prefix_len(&[7], &[]), 1); + assert_eq!(greedy_verified_prefix_len(&[7, 8, 9], &[8, 9]), 3); + assert_eq!(greedy_verified_prefix_len(&[7, 8, 9], &[8, 0]), 2); + assert_eq!(greedy_verified_prefix_len(&[7, 8, 9], &[0, 9]), 1); } #[test] @@ -8572,6 +8840,18 @@ mod tests { assert_eq!(scales, [1.0 / 127.0]); } + #[test] + fn dspark_confidence_quantizes_q8_input_like_ds4() { + let mut weights = [1_u8; 34]; + weights[..2].copy_from_slice(&0x3c00_u16.to_le_bytes()); + let mut features = [0.0; 32]; + features[..2].copy_from_slice(&[1.0, 0.1]); + assert_eq!( + super::dense_dot_bytes(crate::engine::Q8_0, 32, &weights, &features), + (1.0 / 127.0) * 140.0 + ); + } + #[test] fn pro_q4_model_spans_remain_isolated() { assert_eq!( @@ -8974,6 +9254,185 @@ mod tests { run_dspark_target_owned_greedy_cycle(false); } + #[test] + #[ignore = "requires installed 0731/DSpark and DS4SERVER_DSPARK_REFERENCE cycle recording"] + fn dspark_matches_original_summary_cycles() { + use crate::engine::{Model, Rng, sample}; + use crate::model::ModelChoice; + use crate::settings::{ + DiagnosticPreferences, EngineSettings, ExecutionPreferences, ReasoningMode, + SpeculativePreferences, SsdPreferences, SteeringPreferences, + }; + let recording = std::fs::read_to_string( + std::env::var_os("DS4SERVER_DSPARK_REFERENCE").expect("reference recording required"), + ) + .unwrap(); + let events: Vec = recording + .lines() + .map(|line| serde_json::from_str(line).unwrap()) + .collect(); + let prefill = events + .iter() + .find(|v| v["event"] == "reference_prefill" && v["turn"] == 1) + .unwrap(); + let prompt: Vec = serde_json::from_value(prefill["tokens"].clone()).unwrap(); + let cycles: Vec<_> = events + .iter() + .filter(|v| { + v["event"] == "spec_cycle" && v["prompt_tokens"] == prefill["prompt_tokens"] + }) + .take(32) + .collect(); + assert_eq!(cycles.len(), 32, "32 oracle cycles required"); + let settings = EngineSettings { + model: ModelChoice::DeepSeekV4Flash0731, + artifacts: installed_artifacts(ModelChoice::DeepSeekV4Flash0731, true), + context_tokens: 32768, + execution: ExecutionPreferences::default().engine_settings(), + speculative: SpeculativePreferences { + dspark_enabled: true, + ..Default::default() + } + .engine_settings(), + ssd: SsdPreferences::default().engine_settings(), + steering: SteeringPreferences::default().engine_settings(), + diagnostics: DiagnosticPreferences::default().engine_settings(), + }; + super::configure_sources().unwrap(); + let mut executor = super::DeepSeekExecutor::open( + Model::open(&settings).unwrap(), + 32768, + false, + 0, + 100, + settings.speculative, + settings.ssd, + settings.steering, + ) + .unwrap(); + executor.eval(prompt[0]).unwrap(); + let dspark = executor.dspark.as_mut().unwrap(); + dspark.scheduler_skip = 7; + dspark.scheduler_cycles = 3; + dspark.scheduler_lifetime_accepted = 10; + executor.align_prompt(&prompt).unwrap(); + let dspark = executor.dspark.as_ref().unwrap(); + assert_eq!( + ( + dspark.scheduler_skip, + dspark.scheduler_cycles, + dspark.scheduler_lifetime_accepted + ), + (0, 0, 0) + ); + executor.prefill(&prompt[1..], |_| true).unwrap(); + let mut rng = Rng::new(42); + let cancelled = std::sync::atomic::AtomicBool::new(false); + let mut buffers = std::collections::BTreeMap::new(); + let mut reuses = 0; + for cycle in cycles { + let first = sample(executor.logits(), 0.6, 0.95, 0.0, 0, &mut rng); + assert_eq!(i64::from(first), cycle["first"].as_i64().unwrap()); + let expected: Vec = serde_json::from_value(cycle["accepted"].clone()).unwrap(); + let accepted = executor + .eval_speculative_sampled( + first, + 32767 - executor.position(), + ReasoningMode::Low, + 0.6, + 0.95, + 0.0, + 0, + &mut rng, + &cancelled, + ) + .unwrap(); + assert_eq!(accepted, expected, "oracle cycle at {}", cycle["generated"]); + assert_eq!( + u64::from(executor.position()), + cycle["position"].as_u64().unwrap() + ); + let mut remember = |key, buffer: &super::Buffer| { + let pointer = buffer.raw() as usize; + if let Some(previous) = buffers.insert(key, pointer) { + assert_eq!(pointer, previous, "verifier buffer reallocated: {key:?}"); + reuses += 1; + } + }; + if let Some(frontier) = &executor.spare_frontier { + for (slot, layers) in std::iter::once(&frontier.layers) + .chain(executor.spare_prefixes.iter().map(|p| &p.layers)) + .enumerate() + { + for (layer, saved) in layers.iter().enumerate() { + for (kind, state) in + [&saved.compression, &saved.indexer].into_iter().enumerate() + { + if let Some(state) = state { + remember((slot, layer, kind * 2), &state.state_kv); + remember((slot, layer, kind * 2 + 1), &state.state_score); + } + } + } + } + } + let dspark = executor.dspark.as_ref().unwrap(); + for batch in dspark + .spare_batches + .iter() + .chain(dspark.pending_seed.iter().map(|(_, _, batch)| batch)) + { + if batch.output_logits.is_some() { + remember((usize::MAX, batch.rows as usize, 0), &batch.tokens); + } + } + } + assert!(reuses > 0); + assert!(buffers.contains_key(&(usize::MAX, 2, 0))); + assert!(buffers.contains_key(&(usize::MAX, 8, 0))); + let frontier = executor.snapshot_spec_frontier().unwrap(); + let token = super::argmax(executor.logits()); + executor.eval_target(token).unwrap(); + executor.restore_spec_frontier(&frontier).unwrap(); + assert_eq!(executor.position(), frontier.position); + assert_eq!(executor.logits(), frontier.logits); + for (layer, saved) in executor.session.layers.iter().zip(&frontier.layers) { + for (state, saved) in [ + (&layer.compression, &saved.compression), + (&layer.indexer, &saved.indexer), + ] { + if let (Some(state), Some(saved)) = (state, saved) { + assert_eq!(state.rows, saved.rows); + for (live, saved_buffer) in [ + (&state.state_kv, &saved.state_kv), + (&state.state_score, &saved.state_score), + ] { + let mut actual = vec![0; saved.bytes as usize]; + let mut expected = actual.clone(); + live.read(0, &mut actual).unwrap(); + saved_buffer.read(0, &mut expected).unwrap(); + assert_eq!(actual, expected, "speculative rollback bytes differ"); + } + } + } + } + let before = executor.position(); + assert!( + executor + .rewind_speculative_output(before as usize + 1) + .is_err() + ); + assert_eq!(executor.position(), before); + executor + .rewind_speculative_output(before as usize - 1) + .unwrap(); + assert_eq!(executor.position(), before - 1); + assert_eq!(executor.tokens().len(), before as usize - 1); + let dspark = executor.dspark.as_ref().unwrap(); + assert_eq!(dspark.capture_mask, 0); + assert!(dspark.pending_seed.is_none()); + } + #[test] #[ignore = "requires the installed 0731 Flash and checkpoint-specific DSpark GGUF fixtures"] fn ssd_streaming_supports_dspark() { @@ -9172,8 +9631,9 @@ mod tests { glm_mtp_timing: false, keep_vision_loaded: false, dspark: true, - dspark_confidence_threshold: 0.6, - dspark_confidence_threshold_set: false, + // Exercise drafting independently of the fixture's confidence. + dspark_confidence_threshold: 0.0, + dspark_confidence_threshold_set: true, dspark_strict: false, dspark_exact_sampling: true, }, @@ -9198,7 +9658,7 @@ mod tests { let cycle = executor .eval_speculative_sampled( first, - 4, + 16, ReasoningMode::Direct, 0.8, 0.95, @@ -9212,6 +9672,16 @@ mod tests { assert!(executor.dspark.as_ref().unwrap().drafted > 0); assert!(executor.logits().iter().all(|logit| logit.is_finite())); assert!(executor.session.position >= prompt.len() as u32 + cycle.len() as u32); + // Exact verification needs both full distributions, even when the + // ordinary verifier would only read the committed prefix's row. + let verification = executor.eval_batch_tops(&[first, first]).unwrap(); + assert_eq!(verification.logits.len(), 2); + for row in &verification.logits { + assert_eq!(row.len(), executor.model.shape.vocab as usize); + assert!(row.iter().all(|logit| logit.is_finite())); + } + assert_eq!(verification.tops, [argmax(&verification.logits[0])]); + assert_eq!(executor.logits(), verification.logits[1]); } #[test] diff --git a/src/engine/metal/checkpoint.rs b/src/engine/metal/checkpoint.rs index 2a13831..5da0f2e 100644 --- a/src/engine/metal/checkpoint.rs +++ b/src/engine/metal/checkpoint.rs @@ -302,9 +302,7 @@ impl DeepSeekExecutor { self.logits = logits; self.checkpoint_tag = checkpoint_tag; if let Some(dspark) = &mut self.dspark { - dspark.capture_mask = 0; - dspark.cache_start = 0; - dspark.cache_len = 0; + dspark.reset_cache(); } Ok(()) } diff --git a/src/engine/metal/glm.rs b/src/engine/metal/glm.rs index f9c64f0..3c5533e 100644 --- a/src/engine/metal/glm.rs +++ b/src/engine/metal/glm.rs @@ -1008,7 +1008,8 @@ impl GlmExecutor { )?; } std::mem::swap(&mut self.scratch.current, &mut self.scratch.next); - if layer_index + 1 == DECODE_FLUSH_LAYERS { + if glm_decode_flush_layer(layer_index + 1, self.weights.layers.len(), self.ssd.enabled) + { call( unsafe { ds4_gpu_flush_commands() }, "flushing the GLM decode graph", @@ -1419,7 +1420,15 @@ impl GlmExecutor { causal_range: bool, ) -> Result<(), String> { let dsa = layer.dsa()?; - self.encode_batch_attention_raw(batch, dsa, cache, pos, rows, selected_count, causal_range) + self.encode_batch_attention_raw( + batch, + dsa, + cache, + pos, + rows, + selected_count, + if causal_range { pos + rows } else { 0 }, + ) } #[allow(clippy::too_many_arguments)] @@ -1431,7 +1440,7 @@ impl GlmExecutor { pos: u32, rows: u32, selected_count: u32, - causal_range: bool, + dense_limit: u32, ) -> Result<(), String> { let shape = self.model.shape; let map = self.model.main.map_ptr().cast(); @@ -1446,7 +1455,12 @@ impl GlmExecutor { )?; let mut start = 0; while start < rows { - let slice = (rows - start).min(INDEXED_PREFILL_ATTN_SLICE); + let (slice, causal_range) = glm_attention_slice(pos + start, rows - start, dense_limit); + let attention_selected_count = if causal_range { + (pos + rows).min(dense_limit) + } else { + selected_count + }; let q = batch .q .view(u64::from(start) * q_dim * 4, u64::from(slice) * q_dim * 4)?; @@ -1471,7 +1485,7 @@ impl GlmExecutor { cache.kv.raw(), pos + start, slice, - selected_count, + attention_selected_count, self.context, CACHE_F16, shape.heads as u32, @@ -1492,7 +1506,7 @@ impl GlmExecutor { cache.rope.raw(), slice, pos + start, - selected_count, + attention_selected_count, self.context, CACHE_F16, shape.heads as u32, @@ -2007,6 +2021,29 @@ impl GlmExecutor { result } + pub(crate) fn rewind_speculative_output(&mut self, pos: usize) -> Result<(), String> { + if pos > self.tokens.len() { + return Err("GLM emitted token frontier exceeds evaluated tokens".into()); + } + if pos == self.tokens.len() { + return Ok(()); + } + if self.model.shape.model == ModelChoice::Glm53Flash { + return if self.rewind_glm53_mtp(pos)? { + Ok(()) + } else { + Err("GLM could not restore the emitted token frontier".into()) + }; + } + self.tokens.truncate(pos); + if let Some(mtp) = &mut self.mtp { + mtp.pending = None; + mtp.parent = None; + mtp.rollback = None; + } + Ok(()) + } + fn rewind_glm53_mtp(&mut self, pos: usize) -> Result { let Some((start, first)) = self.mtp.as_ref().and_then(|m| m.rollback) else { return Ok(false); @@ -3137,7 +3174,7 @@ impl GlmExecutor { )?; let visible = pos + rows; - let causal = pos < dense_limit; + let causal = visible <= dense_limit; let selected_count = if causal { visible.min(dense_limit) } else { @@ -3250,7 +3287,15 @@ impl GlmExecutor { }, "projecting batched GLM 5.3 low-rank queries", )?; - self.encode_batch_attention_raw(batch, weights, cache, pos, rows, selected_count, causal)?; + self.encode_batch_attention_raw( + batch, + weights, + cache, + pos, + rows, + selected_count, + dense_limit, + )?; glm53_project_rows( &batch.attn_out, weights.output, @@ -3467,7 +3512,7 @@ impl GlmExecutor { "expanding GLM 5.3 FFN hyperconnections", )?; std::mem::swap(&mut self.scratch.hc_current, &mut self.scratch.hc_next); - if ordinal + 1 == DECODE_FLUSH_LAYERS { + if glm_decode_flush_layer(ordinal + 1, self.weights.layers.len(), self.ssd.enabled) { call( unsafe { ds4_gpu_flush_commands() }, "flushing GLM 5.3 decode", @@ -4555,14 +4600,12 @@ impl GlmExecutor { } let pos = self.position(); - let mut rows = - indexed_prefill_rows(pos, remaining.len(), self.model.shape.indexer_top_k as u32); - if self.model.shape.model == ModelChoice::Glm53Flash { - let dense = glm53_dense_limit(self.context, self.ssd.enabled); - if pos < dense { - rows = rows.min((dense - pos) as usize); - } - } + let rows = indexed_prefill_rows( + self.model.shape.model, + pos, + remaining.len(), + self.model.shape.indexer_top_k as u32, + ); let cancelled = self.eval_batch( &remaining[..rows], &mut progress, @@ -4760,13 +4803,7 @@ impl GlmExecutor { let shape = self.model.shape; let map = self.model.main.map_ptr().cast(); let size = self.model.main.len(); - let mut dense_limit = glm53_dense_limit(self.context, self.ssd.enabled); - // DS4 routes the entire verification pair through indexed attention - // when it no longer fits the full-attention span. Normal prefill - // already splits at this boundary before entering the batch function. - if pos < dense_limit && pos + rows > dense_limit { - dense_limit = 0; - } + let dense_limit = glm53_dense_limit(self.context, self.ssd.enabled); // Both HC split/sum and expansion derive their row count from the // output view, not a rows argument. Match DS4's token-sized views even // when a short prefill reuses the full 2048-row workspace. @@ -4930,6 +4967,16 @@ impl GlmExecutor { }, "expanding batched GLM 5.3 attention hyperconnections", )?; + #[cfg(test)] + { + commands = compare_glm53_reference_stages( + commands, + index, + pos, + rows as usize * shape.embd as usize, + &[("attn_out", &batch.attn_out)], + )?; + } glm53_hc_pre( after_attn, &batch.ffn_norm, @@ -4947,6 +4994,16 @@ impl GlmExecutor { size, )?; self.encode_glm53_ffn_batch(batch, layer, index as u32, rows)?; + #[cfg(test)] + { + commands = compare_glm53_reference_stages( + commands, + index, + pos, + rows as usize * shape.embd as usize, + &[("ffn_out", &batch.next)], + )?; + } if let Some(steering) = &self.steering { steering.apply(&batch.next, index as u32, rows, false)?; } @@ -4963,10 +5020,11 @@ impl GlmExecutor { }, "expanding batched GLM 5.3 FFN hyperconnections", )?; + let drain = glm_prefill_flush_layers(rows, read_logits) + && (index + 1).is_multiple_of(PREFILL_DRAIN_LAYERS) + && index + 1 < self.weights.layers.len(); if glm_prefill_flush_layers(rows, read_logits) { - if (index + 1).is_multiple_of(PREFILL_DRAIN_LAYERS) - && index + 1 < self.weights.layers.len() - { + if drain { commands.finish()?; commands = Commands::begin()?; } else { @@ -4989,11 +5047,15 @@ impl GlmExecutor { if let Some(views) = &mut hc_views { views.swap(2, 3); } - let done = u32::try_from( - u64::from(rows) * (index as u64 + 1) / self.weights.layers.len() as u64, - ) - .expect("GLM 5.3 prefill progress is bounded by the batch"); - cancelled |= !progress(pos + done); + // DS4 reports completed work at drains, not every encoded layer. + // An asynchronous flush alone does not advance GPU completion. + if drain { + let done = u32::try_from( + u64::from(rows) * (index as u64 + 1) / self.weights.layers.len() as u64, + ) + .expect("GLM 5.3 prefill progress is bounded by the batch"); + cancelled |= !progress(pos + done); + } } commands.finish()?; @@ -5024,6 +5086,7 @@ impl GlmExecutor { self.logits = logits; } self.tokens.extend_from_slice(tokens); + cancelled |= !progress(self.position()); if let Some(profile) = &self.profile { profile.write()?; } @@ -5513,14 +5576,35 @@ fn glm_prefill_flush_layers(rows: u32, logits_requested: bool) -> bool { logits_requested && rows > 8 } -fn indexed_prefill_rows(pos: u32, remaining: usize, top_k: u32) -> usize { +fn glm_decode_flush_layer(completed: usize, layers: usize, ssd_streaming: bool) -> bool { + // The resident indexed DS4 graph flushes every four layers, except the + // final layer. Static SSD mappings use the expert loader's own boundaries. + !ssd_streaming && completed < layers && completed.is_multiple_of(DECODE_FLUSH_LAYERS) +} + +fn indexed_prefill_rows(model: ModelChoice, pos: u32, remaining: usize, top_k: u32) -> usize { let mut rows = remaining.min(INDEXED_PREFILL_CHUNK as usize); - if pos < top_k { + // GLM 5.3's active DS4 graph keeps full chunks across both the old + // top-k boundary and the dense/sparse boundary. Only attention is sliced. + if model != ModelChoice::Glm53Flash && pos < top_k { rows = rows.min((top_k - pos) as usize); } rows } +fn glm_attention_slice(pos: u32, remaining: u32, dense_limit: u32) -> (u32, bool) { + let causal = pos < dense_limit; + let rows = remaining.min(INDEXED_PREFILL_ATTN_SLICE); + ( + if causal { + rows.min(dense_limit - pos) + } else { + rows + }, + causal, + ) +} + fn full_indexer_layer(shape: super::super::Shape, layer: usize) -> bool { if shape.model == ModelChoice::Glm53Flash { return layer < (shape.layers - shape.nextn) as usize && !glm53_kda_layer(shape, layer); @@ -6432,12 +6516,72 @@ fn f32_project_rows( ) } +// Diagnostic only: no stage reads, extra drains or environment checks in the +// release app. Match original DS4's existing layer0/bootstrap tensor dumps. +#[cfg(test)] +fn compare_glm53_reference_stages( + commands: Commands, + layer: usize, + pos: u32, + values: usize, + tensors: &[(&str, &Buffer)], +) -> Result { + let Some(prefix) = std::env::var_os("DS4SERVER_GLM_STAGE_REFERENCE") else { + return Ok(commands); + }; + let selected_pos = std::env::var("DS4SERVER_GLM_STAGE_POS") + .map_or(Ok(0), |v| v.parse::()) + .map_err(|e| e.to_string())?; + if layer != 0 || pos != selected_pos { + return Ok(commands); + } + commands.finish()?; + for &(name, tensor) in tensors { + let path = format!("{}_{name}-{layer}_pos{pos}.bin", prefix.to_string_lossy()); + let bytes = std::fs::read(&path).map_err(|e| format!("{path}: {e}"))?; + if values == 0 || bytes.len() != values * 4 { + return Err(format!( + "reference tensor geometry mismatch: {path}: {} bytes, expected {}", + bytes.len(), + values * 4 + )); + } + let reference: Vec = bytes + .chunks_exact(4) + .map(|b| f32::from_le_bytes(b.try_into().unwrap())) + .collect(); + let mut actual = vec![0.0; reference.len()]; + tensor.read_f32(&mut actual)?; + if !actual.iter().chain(&reference).all(|v| v.is_finite()) { + return Err(format!("nonfinite stage {name}")); + } + let max_abs = actual + .iter() + .zip(&reference) + .map(|(a, b)| (a - b).abs()) + .fold(0.0_f32, f32::max); + let rms = (actual + .iter() + .zip(&reference) + .map(|(a, b)| f64::from(a - b).powi(2)) + .sum::() + / actual.len() as f64) + .sqrt(); + eprintln!( + "{}", + serde_json::json!({"event":"glm_stage_compare", "name":name, "layer":layer, "pos":pos, "values":actual.len(), "max_abs":max_abs, "rms":rms}) + ); + } + Commands::begin() +} + #[cfg(test)] mod tests { use super::{ - GlmExecutor, argmax, dynamic_expert_budget, full_indexer_layer, glm_prefill_flush_layers, - glm53_dense_compact_prefill, glm53_dense_limit, indexed_prefill_rows, - live_prefix_rewind_target, streaming_token_prefill_eligible, + GlmExecutor, argmax, dynamic_expert_budget, full_indexer_layer, glm_attention_slice, + glm_decode_flush_layer, glm_prefill_flush_layers, glm53_dense_compact_prefill, + glm53_dense_limit, indexed_prefill_rows, live_prefix_rewind_target, + streaming_token_prefill_eligible, }; use crate::engine::{GLM, Model, ReasoningMode}; use crate::model::ModelChoice; @@ -6464,9 +6608,30 @@ mod tests { assert!(!streaming_token_prefill_eligible(8192, 0, 65, 64)); assert!(!streaming_token_prefill_eligible(8192, 8180, 13, 64)); assert!(!streaming_token_prefill_eligible(65_536, 4096, 1, 64)); - assert_eq!(indexed_prefill_rows(0, 5000, 2048), 2048); - assert_eq!(indexed_prefill_rows(2048, 5000, 2048), 2048); - assert_eq!(indexed_prefill_rows(0, 17, 2048), 17); + assert_eq!( + indexed_prefill_rows(ModelChoice::Glm52, 0, 5000, 2048), + 2048 + ); + assert_eq!( + indexed_prefill_rows(ModelChoice::Glm52, 2048, 5000, 2048), + 2048 + ); + assert_eq!(indexed_prefill_rows(ModelChoice::Glm52, 0, 17, 2048), 17); + assert_eq!( + indexed_prefill_rows(ModelChoice::Glm52, 9, 2626, 2048), + 2039 + ); + for pos in [9, 2047, 4095, 8191] { + assert_eq!( + indexed_prefill_rows(ModelChoice::Glm53Flash, pos, 2626, 2048), + 2048 + ); + } + for boundary in [4096, 8192] { + assert_eq!(glm_attention_slice(boundary - 1, 2, boundary), (1, true)); + assert_eq!(glm_attention_slice(boundary, 1, boundary), (1, false)); + assert_eq!(glm_attention_slice(9, 2048, boundary), (2048, true)); + } assert_eq!(glm53_dense_limit(32_768, false), 4096); assert_eq!(glm53_dense_limit(32_768, true), 8192); assert_eq!(glm53_dense_limit(65_536, true), 4096); @@ -6477,6 +6642,212 @@ mod tests { assert!(!glm_prefill_flush_layers(2048, false)); } + #[test] + fn glm_resident_decode_flushes_periodically_without_a_final_or_ssd_flush() { + for layers in [1, 4, 5, 76, 78] { + let actual = (1..=layers) + .filter(|&completed| glm_decode_flush_layer(completed, layers, false)) + .collect::>(); + assert_eq!(actual, (4..layers).step_by(4).collect::>()); + assert!((1..=layers).all(|completed| !glm_decode_flush_layer(completed, layers, true))); + } + } + + #[test] + #[ignore = "installed GLM and DS4SERVER_GLM_REFERENCE_LOG from the standalone reference"] + fn glm53_reference_prompt_tokens_match_shared_runtime() { + let path = std::env::var("DS4SERVER_GLM_REFERENCE_LOG").unwrap(); + let events = std::fs::read_to_string(path) + .unwrap() + .lines() + .map(|line| serde_json::from_str::(line).unwrap()) + .collect::>(); + let start = events + .iter() + .find(|e| e["event"] == "reference_start") + .unwrap(); + assert_eq!(start["family"], "glm"); + let model = Model::open_main( + std::path::Path::new(start["model"].as_str().unwrap()), + ModelChoice::Glm53Flash, + ) + .unwrap(); + let readme = std::fs::read_to_string(start["readme"].as_str().unwrap()).unwrap(); + let prompts = [ + "Reply with exactly: OK".to_owned(), + format!("Give a summary of the following text:\n\n{readme}"), + "Tell me a complete short story about a lighthouse keeper. Do not ask questions.".into(), + "Write a Python function is_prime(n: int) -> bool, followed by five assert examples. No tools.".into(), + ]; + let ids = |v: &serde_json::Value| { + v.as_array() + .unwrap() + .iter() + .map(|v| v.as_i64().unwrap() as i32) + .collect::>() + }; + let mut frontier = Vec::new(); + for (index, prompt) in prompts.iter().enumerate() { + let expected = if index <= 1 { + model.render_prompt("", prompt, ReasoningMode::Low) + } else { + let mut expected = frontier.clone(); + expected.extend(model.tokenizer.encode_continuation( + prompt, + ReasoningMode::Low, + false, + )); + expected + }; + let receipt = events + .iter() + .find(|e| e["event"] == "reference_prefill" && e["turn"] == index) + .unwrap(); + assert_eq!( + ids(&receipt["tokens"]), + expected, + "reference turn {index} prompt differs" + ); + let cached = if index <= 1 { 9 } else { frontier.len() }; + assert_eq!( + receipt["cached_tokens"], cached, + "bootstrap/continuation differs" + ); + let result = events + .iter() + .find(|e| e["event"] == "reference_result" && e["turn"] == index) + .unwrap(); + if index != 0 { + assert_eq!(result["finish_reason"], "stop"); + } + frontier = expected; + frontier.extend(ids(&result["token_ids"])); + } + } + + #[test] + #[ignore = "requires the installed GLM 5.3 Flash Q2 checkpoint and Apple Metal"] + fn glm53_reference_logits_replay_separates_sampling_from_execution() { + use crate::engine::{Rng, sample}; + let events: Vec = + std::fs::read_to_string(std::env::var("DS4SERVER_GLM_REFERENCE_LOG").unwrap()) + .unwrap() + .lines() + .map(|l| serde_json::from_str(l).unwrap()) + .collect(); + let event = |name: &str| { + events + .iter() + .find(|e| e["event"] == name && (name == "reference_start" || e["turn"] == 1)) + .unwrap() + }; + let start = event("reference_start"); + assert_eq!(start["family"], "glm"); + assert_eq!(start["acceleration"], false); + assert_eq!(start["seed"], 42); + let trace = event("reference_logits_trace"); + let path: std::ffi::OsString = serde_json::from_value(trace["path"].clone()).unwrap(); + let bytes = std::fs::read(std::path::PathBuf::from(path)).unwrap(); + let n = trace["vocab"].as_u64().unwrap() as usize; + assert_eq!(bytes.len(), 32 * n * 4); + let values: Vec = bytes + .chunks_exact(4) + .map(|b| f32::from_le_bytes(b.try_into().unwrap())) + .collect(); + let token_ids = |v: &serde_json::Value| { + v.as_array() + .unwrap() + .iter() + .map(|v| v.as_i64().unwrap() as i32) + .collect::>() + }; + let prompt = token_ids(&event("reference_prefill")["tokens"]); + let expected = token_ids(&event("reference_result")["token_ids"]); + assert!(expected.len() >= 32); + // First feed the exact C-produced rows through the production Rust + // sampler, so a model-graph difference cannot masquerade as RNG drift. + let mut host_rng = Rng::new(42); + let host: Vec = values + .chunks_exact(n) + .map(|row| sample(row, 0.6, 0.95, 0.0, 0, &mut host_rng)) + .collect(); + assert_eq!( + host, + expected[..32], + "Rust sampler differs on original DS4 logits" + ); + super::super::configure_sources().unwrap(); + let model = Model::open_main( + std::path::Path::new(start["model"].as_str().unwrap()), + ModelChoice::Glm53Flash, + ) + .unwrap(); + assert_eq!(n, model.shape.vocab as usize); + let mut executor = GlmExecutor::open( + model, + 32768, + false, + EngineSsdSettings { + enabled: false, + cold: false, + cache_experts: 0, + cache_bytes: 0, + full_layers: 0, + full_layers_set: false, + preload_experts: 0, + }, + ) + .unwrap(); + assert_eq!(event("reference_prefill")["cached_tokens"], 9); + for part in [&prompt[..9], &prompt[9..]] { + assert_eq!( + executor + .prefill(part, |pos| { + eprintln!("reference replay prefill {pos}"); + true + }) + .unwrap(), + part.len() + ); + } + let mut gpu_rng = Rng::new(42); + let mut different = Vec::new(); + for (step, row) in values.chunks_exact(n).enumerate() { + let actual = executor.logits(); + let chosen = sample(actual, 0.6, 0.95, 0.0, 0, &mut gpu_rng); + let max_abs = actual + .iter() + .zip(row) + .map(|(a, b)| (a - b).abs()) + .fold(0.0_f32, f32::max); + let rms = (actual + .iter() + .zip(row) + .map(|(a, b)| f64::from(a - b).powi(2)) + .sum::() + / n as f64) + .sqrt(); + eprintln!( + "{}", + serde_json::json!({"event":"glm_logits_compare", "step":step, "expected":expected[step], "actual":chosen, "max_abs":max_abs, "rms":rms}) + ); + if chosen != expected[step] + || !actual + .iter() + .zip(row) + .all(|(a, b)| a.to_bits() == b.to_bits()) + { + different.push(step); + } + // Never diverge the context: both backends see the reference token. + executor.eval(expected[step]).unwrap(); + } + assert!( + different.is_empty(), + "logits or sampled outputs differ on identical history at steps {different:?}" + ); + } + #[test] #[ignore = "requires the installed GLM 5.3 Flash Q2 checkpoint and Apple Metal"] fn glm53_m5_text_performance_gate() { @@ -7171,12 +7542,30 @@ mod tests { .unwrap(); let prepare = |executor: &mut GlmExecutor| { executor.reset().unwrap(); + let mut reported = Vec::new(); executor .prefill(&prompt, |pos| { + reported.push(pos); eprintln!("verify2 prefill {pos}"); true }) .unwrap(); + if prompt.len() <= super::INDEXED_PREFILL_CHUNK as usize { + let layers = executor.weights.layers.len(); + let mut expected = vec![0]; + if prompt.len() > 8 { + expected.extend( + (super::PREFILL_DRAIN_LAYERS..layers) + .step_by(super::PREFILL_DRAIN_LAYERS) + .map(|done| (prompt.len() * done / layers) as u32), + ); + } + expected.push(prompt.len() as u32); + assert_eq!( + reported, expected, + "progress must follow completed GPU work" + ); + } }; let recurrent = |executor: &GlmExecutor| { executor @@ -7230,6 +7619,9 @@ mod tests { eprintln!("partial prefill HC workspace guards passed"); } let before = recurrent(&executor); + assert!(executor.rewind_speculative_output(frontier + 1).is_err()); + executor.rewind_speculative_output(frontier).unwrap(); + assert_eq!(recurrent(&executor), before); let first = argmax(executor.logits()); executor.eval(first).unwrap(); let after_first = recurrent(&executor); @@ -7259,7 +7651,11 @@ mod tests { executor.mtp.as_ref().unwrap().rollback, Some((start, first)) ); - assert!(executor.rewind_glm53_mtp(start as usize + keep).unwrap()); + // The UI/headless consumer uses this same path to remove an + // MTP-returned stop token before retaining the ongoing chat. + executor + .rewind_speculative_output(start as usize + keep) + .unwrap(); assert_eq!(executor.position(), start + keep as u32); assert!( recurrent(&executor) == *if keep == 0 { &before } else { &after_first }, diff --git a/src/engine/metal/gpu.rs b/src/engine/metal/gpu.rs index 03781db..ebbdff5 100644 --- a/src/engine/metal/gpu.rs +++ b/src/engine/metal/gpu.rs @@ -45,6 +45,9 @@ pub(super) struct QwenKernelArgs { pub(super) struct GpuCanarySample { pub(super) scheduled_seconds: f64, pub(super) completed_seconds: f64, + pub(super) gpu_wait_seconds: f64, + pub(super) gpu_interval_seconds: f64, + pub(super) host_return_seconds: f64, } #[derive(Clone, Copy, Default)] diff --git a/src/engine/metal/markov.rs b/src/engine/metal/markov.rs new file mode 100644 index 0000000..6c47abe --- /dev/null +++ b/src/engine/metal/markov.rs @@ -0,0 +1,327 @@ +//! DSpark's CPU Markov head: persistent workers, DS4 row partition/reduction. +use super::{F16, F32, Gguf, Q8_0, Weight, dense_dot_bytes, dense_dot_q8, quantize_q8_activation}; +use memmap2::Mmap; +use std::sync::{Arc, mpsc}; +use std::thread::{self, JoinHandle}; + +type Best = (usize, f32); + +struct Input { + values: Vec, + quantized: Option<(Vec, Vec)>, + logits: Vec, +} + +struct Matrix { + map: Arc, + weight: Weight, + width: usize, + rows: usize, + row_bytes: usize, +} + +impl Matrix { + fn best(&self, input: &Input, start: usize, end: usize) -> Best { + let bytes = &self.map[self.weight.offset as usize..]; + let mut best = (start, -f32::MAX); + for token in start..end { + let row = &bytes[token * self.row_bytes..(token + 1) * self.row_bytes]; + let dot = input.quantized.as_ref().map_or_else( + || dense_dot_bytes(self.weight.kind, self.width, row, &input.values), + |(values, scales)| dense_dot_q8(row, values, scales, self.width), + ); + let score = input.logits[token] + dot; + if score > best.1 { + best = (token, score); + } + } + best + } +} + +struct Worker { + input: Option>>, + result: mpsc::Receiver, + thread: Option>, +} + +impl Drop for Worker { + fn drop(&mut self) { + self.input.take(); + if let Some(thread) = self.thread.take() { + let _ = thread.join(); + } + } +} + +pub(super) struct MarkovPool { + matrix: Arc, + workers: Vec, + chunk: usize, +} + +fn worker_count(online: usize, requested: Option<&str>) -> usize { + requested + .and_then(|value| value.parse::().ok()) + .filter(|&value| value > 0) + .unwrap_or(online.min(12)) + .clamp(1, 32) +} + +impl MarkovPool { + pub(super) fn new(model: &Gguf, weight: Weight) -> Result { + let threads = worker_count( + thread::available_parallelism().map_or(1, std::num::NonZero::get), + std::env::var("DS4_THREADS").ok().as_deref(), + ); + Self::from_map(model.shared_map(), weight, threads) + } + + fn from_map(map: Arc, weight: Weight, threads: usize) -> Result { + let width = usize::try_from(weight.dims[0]).map_err(|_| "Markov width overflow")?; + let rows = usize::try_from(weight.dims[1]).map_err(|_| "Markov rows overflow")?; + if width == 0 || rows == 0 || rows > i32::MAX as usize { + return Err("DSpark dense argmax has invalid dimensions".into()); + } + let row_bytes = match weight.kind { + F32 => width.checked_mul(4), + F16 => width.checked_mul(2), + Q8_0 => width.div_ceil(32).checked_mul(34), + _ => None, + } + .ok_or("unsupported DSpark dense tensor layout")?; + let bytes = row_bytes + .checked_mul(rows) + .ok_or("DSpark dense argmax size overflow")?; + if weight.offset > map.len() as u64 || bytes as u64 > map.len() as u64 - weight.offset { + return Err("DSpark dense argmax is outside the GGUF mapping".into()); + } + let matrix = Arc::new(Matrix { + map, + weight, + width, + rows, + row_bytes, + }); + let mut pool = Self { + matrix, + workers: Vec::new(), + chunk: rows.div_ceil(threads), + }; + for slot in 1..threads { + let matrix = Arc::clone(&pool.matrix); + let start = slot * pool.chunk; + let end = (start + pool.chunk).min(rows); + let (input, jobs) = mpsc::sync_channel::>(1); + let (results, result) = mpsc::channel(); + let thread = thread::Builder::new() + .name(format!("dspark-markov-{slot}")) + .spawn(move || { + while let Ok(input) = jobs.recv() { + let best = matrix.best(&input, start, end); + // Return buffer ownership before signalling completion. + drop(input); + if results.send(best).is_err() { + break; + } + } + }) + .map_err(|error| format!("Cannot start DSpark Markov worker: {error}"))?; + pool.workers.push(Worker { + input: Some(input), + result, + thread: Some(thread), + }); + } + Ok(pool) + } + + pub(super) fn argmax(&mut self, values: &[f32], logits: &mut Vec) -> Result { + if values.len() != self.matrix.width || logits.len() != self.matrix.rows { + return Err("DSpark dense argmax has mismatched dimensions".into()); + } + let input = Arc::new(Input { + values: values.to_vec(), + quantized: (self.matrix.weight.kind == Q8_0).then(|| quantize_q8_activation(values)), + logits: std::mem::take(logits), + }); + let parallel = !self.workers.is_empty() && self.matrix.rows >= 512; + let mut failed = false; + if parallel { + for worker in &self.workers { + failed |= worker + .input + .as_ref() + .expect("worker input open") + .send(Arc::clone(&input)) + .is_err(); + } + } + let mut best = self.matrix.best( + &input, + 0, + if parallel { + self.chunk + } else { + self.matrix.rows + }, + ); + if parallel { + // Join results in slot order: equal scores keep the first token. + // Drain every dispatched job even when another worker failed. + for worker in &self.workers { + match worker.result.recv() { + Ok(candidate) if candidate.1 > best.1 => best = candidate, + Ok(_) => {} + Err(_) => failed = true, + } + } + } + *logits = Arc::try_unwrap(input) + .map_err(|_| "DSpark worker retained input")? + .logits; + if failed { + return Err("DSpark Markov worker failed".into()); + } + Ok(best.0 as i32) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + #[ignore = "requires installed 0731 DSpark support; CPU dispatch diagnostic, not parity acceptance"] + fn installed_markov_worker_dispatch() { + use std::time::Instant; + let artifacts = crate::model::engine_artifacts( + crate::model::ModelChoice::DeepSeekV4Flash0731, + true, + &crate::app::models_path(), + ); + let model = Gguf::open(artifacts.support.as_ref().unwrap()).unwrap(); + let weight = Weight::bind(&model, "mtp.2.markov_head.markov_w2.weight").unwrap(); + let state_weight = Weight::bind(&model, "mtp.2.markov_head.markov_w1.weight").unwrap(); + let values = super::super::dense_row(&model, state_weight, 671).unwrap(); + let mut pool12 = MarkovPool::from_map(model.shared_map(), weight, 12).unwrap(); + let mut pool18 = MarkovPool::from_map(model.shared_map(), weight, 18).unwrap(); + let mut logits = vec![0.0; weight.dims[1] as usize]; + let expected = pool12.argmax(&values, &mut logits).unwrap(); + for round in 0..4 { + let order = if round % 2 == 0 { [0, 1, 2] } else { [2, 1, 0] }; + for mode in order { + let started = Instant::now(); + for _ in 0..128 { + let token = if mode == 0 { + // Same row arithmetic/input ownership, old per-call + // dispatch pattern: isolates worker lifetime/count. + let matrix = &pool18.matrix; + let input = Input { + values: values.clone(), + quantized: Some(quantize_q8_activation(&values)), + logits: std::mem::take(&mut logits), + }; + let chunk = matrix.rows.div_ceil(18); + let best = thread::scope(|scope| { + let handles = (0..matrix.rows) + .step_by(chunk) + .map(|start| { + let input = &input; + scope.spawn(move || { + matrix.best(input, start, (start + chunk).min(matrix.rows)) + }) + }) + .collect::>(); + handles.into_iter().map(|h| h.join().unwrap()).fold( + (0, -f32::MAX), + |best, candidate| { + if candidate.1 > best.1 { + candidate + } else { + best + } + }, + ) + }); + logits = input.logits; + best.0 as i32 + } else if mode == 1 { + pool12.argmax(&values, &mut logits).unwrap() + } else { + pool18.argmax(&values, &mut logits).unwrap() + }; + assert_eq!(token, expected); + } + println!( + "markov_dispatch round={round} mode={} calls=128 elapsed_ms={:.3}", + ["scoped18", "persistent12", "persistent18"][mode], + started.elapsed().as_secs_f64() * 1000.0 + ); + } + } + } + + #[test] + fn markov_workers_reuse_threads_and_buffers_with_ordered_ties() { + assert_eq!(worker_count(18, None), 12); + assert_eq!(worker_count(4, Some("0")), 4); + assert_eq!(worker_count(18, Some("99")), 32); + assert_eq!(worker_count(18, Some("1")), 1); + let rows = 1025; + let mut map = memmap2::MmapMut::map_anon(rows * 4).unwrap(); + for (row, bytes) in map.chunks_exact_mut(4).enumerate() { + bytes.copy_from_slice(&((row % 17) as f32).to_le_bytes()); + } + let map = Arc::new(map.make_read_only().unwrap()); + let weight = Weight { + offset: 0, + kind: F32, + bytes: (rows * 4) as u64, + dims: [1, rows as u64, 1], + }; + for threads in [1, 4, 12] { + let mut pool = MarkovPool::from_map(Arc::clone(&map), weight, threads).unwrap(); + let ids = pool + .workers + .iter() + .map(|w| w.thread.as_ref().unwrap().thread().id()) + .collect::>(); + let mut logits = vec![0.0; rows]; + let pointer = logits.as_ptr(); + assert_eq!(pool.argmax(&[2.0], &mut logits).unwrap(), 16); + logits[rows - 1] = 1000.0; + assert_eq!(pool.argmax(&[2.0], &mut logits).unwrap(), (rows - 1) as i32); + assert_eq!(logits.as_ptr(), pointer); + assert!(pool.argmax(&[], &mut logits).is_err()); + assert_eq!(logits.len(), rows); + assert_eq!( + ids, + pool.workers + .iter() + .map(|w| w.thread.as_ref().unwrap().thread().id()) + .collect::>() + ); + if let Some(worker) = pool.workers.first_mut() { + let (closed, receiver) = mpsc::sync_channel(1); + drop(receiver); + worker.input.replace(closed); + worker.thread.take().unwrap().join().unwrap(); + assert!(pool.argmax(&[2.0], &mut logits).is_err()); + assert_eq!(logits.as_ptr(), pointer); + assert_eq!(logits[rows - 1], 1000.0); + } + } + assert!( + MarkovPool::from_map( + map, + Weight { + offset: 1, + ..weight + }, + 1 + ) + .is_err() + ); + } +} diff --git a/src/engine/metal/qwen_mtplx/request.rs b/src/engine/metal/qwen_mtplx/request.rs index cfc3ae1..b3408d9 100644 --- a/src/engine/metal/qwen_mtplx/request.rs +++ b/src/engine/metal/qwen_mtplx/request.rs @@ -357,8 +357,8 @@ fn qwen_request_progress_uses_ui_frontiers_and_decode_timer() { let timing = PromptTiming { evaluated_tokens: 20, eval_seconds: 0.125, - mtp_history_seconds: 0.025, - restore_seconds: 0.010, + mtp_history_seconds: Some(0.025), + restore_seconds: Some(0.010), }; assert_eq!(metrics.snapshot().prompt_timing, None); metrics.set_prompt_timing(Some(timing)); @@ -399,7 +399,7 @@ fn qwen_request_progress_uses_ui_frontiers_and_decode_timer() { let exact_hit = PromptTiming { evaluated_tokens: 0, eval_seconds: 0.0, - mtp_history_seconds: 0.0, + mtp_history_seconds: Some(0.0), ..timing }; metrics.set_prompt_timing(Some(exact_hit)); diff --git a/src/engine/metal/qwen_mtplx/session_cache.rs b/src/engine/metal/qwen_mtplx/session_cache.rs index bcbcb2e..bb7bd83 100644 --- a/src/engine/metal/qwen_mtplx/session_cache.rs +++ b/src/engine/metal/qwen_mtplx/session_cache.rs @@ -1163,8 +1163,8 @@ impl Execution { observe(TurnProgress::PromptReady(crate::metrics::PromptTiming { evaluated_tokens: prompt.len() - cached, eval_seconds: prepared.eval_seconds + history_seconds, - mtp_history_seconds: history_seconds, - restore_seconds, + mtp_history_seconds: Some(history_seconds), + restore_seconds: Some(restore_seconds), }))?; // The bank owns this boundary before decode starts, including when a // subsequent callback aborts the request. Keeping it in this stack diff --git a/src/engine/metal/qwen_mtplx_tests.rs b/src/engine/metal/qwen_mtplx_tests.rs index e631695..db1a076 100644 --- a/src/engine/metal/qwen_mtplx_tests.rs +++ b/src/engine/metal/qwen_mtplx_tests.rs @@ -78,7 +78,8 @@ fn mtplx_canonical_source_bodies_preserve_pinned_hashes() { .split("// Runtime unit SHA256: ") .skip(1) .collect::>(); - assert_eq!(runtime_units.len(), 22); + // tools/mtplx-kernel-source.py --check verifies this complete pinned export. + assert_eq!(runtime_units.len(), 26); for unit in runtime_units { let expected = unit.lines().next().unwrap(); let body = unit diff --git a/src/engine/tokenizer.rs b/src/engine/tokenizer.rs index 9aba847..b7c29d0 100644 --- a/src/engine/tokenizer.rs +++ b/src/engine/tokenizer.rs @@ -1092,6 +1092,108 @@ fn unicode_punctuation(cp: u32) -> bool { mod tests { use super::*; + #[test] + #[ignore = "CPU-only; requires DS4SERVER_CHAT_REFERENCE and its installed GGUF"] + fn ds4_chat_matches_original_session_tokens() { + use super::super::{Model, ModelRef, conversation_tag, render_text_prompt}; + use crate::model::ModelChoice; + let reference = std::env::var_os("DS4SERVER_CHAT_REFERENCE").unwrap(); + let fixture: serde_json::Value = + serde_json::from_slice(&fs::read(reference).unwrap()).unwrap(); + let choice = match fixture["family"].as_str().unwrap() { + "deepseek" => ModelChoice::DeepSeekV4Flash0731, + "glm" => ModelChoice::Glm53Flash, + other => panic!("unsupported reference family: {other}"), + }; + let model = + Model::open_main(Path::new(fixture["model"].as_str().unwrap()), choice).unwrap(); + let tokenizer = &model.tokenizer; + let cases = fixture["cases"].as_array().unwrap(); + assert_eq!(cases.len(), 3); + let mut previous = Vec::new(); + for (index, case) in cases.iter().enumerate() { + let prompt = case["prompt"].as_str().unwrap(); + let expected: Vec = serde_json::from_value(case["tokens"].clone()).unwrap(); + let actual = if index == 0 { + tokenizer.encode_chat("", prompt, ReasoningMode::Low) + } else { + let mut tokens = previous; + tokens.extend(tokenizer.encode_continuation(prompt, ReasoningMode::Low, false)); + tokens + }; + if actual != expected { + let first = actual + .iter() + .zip(&expected) + .position(|(a, b)| a != b) + .unwrap_or(actual.len().min(expected.len())); + let context = |tokens: &[i32]| { + tokens[first.saturating_sub(4)..tokens.len().min(first + 8)] + .iter() + .map(|&id| { + ( + id, + String::from_utf8_lossy(&tokenizer.token_bytes(id).unwrap()) + .into_owned(), + ) + }) + .collect::>() + }; + panic!( + "turn {} first mismatch {first}, lengths {}/{}; Rust {:?}; DS4 {:?}", + index + 1, + actual.len(), + expected.len(), + context(&actual), + context(&expected) + ); + } + if index == 0 { + let settings = crate::settings::TurnSettings { + kv_cache: crate::settings::KvCachePreferences::default().settings(), + context_tokens: 32768, + max_generated_tokens: i32::MAX, + system_prompt: String::new(), + temperature: 0.6, + top_p: 0.95, + min_p: 0.0, + top_k: 0, + stops: vec![], + seed: Some(42), + reasoning_mode: ReasoningMode::Low, + }; + let messages = [ChatTurn { + user: true, + tool: false, + system: false, + skip_previous_eos: false, + reasoning: None, + reasoning_complete: true, + content: prompt.to_owned(), + }]; + let bootstrap = model.render_history("", &[], settings.reasoning_mode); + let rendered = render_text_prompt( + ModelRef::Gguf(&model), + &bootstrap, + conversation_tag("", settings.reasoning_mode, &[]), + &messages, + &settings, + ); + assert!( + rendered == expected, + "shared bootstrap renderer differs: {}/{} tokens; starts {:?}/{:?}", + rendered.len(), + expected.len(), + &rendered[..4], + &expected[..4] + ); + } + previous = expected; + let reply: Vec = serde_json::from_value(case["reply_tokens"].clone()).unwrap(); + previous.extend(reply); + } + } + #[test] #[ignore = "requires local tokenizer and tools/qwen-chat-reference.py output; no inference"] fn qwen_plain_chat_matches_mtplx_server_tokens() { diff --git a/src/main.rs b/src/main.rs index 7c818dd..9ead797 100644 --- a/src/main.rs +++ b/src/main.rs @@ -35,6 +35,14 @@ use app::{App, Message, app_icon, app_theme}; use iced::{Point, Size, window}; fn main() -> iced::Result { + #[cfg(target_os = "macos")] + if std::env::args().nth(1).as_deref() == Some("gpu-canary") { + if let Err(error) = model_eval::run_external_canary() { + eprintln!("GPU canary failed: {error}"); + std::process::exit(1); + } + return Ok(()); + } if std::env::args().nth(1).as_deref() == Some("validate-a2ui") { if let Err(error) = a2ui_validation::run(std::env::args().skip(2)) { eprintln!("A2UI validation failed: {error}"); diff --git a/src/metrics.rs b/src/metrics.rs index 690a608..79c2a58 100644 --- a/src/metrics.rs +++ b/src/metrics.rs @@ -87,8 +87,9 @@ pub(crate) struct SsdStats { pub(crate) struct PromptTiming { pub(crate) evaluated_tokens: usize, pub(crate) eval_seconds: f64, - pub(crate) mtp_history_seconds: f64, - pub(crate) restore_seconds: f64, + /// Optional decomposition: None means not measured separately, not zero. + pub(crate) mtp_history_seconds: Option, + pub(crate) restore_seconds: Option, } #[derive(Clone, Debug, Default)] @@ -1118,6 +1119,26 @@ fn cache_usage(root: &Path) -> CacheUsage { mod tests { use super::*; + #[test] + fn prompt_timing_keeps_unmeasured_components_distinct_from_zero() { + let metrics = Metrics::new(Path::new("/nonexistent-prompt-timing")); + let timing = PromptTiming { + evaluated_tokens: 20, + eval_seconds: 0.25, + mtp_history_seconds: None, + restore_seconds: None, + }; + metrics.set_prompt_timing(Some(timing)); + metrics.kv_write_finished(Duration::from_secs(2), false); + assert_eq!(metrics.snapshot().prompt_timing, Some(timing)); + let json = serde_json::to_value(timing).unwrap(); + assert_eq!(json["eval_seconds"], 0.25); + assert!(json["mtp_history_seconds"].is_null()); + assert!(json["restore_seconds"].is_null()); + metrics.request_started(WorkSource::LocalChat); + assert_eq!(metrics.snapshot().prompt_timing, None); + } + #[test] fn decode_timer_excludes_later_work_and_resets_between_requests() { let metrics = Metrics::new(Path::new("/path/that/does/not/exist")); diff --git a/src/model_eval.rs b/src/model_eval.rs index f0cd0a8..0353789 100644 --- a/src/model_eval.rs +++ b/src/model_eval.rs @@ -317,6 +317,10 @@ impl EvaluationPhase { Self::Preparing => "preparing", } } + + fn from_label(label: &str) -> Option { + Self::ALL.into_iter().find(|phase| phase.label() == label) + } } enum CanaryEvent { @@ -330,6 +334,9 @@ enum CanaryEvent { completion_phase: EvaluationPhase, scheduled_ms: f64, completed_ms: f64, + gpu_wait_ms: Option, + gpu_interval_ms: Option, + host_return_ms: Option, error: Option, }, } @@ -339,6 +346,9 @@ struct CanaryMeasurement { completion_phase: EvaluationPhase, scheduled_ms: f64, completed_ms: f64, + gpu_wait_ms: Option, + gpu_interval_ms: Option, + host_return_ms: Option, } struct GpuCanaryMonitor { @@ -377,13 +387,30 @@ impl GpuCanaryMonitor { break; } let result = gpu_canary_probe(); - let (scheduled_ms, completed_ms, error) = match result { + let ( + scheduled_ms, + completed_ms, + gpu_wait_ms, + gpu_interval_ms, + host_return_ms, + error, + ) = match result { Ok(sample) => ( sample.scheduled_seconds * 1_000.0, sample.completed_seconds * 1_000.0, + sample.gpu_wait_seconds.map(|v| v * 1_000.0), + sample.gpu_interval_seconds.map(|v| v * 1_000.0), + sample.host_return_seconds.map(|v| v * 1_000.0), None, ), - Err(error) => (0.0, at.elapsed().as_secs_f64() * 1_000.0, Some(error)), + Err(error) => ( + 0.0, + at.elapsed().as_secs_f64() * 1_000.0, + None, + None, + None, + Some(error), + ), }; if sender .send(CanaryEvent::Completed { @@ -392,6 +419,9 @@ impl GpuCanaryMonitor { completion_phase: EvaluationPhase::from_u8(phase.load(Ordering::Relaxed)), scheduled_ms, completed_ms, + gpu_wait_ms, + gpu_interval_ms, + host_return_ms, error, }) .is_err() @@ -447,6 +477,9 @@ impl GpuCanaryMonitor { completion_phase, scheduled_ms, completed_ms, + gpu_wait_ms, + gpu_interval_ms, + host_return_ms, error, } => { let probe_started = self.in_flight.take().map(|(at, _)| at); @@ -476,6 +509,9 @@ impl GpuCanaryMonitor { completion_phase, scheduled_ms, completed_ms, + gpu_wait_ms, + gpu_interval_ms, + host_return_ms, }); eprintln!( "{}", @@ -489,6 +525,9 @@ impl GpuCanaryMonitor { "ok": true, "scheduled_ms": scheduled_ms, "completed_ms": completed_ms, + "gpu_wait_ms": gpu_wait_ms, + "gpu_interval_ms": gpu_interval_ms, + "host_return_ms": host_return_ms, }) ); } @@ -557,6 +596,10 @@ impl GpuCanaryMonitor { "samples": completed.len(), "scheduled_ms": latency_summary(scheduled), "completed_ms": latency_summary(completed), + "gpu_timing_samples": self.measurements.iter().filter(|s| s.phase == phase && s.gpu_interval_ms.is_some()).count(), + "gpu_wait_ms": latency_summary(self.measurements.iter().filter(|s| s.phase == phase).filter_map(|s| s.gpu_wait_ms).collect()), + "gpu_interval_ms": latency_summary(self.measurements.iter().filter(|s| s.phase == phase).filter_map(|s| s.gpu_interval_ms).collect()), + "host_return_ms": latency_summary(self.measurements.iter().filter(|s| s.phase == phase).filter_map(|s| s.host_return_ms).collect()), }), ) }) @@ -599,6 +642,51 @@ fn latency_summary(mut values: Vec) -> Value { }) } +/// The same optional probe as the UI/harness, isolated from a reference engine. +/// Stdin contains phase labels, one per line; EOF ends the observation. No model +/// is loaded and no application settings are read or changed. +pub(crate) fn run_external_canary() -> Result<(), String> { + let started = Instant::now(); + let phase = Arc::new(AtomicU8::new(EvaluationPhase::Startup as u8)); + let (send, receive) = mpsc::channel(); + thread::spawn(move || { + for line in io::stdin().lock().lines() { + if send.send(line.map_err(|e| e.to_string())).is_err() { + break; + } + } + }); + let mut monitor = GpuCanaryMonitor::start(Arc::clone(&phase)); + let mut ready = false; + let result = loop { + monitor.poll(started); + if monitor.failures != 0 || monitor.max_in_flight >= BEACHBALL_THRESHOLD { + break Err("probe failed or exceeded the two-second stall limit".into()); + } + if !ready && !monitor.measurements.is_empty() { + if let Err(error) = write_json(json!({"event":"gpu_canary_ready"})) { + break Err(error); + } + ready = true; + } + match receive.recv_timeout(Duration::from_millis(10)) { + Ok(Ok(label)) => match EvaluationPhase::from_label(&label) { + Some(next) => phase.store(next as u8, Ordering::Relaxed), + None => break Err(format!("unknown canary phase {label:?}")), + }, + Ok(Err(error)) => break Err(error), + Err(RecvTimeoutError::Disconnected) => break Ok(()), + Err(RecvTimeoutError::Timeout) => {} + } + }; + monitor.finish(started, BEACHBALL_THRESHOLD); + write_json(monitor.summary())?; + if monitor.failures != 0 || monitor.max_in_flight >= BEACHBALL_THRESHOLD { + return Err("probe failed or exceeded the two-second stall limit".into()); + } + result +} + /// Full-model diagnostic tests reuse the evaluator's probe and phase accounting. /// This is a safety run, not a clean throughput measurement. A stalled/failed /// probe terminates this test worker; the external supervisor reaps its group. @@ -641,7 +729,7 @@ pub(crate) fn run_supervisor(args: impl IntoIterator) -> Result .args(&args) .stdin(Stdio::null()) .stdout(Stdio::piped()) - .stderr(Stdio::inherit()) + .stderr(Stdio::piped()) .spawn() .map_err(|error| format!("could not start evaluator worker: {error}"))?; let progress = Arc::new(Mutex::new(Progress::new())); @@ -651,6 +739,11 @@ pub(crate) fn run_supervisor(args: impl IntoIterator) -> Result .take() .ok_or("evaluator worker stdout is unavailable")?; let output_thread = thread::spawn(move || forward_worker_output(output, output_progress)); + let diagnostics = child + .stderr + .take() + .ok_or("evaluator worker stderr is unavailable")?; + let diagnostics_thread = thread::spawn(move || forward_worker_diagnostics(diagnostics)); let mut next_sample = started; let mut previous = None; let mut maximum = ResourceUsage::default(); @@ -664,6 +757,9 @@ pub(crate) fn run_supervisor(args: impl IntoIterator) -> Result output_thread .join() .map_err(|_| "evaluator output reader panicked".to_owned())??; + diagnostics_thread + .join() + .map_err(|_| "evaluator diagnostics reader panicked".to_owned())??; let progress = *progress .lock() .map_err(|_| "evaluator progress monitor was poisoned")?; @@ -774,6 +870,19 @@ impl Progress { } } +fn forward_worker_diagnostics(output: impl io::Read) -> Result<(), String> { + // The parent and worker must not serialize JSON fragments to the same fd. + // Emit complete child records under the parent's shared stderr lock, also + // used by resource_sample. Diagnostics never extend inference deadlines. + for line in BufReader::new(output).lines() { + let line = + line.map_err(|error| format!("could not read evaluator diagnostics: {error}"))?; + writeln!(io::stderr().lock(), "{line}") + .map_err(|error| format!("could not forward evaluator diagnostics: {error}"))?; + } + Ok(()) +} + fn forward_worker_output( output: impl io::Read, progress: Arc>, @@ -1597,6 +1706,56 @@ fn physical_memory_bytes() -> u64 { mod tests { use super::*; + #[test] + fn external_canary_protocol_accepts_only_known_phases() { + for phase in EvaluationPhase::ALL { + assert_eq!(EvaluationPhase::from_label(phase.label()), Some(phase)); + } + for invalid in ["", "Prefill", "unknown", "decode\nloading"] { + assert_eq!(EvaluationPhase::from_label(invalid), None); + } + } + + #[test] + #[ignore = "Metal probe only; DS4SERVER_CANARY_BINARY must point to the current release binary"] + fn external_canary_observes_phases_and_finishes_on_eof() { + let mut child = Command::new(std::env::var("DS4SERVER_CANARY_BINARY").unwrap()) + .arg("gpu-canary") + .stdin(Stdio::piped()) + .stdout(Stdio::piped()) + .spawn() + .unwrap(); + let mut input = child.stdin.take().unwrap(); + let mut output = BufReader::new(child.stdout.take().unwrap()); + let mut ready = String::new(); + output.read_line(&mut ready).unwrap(); + assert_eq!( + serde_json::from_str::(&ready).unwrap()["event"], + "gpu_canary_ready" + ); + for phase in ["prefill", "decode"] { + writeln!(input, "{phase}").unwrap(); + thread::sleep(Duration::from_millis(500)); + } + drop(input); + let mut summary = String::new(); + output.read_line(&mut summary).unwrap(); + assert!(child.wait().unwrap().success()); + let summary: Value = serde_json::from_str(&summary).unwrap(); + assert_eq!(summary["event"], "gpu_canary_summary"); + assert_eq!(summary["failures"], 0); + for phase in ["prefill", "decode"] { + let phase = &summary["phases"][phase]; + assert!(phase["samples"].as_u64().unwrap() > 0); + assert_eq!(phase["gpu_timing_samples"], phase["samples"]); + for field in ["gpu_wait_ms", "gpu_interval_ms", "host_return_ms"] { + let max = phase[field]["max"].as_f64().unwrap(); + assert!(max.is_finite() && max >= 0.0); + assert!(max <= phase["completed_ms"]["max"].as_f64().unwrap() + 0.01); + } + } + } + #[test] fn warmup_timing_separates_loading_without_inventing_missing_durations() { let elapsed = Duration::from_millis(1234); @@ -1631,6 +1790,9 @@ mod tests { completion_phase: EvaluationPhase::Warmup, scheduled_ms: 0.1, completed_ms: 2500.0, + gpu_wait_ms: None, + gpu_interval_ms: None, + host_return_ms: None, error: None, }) .unwrap(); @@ -1639,6 +1801,8 @@ mod tests { assert_eq!(summary["phase_assignment"], "probe_start"); assert_eq!(summary["cross_phase_samples"], 1); assert_eq!(summary["phases"]["loading"]["samples"], 1); + assert_eq!(summary["phases"]["loading"]["gpu_timing_samples"], 0); + assert!(summary["phases"]["loading"]["gpu_interval_ms"]["max"].is_null()); assert_eq!(summary["phases"]["warmup"]["samples"], 0); assert_eq!(summary["beachball_risk_observed"], true); } diff --git a/tests/fixtures/ds4-sampling-ec7642c.json b/tests/fixtures/ds4-sampling-ec7642c.json new file mode 100644 index 0000000..2dd9767 --- /dev/null +++ b/tests/fixtures/ds4-sampling-ec7642c.json @@ -0,0 +1 @@ +{"reference":"antirez/ds4","commit":"ec7642cdd9ec81d01ad4b1fd8f8a3d1511533748","recipe":"n=4: [2,2,0,-2]; otherwise ((i*37%101)-50)/8 in f32","cases":[{"n":4,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0],"rng_after":0},{"n":4,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0],"rng_after":42},{"n":4,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[0,0,1,0,1,1,1,0,1,1,0,1,0,1,0,0,0,0,1,0,1,1,1,0,1,0,1,1,0,0,0,1],"rng_after":7191875387263428708},{"n":4,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[0,1,1,1,1,1,0,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,0,0,1,1,1,1,0,0],"rng_after":10733397791740853619},{"n":4,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":0,"tokens":[0,0,1,0,1,1,1,0,1,1,0,1,0,1,0,0,0,0,1,0,1,1,1,0,1,0,1,1,0,0,0,1],"rng_after":7191875387263428708},{"n":4,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":42,"tokens":[0,1,1,1,1,1,0,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,0,0,1,1,1,1,0,0],"rng_after":10733397791740853619},{"n":4,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":0,"tokens":[0,0,1,0,1,1,1,0,1,1,0,1,0,1,0,0,0,0,1,0,1,1,1,0,1,0,1,1,0,0,0,1],"rng_after":7191875387263428708},{"n":4,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":42,"tokens":[0,1,1,1,1,1,0,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,0,0,1,1,1,1,0,0],"rng_after":10733397791740853619},{"n":4,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0],"rng_after":7191875387263428708},{"n":4,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0],"rng_after":10733397791740853619},{"n":4,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":0,"tokens":[0,0,1,0,1,1,1,0,1,1,0,1,0,1,0,0,0,0,1,0,1,1,1,0,1,0,1,1,0,0,0,1],"rng_after":7191875387263428708},{"n":4,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":42,"tokens":[0,1,1,1,1,1,0,0,0,1,0,0,1,0,0,1,0,0,2,0,0,1,0,0,0,0,1,1,1,1,0,0],"rng_after":10733397791740853619},{"n":4,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":0,"tokens":[0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0],"rng_after":0},{"n":4,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":42,"tokens":[0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0],"rng_after":42},{"n":4,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[0,0,1,1,1,1,1,0,1,1,0,1,0,1,0,0,0,0,1,0,1,1,1,0,1,0,1,1,0,0,1,1],"rng_after":7191875387263428708},{"n":4,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[0,1,1,2,1,1,0,0,0,1,0,0,1,0,0,1,0,0,2,0,0,1,0,0,0,0,1,1,1,1,0,0],"rng_after":10733397791740853619},{"n":64,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":0},{"n":64,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":42},{"n":64,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,60,49,60,19,19,8,60,19,8,30,49,30,27,30,60,30,30,38,30,19,49,49,30,49,60,19,38,30,30,60,19],"rng_after":7191875387263428708},{"n":64,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[60,8,8,27,49,8,30,60,60,49,60,60,38,30,30,19,60,60,16,60,60,49,60,60,60,30,19,49,19,8,30,30],"rng_after":10733397791740853619},{"n":64,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":0,"tokens":[8,30,49,30,38,38,60,30,30,60,8,49,19,60,8,30,19,8,60,19,49,49,49,19,49,30,30,60,8,8,30,30],"rng_after":7191875387263428708},{"n":64,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":42,"tokens":[30,60,60,60,60,60,19,30,30,49,30,30,60,19,8,38,30,30,60,30,30,49,30,30,30,30,38,49,38,60,19,19],"rng_after":10733397791740853619},{"n":64,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":0,"tokens":[30,30,60,30,60,60,60,30,30,60,30,60,30,60,30,30,30,30,60,30,60,60,60,30,60,30,30,60,30,30,30,30],"rng_after":7191875387263428708},{"n":64,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":42,"tokens":[30,60,60,60,60,60,30,30,30,60,30,30,60,30,30,60,30,30,60,30,30,60,30,30,30,30,60,60,60,60,30,30],"rng_after":10733397791740853619},{"n":64,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":7191875387263428708},{"n":64,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":10733397791740853619},{"n":64,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":0,"tokens":[30,60,49,19,19,19,38,60,19,38,30,49,30,57,30,60,30,30,38,30,49,49,49,30,49,60,19,27,30,30,60,19],"rng_after":7191875387263428708},{"n":64,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":42,"tokens":[60,8,8,16,8,38,30,60,60,49,60,60,38,30,30,19,60,60,2,60,60,49,60,60,60,60,19,49,19,8,30,30],"rng_after":10733397791740853619},{"n":64,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":0},{"n":64,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":42},{"n":64,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,60,8,49,49,8,57,19,49,57,30,38,30,5,30,60,30,30,16,60,8,38,38,60,38,19,49,46,30,30,49,49],"rng_after":7191875387263428708},{"n":64,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[60,27,27,35,27,57,60,19,60,38,19,60,16,60,30,8,60,60,13,60,19,38,19,60,19,60,8,38,49,27,60,60],"rng_after":10733397791740853619},{"n":513,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":0},{"n":513,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":42},{"n":513,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[131,464,352,19,322,150,240,191,221,139,30,483,333,27,131,464,333,30,68,161,150,453,281,161,453,191,120,472,30,131,494,221],"rng_after":7191875387263428708},{"n":513,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[464,210,311,431,109,341,60,393,363,453,292,464,371,434,131,150,363,363,147,464,292,453,90,464,464,262,49,281,423,311,434,60],"rng_after":10733397791740853619},{"n":513,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":0,"tokens":[30,161,333,240,281,322,423,210,262,423,8,363,60,464,30,161,60,8,434,120,322,333,352,131,333,210,262,442,8,30,232,262],"rng_after":7191875387263428708},{"n":513,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":42,"tokens":[161,393,393,464,382,423,109,232,150,333,221,161,434,90,30,311,161,150,494,161,221,333,191,161,180,131,292,363,292,393,90,109],"rng_after":10733397791740853619},{"n":513,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":0,"tokens":[30,30,131,30,131,131,131,30,131,131,30,131,30,131,30,30,30,30,131,30,131,131,131,30,131,30,131,131,30,30,30,131],"rng_after":7191875387263428708},{"n":513,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":42,"tokens":[30,131,131,131,131,131,30,30,30,131,30,30,131,30,30,131,30,30,131,30,30,131,30,30,30,30,131,131,131,131,30,30],"rng_after":10733397791740853619},{"n":513,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":7191875387263428708},{"n":513,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":10733397791740853619},{"n":513,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":0,"tokens":[131,464,79,120,49,352,371,292,322,169,30,210,434,188,232,464,333,30,199,161,352,281,8,161,180,292,221,229,30,131,19,322],"rng_after":7191875387263428708},{"n":513,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":42,"tokens":[464,139,240,420,412,371,60,494,363,180,393,464,27,60,131,251,464,464,406,464,393,281,191,90,90,262,150,8,49,341,60,60],"rng_after":10733397791740853619},{"n":513,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":0},{"n":513,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":42},{"n":513,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[232,19,442,79,109,38,390,423,382,188,30,401,161,409,333,19,60,30,218,90,139,169,98,90,68,423,281,450,30,131,453,382],"rng_after":7191875387263428708},{"n":513,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[19,431,158,166,229,390,363,150,393,169,49,19,248,363,232,38,494,494,316,19,150,169,322,19,120,292,311,199,109,158,363,363],"rng_after":10733397791740853619},{"n":4096,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":0},{"n":4096,"temperature":0.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":42},{"n":4096,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[1040,3393,2978,19,2948,1059,2260,1403,1635,1452,30,3210,3161,27,1545,3292,2959,232,68,969,1362,3786,1998,1070,3584,1504,1130,3401,232,939,3625,1736],"rng_after":7191875387263428708},{"n":4096,"temperature":0.6000000238418579,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[3595,2230,2836,3764,917,2462,161,2615,2585,3685,2110,3595,2290,3868,1545,857,3191,2888,652,3292,2413,3988,797,3696,3999,1878,4059,2301,3352,3038,3969,161],"rng_after":10733397791740853619},{"n":4096,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":0,"tokens":[221,1351,2686,2009,2312,2533,3412,1646,2170,3371,8,3019,625,3767,311,1351,595,38,3494,969,2555,2757,2937,988,2757,1646,2129,3625,30,191,1949,2181],"rng_after":7191875387263428708},{"n":4096,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":0.20000000298023224,"seed":42,"tokens":[1373,3191,3240,3868,3139,3423,838,1796,1242,2757,1747,1392,3584,767,292,2514,1343,1272,4078,1343,1777,2776,1564,1403,1452,1130,2443,2959,2353,3262,797,838],"rng_after":10733397791740853619},{"n":4096,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":0,"tokens":[30,30,131,30,131,131,131,30,131,131,30,131,30,131,30,30,30,30,131,30,131,131,131,30,131,30,131,131,30,30,30,131],"rng_after":7191875387263428708},{"n":4096,"temperature":0.6000000238418579,"top_k":3,"top_p":0.5,"min_p":0.10000000149011612,"seed":42,"tokens":[30,131,131,131,131,131,30,30,30,131,30,30,131,30,30,131,30,30,131,30,30,131,30,30,30,30,131,131,131,131,30,30],"rng_after":10733397791740853619},{"n":4096,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":7191875387263428708},{"n":4096,"temperature":0.6000000238418579,"top_k":1,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":10733397791740853619},{"n":4096,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":0,"tokens":[1141,3696,382,928,3958,2372,1886,2009,2544,876,30,1422,3262,3693,1646,3696,3060,232,401,1171,2675,1392,3917,1373,1190,2009,2039,835,232,939,322,2746],"rng_after":7191875387263428708},{"n":4096,"temperature":0.6000000238418579,"top_k":2048,"top_p":0.9990000128746033,"min_p":0.0,"seed":42,"tokens":[3898,947,1654,1632,3442,2189,363,3221,2989,1291,2716,3999,3330,4070,1545,2170,3595,3191,2609,3696,3019,1695,1302,90,494,2181,1261,210,352,1957,60,363],"rng_after":10733397791740853619},{"n":4096,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":0,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":0},{"n":4096,"temperature":0.6000000238418579,"top_k":0,"top_p":1.0,"min_p":1.100000023841858,"seed":42,"tokens":[30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30,30],"rng_after":42},{"n":4096,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":0,"tokens":[1747,120,3371,79,715,240,2410,3655,2604,996,30,2623,969,2934,2454,120,767,333,1632,191,644,674,502,393,472,3655,1796,3177,333,1545,3281,2705],"rng_after":7191875387263428708},{"n":4096,"temperature":1.0,"top_k":0,"top_p":0.949999988079071,"min_p":0.0,"seed":42,"tokens":[524,3562,461,772,1542,2713,2989,1463,3019,674,554,524,1157,2383,2454,3947,4029,3423,2336,120,1160,1078,2645,726,1332,1807,2735,1108,1422,865,2585,2989],"rng_after":10733397791740853619}]} diff --git a/tools/build-ds4-session-reference.sh b/tools/build-ds4-session-reference.sh new file mode 100644 index 0000000..fb312a2 --- /dev/null +++ b/tools/build-ds4-session-reference.sh @@ -0,0 +1,31 @@ +#!/bin/bash +# Standalone oracle only. Never add its C objects to Cargo or the app bundle. +set -euo pipefail +if [[ $# != 2 ]]; then + echo "usage: bash tools/build-ds4-session-reference.sh REFERENCE_DIRECTORY OUTPUT_BINARY" >&2 + exit 2 +fi +reference=$(cd "$1" && pwd) +output=$2 +cd "$(dirname "$0")/.." +[[ $(git -C "$reference" rev-parse HEAD) == ec7642cdd9ec81d01ad4b1fd8f8a3d1511533748 ]] +[[ $(shasum -a 256 "$reference/ds4.h" | cut -d ' ' -f 1) == 1fe2491b0b709222dd41076d35b6eb85684d8c467a691ff4d45534d329f06360 ]] +git -C "$reference" diff --exit-code HEAD -- ds4.h ds4.c ds4_metal.m ds4_image.c ds4_distributed.c ds4_tp.c ds4_ssd.c ds4_layer_pack.c metal +# Refuse stale objects. If necessary, build the authorized reference separately. +make -q -C "$reference" ds4 +json_library=$(cargo build --release --bin test-supervisor -j 4 --message-format=json | + jq -rs '[.[] | select(.reason == "compiler-artifact" and .target.name == "serde_json") | .filenames[] | select(endswith(".rlib"))] | last') +[[ -n $json_library ]] +probe_dir=$(mktemp -d) +trap 'rm -f "$probe_dir/canary.o"; rmdir "$probe_dir"' EXIT +clang -c -O3 -fobjc-arc -ffast-math -mcpu=native -I native/metal \ + native/metal/ds4_canary.m -o "$probe_dir/canary.o" +link=() +for object in ds4 ds4_image ds4_distributed ds4_tp ds4_ssd ds4_metal ds4_layer_pack; do + link+=(-C "link-arg=$reference/$object.o") +done +rustc --edition 2024 -O -D warnings tools/ds4-session-reference.rs \ + --extern "serde_json=$json_library" -L dependency=target/release/deps \ + -C "link-arg=$probe_dir/canary.o" \ + "${link[@]}" -l framework=Foundation -l framework=Metal -l objc -l pthread -l m \ + -o "$output" diff --git a/tools/ds4-session-reference.rs b/tools/ds4-session-reference.rs new file mode 100644 index 0000000..0778c07 --- /dev/null +++ b/tools/ds4-session-reference.rs @@ -0,0 +1,775 @@ +//! Standalone reference benchmark, NEVER linked into DS4Server/Cargo targets. +//! Uses antirez/ds4's public session API, pinned to ec7642c (ds4.h ABI, arm64). +//! Build separately with the original reference objects; run with test-supervisor. +//! Arguments: MODEL FAMILY(glm|deepseek) ACCELERATION(on|off) README [SUPPORT_GGUF] +//! Power 100, Low, context 32768, temp .6/top-p .95/top-k 0/min-p 0/seed 42. +//! Separate 32-token warmup, then one ongoing Summary/Story/Python chat to EOS. +//! Optional DS4_REFERENCE_CANARY=/path/to/ds4-server uses the production probe +//! in a separate process. Leave unset for clean throughput measurements. +use serde_json::json; +use std::ffi::{CStr, CString, c_char, c_void}; +use std::io::{self, BufRead, BufReader, Write}; +use std::mem::{offset_of, size_of}; +use std::process::{Child, Command, Stdio}; +use std::ptr::{null, null_mut}; +use std::sync::{ + Arc, Mutex, + atomic::{AtomicBool, Ordering}, + mpsc, +}; +use std::time::{Duration, Instant}; + +#[path = "sampler-replay-benchmark.rs"] +mod sampler_replay_benchmark; + +// Same native probe object as DS4Server, linked only into this standalone oracle. +#[repr(C)] +#[derive(Default)] +struct CanarySample { + scheduled: f64, + completed: f64, + gpu_wait: f64, + gpu_interval: f64, + host_return: f64, +} +unsafe extern "C" { + fn ds4_gpu_canary_probe(sample: *mut CanarySample) -> i32; + fn ds4_engine_vocab_size(engine: *mut c_void) -> i32; + fn ds4_engine_power(engine: *mut c_void) -> i32; + fn ds4_session_prefill_cap(session: *mut c_void) -> i32; + fn ds4_session_copy_logits(session: *mut c_void, out: *mut f32, cap: i32) -> i32; + fn ds4_sample_logits( + logits: *const f32, + n: i32, + temperature: f32, + top_k: i32, + top_p: f32, + min_p: f32, + rng: *mut u64, + ) -> i32; +} + +// CPU-only oracle fixture: no weights, model initialization or Metal work. +fn sampler_fixture() { + let mut cases = Vec::new(); + for n in [4, 64, 513, 4096] { + let logits: Vec = if n == 4 { + vec![2.0, 2.0, 0.0, -2.0] + } else { + (0..n) + .map(|i| ((i * 37 % 101) as f32 - 50.0) / 8.0) + .collect() + }; + for (temperature, top_k, top_p, min_p) in [ + (0.0, 0, 0.95, 0.0), + (0.6, 0, 0.95, 0.0), + (0.6, 0, 1.0, 0.2), + (0.6, 3, 0.5, 0.1), + (0.6, 1, 0.95, 0.0), + (0.6, 2048, 0.999, 0.0), + (0.6, 0, 1.0, 1.1), + (1.0, 0, 0.95, 0.0), + ] { + for seed in [0, 42] { + let mut rng = seed; + let tokens: Vec = (0..32) + .map(|_| unsafe { + ds4_sample_logits( + logits.as_ptr(), + n, + temperature, + top_k, + top_p, + min_p, + &mut rng, + ) + }) + .collect(); + cases.push(json!({"n":n,"temperature":temperature,"top_k":top_k,"top_p":top_p,"min_p":min_p,"seed":seed,"tokens":tokens,"rng_after":rng})); + } + } + } + emit( + json!({"reference":"antirez/ds4", "commit":"ec7642cdd9ec81d01ad4b1fd8f8a3d1511533748", "recipe":"n=4: [2,2,0,-2]; otherwise ((i*37%101)-50)/8 in f32", "cases":cases}), + ); +} +struct InlineProbe { + phase: Arc>, + stop: Arc, + worker: Option>>, +} +impl InlineProbe { + fn start() -> Result { + assert_eq!(size_of::(), 40); + let phase = Arc::new(Mutex::new("startup")); + let stop = Arc::new(AtomicBool::new(false)); + let worker_phase = Arc::clone(&phase); + let worker_stop = Arc::clone(&stop); + let (ready_send, ready_recv) = mpsc::sync_channel(1); + let worker = std::thread::spawn(move || { + let started = Instant::now(); + let mut first = true; + let mut samples = 0; + loop { + let at = Instant::now(); + let phase = *worker_phase.lock().unwrap(); + let mut sample = CanarySample::default(); + let ok = unsafe { ds4_gpu_canary_probe(&mut sample) } != 0; + let ms = |s: f64| (s >= 0.0).then_some(s * 1000.0); + eprintln!( + "{}", + json!({"event":"gpu_canary_sample", "observer":"in_process", "elapsed_ms":started.elapsed().as_millis(), "started_elapsed_ms":at.duration_since(started).as_millis(), "phase":phase, "completion_phase":*worker_phase.lock().unwrap(), "ok":ok, "scheduled_ms":ms(sample.scheduled), "completed_ms":ms(sample.completed), "gpu_wait_ms":ms(sample.gpu_wait), "gpu_interval_ms":ms(sample.gpu_interval), "host_return_ms":ms(sample.host_return)}) + ); + if first { + let _ = ready_send.send(ok); + first = false; + } + if !ok { + return Err("in-process Metal canary failed".into()); + } + samples += 1; + if worker_stop.load(Ordering::Relaxed) { + emit( + json!({"event":"gpu_canary_summary", "observer":"in_process", "samples":samples, "failures":0}), + ); + return Ok(()); + } + if let Some(wait) = Duration::from_millis(100).checked_sub(at.elapsed()) { + std::thread::sleep(wait); + } + } + }); + let probe = Self { + phase, + stop, + worker: Some(worker), + }; + if ready_recv.recv_timeout(Duration::from_secs(2)) != Ok(true) { + return Err("in-process canary did not become ready".into()); + } + emit(json!({"event":"gpu_canary_ready", "observer":"in_process"})); + Ok(probe) + } + fn finish(&mut self) -> Result<(), String> { + self.stop.store(true, Ordering::Relaxed); + let deadline = Instant::now() + Duration::from_secs(2); + while self.worker.as_ref().is_some_and(|w| !w.is_finished()) { + if Instant::now() >= deadline { + return Err("in-process canary did not finish".into()); + } + std::thread::sleep(Duration::from_millis(10)); + } + if let Some(worker) = self.worker.take() { + worker.join().map_err(|_| "in-process canary panicked")??; + } + Ok(()) + } +} +impl Drop for InlineProbe { + fn drop(&mut self) { + self.stop.store(true, Ordering::Relaxed); + } +} + +struct Probe { + child: Option, + reader: Option>>, + inline: Option, +} +impl Probe { + fn start() -> Result { + let child = std::env::var_os("DS4_REFERENCE_CANARY") + .map(|path| { + Command::new(path) + .arg("gpu-canary") + .stdin(Stdio::piped()) + .stdout(Stdio::piped()) + .spawn() + }) + .transpose() + .map_err(|e| e.to_string())?; + let mut probe = Self { + child, + reader: None, + inline: None, + }; + match std::env::var("DS4_REFERENCE_IN_PROCESS_CANARY").as_deref() { + Ok("1") => probe.inline = Some(InlineProbe::start()?), + Err(std::env::VarError::NotPresent) | Ok("0") => {} + _ => return Err("DS4_REFERENCE_IN_PROCESS_CANARY must be 0 or 1".into()), + } + if let Some(child) = &mut probe.child { + let mut output = BufReader::new(child.stdout.take().ok_or("canary stdout missing")?); + let mut ready = String::new(); + output.read_line(&mut ready).map_err(|e| e.to_string())?; + let ready: serde_json::Value = + serde_json::from_str(&ready).map_err(|e| e.to_string())?; + if ready["event"] != "gpu_canary_ready" { + return Err(format!("canary not ready: {ready}")); + } + emit(ready); + probe.reader = Some(std::thread::spawn(move || { + for line in output.lines() { + let line = line.map_err(|e| e.to_string())?; + emit(serde_json::from_str(&line).map_err(|e| e.to_string())?); + } + Ok(()) + })); + } + Ok(probe) + } + fn phase(&mut self, phase: &'static str) -> Result<(), String> { + if let Some(probe) = &self.inline { + if probe.worker.as_ref().is_some_and(|w| w.is_finished()) { + return Err("in-process canary exited early".into()); + } + *probe.phase.lock().unwrap() = phase; + } + if let Some(child) = &mut self.child { + if let Some(status) = child.try_wait().map_err(|e| e.to_string())? { + return Err(format!("canary exited early: {status}")); + } + writeln!( + child.stdin.as_mut().ok_or("canary stdin closed")?, + "{phase}" + ) + .map_err(|e| e.to_string())?; + } + Ok(()) + } + fn finish(&mut self) -> Result<(), String> { + self.phase("finishing")?; + if let Some(probe) = &mut self.inline { + probe.finish()?; + } + if let Some(child) = &mut self.child { + drop(child.stdin.take()); + let deadline = Instant::now() + Duration::from_secs(3); + loop { + if let Some(status) = child.try_wait().map_err(|e| e.to_string())? { + if let Some(reader) = self.reader.take() { + reader.join().map_err(|_| "canary reader panicked")??; + } + return if status.success() { + Ok(()) + } else { + Err(format!("canary failed: {status}")) + }; + } + if Instant::now() >= deadline { + return Err("canary did not finish".into()); + } + std::thread::sleep(Duration::from_millis(10)); + } + } + Ok(()) + } +} +impl Drop for Probe { + fn drop(&mut self) { + if let Some(child) = &mut self.child { + let _ = child.kill(); + let _ = child.wait(); + } + if let Some(reader) = self.reader.take() { + let _ = reader.join(); + } + } +} + +// Public reference header declarations, not a port of inference code. +#[repr(C)] +#[derive(Default)] +struct Tokens { + v: *mut i32, + len: i32, + cap: i32, +} +#[repr(C)] +#[derive(Default)] +struct Layers { + start: u32, + end: u32, + output: bool, + set: bool, +} +#[repr(C)] +#[derive(Default)] +struct Distributed { + role: i32, + layers: Layers, + listen_host: *const c_char, + listen_port: i32, + coordinator_host: *const c_char, + coordinator_port: i32, + prefill_chunk: u32, + prefill_window: u32, + activation_bits: u32, + replay_check: bool, + debug: bool, +} +#[repr(C)] +#[derive(Default)] +struct Tp { + role: i32, + requested: bool, + listen_host: *const c_char, + listen_port: i32, + leader_host: *const c_char, + leader_port: i32, + transport: i32, + rdma_device: *const c_char, + rdma_gid_index: i32, + rdma_gid_index_set: bool, + glm_token_prefill: bool, + debug_hash: i32, +} +#[repr(C)] +#[derive(Default)] +struct Options { + model: *const c_char, + mtp: *const c_char, + vision: *const c_char, + backend: i32, + threads: i32, + context: i32, + prefill_chunk: u32, + mtp_draft_tokens: i32, + mtp_margin: f32, + dspark_confidence_threshold: f32, + steering: *const c_char, + expert_profile: *const c_char, + steering_attn: f32, + steering_ffn: f32, + power: i32, + cache_experts: u32, + cache_bytes: u64, + full_layers: u32, + preload_experts: u32, + simulated_memory: u64, + warm_weights: bool, + quality: bool, + glm_mtp: bool, + glm_mtp_timing: bool, + dspark: bool, + dspark_strict: bool, + dspark_exact: bool, + confidence_set: bool, + cuda_tp: bool, + ssd: bool, + ssd_cold: bool, + full_layers_set: bool, + inspect: bool, + placement_ctx: i32, + placement_sessions: i32, + share_prefill: bool, + first_token_test: bool, + graph_test: bool, + load_slice: bool, + load_start: u32, + load_end: u32, + load_output: bool, + distributed: Distributed, + tp: Tp, +} +unsafe extern "C" { + fn ds4_engine_open(out: *mut *mut c_void, options: *const Options) -> i32; + fn ds4_engine_close(engine: *mut c_void); + fn ds4_engine_mtp_draft_tokens(engine: *mut c_void) -> i32; + fn ds4_session_create(out: *mut *mut c_void, engine: *mut c_void, ctx: i32) -> i32; + fn ds4_session_free(session: *mut c_void); + fn ds4_session_set_progress( + s: *mut c_void, + f: unsafe extern "C" fn(*mut c_void, *const c_char, i32, i32), + ud: *mut c_void, + ); + fn ds4_session_set_display_progress( + s: *mut c_void, + f: unsafe extern "C" fn(*mut c_void, *const c_char, i32, i32), + ud: *mut c_void, + ); + fn ds4_session_sync(s: *mut c_void, tokens: *const Tokens, err: *mut c_char, len: usize) + -> i32; + fn ds4_session_common_prefix(s: *mut c_void, tokens: *const Tokens) -> i32; + fn ds4_session_sample( + s: *mut c_void, + temp: f32, + top_k: i32, + top_p: f32, + min_p: f32, + rng: *mut u64, + ) -> i32; + fn ds4_session_eval(s: *mut c_void, token: i32, err: *mut c_char, len: usize) -> i32; + fn ds4_session_pos(s: *mut c_void) -> i32; + fn ds4_session_eval_speculative( + s: *mut c_void, + first: i32, + max: i32, + eos: i32, + temp: f32, + top_k: i32, + top_p: f32, + min_p: f32, + rng: *mut u64, + out: *mut i32, + cap: i32, + err: *mut c_char, + len: usize, + ) -> i32; + fn ds4_session_rewind(s: *mut c_void, pos: i32); + fn ds4_tokens_push(tokens: *mut Tokens, token: i32); + fn ds4_tokens_free(tokens: *mut Tokens); + fn ds4_chat_begin(e: *mut c_void, tokens: *mut Tokens); + fn ds4_chat_append_message( + e: *mut c_void, + tokens: *mut Tokens, + role: *const c_char, + text: *const c_char, + ); + fn ds4_chat_append_assistant_prefix(e: *mut c_void, tokens: *mut Tokens, think: i32); + fn ds4_token_text(e: *mut c_void, token: i32, len: *mut usize) -> *mut c_char; + fn ds4_token_eos(e: *mut c_void) -> i32; + fn ds4_token_is_stop_for_think_mode(e: *mut c_void, token: i32, think: i32) -> bool; + fn free(p: *mut c_void); +} +struct Engine(*mut c_void); +impl Drop for Engine { + fn drop(&mut self) { + unsafe { ds4_engine_close(self.0) }; + } +} +struct Session(*mut c_void); +impl Drop for Session { + fn drop(&mut self) { + unsafe { ds4_session_free(self.0) }; + } +} +impl Drop for Tokens { + fn drop(&mut self) { + unsafe { ds4_tokens_free(self) }; + } +} + +fn emit(value: serde_json::Value) { + let mut out = io::stdout().lock(); + if writeln!(out, "{value}").and_then(|()| out.flush()).is_err() { + std::process::exit(1); + } +} +unsafe extern "C" fn progress(_: *mut c_void, event: *const c_char, current: i32, total: i32) { + emit( + json!({"event":"reference_progress", "kind":unsafe { CStr::from_ptr(event) }.to_string_lossy(), "current":current, "total":total}), + ); +} +fn check(code: i32, err: &[c_char]) -> Result<(), String> { + if code == 0 { + Ok(()) + } else { + Err(format!( + "reference error {code}: {}", + unsafe { CStr::from_ptr(err.as_ptr()) }.to_string_lossy() + )) + } +} +fn session(engine: &Engine) -> Result { + let mut s = null_mut(); + if unsafe { ds4_session_create(&mut s, engine.0, 32768) } != 0 || s.is_null() { + return Err("reference session creation failed".into()); + } + unsafe { + ds4_session_set_progress(s, progress, null_mut()); + ds4_session_set_display_progress(s, progress, null_mut()); + } + Ok(Session(s)) +} +fn turn( + engine: &Engine, + session: &Session, + tokens: &mut Tokens, + prompt: &str, + glm: bool, + index: usize, + max: i32, + probe: &mut Probe, +) -> Result<(), String> { + let started = Instant::now(); + probe.phase("prefill")?; + let mut err = [0; 256]; + let mut bootstrap_ms = 0.0; + let prompt = CString::new(prompt).map_err(|e| e.to_string())?; + if tokens.len == 0 { + unsafe { + ds4_chat_begin(engine.0, tokens); + if glm { + // Low is prompt policy; the public CLI aliases it to High. + ds4_chat_append_message( + engine.0, + tokens, + c"system".as_ptr(), + c"Reasoning Effort: Low".as_ptr(), + ); + } + } + // The UI's shared runtime prepares its system-prefix KV separately. + // Reproduce that boundary; a cold full-prompt batch can round differently. + let bootstrap = Instant::now(); + check( + unsafe { ds4_session_sync(session.0, tokens, err.as_mut_ptr(), err.len()) }, + &err, + )?; + bootstrap_ms = bootstrap.elapsed().as_secs_f64() * 1000.0; + } else if !glm { + unsafe { + ds4_tokens_push(tokens, ds4_token_eos(engine.0)); + } + } + unsafe { + ds4_chat_append_message(engine.0, tokens, c"user".as_ptr(), prompt.as_ptr()); + ds4_chat_append_assistant_prefix(engine.0, tokens, 1); + } + let cached = unsafe { ds4_session_common_prefix(session.0, tokens) }; + let input = tokens.len; + emit( + json!({"event":"reference_prefill", "turn":index, "prompt_tokens":input, "cached_tokens":cached, "session_prefill_cap":unsafe { ds4_session_prefill_cap(session.0) }, "tokens":unsafe { std::slice::from_raw_parts(tokens.v, input as usize) }}), + ); + let prefill = Instant::now(); + check( + unsafe { ds4_session_sync(session.0, tokens, err.as_mut_ptr(), err.len()) }, + &err, + )?; + let prefill_ms = prefill.elapsed().as_secs_f64() * 1000.0; + emit(json!({"event":"reference_decode", "turn":index, "prefill_ms":prefill_ms})); + probe.phase("decode")?; + let mut rng = 42_u64; + let mut ids = Vec::new(); + let mut text = Vec::new(); + let decode = Instant::now(); + let mut stop = None; + let trace_cycles = std::env::var_os("DS4_SPEC_CYCLE_TRACE").is_some(); + let mut trace = if index == 1 { + std::env::var_os("DS4_REFERENCE_LOGITS_TRACE").map(|path| { + let n = unsafe { ds4_engine_vocab_size(engine.0) }; + if !(1..=1_000_000).contains(&n) { return Err("invalid vocabulary size".to_owned()); } + let file = std::fs::OpenOptions::new().create_new(true).write(true).open(&path).map_err(|e| e.to_string())?; + emit(json!({"event":"reference_logits_trace", "path":path, "vocab":n, "rows":32, "turn":index, "format":"little-endian-f32"})); + Ok((file, vec![0.0_f32; n as usize])) + }).transpose()? + } else { + None + }; + while ids.len() < max as usize && tokens.len < 32767 { + if ids.len() < 32 + && let Some((file, logits)) = &mut trace + { + if unsafe { + ds4_session_copy_logits(session.0, logits.as_mut_ptr(), logits.len() as i32) + } != logits.len() as i32 + { + return Err("copying reference logits failed".into()); + } + let bytes: Vec = logits.iter().flat_map(|v| v.to_le_bytes()).collect(); + file.write_all(&bytes).map_err(|e| e.to_string())?; + } + let first = unsafe { ds4_session_sample(session.0, 0.6, 0, 0.95, 0.0, &mut rng) }; + if unsafe { ds4_token_is_stop_for_think_mode(engine.0, first, 1) } { + stop = Some(first); + break; + } + let mut accepted = [first; 17]; + let room = (max - ids.len() as i32).min(32767 - tokens.len); + let count = unsafe { + if ds4_engine_mtp_draft_tokens(engine.0) > 1 { + ds4_session_eval_speculative( + session.0, + first, + room, + ds4_token_eos(engine.0), + 0.6, + 0, + 0.95, + 0.0, + &mut rng, + accepted.as_mut_ptr(), + room.min(17), + err.as_mut_ptr(), + err.len(), + ) + } else { + check( + ds4_session_eval(session.0, first, err.as_mut_ptr(), err.len()), + &err, + )?; + 1 + } + }; + if count <= 0 || count > 17 { + return Err(format!( + "invalid reference verify result {count}: {}", + unsafe { CStr::from_ptr(err.as_ptr()) }.to_string_lossy() + )); + } + if trace_cycles { + emit(json!({"event":"spec_cycle", "prompt_tokens":input, + "generated":ids.len(), "first":first, "accepted":&accepted[..count as usize], + "position":unsafe { ds4_session_pos(session.0) }})); + } + for &token in &accepted[..count as usize] { + if unsafe { ds4_token_is_stop_for_think_mode(engine.0, token, 1) } { + stop = Some(token); + break; + } + unsafe { + ds4_tokens_push(tokens, token); + } + ids.push(token); + let mut len = 0; + let piece = unsafe { ds4_token_text(engine.0, token, &mut len) }; + if piece.is_null() { + return Err("null reference token text".into()); + } + text.extend_from_slice(unsafe { std::slice::from_raw_parts(piece.cast::(), len) }); + unsafe { + free(piece.cast()); + } + } + emit(json!({"event":"reference_tokens", "turn":index, "generated":ids.len()})); + if stop.is_some() { + break; + } + } + let decode_ms = decode.elapsed().as_secs_f64() * 1000.0; + probe.phase("finishing")?; + // Keep exactly the visible token frontier, excluding an MTP-returned EOS. + unsafe { + ds4_session_rewind(session.0, tokens.len); + } + emit( + json!({"event":"reference_result", "turn":index, "prompt_tokens":input, "cached_tokens":cached, "completion_tokens":ids.len(), "token_ids":ids, "text":String::from_utf8(text).map_err(|e| e.to_string())?, "stop_token":stop, "finish_reason":if stop.is_some(){"stop"}else{"length"}, "bootstrap_ms":bootstrap_ms, "prefill_ms":prefill_ms, "decode_ms":decode_ms, "decode_tokens_per_second":ids.len() as f64 * 1000.0 / decode_ms, "elapsed_ms":started.elapsed().as_secs_f64()*1000.0}), + ); + if index != 0 && stop.is_none() { + return Err("measured turn exhausted context without EOS".into()); + } + Ok(()) +} +fn run() -> Result<(), String> { + // Checked against clang -fdump-record-layouts-complete for the pinned header. + assert_eq!( + ( + size_of::(), + offset_of!(Options, distributed), + offset_of!(Options, tp) + ), + (280, 152, 216) + ); + let args = std::env::args().skip(1).collect::>(); + if args.first().is_some_and(|a| a == "--sampler-benchmark") { + if args.len() != 3 { + return Err("usage: reference --sampler-benchmark F32_RECORDING VOCAB".into()); + } + let vocab = args[2].parse::().map_err(|e| e.to_string())?; + let mut rng = 42_u64; + emit(sampler_replay_benchmark::run( + std::path::Path::new(&args[1]), + vocab, + |row| unsafe { + ds4_sample_logits(row.as_ptr(), row.len() as i32, 0.6, 0, 0.95, 0.0, &mut rng) + }, + )?); + return Ok(()); + } + if args == ["--sampler-fixture"] { + sampler_fixture(); + return Ok(()); + } + if args == ["--canary-self-test"] { + let mut probe = InlineProbe::start()?; + for phase in ["prefill", "decode"] { + *probe.phase.lock().unwrap() = phase; + std::thread::sleep(Duration::from_millis(300)); + } + return probe.finish(); + } + if !(4..=5).contains(&args.len()) + || !["glm", "deepseek"].contains(&args[1].as_str()) + || !["on", "off"].contains(&args[2].as_str()) + { + return Err("usage: reference MODEL glm|deepseek on|off README [SUPPORT_GGUF]".into()); + } + for file in [Some(&args[0]), args.get(4)].into_iter().flatten() { + if !std::path::Path::new(file).is_file() { + return Err(format!("installed artifact missing: {file}")); + } + } + let readme = std::fs::read_to_string(&args[3]).map_err(|e| e.to_string())?; + let model = CString::new(args[0].as_str()).map_err(|e| e.to_string())?; + let support = args + .get(4) + .map(|s| CString::new(s.as_str())) + .transpose() + .map_err(|e| e.to_string())?; + let glm = args[1] == "glm"; + let enabled = args[2] == "on"; + if enabled && std::env::var_os("DS4_REFERENCE_LOGITS_TRACE").is_some() { + return Err("logits tracing requires acceleration off".into()); + } + if !glm && enabled && support.is_none() { + return Err("DSpark requires installed support GGUF".into()); + } + let options = Options { + model: model.as_ptr(), + mtp: support.as_ref().map_or(null(), |s| s.as_ptr()), + context: 32768, + prefill_chunk: 0, + mtp_draft_tokens: 1, + mtp_margin: 3.0, + dspark_confidence_threshold: 0.6, + power: 100, + glm_mtp: glm && enabled, + dspark: !glm && enabled, + ..Options::default() + }; + let mut probe = Probe::start()?; + probe.phase("loading")?; + emit( + json!({"event":"reference_start", "family":args[1], "acceleration":enabled, "power_percent":100, "reasoning":"low", "model":args[0], "readme":args[3], "context":32768, "prefill_chunk":options.prefill_chunk, "temperature":0.6_f32, "top_p":0.95_f32, "min_p":0, "top_k":0, "seed":42, "canary":probe.child.is_some() || probe.inline.is_some(), "canary_in_process":probe.inline.is_some()}), + ); + let mut e = null_mut(); + if unsafe { ds4_engine_open(&mut e, &options) } != 0 || e.is_null() { + return Err("reference engine load failed".into()); + } + let engine = Engine(e); + let power = unsafe { ds4_engine_power(engine.0) }; + if power != 100 { + return Err(format!("reference engine power is {power}, expected 100")); + } + emit(json!({"event":"reference_engine", "power_percent":power})); + { + let warmup = session(&engine)?; + turn( + &engine, + &warmup, + &mut Tokens::default(), + "Reply with exactly: OK", + glm, + 0, + 32, + &mut probe, + )?; + } + let measured = session(&engine)?; + let mut tokens = Tokens::default(); + for (index, prompt) in [format!("Give a summary of the following text:\n\n{readme}"), "Tell me a complete short story about a lighthouse keeper. Do not ask questions.".into(), "Write a Python function is_prime(n: int) -> bool, followed by five assert examples. No tools.".into()].iter().enumerate() { + turn(&engine, &measured, &mut tokens, prompt, glm, index + 1, i32::MAX, &mut probe)?; + } + drop(tokens); + drop(measured); + drop(engine); + probe.finish()?; + Ok(()) +} +fn main() { + if let Err(error) = run() { + eprintln!("reference failed: {error}"); + std::process::exit(1); + } +} diff --git a/tools/sampler-replay-benchmark.rs b/tools/sampler-replay-benchmark.rs new file mode 100644 index 0000000..5dcc966 --- /dev/null +++ b/tools/sampler-replay-benchmark.rs @@ -0,0 +1,39 @@ +// CPU-only diagnostic shared by the standalone DS4 oracle and Rust tests. +// The input is the existing 32-row, little-endian F32 logit recording. +pub fn run( + path: &std::path::Path, + vocab: usize, + mut sample: impl FnMut(&[f32]) -> i32, +) -> Result { + if !(1..=1_000_000).contains(&vocab) { + return Err("invalid sampler vocabulary size".into()); + } + let expected = 32 * vocab * 4; + if std::fs::metadata(path).map_err(|e| e.to_string())?.len() != expected as u64 { + return Err("sampler recording must contain exactly 32 F32 rows".into()); + } + let bytes = std::fs::read(path).map_err(|e| e.to_string())?; + if bytes.len() != expected { + return Err("sampler recording changed while reading".into()); + } + let values = bytes + .chunks_exact(4) + .map(|b| f32::from_le_bytes(b.try_into().unwrap())) + .collect::>(); + for row in values.chunks_exact(vocab) { + std::hint::black_box(sample(row)); + } + let mut batch_ms = Vec::with_capacity(16); + let mut token_ids = Vec::with_capacity(512); + for _ in 0..16 { + let started = std::time::Instant::now(); + for row in values.chunks_exact(vocab) { + token_ids.push(std::hint::black_box(sample(row))); + } + batch_ms.push(started.elapsed().as_secs_f64() * 1000.0); + } + Ok(serde_json::json!({ + "event":"sampler_replay_benchmark", "vocab":vocab, "rows":32, + "warmup_draws":32, "draws":512, "batch_ms":batch_ms, "token_ids":token_ids, + })) +} diff --git a/tools/test-supervisor.rs b/tools/test-supervisor.rs index a3fa868..a879eea 100644 --- a/tools/test-supervisor.rs +++ b/tools/test-supervisor.rs @@ -59,6 +59,8 @@ mod supervisor { value["event"].as_str(), Some( "gpu_canary_sample" + | "gpu_canary_ready" + | "gpu_canary_stall" | "gpu_canary_summary" | "reference_canary_summary" | "test_resource_sample" @@ -80,6 +82,31 @@ mod supervisor { progressed } + // Keep complete records together when the watchdog writes resource samples + // to the same stream. Pipe reads can split a single worker JSON record. + fn forward_records( + output: &mut impl Write, + pending: &mut Vec, + bytes: &[u8], + eof: bool, + ) -> io::Result<()> { + pending.extend_from_slice(bytes); + let end = if eof || pending.len() > 64 * 1024 { + pending.len() + } else { + pending + .iter() + .rposition(|&b| b == b'\n') + .map_or(0, |i| i + 1) + }; + if end > 0 { + output.write_all(&pending[..end])?; + output.flush()?; + pending.drain(..end); + } + Ok(()) + } + fn forward( mut input: impl Read + Send + 'static, mut output: impl Write + Send + 'static, @@ -88,15 +115,17 @@ mod supervisor { thread::spawn(move || { let mut bytes = [0; 4096]; let mut pending = Vec::new(); + let mut records = Vec::new(); loop { match input.read(&mut bytes) { - Ok(0) => return, + Ok(0) => { + if forward_records(&mut output, &mut records, &[], true).is_err() { + progress.lock().unwrap().failed = true; + } + return; + } Ok(n) => { - if output - .write_all(&bytes[..n]) - .and_then(|()| output.flush()) - .is_err() - { + if forward_records(&mut output, &mut records, &bytes[..n], false).is_err() { progress.lock().unwrap().failed = true; return; } @@ -257,6 +286,29 @@ mod supervisor { mod tests { use super::*; + #[test] + fn split_records_cannot_be_interleaved_with_watchdog_samples() { + let mut output = Vec::new(); + let mut pending = Vec::new(); + forward_records( + &mut output, + &mut pending, + b"{\"event\":\"gpu_canary_", + false, + ) + .unwrap(); + assert!(output.is_empty()); + output.extend_from_slice(b"{\"event\":\"test_resource_sample\"}\n"); + forward_records(&mut output, &mut pending, b"sample\"}\npartial", false).unwrap(); + for line in output.split(|&b| b == b'\n').filter(|l| !l.is_empty()) { + serde_json::from_slice::(line).unwrap(); + } + assert_eq!(pending, b"partial"); + forward_records(&mut output, &mut pending, &[], true).unwrap(); + assert!(output.ends_with(b"partial")); + assert!(pending.is_empty()); + } + #[test] fn direct_worker_arguments_are_not_rust_test_arguments() { let args = |a: &[&str]| a.iter().map(std::ffi::OsString::from).collect::>();