diff --git a/.agents/skills/benchmark-model/SKILL.md b/.agents/skills/benchmark-model/SKILL.md index eccdf26078..b1ea73261c 100644 --- a/.agents/skills/benchmark-model/SKILL.md +++ b/.agents/skills/benchmark-model/SKILL.md @@ -1,6 +1,6 @@ --- name: benchmark-model -description: Measure every retained number in docs/models//BENCHMARKS.md for one Gufo model on Strix Halo gfx1151 over HTTP, run the same workload against the reference project's OpenAI-compatible server (llama.cpp, audio.cpp, ...), and report the gain per cell. +description: Measure every retained number in a Gufo model's BENCHMARKS.md on Strix Halo gfx1151 over HTTP, run the same workload against the reference project's OpenAI-compatible server (llama.cpp, audio.cpp, ...), and report the gain per cell. metadata: origin: gufo --- @@ -21,6 +21,12 @@ paths remain valid. Prefer one warmed sample per point, run only the requested sweeps, and repeat only to investigate a concrete discrepancy. Prefix setup at 64K/128K dominates runtime; avoid rebuilding it for unrelated checks. +Record the power profile and starting CPU/GPU temperatures. Paired runs need +the same profile and comparable thermal/warmup conditions. On this host, +check `/sys/firmware/acpi/platform_profile`: Balanced and Performance can +differ materially with identical binaries. Monitor the profile throughout +long sweeps and retain their warmup/cache history. + Start long runs in the background with output redirected to a log. Inspect progress and process status at short intervals while continuing independent work; do not leave a run unreviewed. Before starting, establish diff --git a/README.md b/README.md index fb4f763688..3f6394d5dd 100644 --- a/README.md +++ b/README.md @@ -22,7 +22,7 @@ All model documentation lives under [docs/models](docs/models/README.md): | Model | Inference modes | Hugging Face weights | Benchmarks | Quality | | --- | --- | --- | --- | --- | | [Qwen3.8 27B](docs/models/qwen3.8-27b/README.md) | Q4/Q8, images, AR, DFlash2 | Unsloth [Q4_K_XL](https://huggingface.co/unsloth/Qwen3.8-27B-GGUF/blob/4ca720788d1e01f1bff70c033e0d0028fd02e502/Qwen3.8-27B-UD-Q4_K_XL.gguf) / [Q8_K_XL](https://huggingface.co/unsloth/Qwen3.8-27B-GGUF/blob/4ca720788d1e01f1bff70c033e0d0028fd02e502/Qwen3.8-27B-UD-Q8_K_XL.gguf) · [DFlash2 Q4_K_M](https://huggingface.co/z-lab/Qwen3.8-27B-DFlash2-GGUF/blob/2d9571f8ce46e151f61c6499c99dee6079e1d610/Qwen3.8-27B-DFlash2-Q4_K_M.gguf) | Q4: **656.33 tok/s pp**; up to **70.56 tok/s tg** single user and **123.00 aggregated tok/s** on 8 concurrent requests with DFlash2 · [Benchmarks](docs/models/qwen3.8-27b/BENCHMARKS.md) | [Quality](docs/models/qwen3.8-27b/QUALITY.md) | -| [Qwen3.8 Flash-Next](docs/models/qwen3.8-flash-next/README.md) | Q4, images, AR, MTP | Unsloth [Q4_K_XL](https://huggingface.co/unsloth/Qwen3.8-Flash-Next-GGUF/tree/38bb39ee97821de2c9009abb7e93950eec396e66/UD-Q4_K_XL) · [MTP Q8_0](https://huggingface.co/unsloth/Qwen3.8-Flash-Next-GGUF/blob/38bb39ee97821de2c9009abb7e93950eec396e66/MTP/mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf) | **1,639.03 tok/s pp**; up to **60.36 tok/s tg** single user and **157.22 aggregated tok/s** on 8 concurrent requests with MTP · [Benchmarks](docs/models/qwen3.8-flash-next/BENCHMARKS.md) | [Quality](docs/models/qwen3.8-flash-next/QUALITY.md) | +| [Qwen3.8 Flash-Next](docs/models/qwen3.8-flash-next/README.md) | Q4, images, AR, MTP | Unsloth [Q4_K_XL](https://huggingface.co/unsloth/Qwen3.8-Flash-Next-GGUF/tree/38bb39ee97821de2c9009abb7e93950eec396e66/UD-Q4_K_XL) · [MTP Q8_0](https://huggingface.co/unsloth/Qwen3.8-Flash-Next-GGUF/blob/38bb39ee97821de2c9009abb7e93950eec396e66/MTP/mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf) | **1,644.68 tok/s pp**; up to **59.02 tok/s tg** single user and **157.22 aggregated tok/s** on 8 concurrent requests with MTP · [Benchmarks](docs/models/qwen3.8-flash-next/BENCHMARKS.md) | [Quality](docs/models/qwen3.8-flash-next/QUALITY.md) | | [DeepSeek V4 Flash](docs/models/deepseek-v4-flash/README.md) | AR, DSpark | [antirez Flash 0731 IQ2XXS](https://huggingface.co/antirez/deepseek-v4-gguf/blob/1cd7b564460821938add0475a60b942c409295e0/DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix-0731.gguf) · [DSpark](https://huggingface.co/antirez/deepseek-v4-gguf/blob/e7f04037032990db0346398d249baf9fb9df1ccc/DeepSeek-V4-Flash-DSpark-support-0731.gguf) | **484.62 tok/s pp**; up to **26.62 tok/s tg** single user and **54.74 aggregated tok/s** on 8 concurrent requests with DSpark · [Benchmarks](docs/models/deepseek-v4-flash/BENCHMARKS.md) | [Quality](docs/models/deepseek-v4-flash/QUALITY.md) | | [Qwen3-ASR 1.7B](docs/models/qwen3-asr/README.md) | Speech recognition | [BF16](https://huggingface.co/Qwen/Qwen3-ASR-1.7B/tree/7278e1e70fe206f11671096ffdd38061171dd6e5) | **15.27× realtime** · [Benchmarks](docs/models/qwen3-asr/BENCHMARKS.md) | [Quality](docs/models/qwen3-asr/QUALITY.md) | | [Qwen3-TTS 1.7B](docs/models/qwen3-tts/README.md) | Speech synthesis and voice cloning | BF16 [CustomVoice](https://huggingface.co/Qwen/Qwen3-TTS-12Hz-1.7B-CustomVoice) / [VoiceDesign](https://huggingface.co/Qwen/Qwen3-TTS-12Hz-1.7B-VoiceDesign) / [Base](https://huggingface.co/Qwen/Qwen3-TTS-12Hz-1.7B-Base) | Up to **2.54× realtime**; **201 ms** to first audio (CustomVoice) · [Benchmarks](docs/models/qwen3-tts/BENCHMARKS.md) | [Quality](docs/models/qwen3-tts/QUALITY.md) | diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index d9925673c0..72c9eac4ca 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -134,10 +134,10 @@ end-to-end model measurements. Collect profiler traces separately and use ## Reporting and artifact retention Record model/engine revisions, artifact hashes, hardware/software fingerprint, -input identity, context, output count, sampling, concurrency, cache state, -oracle identity, and the exact timed scope. A numerical mismatch or device -failure is not a passing result. Add repetitions when timing noise or a -suspected regression requires them. +power profile and CPU/GPU temperatures, input identity, context, output count, +sampling, concurrency, cache state, oracle identity, and the exact timed scope. +A numerical mismatch or device failure is not a passing result. Add repetitions +when timing noise or a suspected regression requires them. Checked-in JSON retains fixtures, model provenance, controller calibration, and bounded result summaries. Serving summaries preserve aggregate metrics, diff --git a/docs/models/qwen3.8-flash-next/BENCHMARKS.md b/docs/models/qwen3.8-flash-next/BENCHMARKS.md index 49cb24c128..bd46ca9fed 100644 --- a/docs/models/qwen3.8-flash-next/BENCHMARKS.md +++ b/docs/models/qwen3.8-flash-next/BENCHMARKS.md @@ -2,8 +2,8 @@ AMD Strix Halo `gfx1151`, 128 GB unified memory. Unsloth `UD-Q4_K_XL` target and shared-Q8_0 MTP sidecar; Gufo uses adaptive MTP. HTTP, greedy, thinking off. -Gufo single-user pp/tg: October 6, 2026 (`4c5a00d2`); concurrency, loading, -memory and reference results: September 22–23. +Gufo single-user pp/tg: October 7, 2026 (`def2ed1e`), Performance power profile. +Concurrency, loading, memory and reference results: September 22–23. llama.cpp uses `b11069` for AR and `6fcaa16f` for MTP. Positive gain favors Gufo. @@ -17,14 +17,14 @@ Context capacities differ between engines; see the measurement details. | Flash-Next Q4 AR
Depth (tokens) | Gufo pp (tok/s) | llama.cpp pp (tok/s) | Gain | Gufo tg (tok/s) | llama.cpp tg (tok/s) | Gain | | ---: | ---: | ---: | ---: | ---: | ---: | ---: | -| 0 | 1639.03 | 489.59 | +234.8% | 25.96 | 22.20 | +16.9% | -| 4,096 | 1564.36 | 455.24 | +243.6% | 25.96 | 21.21 | +22.4% | -| 8,192 | 1557.47 | 428.85 | +263.2% | 25.95 | 20.40 | +27.2% | -| 12,288 | 1553.60 | 400.21 | +288.2% | 25.93 | 19.64 | +32.0% | -| 16,384 | 1553.97 | 375.89 | +313.4% | 25.89 | 18.95 | +36.6% | -| 32,768 | 1540.25 | 301.31 | +411.2% | 25.75 | 16.54 | +55.7% | -| 65,536 | 1459.69 | 221.94 | +557.7% | 25.50 | 11.62 | +119.4% | -| 131,072 | 1444.79 | 144.78 | +897.9% | 24.90 | 7.98 | +212.0% | +| 0 | 1644.68 | 489.59 | +235.9% | 25.99 | 22.20 | +17.1% | +| 4,096 | 1518.57 | 455.24 | +233.6% | 25.73 | 21.21 | +21.3% | +| 8,192 | 1567.32 | 428.85 | +265.5% | 25.99 | 20.40 | +27.4% | +| 12,288 | 1560.97 | 400.21 | +290.0% | 25.96 | 19.64 | +32.2% | +| 16,384 | 1556.02 | 375.89 | +314.0% | 25.93 | 18.95 | +36.8% | +| 32,768 | 1546.81 | 301.31 | +413.4% | 25.80 | 16.54 | +56.0% | +| 65,536 | 1425.49 | 221.94 | +542.3% | 25.43 | 11.62 | +118.8% | +| 131,072 | 1446.10 | 144.78 | +898.8% | 24.97 | 7.98 | +212.9% | ![Single user, autoregressive](artifacts/charts/single-ar.svg) @@ -37,14 +37,14 @@ text, including Gufo predictor catch-up. | Flash-Next Q4 MTP
Depth (tokens) | Gufo pp (tok/s) | llama.cpp pp (tok/s) | Gain pp | Gufo tg mixed (tok/s) | llama.cpp tg mixed (tok/s) | Gain mixed | Gufo tg repetitive (tok/s) | llama.cpp tg repetitive (tok/s) | Gain repetitive | | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | -| 0 | 1584.63 | 468.76 | +238.0% | 31.33 | 31.73 | -1.3% | 60.36 | 48.12 | +25.4% | -| 4,096 | 1528.90 | 423.79 | +260.8% | 32.83 | 34.48 | -4.8% | 57.61 | 45.71 | +26.0% | -| 8,192 | 1528.87 | 394.98 | +287.1% | 32.41 | 32.20 | +0.7% | 49.22 | 44.31 | +11.1% | -| 12,288 | 1511.31 | 365.91 | +313.0% | 34.72 | 32.32 | +7.4% | 42.17 | 43.53 | -3.1% | -| 16,384 | 1503.69 | 341.98 | +339.7% | 35.21 | 32.07 | +9.8% | 50.56 | 43.66 | +15.8% | -| 32,768 | 1501.78 | 277.61 | +441.0% | 34.21 | 25.41 | +34.6% | 44.99 | 39.28 | +14.5% | -| 65,536 | 1415.89 | 208.04 | +580.6% | 33.50 | 19.65 | +70.5% | 44.45 | 27.75 | +60.2% | -| 131,072 | 1403.96 | 135.42 | +936.7% | 33.10 | 14.63 | +126.2% | 44.34 | 20.96 | +111.5% | +| 0 | 1560.11 | 468.76 | +232.8% | 30.16 | 31.73 | -4.9% | 59.02 | 48.12 | +22.7% | +| 4,096 | 1529.17 | 423.79 | +260.8% | 32.86 | 34.48 | -4.7% | 57.58 | 45.71 | +26.0% | +| 8,192 | 1525.11 | 394.98 | +286.1% | 32.47 | 32.20 | +0.8% | 49.26 | 44.31 | +11.2% | +| 12,288 | 1517.04 | 365.91 | +314.6% | 34.78 | 32.32 | +7.6% | 42.19 | 43.53 | -3.1% | +| 16,384 | 1502.95 | 341.98 | +339.5% | 35.27 | 32.07 | +10.0% | 50.60 | 43.66 | +15.9% | +| 32,768 | 1503.08 | 277.61 | +441.4% | 34.27 | 25.41 | +34.9% | 44.98 | 39.28 | +14.5% | +| 65,536 | 1390.03 | 208.04 | +568.2% | 33.51 | 19.65 | +70.5% | 43.12 | 27.75 | +55.4% | +| 131,072 | 1408.06 | 135.42 | +939.8% | 33.18 | 14.63 | +126.8% | 44.43 | 20.96 | +112.0% | ![Single user, MTP](artifacts/charts/single-mtp.svg) diff --git a/docs/models/qwen3.8-flash-next/EXPERIMENTS.md b/docs/models/qwen3.8-flash-next/EXPERIMENTS.md index 7ca5af2a40..6cfff14603 100644 --- a/docs/models/qwen3.8-flash-next/EXPERIMENTS.md +++ b/docs/models/qwen3.8-flash-next/EXPERIMENTS.md @@ -59,6 +59,9 @@ | Live final frontier and async prompt snapshots | Retained; immutable branch snapshot, worker capture, bounded persistence outside the lookup lock. | | Thread-local decode graph capture | Retained; independent snapshot workers no longer invalidate a peer's capture. Snapshot bytes and captured/replayed logits are exact; no request serialization added. | | Chunk-equivalent projections/attention | Retained with exact continued-image/cache/full-logit gates; one-token tails keep prefill arithmetic. | +| Dense split at the indexer budget | Retained; pre-budget queries always take the dense tiles, so prefill logits no longer depend on the prefill batch. Before: batches 4096 and 8192 differed from 2048 on 4/4 repository documents; after: byte-identical decode dumps at 2048/4096/8192. | +| 4096-token prefill chunks | Retained together with 128 readers and the next-chunk prefetch; exact given the dense split (byte-identical dumps on 4/4 repository documents, identical served greedy outputs). Served repository prompts on two gfx1151 machines, ABBA, 8+ samples per prompt: 24K–102K-token prompts +6–11% (median per prompt), 7K–8K prompts −0.1% to +0.6% in dedicated rechecks; pp4096 +5.3%/+5.4% at d0 and +5.8%/+5.9% at d32K; TG unchanged. Peak memory at 262K +1.3 GiB. Prefill beside an active decode stays at `--prefill-chunk` (512); two-session gaps unchanged. | +| Next-chunk n-gram prefetch | Retained; a prefill chunk reads the next chunk's rows once its own have arrived, and a batch uses them only with identical tokens and history, so output is unchanged. With 4096-token chunks the gather then overlaps a whole chunk's forward instead of only layer 0. Adds one row buffer (≤ 40 MiB pageable, allocated on first use). | | More Q8 vocabulary rows/block | Rejected: no C1 improvement. | | Private shared-expert F16 rows | Retained; the shared expert's SwiGLU rows for its F16 down projection get their own buffer instead of overwriting the routed experts' narrowed token rows, which removes one full-batch narrowing pass per layer. Every GEMM reads identical bytes; greedy output, chunk-boundary logits and the C2–C8 batch checks are exact. Interleaved Nix A/B at d0: 1559.6 → 1568.1 tok/s (+0.6%); the profiled kernel time fell 1362.8 → 1333.1 ms. | | 256-row routed down-projection blocks | Rejected: exact output, Q5_1 unchanged and Q8_0 slower. The routed kernels already stream expert weights at roughly two thirds of DRAM bandwidth, so per-block prologue/epilogue and activation restaging were not the bound. | @@ -70,7 +73,7 @@ | Whole-tile carry across windows | Rejected: same output, but 16.5 KiB of LDS cost a resident block per CU and slowed d2K eight-row verification 8.5%. | | GPU greedy penalties and linear CPU anchor selection | Retained; exact FP64 penalties, unchanged proposals and snapshots. Short heat-pump tg400: 32.12 → 33.80 tok/s; C2/4/6/8 improve 7.6/8.5/16.3/14.2%. Unpenalized control unchanged. [Evidence](artifacts/penalty-verification.json). | | Prefill projection, attention and indexer kernels | Retained; bit-identical logit dumps and greedy hashes, plus a bitwise GEMM sweep against hipBLASLt for every routed n ≤ 4096. pp4096 +1.4% at d0 and +2.5–2.8% at 32K/115K (interleaved ABBA, `-r 6`, two machines); repository prompts 6.5K–102K +3.2% median; tg unchanged. The indexer projections use the own kernel only when hipBLASLt selects algorithm 4438 and otherwise stay on the library. | -| 128 n-gram readers | Rejected: first gather 179 → 140 ms, but no end-to-end change on repository prompts on two different NVMe drives (ABBA medians within 1%). | +| 128 n-gram readers | Retained with 4096-token chunks, which double each gather; readers block in `pread`, so the pool sets the queue depth. Readers are woken per read, so small gathers do not wake the pool; 96 more threads. | Separate d32K pp2048 profiling attributes 29.1% of kernel time to MoE, 34.9% to dense projections and 12.6% to attention/indexing. Final-tile catch-up diff --git a/docs/models/qwen3.8-flash-next/QUALITY.md b/docs/models/qwen3.8-flash-next/QUALITY.md index 8b81f5bf2e..dd40e8a21c 100644 --- a/docs/models/qwen3.8-flash-next/QUALITY.md +++ b/docs/models/qwen3.8-flash-next/QUALITY.md @@ -69,8 +69,9 @@ weights. [Vision reproduction](../qwen3.8-27b/QUALITY.md#vision). ## Benchmark method -Gufo single-user pp/tg refreshed October 6, 2026 (`4c5a00d2`); reference, -concurrency, loading and memory measurements retain September 22–23 provenance. +Gufo single-user pp/tg refreshed October 7, 2026 (`def2ed1e`), Performance power +profile. Reference, concurrency, loading and memory retain September 22–23 +provenance. One warmed sample per point, greedy, thinking off, penalties disabled. Single-user uses pp2048/tg128; MTP pp is the maximum across mixed/repetitive workloads. Gufo capacity is 133760; reference capacity is 35456 through 32K, diff --git a/docs/models/qwen3.8-flash-next/artifacts/charts/single-ar.svg b/docs/models/qwen3.8-flash-next/artifacts/charts/single-ar.svg index c8ab818e38..2fa713d8d0 100644 --- a/docs/models/qwen3.8-flash-next/artifacts/charts/single-ar.svg +++ b/docs/models/qwen3.8-flash-next/artifacts/charts/single-ar.svg @@ -144,107 +144,107 @@ L -3.5 0 - - + - 250 + 250 - - + - 500 + 500 - - + - 750 + 750 - - + - 1000 + 1000 - - + - 1250 + 1250 - - + - 1500 + 1500 - - + - 1750 + 1750 @@ -253,13 +253,13 @@ L 286.549 47.326945 - - - - - - - + + + + + + + - - - - - - - - - + + + + + + + + @@ -468,92 +468,92 @@ L 563.96275 191.679297 - - + - 5 + 5 - - + - 10 + 10 - - + - 15 + 15 - - + - 20 + 20 - - + - 25 + 25 - - + - 30 + 30 @@ -562,44 +562,44 @@ L 563.96275 35.440357 - - - - - - - + + + + + + + - - - - - - - - - + + + + + + + + diff --git a/docs/models/qwen3.8-flash-next/artifacts/charts/single-mtp.svg b/docs/models/qwen3.8-flash-next/artifacts/charts/single-mtp.svg index 0e1e50417a..fb12bde782 100644 --- a/docs/models/qwen3.8-flash-next/artifacts/charts/single-mtp.svg +++ b/docs/models/qwen3.8-flash-next/artifacts/charts/single-mtp.svg @@ -144,107 +144,107 @@ L -3.5 0 - - + - 250 + 250 - - + - 500 + 500 - - + - 750 + 750 - - + - 1000 + 1000 - - + - 1250 + 1250 - - + - 1500 + 1500 - - + - 1750 + 1750 @@ -253,13 +253,13 @@ L 289.806481 55.445415 - - - - - - - + + + + + + + - - - - - - - - - + + + + + + + + @@ -471,62 +471,62 @@ L 570.748981 191.679297 - - + - 10 + 10 - - + - 20 + 20 - - + - 30 + 30 - - + - 40 + 40 @@ -534,45 +534,45 @@ L 570.748981 51.537099 - - - - - + + + + - - - + + + - - - - - - - - - + + + + + + + + @@ -727,107 +727,107 @@ L 851.691481 191.679297 - - + - 10 + 10 - - + - 20 + 20 - - + - 30 + 30 - - + - 40 + 40 - - + - 50 + 50 - - + - 60 + 60 - - + - 70 + 70 @@ -836,44 +836,44 @@ L 851.691481 48.61747 - - - - - - - + + + + + + + - - - - - - - - - + + + + + + + + diff --git a/docs/models/qwen3.8-flash-next/artifacts/model-identities.json b/docs/models/qwen3.8-flash-next/artifacts/model-identities.json index ebb7ae6c5b..bdc43b223f 100644 --- a/docs/models/qwen3.8-flash-next/artifacts/model-identities.json +++ b/docs/models/qwen3.8-flash-next/artifacts/model-identities.json @@ -57,5 +57,32 @@ "measurement": "One warmed valid sample per point, no repetitions. All 24 requests generated 128 tokens and passed the documented cache/prefill count tolerances. Token-count calibration attempts are retained in raw logs and excluded from the tables.", "quality": "The d0 mixed MTP completion matches the fresh AR sweep; repetitive d0 matches the retained fresh AR control. Full-vocabulary pruning and sampled snapshot replay are qualified separately in prefill-final-rows.json, including mixed text, images and 128K. Deep HTTP AR/MTP frontiers differ by one token, so their free-running completion hashes are not an exact-parity test.", "retained": "Reference, concurrency, loading and memory results retain their September 22-23 measurements." + }, + "benchmarks-2026-10-07": { + "scope": "Gufo single-user AR and MTP mixed/repetitive HTTP pp2048/tg128 at 0, 4096, 8192, 12288, 16384, 32768, 65536 and 131072 cached tokens.", + "gufo": { + "revision": "def2ed1e904f912e69550344aa7c13a99f535d0e", + "binarySha256": "1ea59e9bfb214b5cdd4d4b6bcc896058df43a06f9c5e180086ea6a8d4301ca58", + "flakeLockSha256": "faabad820ec09caab7351bf48910cec895a38cf1766e098e3065aea86f405ed6", + "build": "Production CMake release preset (RelWithDebInfo), pinned Nix toolchain." + }, + "conditions": { + "powerProfile": "performance", + "profileVerifiedThroughoutSweep": true, + "temperatureScope": "Entire ordered AR, mixed MTP and repetitive MTP sweep.", + "sweepStartGpuEdgeC": 48.0, + "sweepStartCpuTctlC": 54.0, + "gpuEdgeRangeC": [ + 48.0, + 85.0 + ], + "cpuTctlRangeC": [ + 55.25, + 96.125 + ] + }, + "measurement": "One warmed valid sample per point. Ordered full depth sweep; cache/prefill count calibration attempts are retained in private raw logs.", + "quality": "All 24 requests generate 128 tokens and retain the October 6 prompt/cache/output counts, completion hashes and speculative counts. This checks consistency, not original unquantized-model accuracy.", + "retained": "Reference, concurrency, loading and memory retain September 22-23 measurements." } } diff --git a/docs/models/qwen3.8-flash-next/artifacts/single-ar-gufo.json b/docs/models/qwen3.8-flash-next/artifacts/single-ar-gufo.json index 8f2da9cff8..668b880794 100644 --- a/docs/models/qwen3.8-flash-next/artifacts/single-ar-gufo.json +++ b/docs/models/qwen3.8-flash-next/artifacts/single-ar-gufo.json @@ -1,5 +1,9 @@ { "artifactType": "model-bench-table", + "conditions": { + "powerProfile": "performance", + "profileVerifiedThroughoutSweep": true + }, "fingerprint": { "canonical": { "cppStandard": "C++20", @@ -14,15 +18,15 @@ "gpuName": "AMD Radeon 8060S Graphics", "gpuPciId": "1002:1586", "kernelRelease": "7.1.8", - "memoryTotalBytes": 134309883904, + "memoryTotalBytes": 134309888000, "memoryType": "LPDDR5X (Unified)", "rocmVersion": "7.2.3", "schemaVersion": "1.0.0" }, - "fingerprintId": "46337fb568b92fb76472ec90a0af584eb7d7ae401f9878d8d18b09762e7c4caa", + "fingerprintId": "4dcf1077494084b2d72ecf90594aad3c844b087618d3db6af2cd796f06cb1c36", "schemaVersion": "1.0.0" }, - "measuredOn": "2026-10-06", + "measuredOn": "2026-10-07", "mode": null, "model": { "id": "Qwen3.8-Flash-Next", @@ -30,16 +34,12 @@ }, "notes": [ "pp2048/tg128; depth is a cached conversation prefix of synthetic text (8-token reply); measured turn task is recorded per row; tolerance max(32, 0.500%) on cache_n and prompt_n; actual counts per sample in rows", - "synthetic text: 1.161 tokens/word, template overhead 12 tokens", - "One warmed valid measurement per depth, without repetitions; out-of-tolerance token-count calibration attempts remain in the raw logs.", - "The immutable Nix release binary was built before committing the prefill patch; its production sources match eecb008a and 4c5a00d2. Only tests and documentation changed between those revisions.", - "Reference, concurrency, loading and memory measurements were not rerun. Deep AR/MTP HTTP cache frontiers differ by one token; exact numerical replay is qualified separately with identical prefill boundaries in prefill-final-rows.json." + "synthetic text: 1.161 tokens/word, template overhead 12 tokens" ], - "rawArtifactSha256": "e18aab305722ed191caa6a06b9d412e66648f3c8628e236f9254d40ee2f01262", "referenceName": "llama.cpp", "rows": { "0": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 0, @@ -50,15 +50,15 @@ "prompt_n": 2040 } ], - "pp": 1639.03, + "pp": 1644.68, "pp_sd": null, "samples": 1, - "tg": 25.96, + "tg": 25.99, "tg_sd": null, "workload": "prose" }, "12288": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 12249, @@ -69,15 +69,15 @@ "prompt_n": 2047 } ], - "pp": 1553.6, + "pp": 1560.97, "pp_sd": null, "samples": 1, - "tg": 25.93, + "tg": 25.96, "tg_sd": null, "workload": "prose" }, "131072": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 130932, @@ -88,15 +88,15 @@ "prompt_n": 2047 } ], - "pp": 1444.79, + "pp": 1446.1, "pp_sd": null, "samples": 1, - "tg": 24.9, + "tg": 24.97, "tg_sd": null, "workload": "prose" }, "16384": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 16324, @@ -107,15 +107,15 @@ "prompt_n": 2047 } ], - "pp": 1553.97, + "pp": 1556.02, "pp_sd": null, "samples": 1, - "tg": 25.89, + "tg": 25.93, "tg_sd": null, "workload": "prose" }, "32768": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 32718, @@ -126,15 +126,15 @@ "prompt_n": 2029 } ], - "pp": 1540.25, + "pp": 1546.81, "pp_sd": null, "samples": 1, - "tg": 25.75, + "tg": 25.8, "tg_sd": null, "workload": "prose" }, "4096": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 4095, @@ -145,15 +145,15 @@ "prompt_n": 2043 } ], - "pp": 1564.36, + "pp": 1518.57, "pp_sd": null, "samples": 1, - "tg": 25.96, + "tg": 25.73, "tg_sd": null, "workload": "prose" }, "65536": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 65447, @@ -164,15 +164,15 @@ "prompt_n": 2058 } ], - "pp": 1459.69, + "pp": 1425.49, "pp_sd": null, "samples": 1, - "tg": 25.5, + "tg": 25.43, "tg_sd": null, "workload": "prose" }, "8192": { - "command": "gufo serve --port 57693 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", + "command": "fix-gufo serve --port 58131 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8", "counts": [ { "cache_n": 8184, @@ -183,10 +183,10 @@ "prompt_n": 2048 } ], - "pp": 1557.47, + "pp": 1567.32, "pp_sd": null, "samples": 1, - "tg": 25.95, + "tg": 25.99, "tg_sd": null, "workload": "prose" } @@ -194,10 +194,10 @@ "schemaVersion": 1, "server": { "command": [ - "gufo", + "fix-gufo", "serve", "--port", - "57693", + "58131", "--sessions", "1", "llm", @@ -214,12 +214,10 @@ ] }, "source": { - "binaryRevision": "eecb008a909ca299c13d1ef773765c7e2e24d395", - "binarySha256": "2ebf7a2cf22eae84b0c78bf87bc39adda3336fc59f9ecb4e3f1b7c06227571d4", - "buildMode": "nix-release", + "binarySha256": "1ea59e9bfb214b5cdd4d4b6bcc896058df43a06f9c5e180086ea6a8d4301ca58", + "buildMode": "cmake-release-preset", "dirty": false, - "flakeLockSha256": "faabad820ec09caab7351bf48910cec895a38cf1766e098e3065aea86f405ed6", - "revision": "4c5a00d2d609846cb73e79aa33de324624aa9c2c" + "revision": "def2ed1e904f912e69550344aa7c13a99f535d0e" }, "table": "single-ar", "target": "gufo" diff --git a/docs/models/qwen3.8-flash-next/artifacts/single-mtp-gufo.json b/docs/models/qwen3.8-flash-next/artifacts/single-mtp-gufo.json index 3b5fd6157c..d1a82abb62 100644 --- a/docs/models/qwen3.8-flash-next/artifacts/single-mtp-gufo.json +++ b/docs/models/qwen3.8-flash-next/artifacts/single-mtp-gufo.json @@ -1,5 +1,9 @@ { "artifactType": "model-bench-table", + "conditions": { + "powerProfile": "performance", + "profileVerifiedThroughoutSweep": true + }, "fingerprint": { "canonical": { "cppStandard": "C++20", @@ -14,15 +18,15 @@ "gpuName": "AMD Radeon 8060S Graphics", "gpuPciId": "1002:1586", "kernelRelease": "7.1.8", - "memoryTotalBytes": 134309883904, + "memoryTotalBytes": 134309888000, "memoryType": "LPDDR5X (Unified)", "rocmVersion": "7.2.3", "schemaVersion": "1.0.0" }, - "fingerprintId": "46337fb568b92fb76472ec90a0af584eb7d7ae401f9878d8d18b09762e7c4caa", + "fingerprintId": "4dcf1077494084b2d72ecf90594aad3c844b087618d3db6af2cd796f06cb1c36", "schemaVersion": "1.0.0" }, - "measuredOn": "2026-10-06", + "measuredOn": "2026-10-07", "mode": "mtp", "model": { "id": "Qwen3.8-Flash-Next", @@ -30,12 +34,8 @@ }, "notes": [ "pp2048/tg128; depth is a cached conversation prefix of synthetic text (8-token reply); measured turn task is recorded per row; tolerance max(32, 0.500%) on cache_n and prompt_n; actual counts per sample in rows", - "synthetic text: 1.161 tokens/word, template overhead 12 tokens", - "One warmed valid measurement per depth, without repetitions; out-of-tolerance token-count calibration attempts remain in the raw logs.", - "The immutable Nix release binary was built before committing the prefill patch; its production sources match eecb008a and 4c5a00d2. Only tests and documentation changed between those revisions.", - "Reference, concurrency, loading and memory measurements were not rerun. Deep AR/MTP HTTP cache frontiers differ by one token; exact numerical replay is qualified separately with identical prefill boundaries in prefill-final-rows.json." + "synthetic text: 1.161 tokens/word, template overhead 12 tokens" ], - "rawArtifactSha256": "1def651097763e335a9d036dbe664c18bf3fbdd01d148f3f6bdd2fd7e4e4024b", "referenceName": "llama.cpp", "rows": { "0": { @@ -43,7 +43,7 @@ "acceptance_sd": null, "accepted_per_step": 0.66, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 0, @@ -54,10 +54,10 @@ "prompt_n": 2040 } ], - "pp": 1584.63, + "pp": 1560.11, "pp_sd": null, "samples": 1, - "tg": 31.33, + "tg": 30.16, "tg_sd": null, "workload": "prose" }, @@ -66,7 +66,7 @@ "acceptance_sd": null, "accepted_per_step": 0.86, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 12250, @@ -77,10 +77,10 @@ "prompt_n": 2046 } ], - "pp": 1503.9, + "pp": 1511.22, "pp_sd": null, "samples": 1, - "tg": 34.72, + "tg": 34.78, "tg_sd": null, "workload": "prose" }, @@ -89,7 +89,7 @@ "acceptance_sd": null, "accepted_per_step": 1.03, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 130933, @@ -100,10 +100,10 @@ "prompt_n": 2046 } ], - "pp": 1403.96, + "pp": 1408.06, "pp_sd": null, "samples": 1, - "tg": 33.1, + "tg": 33.18, "tg_sd": null, "workload": "prose" }, @@ -112,7 +112,7 @@ "acceptance_sd": null, "accepted_per_step": 0.91, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 16325, @@ -123,10 +123,10 @@ "prompt_n": 2046 } ], - "pp": 1492.72, + "pp": 1499.47, "pp_sd": null, "samples": 1, - "tg": 35.21, + "tg": 35.27, "tg_sd": null, "workload": "prose" }, @@ -135,7 +135,7 @@ "acceptance_sd": null, "accepted_per_step": 1.0, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 32719, @@ -146,10 +146,10 @@ "prompt_n": 2028 } ], - "pp": 1489.54, + "pp": 1496.75, "pp_sd": null, "samples": 1, - "tg": 34.21, + "tg": 34.27, "tg_sd": null, "workload": "prose" }, @@ -158,7 +158,7 @@ "acceptance_sd": null, "accepted_per_step": 0.75, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 4096, @@ -169,10 +169,10 @@ "prompt_n": 2042 } ], - "pp": 1522.58, + "pp": 1521.72, "pp_sd": null, "samples": 1, - "tg": 32.83, + "tg": 32.86, "tg_sd": null, "workload": "prose" }, @@ -181,7 +181,7 @@ "acceptance_sd": null, "accepted_per_step": 0.88, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 65448, @@ -192,10 +192,10 @@ "prompt_n": 2057 } ], - "pp": 1415.89, + "pp": 1390.03, "pp_sd": null, "samples": 1, - "tg": 33.5, + "tg": 33.51, "tg_sd": null, "workload": "prose" }, @@ -204,7 +204,7 @@ "acceptance_sd": null, "accepted_per_step": 0.68, "accepted_per_step_sd": null, - "command": "gufo serve --port 45747 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 46959 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 8185, @@ -215,10 +215,10 @@ "prompt_n": 2047 } ], - "pp": 1519.38, + "pp": 1525.11, "pp_sd": null, "samples": 1, - "tg": 32.41, + "tg": 32.47, "tg_sd": null, "workload": "prose" } @@ -226,10 +226,10 @@ "schemaVersion": 1, "server": { "command": [ - "gufo", + "fix-gufo", "serve", "--port", - "45747", + "46959", "--sessions", "1", "llm", @@ -250,12 +250,10 @@ ] }, "source": { - "binaryRevision": "eecb008a909ca299c13d1ef773765c7e2e24d395", - "binarySha256": "2ebf7a2cf22eae84b0c78bf87bc39adda3336fc59f9ecb4e3f1b7c06227571d4", - "buildMode": "nix-release", + "binarySha256": "1ea59e9bfb214b5cdd4d4b6bcc896058df43a06f9c5e180086ea6a8d4301ca58", + "buildMode": "cmake-release-preset", "dirty": false, - "flakeLockSha256": "faabad820ec09caab7351bf48910cec895a38cf1766e098e3065aea86f405ed6", - "revision": "4c5a00d2d609846cb73e79aa33de324624aa9c2c" + "revision": "def2ed1e904f912e69550344aa7c13a99f535d0e" }, "table": "single-mtp", "target": "gufo" diff --git a/docs/models/qwen3.8-flash-next/artifacts/single-mtp-repetition-gufo.json b/docs/models/qwen3.8-flash-next/artifacts/single-mtp-repetition-gufo.json index 8aee40af48..fc29dd5670 100644 --- a/docs/models/qwen3.8-flash-next/artifacts/single-mtp-repetition-gufo.json +++ b/docs/models/qwen3.8-flash-next/artifacts/single-mtp-repetition-gufo.json @@ -1,5 +1,9 @@ { "artifactType": "model-bench-table", + "conditions": { + "powerProfile": "performance", + "profileVerifiedThroughoutSweep": true + }, "fingerprint": { "canonical": { "cppStandard": "C++20", @@ -14,15 +18,15 @@ "gpuName": "AMD Radeon 8060S Graphics", "gpuPciId": "1002:1586", "kernelRelease": "7.1.8", - "memoryTotalBytes": 134309883904, + "memoryTotalBytes": 134309888000, "memoryType": "LPDDR5X (Unified)", "rocmVersion": "7.2.3", "schemaVersion": "1.0.0" }, - "fingerprintId": "46337fb568b92fb76472ec90a0af584eb7d7ae401f9878d8d18b09762e7c4caa", + "fingerprintId": "4dcf1077494084b2d72ecf90594aad3c844b087618d3db6af2cd796f06cb1c36", "schemaVersion": "1.0.0" }, - "measuredOn": "2026-10-06", + "measuredOn": "2026-10-07", "mode": "mtp", "model": { "id": "Qwen3.8-Flash-Next", @@ -30,12 +34,8 @@ }, "notes": [ "pp2048/tg128; depth is a cached conversation prefix of synthetic text (8-token reply); measured turn task is recorded per row; tolerance max(32, 0.500%) on cache_n and prompt_n; actual counts per sample in rows", - "synthetic text: 1.160 tokens/word, template overhead 12 tokens", - "One warmed valid measurement per depth, without repetitions; out-of-tolerance token-count calibration attempts remain in the raw logs.", - "The immutable Nix release binary was built before committing the prefill patch; its production sources match eecb008a and 4c5a00d2. Only tests and documentation changed between those revisions.", - "Reference, concurrency, loading and memory measurements were not rerun. Deep AR/MTP HTTP cache frontiers differ by one token; exact numerical replay is qualified separately with identical prefill boundaries in prefill-final-rows.json." + "synthetic text: 1.160 tokens/word, template overhead 12 tokens" ], - "rawArtifactSha256": "98ac280c411d053f04c6f35dd03262a6dc38a97768e9a91ada878b65260e0cb8", "referenceName": "llama.cpp", "rows": { "0": { @@ -43,7 +43,7 @@ "acceptance_sd": null, "accepted_per_step": 3.74, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 0, @@ -54,10 +54,10 @@ "prompt_n": 2034 } ], - "pp": 1577.11, + "pp": 1521.3, "pp_sd": null, "samples": 1, - "tg": 60.36, + "tg": 59.02, "tg_sd": null, "workload": "repetition" }, @@ -66,7 +66,7 @@ "acceptance_sd": null, "accepted_per_step": 1.37, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 12256, @@ -77,10 +77,10 @@ "prompt_n": 2043 } ], - "pp": 1511.31, + "pp": 1517.04, "pp_sd": null, "samples": 1, - "tg": 42.17, + "tg": 42.19, "tg_sd": null, "workload": "repetition" }, @@ -89,7 +89,7 @@ "acceptance_sd": null, "accepted_per_step": 2.28, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 131000, @@ -100,10 +100,10 @@ "prompt_n": 2040 } ], - "pp": 1349.84, + "pp": 1353.33, "pp_sd": null, "samples": 1, - "tg": 44.34, + "tg": 44.43, "tg_sd": null, "workload": "repetition" }, @@ -112,7 +112,7 @@ "acceptance_sd": null, "accepted_per_step": 2.37, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 16332, @@ -123,10 +123,10 @@ "prompt_n": 2041 } ], - "pp": 1503.69, + "pp": 1502.95, "pp_sd": null, "samples": 1, - "tg": 50.56, + "tg": 50.6, "tg_sd": null, "workload": "repetition" }, @@ -135,7 +135,7 @@ "acceptance_sd": null, "accepted_per_step": 1.78, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 32734, @@ -146,10 +146,10 @@ "prompt_n": 2022 } ], - "pp": 1501.78, + "pp": 1503.08, "pp_sd": null, "samples": 1, - "tg": 44.99, + "tg": 44.98, "tg_sd": null, "workload": "repetition" }, @@ -158,7 +158,7 @@ "acceptance_sd": null, "accepted_per_step": 3.57, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 4096, @@ -169,10 +169,10 @@ "prompt_n": 2038 } ], - "pp": 1528.9, + "pp": 1529.17, "pp_sd": null, "samples": 1, - "tg": 57.61, + "tg": 57.58, "tg_sd": null, "workload": "repetition" }, @@ -181,7 +181,7 @@ "acceptance_sd": null, "accepted_per_step": 1.84, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 65481, @@ -192,10 +192,10 @@ "prompt_n": 2052 } ], - "pp": 1375.45, + "pp": 1349.8, "pp_sd": null, "samples": 1, - "tg": 44.45, + "tg": 43.12, "tg_sd": null, "workload": "repetition" }, @@ -204,7 +204,7 @@ "acceptance_sd": null, "accepted_per_step": 2.28, "accepted_per_step_sd": null, - "command": "gufo serve --port 37661 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", + "command": "fix-gufo serve --port 49787 --sessions 1 llm --model Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf --context 133760 --served-model-name bench --think off --max-pending-per-client 8 --speculative mtp --mtp-model mtp-Qwen3.8-Flash-Next-shared-Q8_0.gguf", "counts": [ { "cache_n": 8189, @@ -215,10 +215,10 @@ "prompt_n": 2042 } ], - "pp": 1528.87, + "pp": 1524.78, "pp_sd": null, "samples": 1, - "tg": 49.22, + "tg": 49.26, "tg_sd": null, "workload": "repetition" } @@ -226,10 +226,10 @@ "schemaVersion": 1, "server": { "command": [ - "gufo", + "fix-gufo", "serve", "--port", - "37661", + "49787", "--sessions", "1", "llm", @@ -250,12 +250,10 @@ ] }, "source": { - "binaryRevision": "eecb008a909ca299c13d1ef773765c7e2e24d395", - "binarySha256": "2ebf7a2cf22eae84b0c78bf87bc39adda3336fc59f9ecb4e3f1b7c06227571d4", - "buildMode": "nix-release", + "binarySha256": "1ea59e9bfb214b5cdd4d4b6bcc896058df43a06f9c5e180086ea6a8d4301ca58", + "buildMode": "cmake-release-preset", "dirty": false, - "flakeLockSha256": "faabad820ec09caab7351bf48910cec895a38cf1766e098e3065aea86f405ed6", - "revision": "4c5a00d2d609846cb73e79aa33de324624aa9c2c" + "revision": "def2ed1e904f912e69550344aa7c13a99f535d0e" }, "table": "single-mtp-repetition", "target": "gufo" diff --git a/src/cli/serve/inference_backend.cpp b/src/cli/serve/inference_backend.cpp index 0ea279aa2e..351a41d2b4 100644 --- a/src/cli/serve/inference_backend.cpp +++ b/src/cli/serve/inference_backend.cpp @@ -2579,8 +2579,17 @@ class QwenFlashNextTextRunner final : public HipTextModelRunner { : model_->PrefillCapacity()}); const std::size_t next_position = offset + consumed; const auto prefix = QwenFlashNextEngineTokens(prompt.first(next_position)); + // The next step most likely takes the same budget; its n-gram rows are + // read during this one and used only if it does. + const std::size_t after = prompt.size() - next_position; + const auto next = QwenFlashNextEngineTokens(prompt.subspan( + next_position, + std::min({max_input_tokens, after, + after <= model_->PrefillThroughCapacity() + ? model_->PrefillThroughCapacity() + : model_->PrefillCapacity()}))); std::string error; - if (!qfn.session().Sync(prefix, &error)) { + if (!qfn.session().Sync(prefix, &error, next)) { qfn.set_position(0); throw std::runtime_error("Qwen3.8-Flash-Next prefill failed: " + error); } diff --git a/src/cli/serve/text_generation_scheduler.cpp b/src/cli/serve/text_generation_scheduler.cpp index a68b4b9c10..6b0152b001 100644 --- a/src/cli/serve/text_generation_scheduler.cpp +++ b/src/cli/serve/text_generation_scheduler.cpp @@ -1074,7 +1074,15 @@ struct TextGenerationScheduler::Impl { request->decode_due = true; decoding.push_back(std::move(request)); } else { - prefilling.push_back(std::move(request)); + // The last-served prefill just moved to the back of its round. + // Give an arrival its first turn before repeating that work, while + // retaining the order of peers that have already been waiting. + auto next = prefilling.end(); + if (!prefilling.empty() && + prefilling.back()->id == last_prefill_request_id) { + --next; + } + prefilling.insert(next, std::move(request)); } } catch (...) { CompleteFailure(request, std::current_exception()); @@ -1083,7 +1091,7 @@ struct TextGenerationScheduler::Impl { } void StepPrefill(const std::shared_ptr& request, - bool decoder_runnable, bool snapshot_pending = false) { + bool decoder_runnable, bool peer_waiting = false) { try { if (CompleteIfStopped(request)) { return; @@ -1091,11 +1099,11 @@ struct TextGenerationScheduler::Impl { request->phase.store(TextRequestPhase::kPrefilling, std::memory_order_release); - // A published first token is also latency-sensitive while its frozen - // prompt is being captured. Bound other prefill work so we can poll - // that capture promptly. Spare slots alone do not change lone prefill. - const bool bounded_prefill = incremental_prefill_supported && - (decoder_runnable || snapshot_pending); + last_prefill_request_id = request->id; + // Bound work for short waiting prefills and captures as well as decoders. + // Long prefills alone retain the model's efficient chunk size. + const bool bounded_prefill = + incremental_prefill_supported && (decoder_runnable || peer_waiting); const std::size_t budget = bounded_prefill ? prefill_policy.decode_active_tokens : request->runner_request.prompt_tokens(); @@ -1786,6 +1794,9 @@ struct TextGenerationScheduler::Impl { const bool assemble_initial_batch = preparing_multi_token_batch && consecutive_active_prefill_chunks == 0 && + prefilling.front()->runner_request.prompt_tokens() - + prefilling.front()->runner_request.prefill_position() <= + prefill_policy.decode_active_tokens && std::all_of(decoding.begin(), decoding.end(), [](const auto& request) { return request->result.tokens.empty(); @@ -1841,7 +1852,13 @@ struct TextGenerationScheduler::Impl { auto request = std::move(prefilling.front()); prefilling.pop_front(); - StepPrefill(request, false, !capturing.empty()); + const bool short_peer_waiting = std::any_of( + prefilling.begin(), prefilling.end(), [&](const auto& peer) { + return peer->runner_request.prompt_tokens() - + peer->runner_request.prefill_position() <= + prefill_policy.decode_active_tokens; + }); + StepPrefill(request, false, !capturing.empty() || short_peer_waiting); if (!IsTerminal(request)) { if (request->runner_request.SnapshotPending()) { capturing.push_back(std::move(request)); @@ -1880,6 +1897,7 @@ struct TextGenerationScheduler::Impl { TextGenerationErrorCode::kDeviceLost, kDeviceLostMessage)); std::vector> device_lost_requests; std::size_t consecutive_active_prefill_chunks{0}; + std::uint64_t last_prefill_request_id{0}; std::atomic next_request_id{1}; std::jthread worker; }; diff --git a/src/models/qwen38_flash_next/engine.cpp b/src/models/qwen38_flash_next/engine.cpp index 0a9f423f47..2fb27d3209 100644 --- a/src/models/qwen38_flash_next/engine.cpp +++ b/src/models/qwen38_flash_next/engine.cpp @@ -20,9 +20,12 @@ namespace gufo::models::qwen38_flash_next { namespace { -// gfx1151 pp4096 at depths 0/4096: the 512/1024/2048/4096 sweep favored -// 2048; larger chunks used more scratch without improving throughput. -constexpr std::uint32_t kPrefillChunkTokens = 2048; +// gfx1151, served repository prompts: 4096-token chunks prefill 24K-102K +// tokens 6-11% faster than 2048 (pp4096 +5% at d0, +6% at d32K). Each +// chunk's n-gram gather doubles; next-chunk prefetch overlaps that read. +// Prefill logits do not depend on the chunk size +// (pre-budget queries always take the dense attention tiles). +constexpr std::uint32_t kPrefillChunkTokens = 4096; // A prompt ending this close past a chunk finishes in that chunk: a separate // tail pass costs about as much as this many more chunk rows. constexpr std::uint32_t kPrefillTailTokens = 128; @@ -577,7 +580,7 @@ bool Session::DraftCatchUpBatch(std::span requests, bool Session::Feed(std::span tokens, std::string* error_msg, bool prefill, std::uint32_t boundary, std::unique_ptr* checkpoint, - double* capture_ms) { + double* capture_ms, std::span after) { rocm::Executor& exec = *model_->executor_; for (std::size_t off = 0; off < tokens.size();) { const auto remaining = tokens.size() - off; @@ -601,8 +604,14 @@ bool Session::Feed(std::span tokens, std::string* error_msg, if (!capture) return false; } + // The chunk this loop, or the following Sync, takes next. + const auto rest = off + n < tokens.size() ? tokens.subspan(off + n) : after; + const auto next = rest.first(std::min( + rest.size() <= exec.max_batch() ? exec.max_batch() + : model_->PrefillCapacity(), + rest.size())); if (!exec.Forward(*session_, chunk, 1, logits_.data(), mode, error_msg, - capture.get())) { + capture.get(), next)) { return false; } const auto kept = capture && capture->tokens < n ? n - capture->tokens : n; @@ -629,9 +638,9 @@ bool Session::Feed(std::span tokens, std::string* error_msg, return true; } -bool Session::Sync(std::span prompt, - std::string* error_msg) { - return SyncImpl(prompt, error_msg, 0, nullptr); +bool Session::Sync(std::span prompt, std::string* error_msg, + std::span next) { + return SyncImpl(prompt, error_msg, 0, nullptr, nullptr, next); } bool Session::SyncThrough(std::span prompt, @@ -653,7 +662,7 @@ bool Session::SyncThrough(std::span prompt, bool Session::SyncImpl(std::span prompt, std::string* error_msg, std::uint32_t boundary, std::unique_ptr* checkpoint, - double* capture_ms) { + double* capture_ms, std::span next) { if (prompt.empty()) { AssignError(error_msg, "prompt is empty"); return false; @@ -680,7 +689,7 @@ bool Session::SyncImpl(std::span prompt, } valid_ = false; const bool ok = Feed(prompt.subspan(common), error_msg, true, boundary, - checkpoint, capture_ms); + checkpoint, capture_ms, next); valid_ = ok; return ok; } diff --git a/src/models/qwen38_flash_next/engine.hpp b/src/models/qwen38_flash_next/engine.hpp index 695443ada7..37df2f9c7d 100644 --- a/src/models/qwen38_flash_next/engine.hpp +++ b/src/models/qwen38_flash_next/engine.hpp @@ -125,9 +125,12 @@ class Session final { Session& operator=(const Session&) = delete; /// Makes the session state equal to `prompt`: keeps the longest common - /// prefix with the current tokens, reprocesses the rest. + /// prefix with the current tokens, reprocesses the rest. `next` may name + /// the tokens the following Sync will append, so their n-gram rows can be + /// read during this one; it never changes the result. [[nodiscard]] bool Sync(std::span prompt, - std::string* error_msg = nullptr); + std::string* error_msg = nullptr, + std::span next = {}); /// Prefills through a stable text boundary in the same forward, returning /// a checkpoint of that boundary while the live session reaches prompt. [[nodiscard]] bool SyncThrough(std::span prompt, @@ -234,11 +237,13 @@ class Session final { bool Feed(std::span tokens, std::string* error_msg, bool prefill = false, std::uint32_t boundary = 0, std::unique_ptr* checkpoint = nullptr, - double* capture_ms = nullptr); + double* capture_ms = nullptr, + std::span after = {}); bool SyncImpl(std::span prompt, std::string* error_msg, std::uint32_t boundary, std::unique_ptr* checkpoint, - double* capture_ms = nullptr); + double* capture_ms = nullptr, + std::span next = {}); std::unique_ptr SaveSnapshotImpl( const rocm::Session& state, std::span tokens, std::span logits, std::uint32_t hidden_rows, diff --git a/src/models/qwen38_flash_next/kernels/rocm/batch.cpp b/src/models/qwen38_flash_next/kernels/rocm/batch.cpp index a8cca29b87..935266bb78 100644 --- a/src/models/qwen38_flash_next/kernels/rocm/batch.cpp +++ b/src/models/qwen38_flash_next/kernels/rocm/batch.cpp @@ -723,6 +723,7 @@ bool Executor::ForwardBatch(std::span items, if (ple_pending_ && !WaitPle(error)) { return false; } + FinishPrefetch(); batch_rows_ = 0; for (std::size_t i = 0; i < items.size(); ++i) { const auto& item = items[i]; diff --git a/src/models/qwen38_flash_next/kernels/rocm/executor.cpp b/src/models/qwen38_flash_next/kernels/rocm/executor.cpp index 77ebf3d701..8ae7902456 100644 --- a/src/models/qwen38_flash_next/kernels/rocm/executor.cpp +++ b/src/models/qwen38_flash_next/kernels/rocm/executor.cpp @@ -840,7 +840,7 @@ void Session::Reset() { Executor::~Executor() { // Readers write pinned staging memory; drain them before freeing it, // including when a forward failed before reaching PLE. - if (ple_pending_) { + if (ple_pending_ || prefetch_.pending) { (void)ngram_->WaitRead(); } (void)hipFree(batch_logits_); @@ -928,33 +928,63 @@ std::unique_ptr Executor::Create(const DeviceModel& model, s.xn_half = Alloc<__half>(a, T * hc_dim, error_msg); s.xn_q8t = Alloc(a, Q8TiledBytes(T, hc_dim), error_msg); s.lo = f32(T * c.hc_low_rank); - s.hc_gate = f32(T * hc_dim); s.mixed = f32(T * hidden); s.inject = f32(T * c.hc_count * HcInjectParts(hidden)); s.block_out = f32(T * hidden); + const std::uint32_t max_blocks = + (c.context_length + c.compress_ratio - 1) / c.compress_ratio; + e->mask_words_ = (max_blocks + 31) / 32; + const std::size_t linear_projection = std::max( + c.SsmConvChannels() + c.SsmValueDim(), c.ple_layer >= 0 ? hc_dim : 0); + const std::size_t linear_floats = + T * (linear_projection + 2 * c.ssm_num_v_heads + 2 * c.SsmKeyDim() + + 2 * c.SsmValueDim()) + + (T + c.ssm_conv_kernel) * c.SsmConvChannels(); + const std::size_t attention_floats = + T * (5 * c.AttentionQDim() + 4 * c.AttentionKvDim() + + c.indexer_heads * c.indexer_head_dim + c.indexer_head_dim + + e->mask_words_); + const std::size_t expert_floats = slots * (2 * c.expert_ff + hidden); + // Linear/full attention and routed experts run sequentially on stream_. + // Their intermediates die before the next stage, including batched rows, + // MTP catch-up and the deferred MoE epilogue consumed by Combine. Keep + // stable addresses for graph replay without reserving all three stages. + float* const stage = + f32(std::max({linear_floats, attention_floats, expert_floats, + T * (linear_projection + hc_dim)})); + if (stage == nullptr) { + return nullptr; + } + // Mixer gates die before attention/experts start. PLE also uses this + // buffer as its key, so keep it beyond PLE's gated output at stage. + s.hc_gate = stage + T * linear_projection; + float* cursor = stage; + const auto take = [&](std::size_t count) { + float* result = cursor; + cursor += count; + return result; + }; // Stacked and separate SSM projections are mutually exclusive. MTP's // projected embedding is consumed before attention and reuses this space. - s.qkvz = f32(T * std::max({c.SsmConvChannels() + c.SsmValueDim(), - c.ple_layer >= 0 ? hc_dim : 0})); + s.qkvz = take(T * linear_projection); s.qkv = s.qkvz; s.z = s.qkvz != nullptr ? s.qkvz + T * c.SsmConvChannels() : nullptr; - s.alpha_beta = f32(T * 2 * c.ssm_num_v_heads); - s.conv_scratch = f32((T + c.ssm_conv_kernel) * c.SsmConvChannels()); - s.qn = f32(T * c.SsmKeyDim()); - s.kn = f32(T * c.SsmKeyDim()); - s.gdn_raw = f32(T * c.SsmValueDim()); - s.gdn_out = f32(T * c.SsmValueDim()); - s.qg = f32(T * (2 * c.AttentionQDim() + 2 * c.AttentionKvDim())); - s.q = f32(T * c.AttentionQDim()); - s.attn_gate = f32(T * c.AttentionQDim()); - s.k = f32(T * c.AttentionKvDim()); - s.v = f32(T * c.AttentionKvDim()); - s.iq = f32(T * c.indexer_heads * c.indexer_head_dim); - s.ik = f32(T * c.indexer_head_dim); - const std::uint32_t max_blocks = - (c.context_length + c.compress_ratio - 1) / c.compress_ratio; - e->mask_words_ = (max_blocks + 31) / 32; - s.mask = Alloc(a, T * e->mask_words_, error_msg); + s.alpha_beta = take(T * 2 * c.ssm_num_v_heads); + s.conv_scratch = take((T + c.ssm_conv_kernel) * c.SsmConvChannels()); + s.qn = take(T * c.SsmKeyDim()); + s.kn = take(T * c.SsmKeyDim()); + s.gdn_raw = take(T * c.SsmValueDim()); + s.gdn_out = take(T * c.SsmValueDim()); + cursor = stage; + s.qg = take(T * (2 * c.AttentionQDim() + 2 * c.AttentionKvDim())); + s.q = take(T * c.AttentionQDim()); + s.attn_gate = take(T * c.AttentionQDim()); + s.k = take(T * c.AttentionKvDim()); + s.v = take(T * c.AttentionKvDim()); + s.iq = take(T * c.indexer_heads * c.indexer_head_dim); + s.ik = take(T * c.indexer_head_dim); + s.mask = reinterpret_cast(take(T * e->mask_words_)); + s.ctx = take(T * c.AttentionQDim()); // Keep score/selection traffic near the device cache size. Captured // verification still needs all its rows at the maximum context. const std::size_t score_stride = std::size_t{e->mask_words_} * 32; @@ -962,7 +992,6 @@ std::unique_ptr Executor::Create(const DeviceModel& model, std::max(kVecBatch * score_stride, std::min(std::size_t{5 * 1024 * 1024}, 512 * score_stride)); s.scores = f32(e->select_score_floats_); - s.ctx = f32(T * c.AttentionQDim()); s.attn_partials = f32(static_cast(kVecBatch) * c.num_heads * kAttnSplits * (c.head_dim + 2)); if (c.ple_layer >= 0) { @@ -999,9 +1028,10 @@ std::unique_ptr Executor::Create(const DeviceModel& model, Alloc(a, 3 * RoutedTileCapacity(slots, c), error_msg); } s.weights = f32(slots); - s.gate_e = f32(slots * c.expert_ff); - s.up_e = f32(slots * c.expert_ff); - s.down_e = f32(slots * hidden); + cursor = stage; + s.gate_e = take(slots * c.expert_ff); + s.up_e = take(slots * c.expert_ff); + s.down_e = take(slots * hidden); s.shexp_gate = f32(T * c.shared_expert_ff); s.shexp_up = f32(T * c.shared_expert_ff); s.shexp_out = f32(T * hidden); @@ -1812,6 +1842,16 @@ bool Executor::PleFetch(Session& session, std::span tokens, if (ple_pending_ && !WaitPle(error_msg)) { return false; } + if ((prefetch_.pending || prefetch_.ready) && !speculative && + session.ngram_ == prefetch_.before && + std::ranges::equal(tokens, prefetch_.tokens)) { + // These are this batch's rows: WaitPle collects them. + session.ngram_ = prefetch_.after; + prefetch_.claimed = true; + ple_pending_ = true; + return true; + } + FinishPrefetch(); const Config& c = config(); const auto n = static_cast(tokens.size()); // Only proper prefixes need snapshots; the full batch keeps its live state. @@ -1839,7 +1879,20 @@ bool Executor::PleFetch(Session& session, std::span tokens, } bool Executor::WaitPle(std::string* error_msg) const { - const bool ok = ple_pending_ && ngram_->WaitRead(); + bool ok = ple_pending_; + if (ok && prefetch_.claimed) { + FinishPrefetch(); + ok = prefetch_.ready; + if (ok) { + std::copy_n(prefetch_.rows.data(), + prefetch_.tokens.size() * config().PleEmbeddingDim(), + host_emb_); + } + prefetch_.ready = false; + prefetch_.claimed = false; + } else if (ok) { + ok = ngram_->WaitRead(); + } ple_pending_ = false; if (!ok) { AssignError(error_msg, "n-gram table read failed"); @@ -1847,6 +1900,37 @@ bool Executor::WaitPle(std::string* error_msg) const { return ok; } +void Executor::PrefetchPle(const Session& session, + std::span next) const { + const Config& c = config(); + if (ple_pending_ || prefetch_.pending || next.empty() || + next.size() > options_.max_batch) { + return; + } + // Hashed from the history after this batch, which PleFetch already + // advanced; the session keeps it until the next batch claims the rows. + prefetch_.ready = false; + prefetch_.before = session.ngram_; + prefetch_.after = session.ngram_; + const std::span rows(host_rows_.data(), + next.size() * c.ple_heads); + HashNgramRows(c, prefetch_.after, next, rows); + prefetch_.tokens.assign(next.begin(), next.end()); + const std::size_t count = next.size() * c.PleEmbeddingDim(); + if (prefetch_.rows.size() < count) { + prefetch_.rows.resize(count); + } + prefetch_.pending = + ngram_->StartRead(rows, std::span(prefetch_.rows.data(), count)); +} + +void Executor::FinishPrefetch() const { + if (prefetch_.pending) { + prefetch_.ready = ngram_->WaitRead(); + prefetch_.pending = false; + } +} + bool Executor::Ple(const DeviceLayer& l, Session& session, std::uint32_t n, float* res, bool speculative, std::string* error_msg, bool embeddings_ready, PrefillCheckpoint* checkpoint) const { @@ -1860,6 +1944,10 @@ bool Executor::Ple(const DeviceLayer& l, Session& session, std::uint32_t n, "n-gram upload", error_msg))) { return false; } + if (!embeddings_ready && !prefetch_next_.empty()) { + PrefetchPle(session, prefetch_next_); + prefetch_next_ = {}; + } const std::uint32_t hc_dim = c.HcDim(); Q8Input emb; if (!Quantize(s_.ple_emb, n, c.PleEmbeddingDim(), &emb, error_msg) || @@ -2114,36 +2202,46 @@ bool Executor::Attention(const DeviceLayer& l, Session::AttentionState& s, "partial attention output initialization", error_msg)) { return false; } + // Queries before the token budget always take the dense tiles, also in a + // batch that crosses it, so each position's key sweep (and its rounding) + // does not depend on where the prefill chunks fall. The sparse part goes + // first: it is the launch that can be refused. + const auto attend = [&](std::uint32_t first, std::uint32_t rows, bool last) { + const std::uint32_t pos = start_pos + first; + const std::uint32_t dense_rows = + mask != nullptr && !last && pos < c.indexer_top_k + ? std::min(rows, c.indexer_top_k - pos) + : 0; + const auto at = [&](std::uint32_t row) { + return std::size_t{first + row} * c.AttentionQDim(); + }; + return (dense_rows == rows || + WmmaCausalAttention( + s_.q + at(dense_rows), s_.attn_gate + at(dense_rows), s.k_cache, + s.v_cache, + mask ? mask + std::size_t{first + dense_rows} * mask_words_ + : nullptr, + mask_words_, s_.ctx + at(dense_rows), rows - dense_rows, + pos + dense_rows, c.num_heads, c.num_kv_heads, c.head_dim, + c.compress_ratio, stream_, last)) && + (dense_rows == 0 || + WmmaCausalAttention(s_.q + at(0), s_.attn_gate + at(0), s.k_cache, + s.v_cache, nullptr, mask_words_, s_.ctx + at(0), + dense_rows, pos, c.num_heads, c.num_kv_heads, + c.head_dim, c.compress_ratio, stream_, false)); + }; if (checkpoint_tokens != 0 && checkpoint_tokens < n_tokens) { // Sparse tiles compact the union of their queries' selected keys. // Keep the checkpoint's query grouping identical to a pass ending there. - const auto prefix_sparse = - c.compress_ratio > 0 && start_pos + checkpoint_tokens > c.indexer_top_k; - const auto tail = n_tokens - checkpoint_tokens; - const auto offset = std::size_t{checkpoint_tokens} * c.AttentionQDim(); - if (!WmmaCausalAttention(s_.q, s_.attn_gate, s.k_cache, s.v_cache, - prefix_sparse ? mask : nullptr, mask_words_, - s_.ctx, checkpoint_tokens, start_pos, c.num_heads, - c.num_kv_heads, c.head_dim, c.compress_ratio, - stream_) || - !WmmaCausalAttention( - s_.q + offset, s_.attn_gate + offset, s.k_cache, s.v_cache, - mask ? mask + std::size_t{checkpoint_tokens} * mask_words_ - : nullptr, - mask_words_, s_.ctx + offset, tail, start_pos + checkpoint_tokens, - c.num_heads, c.num_kv_heads, c.head_dim, c.compress_ratio, - stream_)) { + if (!attend(0, checkpoint_tokens, false) || + !attend(checkpoint_tokens, n_tokens - checkpoint_tokens, false)) { AssignError(error_msg, "checkpoint attention geometry is unsupported"); return false; } return !project_output || Dense(l.attn_out, s_.ctx, out, n_tokens, error_msg); } - if (MatrixRows(n_tokens) && - WmmaCausalAttention(s_.q, s_.attn_gate, s.k_cache, s.v_cache, mask, - mask_words_, s_.ctx, n_tokens, start_pos, c.num_heads, - c.num_kv_heads, c.head_dim, c.compress_ratio, stream_, - last_only)) { + if (MatrixRows(n_tokens) && attend(0, n_tokens, last_only)) { return !project_output || Dense(l.attn_out, s_.ctx, out, n_tokens, error_msg); } @@ -2443,8 +2541,16 @@ bool Executor::Run(Session& session, std::uint64_t key, bool graph, bool Executor::Forward(Session& session, std::span tokens, std::uint32_t n_logits, float* logits, ForwardMode mode, - std::string* error_msg, - PrefillCheckpoint* checkpoint) const { + std::string* error_msg, PrefillCheckpoint* checkpoint, + std::span next) const { + // Only this call's PLE layer may start the prefetch of `next`. + struct NextHint { + std::span& hint; + ~NextHint() { hint = {}; } + } next_hint{prefetch_next_}; + if (mode == ForwardMode::kPrefill) { + prefetch_next_ = next; + } const bool speculative = mode == ForwardMode::kVerify; PrefillPhase phase(mode == ForwardMode::kPrefill); selected_logits_ = nullptr; diff --git a/src/models/qwen38_flash_next/kernels/rocm/executor.hpp b/src/models/qwen38_flash_next/kernels/rocm/executor.hpp index b7bbcad441..2539dfb909 100644 --- a/src/models/qwen38_flash_next/kernels/rocm/executor.hpp +++ b/src/models/qwen38_flash_next/kernels/rocm/executor.hpp @@ -258,12 +258,16 @@ class Executor { /// verification. The final wide residual of those tokens stays on /// the device for MtpForward. kVerify permits Rollback (at most /// max_speculative rows). kPrefill uses consistent prompt arithmetic at - /// every chunk width. + /// every chunk width. A kPrefill call may name the `next` chunk of the same + /// prompt: once its own n-gram rows have arrived it starts reading the + /// next chunk's, so that read overlaps this chunk's layers. The rows are + /// used only if the next batch has exactly those tokens. [[nodiscard]] bool Forward(Session& session, std::span tokens, std::uint32_t n_logits, float* logits, ForwardMode mode, std::string* error_msg, - PrefillCheckpoint* checkpoint = nullptr) const; + PrefillCheckpoint* checkpoint = nullptr, + std::span next = {}) const; struct BatchItem { Session* session; @@ -454,6 +458,12 @@ class Executor { bool PleFetch(Session& s, std::span tokens, bool speculative, std::string* error_msg) const; bool WaitPle(std::string* error_msg) const; + /// Starts reading `next`'s rows into the prefetch buffer. Best effort: a + /// prefetch that cannot start is skipped. + void PrefetchPle(const Session& s, std::span next) const; + /// Waits for a prefetch in flight before another read; its rows stay + /// claimable. + void FinishPrefetch() const; bool Ple(const DeviceLayer& l, Session& s, std::uint32_t n_tokens, float* res, bool speculative, std::string* error_msg, bool embeddings_ready = false, @@ -631,6 +641,22 @@ class Executor { float* host_emb_{nullptr}; mutable std::vector host_rows_; mutable bool ple_pending_{false}; + /// The rows a prefill chunk reads are a pure function of its tokens and + /// the n-gram history before them, so prefetched rows are claimed only by + /// a batch with exactly those tokens and that history, whichever session + /// it belongs to. A claim copies them into host_emb_, the only buffer + /// uploads (and captured graphs) read. + struct PlePrefetch { + bool pending{false}; ///< read in flight + bool ready{false}; ///< rows complete in `rows` + bool claimed{false}; ///< WaitPle collects them for the current batch + NgramHistory before; + NgramHistory after; + std::vector tokens; + std::vector rows; + }; + mutable PlePrefetch prefetch_; + mutable std::span prefetch_next_; // Pinned host staging the launched (or captured) work reads and writes. Session::Control* control_host_{nullptr}; std::int32_t* tokens_host_{nullptr}; diff --git a/src/models/qwen38_flash_next/ngram.cpp b/src/models/qwen38_flash_next/ngram.cpp index cf20456bb7..8412c80e4e 100644 --- a/src/models/qwen38_flash_next/ngram.cpp +++ b/src/models/qwen38_flash_next/ngram.cpp @@ -15,7 +15,8 @@ namespace gufo::models::qwen38_flash_next { namespace { constexpr std::size_t kPage = 4096; -// Keep enough direct reads outstanding while layer 0 runs. +// Next-chunk prefetch hides large gathers behind the current transformer +// pass. More readers increased CPU work without improving served prefill. constexpr std::size_t kWorkers = 32; constexpr std::size_t kReadBatch = 8; constexpr std::size_t kBatchJobs = 1024; @@ -109,9 +110,7 @@ std::unique_ptr NgramTable::Open( } return nullptr; } - const std::size_t workers = std::min( - kWorkers, static_cast(std::thread::hardware_concurrency())); - for (std::size_t i = 0; i < std::max(1, workers); ++i) { + for (std::size_t i = 0; i < kWorkers; ++i) { t->workers_.emplace_back([raw = t.get()] { raw->Worker(); }); } return t; @@ -273,8 +272,13 @@ bool NgramTable::StartRead(std::span rows, pending_ = jobs_.size(); failed_ = false; active_ = true; - if (pending_ != 0) { + // A decode gather has a handful of reads: wake one reader per read, not + // the whole pool. Readers keep taking jobs until the queue is empty. + if (pending_ >= workers_.size()) { wake_.notify_all(); + } else { + for (std::size_t i = 0; i < pending_; ++i) + wake_.notify_one(); } return true; } diff --git a/src/models/qwen38_flash_next/ngram.hpp b/src/models/qwen38_flash_next/ngram.hpp index 7f97373090..188975bf99 100644 --- a/src/models/qwen38_flash_next/ngram.hpp +++ b/src/models/qwen38_flash_next/ngram.hpp @@ -25,6 +25,7 @@ struct NgramHistory { std::array prev{kNone, kNone}; void Reset() noexcept { prev.fill(kNone); } + bool operator==(const NgramHistory&) const = default; }; /// Computes the table rows every token gathers: `ple_heads` rows per token, diff --git a/tests/cli/text_generation_scheduler_test.cpp b/tests/cli/text_generation_scheduler_test.cpp index ed5db55936..a84a1ae897 100644 --- a/tests/cli/text_generation_scheduler_test.cpp +++ b/tests/cli/text_generation_scheduler_test.cpp @@ -1165,6 +1165,83 @@ void TestMultiResidentPrefillUsesBoundedWorkUnits() { } } +void TestArrivalDoesNotWaitForAnotherLongChunk() { + for (const bool speculative : {false, true}) { + auto control = std::make_shared(); + control->multi_token_decode = speculative; + control->batched_multi_token_decode = speculative; + control->supports_batched_advance = speculative; + control->prefill_capacity = 4; + control->block_prefill_label = 1; + auto scheduler = MakeScheduler(control, 2, {.decode_active_tokens = 2}); + auto long_request = scheduler->Submit( + {1, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20}, 4, 0.0F); + control->WaitForPrefill(1); + auto short_request = scheduler->Submit({2}, 1, 0.0F); + control->ReleasePrefill(); + Expect(short_request.Wait().tokens == ExpectedTokens(2, 1), + "the arrival produces its independent first token"); + Expect(long_request.Wait().tokens == ExpectedTokens(1, 4), + "the interrupted prefill retains its trajectory"); + const auto events = control->Events(); + const auto second_long = EventIndex(events, EventKind::kPrefill, 1, 1); + Expect(EventIndex(events, EventKind::kPrefill, 2) < second_long, + "a newcomer prefills before another chunk of the previous request"); + Expect(EventIndex(events, EventKind::kAdvance, 2) < second_long, + "a ready short request does not wait to batch with a distant peer"); + } +} + +void TestLongPrefillsKeepWideFairTurns() { + auto control = std::make_shared(); + control->prefill_capacity = 4; + control->block_prefill_label = 1; + auto scheduler = MakeScheduler(control, 2, {.decode_active_tokens = 2}); + auto first = scheduler->Submit( + {1, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20}, 2, 0.0F); + control->WaitForPrefill(1); + auto second = scheduler->Submit( + {2, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30}, 2, 0.0F); + control->ReleasePrefill(); + Expect(first.Wait().tokens == ExpectedTokens(1, 2) && + second.Wait().tokens == ExpectedTokens(2, 2), + "interleaved prefills preserve both outputs"); + const auto events = control->Events(); + const auto a = EventIndex(events, EventKind::kPrefill, 1, 1); + const auto b = EventIndex(events, EventKind::kPrefill, 2); + const auto next_b = EventIndex(events, EventKind::kPrefill, 2, 1); + Expect(b < a && a < next_b, + "the newcomer gets one turn, then the older prefill progresses"); + Expect( + events[a].count == 4 && events[b].count == 4 && events[next_b].count == 4, + "long waiting prefills keep wide chunks before any decoder is ready"); +} + +void TestShortArrivalBehindWaitingPrefill() { + auto control = std::make_shared(); + control->prefill_capacity = 4; + control->block_prefill_label = 1; + auto scheduler = MakeScheduler(control, 3, {.decode_active_tokens = 2}); + auto first = scheduler->Submit( + {1, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20}, 2, 0.0F); + control->WaitForPrefill(1); + auto second = scheduler->Submit({2, 20, 21, 22, 23, 24, 25, 26}, 2, 0.0F); + auto short_request = scheduler->Submit({3}, 1, 0.0F); + control->ReleasePrefill(); + Expect(first.Wait().tokens == ExpectedTokens(1, 2) && + second.Wait().tokens == ExpectedTokens(2, 2) && + short_request.Wait().tokens == ExpectedTokens(3, 1), + "three independent requests preserve their outputs"); + const auto events = control->Events(); + const auto second_start = EventIndex(events, EventKind::kPrefill, 2); + Expect(events[second_start].count == 2, + "a long peer gets bounded work ahead of a waiting short request"); + Expect(second_start < EventIndex(events, EventKind::kPrefill, 3) && + EventIndex(events, EventKind::kAdvance, 3) < + EventIndex(events, EventKind::kPrefill, 1, 1), + "the short newcomer gets its first token without repeating the round"); +} + void TestDecodeActivePrefillIsBounded(bool multi_token, bool batched) { auto control = std::make_shared(); control->multi_token_decode = multi_token; @@ -2775,6 +2852,9 @@ int main() { TestBatchFailureIsolation(); TestBatchDeviceLossFailsSuccessfulPeers(); TestMultiResidentPrefillUsesBoundedWorkUnits(); + TestArrivalDoesNotWaitForAnotherLongChunk(); + TestLongPrefillsKeepWideFairTurns(); + TestShortArrivalBehindWaitingPrefill(); for (const bool multi_token : {false, true}) { for (const bool batched : {false, true}) { if (batched && !multi_token) diff --git a/tests/functional/README.md b/tests/functional/README.md index ad322afd04..9afb6fd81d 100644 --- a/tests/functional/README.md +++ b/tests/functional/README.md @@ -96,6 +96,7 @@ draft limit for this suite. Audio and image/video generation have separate tests | `batch` | Independent requests across Chat, Responses and Completions; sessions 1–8 | | `progress` | Opt-in progress on all text endpoints; output/sampling equality, limits, stops, images, batching and cancel/resume | | `stream-start` | Plain streams on all text endpoints send headers before a cold prefill completes; a stream queued behind every session sends them after the five-second bound | +| `prefill-scheduling` | Short arrival during a cold prefill: work-aligned arrival, independent per-request timings, unchanged output and no repeated prefill. Requires sessions ≥2 and context ≥16384; not in `all` | | `long-context` | Longer multi-turn recall, endpoint switching, sampled JSON and cancellation | | `metrics` | Live slots, Prometheus cache/time/draft counters, uncached work, endpoint totals, queueing and cancellation | | `cache` | Interrupted text/thinking/tool/image histories, ordinary and legacy tool names, RAM and disk restart; disk checkpoint spacing for a growing conversation and a branch restored after restart | diff --git a/tests/functional/openai_sdk.py b/tests/functional/openai_sdk.py index 9468bf41f6..3bf68cf51a 100644 --- a/tests/functional/openai_sdk.py +++ b/tests/functional/openai_sdk.py @@ -39,6 +39,7 @@ from cache_growth import check_cache_growth from cache_depth import check_cache_depth from cache_rotation import check_cache_rotation +from prefill_scheduling import check_prefill_scheduling class CompletionStreamChoice(CompletionChoice): @@ -2525,7 +2526,7 @@ def usage(snapshot): SDK_SUITES = ("discovery", "responses", "stops", "conversation", "image-inputs", "image-count", "structured", "structured-limits", "tool-reasoning", "reasoning-separator", "tools", "auto-tools", "tool-edges", "tool-agent", "tool-agent-loop", "tool-history", "tool-untyped", "tool-mixed", "tool-native-schemas", "tool-native-types", "tool-schema-edges", "sampling-defaults", "sampling-ranges", "batch", - "long-context", "state-edges", "progress", "stream-start", "metrics", "cache-edits", "cache-growth", "cache-depth", "cache-rotation", "cache-concurrency", "cache-shared-prefix", + "long-context", "state-edges", "progress", "stream-start", "prefill-scheduling", "metrics", "cache-edits", "cache-growth", "cache-depth", "cache-rotation", "cache-concurrency", "cache-shared-prefix", "cache-bridge", "system-injection") @@ -2657,6 +2658,8 @@ async def async_local_only(request): args.allow_missing_progress), "stream-start": lambda: check_stream_start( client, args.model, checks, args.concurrency, args.context), + "prefill-scheduling": lambda: check_prefill_scheduling( + client, args.model, checks, chat_result, args.concurrency, args.context), "metrics": lambda: check_server_metrics(client, args.model, checks, args.concurrency, args.context, args.speculative), "cache-edits": lambda: check_cache_edits(client, args.model, checks, chat_result), @@ -2674,7 +2677,7 @@ async def async_local_only(request): "system-injection": lambda: check_system_injection( client, args.model, checks, chat_result), } - selected = ([name for name in suites if name not in ("tool-native-types", "cache-bridge") + selected = ([name for name in suites if name not in ("tool-native-types", "cache-bridge", "prefill-scheduling") and (name not in ("image-inputs", "image-count") or args.vision)] if args.suite == "all" else ["native-tools", "auto-tools"] if args.suite == "tools" else [args.suite]) diff --git a/tests/functional/prefill_scheduling.py b/tests/functional/prefill_scheduling.py new file mode 100644 index 0000000000..499a33b719 --- /dev/null +++ b/tests/functional/prefill_scheduling.py @@ -0,0 +1,71 @@ +"""A small request must make progress beside a longer cold prefill.""" + +from concurrent.futures import ThreadPoolExecutor +from copy import deepcopy +import threading + +from progress import ProgressTrace +from stream_start import cold_prompt + + +def check_prefill_scheduling(client, model, checks, chat_result, width, context): + assert width >= 2 and context >= 16384, "use --sessions 2 --context 32768 or larger" + common = dict(model=model, temperature=0, seed=31, top_p=1, + max_completion_tokens=8, presence_penalty=0, frequency_penalty=0, + reasoning_effort="none", extra_body={ + "top_k": 0, "min_p": 0, "repeat_penalty": 1, + "cache_prompt": False, "return_progress": True, + "chat_template_kwargs": {"enable_thinking": False}}) + short = {**deepcopy(common), "messages": [ + {"role": "user", "content": "What is 2 + 2? Reply only the digit."}]} + control = chat_result(client, short, True) + checks["prefill_short_control"] = control + + # The first positive progress event aligns arrival by work rather than + # wall time. Keep a real multi-chunk prefill, without filling the context. + long = {**deepcopy(common), "max_completion_tokens": 1, "messages": [ + {"role": "user", "content": cold_prompt("prefill_scheduling", min(context, 65536))}]} + ready = threading.Event() + lock = threading.Lock() + trace = ProgressTrace(True) + at_arrival = None + at_output = None + + def long_progress(chunk): + nonlocal at_arrival + with lock: + trace(chunk) + if trace.updates: + update = trace.updates[-1] + if 0 < update["processed"] < update["total"] and at_arrival is None: + at_arrival = dict(update) + ready.set() + + def short_progress(chunk): + nonlocal at_output + if any(choice.delta.content for choice in chunk.choices): + with lock: + if at_output is None: + at_output = dict(trace.updates[-1]) + + with ThreadPoolExecutor(max_workers=1) as pool: + pending = pool.submit(chat_result, client, long, True, long_progress) + assert ready.wait(120), "long prompt did not publish intermediate progress" + result = chat_result(client, short, True, short_progress) + completed = pending.result(timeout=180) + + # Mark the cohort on its owner thread; timing records still retain each + # request and phase separately. Faster peers cannot hide a slow request. + checks["prefill_overlap"] = { + "short": result, "long": completed, "arrival": at_arrival, "first_output": at_output} + for key in ("text", "reasoning", "tools", "finish"): + assert result[key] == control[key], (key, control, result) + assert result["text"].strip().rstrip(".") == "4", result + assert at_output and at_output["processed"] < at_output["total"], ( + "short output waited for the entire peer prefill", at_output) + for row in (control, result, completed): + usage = row["usage"] + assert usage["cached_tokens"] == 0, usage + assert usage["gufo"]["prefill_tokens"] == usage["prompt_tokens"], usage + assert completed["usage"]["completion_tokens"] == 1, completed + assert trace.updates[-1]["processed"] == trace.updates[-1]["total"], trace.updates diff --git a/tests/functional/run.py b/tests/functional/run.py index bff7052df7..9bfa55f4e9 100644 --- a/tests/functional/run.py +++ b/tests/functional/run.py @@ -31,7 +31,7 @@ SUITES = ("discovery", "responses", "stops", "conversation", "image-inputs", "image-count", "structured", "structured-limits", "tool-reasoning", "reasoning-separator", "tools", "auto-tools", "tool-edges", "tool-agent", "tool-agent-loop", "tool-history", "tool-untyped", "tool-mixed", "tool-native-schemas", "tool-native-types", "tool-schema-edges", "sampling-defaults", "sampling-ranges", "batch", - "long-context", "state-edges", "progress", "stream-start", "metrics", "cache-edits", "cache-growth", "cache-depth", "cache-rotation", "cache-concurrency", "cache-shared-prefix", "cache-bridge", "system-injection", "cache") + "long-context", "state-edges", "progress", "stream-start", "prefill-scheduling", "metrics", "cache-edits", "cache-growth", "cache-depth", "cache-rotation", "cache-concurrency", "cache-shared-prefix", "cache-bridge", "system-injection", "cache") SAMPLING = { "--temperature": ("temperature", float), "--top-p": ("top_p", float), "--top-k": ("top_k", int), "--min-p": ("min_p", float), @@ -48,7 +48,7 @@ def provenance(): source = hashlib.sha256() - for name in ("run.py", "metrics.py", "progress.py", "stream_start.py", "server_metrics.py", "openai_sdk.py", "continuation.py", + for name in ("run.py", "metrics.py", "progress.py", "stream_start.py", "prefill_scheduling.py", "server_metrics.py", "openai_sdk.py", "continuation.py", "tool_reasoning.py", "tool_agent.py", "tool_native.py", "discovery.py", "image_inputs.py", "cache_edits.py", "cache_growth.py", "cache_depth.py", "cache_rotation.py", "cache_concurrency.py", "cache_shared_prefix.py", "cache_bridge.py", "system_injection.py", "cache_disk_spacing.py"): source.update((TESTS / name).read_bytes()) @@ -246,7 +246,7 @@ def main(): if "all" in selected: if len(selected) != 1: parser.error("all cannot be combined with other suites") - selected = [suite for suite in SUITES if suite not in ("auto-tools", "tool-native-types", "cache-bridge") + selected = [suite for suite in SUITES if suite not in ("auto-tools", "tool-native-types", "cache-bridge", "prefill-scheduling") and (suite not in ("image-inputs", "image-count") or option(command, "--mmproj") is not None)] selected = list(dict.fromkeys(selected)) diff --git a/tests/models/qwen38_flash_next/gpu_probe.cpp b/tests/models/qwen38_flash_next/gpu_probe.cpp index 8c0e3783fd..85467370fd 100644 --- a/tests/models/qwen38_flash_next/gpu_probe.cpp +++ b/tests/models/qwen38_flash_next/gpu_probe.cpp @@ -321,10 +321,15 @@ int main(int argc, char** argv) { const auto rows = static_cast( std::min(n, options.max_logit_rows)); std::vector out(static_cast(rows) * c.vocab_size); + // Name the following chunk as Session::Feed does, so its n-gram read + // overlaps this one. + const std::size_t next = + std::min(batch, tokens.size() - off - n); if (!executor->Forward( *session, std::span(tokens.data() + off, n), - rows, out.data(), q::rocm::Executor::ForwardMode::kPrefill, - &error)) { + rows, out.data(), q::rocm::Executor::ForwardMode::kPrefill, &error, + nullptr, + std::span(tokens.data() + off + n, next))) { std::fprintf(stderr, "prefill failed: %s\n", error.c_str()); return 1; }