From 5e575f75490200a0bf82afd1d0ac6226237b5352 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 18:29:12 +0000 Subject: [PATCH 1/2] =?UTF-8?q?spike(quant):=20MXFP4-FLASH-OCCUPANCY=20?= =?UTF-8?q?=E2=80=94=20open=20the=20owed=20ncu=20diff;=20free=20nsys=20re-?= =?UTF-8?q?read=20reframes=20the=20premise?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Opens row/QUANT-CT-MXFP4-FLASH-OCCUPANCY (helper) for the ncu diff #69 owed after the box OOM-reboot lost the vLLM-side capture. A free re-read of the existing c8 nsys CSVs (gpu_trace_c8_dflt vs vllm_offline_trace_c8) reframes the premise BEFORE any GPU work: - ours-vs-vLLM c8 decode flash_fwd_splitkv grid is IDENTICAL: 1x3x64 (num_splits=3, gridZ=64=batch*kv_heads) — same GQA-pack, same split heuristic. - ours uses FEWER registers (216) than vLLM (241) => ours has MORE occupancy headroom, not less. The #69 "8.3% occupancy, register-limited" figure was the batch-1 short-context num_splits=1 kernel (Split=false); the prior ncu pair was context-MISMATCHED (ours ~5-token cli prompt vs vLLM lens=1024). - so at the REAL failing c8 kernel occupancy is NOT vLLM's advantage; the +18us (offline; +12.5 online) residual is codegen (our native sm_121a SASS vs vLLM's sm_80-PTX driver-JIT SASS) or L2 — the matched-workload ncu will disambiguate. Spec .agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md carries the spike contract (W1 matched-c8 ncu, W2 mirror-first codegen lever, W3 binding+verdict). STATUS/BENCHMARKS keyed rows updated (spike-open); STATUS ratchet lowered 3. check-fusion-consistency (minimax_h3_video_vae_device) is a pre-existing red, untouched by this change. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .../qwen3-mxfp4-flash-occupancy-2026-08-06.md | 59 +++++++++++++++++++ docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- scripts/check-public-doc-tables.py | 2 +- 4 files changed, 62 insertions(+), 3 deletions(-) create mode 100644 .agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md diff --git a/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md b/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md new file mode 100644 index 000000000..1e120618c --- /dev/null +++ b/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md @@ -0,0 +1,59 @@ +# QUANT-CT-MXFP4-FLASH-OCCUPANCY — the owed ours-vs-vLLM flash decode ncu diff + the occupancy/L2/codegen lever + + +Row: `QUANT-CT-MXFP4-FLASH-OCCUPANCY` (helper, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`). +Base: `origin/main` `f7a1e322`. Vehicle: `Yi30/Qwen3-8B-MXFP4` (dense +`Qwen3ForCausalLM`, W4A16 Marlin keep-quant); oracle arm +`VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`. GB10 sm_121a. + +## Why this row exists + +`#68/#69` (FLASH-AUDIT) closed the compile lens: our flash source is byte-identical +to vLLM's pinned `2c839c33`; `-use_fast_math` was TRIED and REJECTED (measured ++21 us/call regression — the 246→255 reg bump cuts occupancy on a latency-bound +kernel). Fresh same-tool nsys: ours 168.8 vs vLLM 156.3 us/call at c8 (+12.5, ++450 us/step). `#69` characterized OURS (occ 8.3%, L2 53%, ~38% smem-scoreboard + +~37% CTA-barrier stalls) but the vLLM-side ncu was LOST to a box OOM-reboot — the +precise ours-vs-vLLM ncu diff is OWED. This row runs that diff on an idle box and +takes whichever lever it names. + +## W0 free finding (from existing c8 nsys CSVs, no GPU) — the premise is REFRAMED + +`gpu_trace_c8_dflt` (ours) and `vllm_offline_trace_c8` (vLLM), flash_fwd_splitkv +non-combine, c8 dominant decode: + +| engine | grid (GrdX×Y×Z) | blk | Reg/Trd | mean us/call | combine | +|---|---|---|---|---|---| +| OURS | 1×3×64 | 128 | **216** | **174.2** | 64×1×1 reg46 3.7us | +| vLLM | 1×3×64 | 128 | **241** | **155.8** | 64×1×1 reg64 3.9us | + +The grid is **IDENTICAL** (num_splits=3, gridZ=64=batch×kv_heads — same GQA-pack, +same split heuristic), and **ours uses FEWER registers (216 < 241)** → ours has +MORE occupancy headroom, not less. The `#69` "8.3% occupancy, register-limited" +number was the batch-1 short-context num_splits=1 kernel (a DIFFERENT regime); the +prior ncu pair was context-MISMATCHED (ours ~5-token cli prompt vs vLLM lens=1024). +So at the real c8 decode kernel, occupancy is NOT the vLLM advantage. The residual +is codegen (our native sm_121a SASS vs vLLM's sm_80-PTX driver-JIT SASS) or L2. + +## Work items + +- **W1** — matched-workload ncu diff on BOTH engines at c8 (grid 1×3×64): + OURS `vllm-bench --input-len 1024 --output-len 128 --concurrency 8`; vLLM + `vllm_offline_decode.py NPROMPTS=8`. Full section set (LaunchStats, Occupancy, + SOL, MemoryWorkload/L2/DRAM, WarpStateStats, SchedulerStats). Table the diff; + disambiguate codegen (SM throughput / instruction stats differ) vs L2 (memory + section differs). +- **W2** — the lever the diff names, mirror-first. Leading candidate: compile our + flash TUs for `80-virtual` (compute_80 PTX → driver JIT to sm_121), MIRRORING + vLLM's exact build path, if the diff shows codegen. A build-flag change like + fast-math — runtime arbitrates; keep the `#69` lesson (measure, never assume). + Byte-exact-first; near-tie razor + full battery if the reduction order shifts. +- **W3** — binding c1..c8 ×3 production defaults vs `1.020/0.962/0.966/0.969`; + THE PARITY VERDICT (≥1.0 every axis ⇒ MXFP4 DONE; short ⇒ honest residual map). + +## Gates + +Byte-exact razor: #44 MXFP4-8B smoke 3/3 token-exact + coherent. If reduction +order shifts: SACRED 0.6B/4B distributional + 32B strict, async, memcheck, eager ++ graphed. Box safety: BOTH flock locks, free -g ≥ 90, worker STOPPED, tmux + +done-markers, sequential arms, single-load steady-state. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 611b9ff49..80480f7d5 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -303,7 +303,7 @@ built on it rather than keeping the flattering one. | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | | MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) but frame is a non-scene patch-grid | Render coherence ROOT-CAUSED (#70): VAE fine, DiT latent spatially white. fp4 speed CLOSED. Detail: benchmark-record + spec §8 | -| MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** banks +951us marlin. `FUSED-GLUE` W0: glue-into-marlin REFUTED (vLLM doesn't fuse into extern marlin); `FLASH-AUDIT` #68: c8 flash occupancy/L2-bound; `-use_fast_math` regresses, rejected | +| MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us marlin). glue-into-marlin + `-use_fast_math` both REFUTED (#68). `FLASH-OCCUPANCY` spike: c8 decode-flash grid ours==vLLM 1×3×64, ours **216<241 reg** ⇒ codegen/L2, not occupancy (owed ncu diff) | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx | | Ampere consumer (`sm_86`, RTX 3090 class) | **No number owed; no such board here.** 2026-08-06 build-verify: 7/7 FA2 TUs 0-warn, real `sm_86` SASS. [Detail](../.agents/benchmark-record.md) | External RTX 3090 report. Floor is llama.cpp on that card (GGUF, not our Blackwell-only NVFP4 grid) | diff --git a/docs/STATUS.md b/docs/STATUS.md index 76ef187f5..ab6223cab 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -58,7 +58,7 @@ token-for-token correctness against the pinned oracle. |---|---|---| | Qwen3.6-27B (NVFP4) text generation | Correctness-complete, at/above vLLM speed | Token-exact greedy on GB10; beats vLLM 0.25.0 total throughput at every concurrency (1.007-1.045x), effective parity 115/124 axes | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; 3-rep grid 0.93-1.03x. Async batch-1 token-0 degeneration FIXED: `VT_ASYNC_DEVICE_MIRROR` default ON | Token-exact SYNC+ASYNC (RED→GREEN); c16 0.93x; `VT_ASYNC_EXECUTOR` Option A (H2D out of capture) GREEN+RED but A/B NEUTRAL → OFF; c16 residual is prefill glue | -| Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA group-swap FLIPPED DEFAULT-ON — binding c2-c8 →0.92-0.94x, still <1.0x. state `KERNEL-FA2-GQA-SWAP-FLIP`. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | +| Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-OCCUPANCY` spike: c8 flash grid ours==vLLM, 216<241 reg ⇒ codegen/L2. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | | Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete, speed-pending | Revalidated after merging current upstream: local throughput is unchanged at 0.99997x its prior run; against the freshly measured pinned oracle it is 0.9971x. TTFT 0.7719x and host PSS 0.3127x pass; TPOT/ITL 1.1244x and VRAM 1.0014x remain open. Direct ON/OFF outputs remain 128/128 identical | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 01587c75b..7c175f8f4 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -321,7 +321,7 @@ def features_errors(text: str) -> list[str]: # Lowering these numbers as the page is compacted is the gate closing. STATUS = ROOT / "docs/STATUS.md" STATUS_RATCHET = { - "chars": 284081, + "chars": 284078, "h2_sections": 11, "long_paragraphs": 89, "oversized_cells": 47, From 5ee04b1ab370afcae6c4dc1a21af123522e67c3c Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 6 Aug 2026 19:49:34 +0000 Subject: [PATCH 2/2] =?UTF-8?q?row(quant):=20MXFP4-FLASH-OCCUPANCY=20CLOSE?= =?UTF-8?q?D=20=E2=80=94=20matched-c8=20ncu,=20occupancy=20IDENTICAL,=20ga?= =?UTF-8?q?p=20is=20an=20irreducible-for-us=20ptxas=20SASS-quality=20diffe?= =?UTF-8?q?rence;=20no=20lever=20(#75)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Runs the #69-owed ours-vs-vLLM flash decode ncu diff to a MEASURED verdict on an idle box, and CORRECTS #69's "occupancy/L2-bound" framing. W1 (matched-c8 ncu, both engines, grid 1x3x64, full section set): - vLLM runs at the SAME 8.33% occupancy as ours — both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB kSmemSize (the 216-vs-241 reg delta is moot). - L2 hit ~1% on BOTH (KV streamed from DRAM, no adjacency/warmth to win), identical short-scoreboard stall structure. Only measured diff: +13% instr. - The prior #69 "8.3% occupancy" was the batch-1 short-context num_splits=1 kernel; the prior ncu pair was context-mismatched. Now measured on the real c8 kernel with everything else held identical. W2 (mirror-first levers, all refuted DECISIVELY by building vLLM's exact recipe): - arch-mirror (compute_80 PTX driver-JIT): neutral (no change from native). - compute_80 + -use_fast_math (vLLM's EXACT recipe): reproduces vLLM's SASS profile EXACTLY (241 reg, 17,008 instr vs vLLM 241/17,020) and is STILL ~167us = the same as native ~165us and ~10us slower than vLLM's ~157us. So matching vLLM's arch, fast-math, register AND instruction count does NOT close the gap: the kernel is not instruction-bound; the residual is vLLM's wheel-ptxas SASS-scheduling quality, un-reachable from nvcc 13.0. (native+fast-math REGRESSES to 189.8us per #68; compute_80+fast-math is neutral.) W3 VERDICT: MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged 1.020/0.962/0.966/0.969). NO lever exists on our stack — occupancy, L2, num_splits, reg count, instruction count, __launch_bounds__, arch-mirror and fast-math ALL refuted by measurement. The flash term (+450us/step, ~40% of the c8 residual) is an irreducible-for-us ptxas quality gap. NO default flip owed; no functional code ships. CMakeLists NOTE + benchmark-record #75 record the closed levers so they are not re-tried. No byte-exact razor / SACRED battery owed (nothing shipped to gate). check-fusion-consistency (minimax_h3_video_vae_device) is a pre-existing red, untouched by this change. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode] --- .agents/NOW.md | 2 +- .agents/benchmark-record.md | 40 ++++++++++++++++++ .../qwen3-mxfp4-flash-occupancy-2026-08-06.md | 20 ++++++++- .agents/state.md | 42 +++++++++++++++++++ CMakeLists.txt | 30 +++++++++---- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- scripts/check-public-doc-tables.py | 2 +- 8 files changed, 126 insertions(+), 14 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 4b93fbf43..63c53cb73 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -22,7 +22,7 @@ checkpoint on `upstream/main` at `59674cf1d`. | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **e2e RUNS** (bf16-resident §13): 13/13·656. Token gate **NEAR-TIE 106/128** | device GDN/MLA islands; 1.59 tok/s; default OFF | | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) | | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` | -| MXFP4 parity | **c1 1.020, c2-c8 0.962-0.969, mem 2.63x**. `FLASH-AUDIT` #68: c8 flash +12.5us/call = occupancy/L2-bound; `-use_fast_math` TRIED→REGRESSES, rejected | real lever: flash occupancy/L2; no flip | +| MXFP4 parity | c1 1.020, c2-c8 0.962-0.969 (BELOW-FLOOR). `FLASH-OCCUPANCY` #75 CLOSED: matched-c8 ncu, occupancy IDENTICAL 8.33%; matched vLLM's exact reg+instr, STILL +10us = ptxas gap, no lever/flip (#75) | c8: flash ptxas ~40% + glue ~18% | | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): async mirror on classic dense Qwen3; SACRED 184/184 | Residual: sibling scope one-liner | | CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Op-dispatch profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn**. **G5 not next** | Parakeet encoder; attn dtype hoist | | Supported-models list (`row/DOCS-SUPPORTED-MODELS-MATRIX`) | **DRAFT PR**: FEATURES per-arch table CI-bound to registry (30 archs) | Reviewer merge | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 75e2539c8..4b4d24b61 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -19,6 +19,46 @@ from relative link targets repointed for this file's location. # Benchmarks +## QUANT-CT-MXFP4-FLASH-OCCUPANCY — the owed ours-vs-vLLM flash decode ncu diff: occupancy is IDENTICAL (8.33%, smem-limited); the flash gap is a COMPILER-CODEGEN difference (vLLM's wheel SASS runs 13% fewer instructions on the byte-identical kernel), NOT occupancy/L2/arch/fast-math — no lever exists on our stack (2026-08-06, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`, base `f7a1e322`, GB10 sm_121a, PR #75) + +#69 closed the compile lens (flash source byte-identical to vLLM's `2c839c33`; `-use_fast_math` REJECTED, +21us/call) and OWED the ours-vs-vLLM ncu diff (vLLM side lost to a box OOM-reboot). This row runs that diff to a MEASURED verdict on an idle box and re-frames the #69 premise. **Vehicle** `Yi30/Qwen3-8B-MXFP4` (dense W4A16 Marlin), oracle arm `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`. + +**W0 (free, existing c8 nsys CSVs) — the #69 "8.3% occupancy, register-limited" figure was the WRONG kernel.** The prior ncu pair was context-MISMATCHED: ours ran `vllm-cli` with a ~5-token prompt (context ~5 → `num_splits=1`, Split=false — a different kernel instance), vLLM ran `lens=[1024]` (Split=true). At the REAL c8 decode kernel (`gpu_trace_c8_dflt` vs `vllm_offline_trace_c8`) both run the IDENTICAL grid `1x3x64` (num_splits=3, gridZ=64=batch×kv_heads) and ours uses FEWER registers (216 vs 241) → more occupancy headroom, not less. + +**W1 — matched-c8 ncu diff (both engines, grid `1x3x64`, full section set; ours = `vllm-bench --input-len 1024 --output-len 128 -c 8`; vLLM = offline NPROMPTS=8, `enforce_eager` to dodge the FULL-cudagraph node-profiling driver-resource conflict; vLLM `gpu_memory_utilization` dropped 0.5→0.15 after full-section replay + the 0.5-util reservation OOM-rebooted the 119 GiB unified pool once — the SAME #69 failure mode, now root-caused: ncu kernel-replay on the vLLM process is memory-heavy):** + +| metric | OURS (native sm_121a) | vLLM (sm_80-PTX driver-JIT wheel) | +|---|---|---| +| grid (X×Y×Z) | **192 (1×3×64)** | **192 (1×3×64)** — identical | +| block / dyn smem per block | 128 / **81.92 KB** | 128 / **81.92 KB** — identical | +| Block Limit Shared Mem | **1** | **1** (both smem-limited to 1 CTA/SM) | +| **Achieved Occupancy** | **8.33%** | **8.33%** — IDENTICAL | +| Registers / thread | 216 | 241 | +| **Executed Instructions** | **3,695,180** (19,245/sched) | **3,267,840** (17,020/sched) — ours **+13.1%** | +| Duration (ncu, median) | ~165 us | ~157 us | +| L2 hit / L1 hit | 1.03% / 0.39% | 1.23% / 1.07% | +| SM / Memory throughput | 13.5% / 10.4% | 14.2% / 11.0% | +| Active warps/sched (eligible) | ~1.0 (0.06) | ~1.0 (0.06) | +| stall: short-scoreboard (LDS→MMA) | ~63% of ~16.6 cyc | ~58% of ~16.9 cyc — same KIND | + +**Answer to the owed W1 question: YES — vLLM's instance runs at the SAME 8.33% occupancy.** Both are SMEM-limited to 1 CTA/SM by the byte-identical 81.92 KB `kSmemSize` (the 216-vs-241 reg delta is MOOT — smem binds first). Occupancy is refuted as the lever. L2-warmth/adjacency is refuted: both stream KV from DRAM at ~1% L2 hit (c8 KV working set ≫ L2 → nothing a preceding kernel leaves stays hot), so the flash kernel's own memory profile is context-independent. num_splits/grid/GQA-pack are identical. The ONLY measured difference is **executed instructions: ours +13.1%** (the same +11.6% scalar bloat #69's cuobjdump saw on the non-split kernel, now on the actual c8 decode kernel with EVERYTHING ELSE held identical). That made instruction-count the leading CODEGEN hypothesis — but W2 tests it directly and **REFUTES it as the cause**: matching vLLM's exact instruction count does NOT close the gap. + +**W2 — mirror-first lever exploration (controlled ncu on the SAME c8 kernel; flash TUs recompiled per cell into an isolated build):** + +| flash-TU build variant | reg | executed instr | ncu duration | vs vLLM 3.27M/157us | +|---|---|---|---|---| +| A native sm_121a (main default) | 216 | 3.70M (19,245/sched) | ~165 us | baseline | +| B compute_80 PTX driver-JIT (mirror vLLM's arch) | 236 | 3.69M (19,212/sched) | ~165 us | **NEGATIVE — no change** | +| **C** compute_80 + `-use_fast_math` (vLLM's EXACT recipe) | **241** | **3,267,072 (17,008/sched)** | **~167 us (161-170)** | **matches vLLM reg+instr EXACTLY, STILL ~10us slower** | +| #68 native sm_121a + `-use_fast_math` | 255 | (fewer) | 189.8 us | REGRESSED (prior row) | +| vLLM `_vllm_fa2_C` (sm_80+PTX+fast-math wheel) | 241 | 3,267,840 (17,020/sched) | ~157 us | the target | + +**The decisive cell is C.** Arch-mirror alone (B) is NEUTRAL (compute_80 PTX driver-JIT = native instr/duration). Adding fast-math (C) reproduces vLLM's SASS profile EXACTLY — 241 reg and 17,008 instr vs vLLM's 241/17,020, a −13% instruction drop from A's 19,245 — and YET C is ~167 us, essentially the SAME as A/B's ~165 us and still ~10 us SLOWER than vLLM's ~157 us. So the −13% instruction reduction does NOT translate to speed (the kernel is NOT instruction-bound — our toolchain sits at ~165 us regardless of 19,245 or 17,008 instr), and matching vLLM's exact register+instruction profile still leaves the gap. The residual is the specific SASS instruction-SCHEDULING/selection quality of vLLM's wheel `ptxas`, applied to the byte-identical source — which our nvcc-13.0 `ptxas` does not reproduce even when told to emit the same arch, fast-math, register count and instruction count. (Sub-finding: native+fast-math REGRESSES to 189.8 us (#68) but compute_80+fast-math is neutral at ~167 us — the arch target changes how fast-math's SASS lands, but neither is a win.) Build variants were made by swapping `--generate-code` and adding `-use_fast_math` in the exact ninja compile command, then ar+link'ing `libvllm.a`/binaries manually (CMake has no per-source `CUDA_ARCHITECTURES` property — `set_source_files_properties(... CUDA_ARCHITECTURES ...)` silently no-ops; each variant verified via `cuobjdump` and an ncu regs/instr read). + +**W3 — THE PARITY VERDICT: MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged: c1 1.020 / c2 0.962 / c4 0.966 / c8 0.969 — nothing landed to move it).** Every candidate the diff could name is refuted by measurement: occupancy (identical 8.33%, smem-bound), L2-warmth (~1% both, KV ≫ L2), num_splits/grid (identical), register count (matched vLLM's 241 in C), instruction count (matched vLLM's 17k in C — and it did NOT help: the kernel is not instruction-bound), `__launch_bounds__`/reg-pressure (moot — smem-bound), arch-mirror (neutral), fast-math (regresses on native, neutral on compute_80). The flash decode term (+12.5 us/call ≈ +450 us/step, ~40% of the c8 residual after the dense-direct marlin banked the rest) is an IRREDUCIBLE (for us) **ptxas SASS-scheduling-quality gap**: on the byte-identical source, at identical grid/occupancy/smem/L2/stall-structure AND matched register+instruction counts, vLLM's wheel-`ptxas` schedules the SAME work ~10 us/call tighter than our nvcc-13.0 `ptxas`. Closing it requires vLLM's exact build `ptxas`, not something reachable by source, arch, flag, register or instruction control on our stack — so it is NOT a code, kernel, occupancy, or scheduling lever available to us. Honest c8 residual map: ~0.969 = flash ptxas-quality ~40% (this row, un-closable-by-us) + glue ~18% (portable-fusion, numerics-delicate, sub-parity) + residual-marlin/host the remainder. NO default flip owed (no throughput win). The CMakeLists NOTE records why -use_fast_math and the arch-mirror are NOT used, so they are not re-tried. + +Evidence: `dgx:~/mxfp4-nsys/{ours_flash_c8_ncu,vllm_flash_c8_ncu,buildB_flash_c8_ncu,buildC_flash_c8_ncu}.ncu-rep`; drivers `{ncu_flash_c8,ncu_vllm_eager2,perf_ab,buildc_ncu}.sh` + `vllm_offline_decode8_eager015.py`. Box OOM-rebooted once (util-0.5 vLLM ncu-replay), recovered, left clean. + ## QUANT-CT-MXFP4-FLASH-AUDIT — `--use_fast_math` on the FA2 TUs REJECTED (measured flash REGRESSION); the flash decode gap vs vLLM is a runtime memory/occupancy effect, not the SASS instruction count (2026-08-06, `row/QUANT-CT-MXFP4-FLASH-AUDIT`, base `origin/main` `4ce9fb74`, GB10 sm_121a, PR #69) The #67 refutation re-attributed the c2-c8 MXFP4 residual FLASH-dominant and OWED a diff --git a/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md b/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md index 1e120618c..acff2a33f 100644 --- a/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md +++ b/.agents/specs/qwen3-mxfp4-flash-occupancy-2026-08-06.md @@ -1,6 +1,6 @@ # QUANT-CT-MXFP4-FLASH-OCCUPANCY — the owed ours-vs-vLLM flash decode ncu diff + the occupancy/L2/codegen lever - + Row: `QUANT-CT-MXFP4-FLASH-OCCUPANCY` (helper, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`). Base: `origin/main` `f7a1e322`. Vehicle: `Yi30/Qwen3-8B-MXFP4` (dense `Qwen3ForCausalLM`, W4A16 Marlin keep-quant); oracle arm @@ -57,3 +57,21 @@ Byte-exact razor: #44 MXFP4-8B smoke 3/3 token-exact + coherent. If reduction order shifts: SACRED 0.6B/4B distributional + 32B strict, async, memcheck, eager + graphed. Box safety: BOTH flock locks, free -g ≥ 90, worker STOPPED, tmux + done-markers, sequential arms, single-load steady-state. + +## CLOSED — verdict (2026-08-06) + +W1 DONE (matched-c8 ncu, both engines, grid 1x3x64): **vLLM runs at the SAME +8.33% occupancy** (both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB +smem), L2 ~1% both, identical stall structure. The ONLY measured diff was +13% +executed instructions — the CODEGEN hypothesis. W2 tested every mirror-first +lever and REFUTED all of them: arch-mirror (compute_80 PTX driver-JIT) is +neutral; and building vLLM's EXACT recipe (compute_80 + `-use_fast_math`) +reproduces vLLM's SASS profile EXACTLY (241 reg, 17,008 instr vs vLLM 241/17,020) +and is STILL ~167 us, ~10 us slower than vLLM's ~157 us. So matching vLLM's arch, +fast-math, register AND instruction count does NOT close the gap: the kernel is +not instruction-bound, and the residual is vLLM's wheel-`ptxas` SASS-scheduling +quality, un-reachable from our nvcc 13.0. W3 VERDICT: MXFP4 stays BELOW-FLOOR +(binding unchanged 1.020/0.962/0.966/0.969), NO lever exists on our stack, NO +default flip owed. No functional code shipped; a CMakeLists NOTE + benchmark +record (#75) capture the closed levers so they are not re-tried. NO byte-exact +razor / SACRED battery owed (nothing shipped to gate). diff --git a/.agents/state.md b/.agents/state.md index 7c86e3a82..a9913269d 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -39537,3 +39537,45 @@ are committed on the row for the next session. fp4 speed path stays CLOSED. Box left clean (GPU idle, both locks free, worker down, ~39 G ckpt cached at `dgx:~/h3fp4/ckpt`, diagnostic latents at `dgx:~/h3fp4/diag`). Evidence: `dgx:~/h3fp4/{diag,rt_out,out_small,out_cond}`; PR #70. + +## 2026-08-06T21:45 - QUANT-CT-MXFP4-FLASH-OCCUPANCY: the owed matched-c8 flash ncu diff - occupancy is IDENTICAL (8.33%), the gap is an irreducible-for-us ptxas SASS-scheduling quality difference (matched vLLM's exact reg+instr, still +10us), NO lever on our stack + + + +Ran the #69-owed ours-vs-vLLM flash decode ncu diff to a MEASURED verdict on an +idle box (PR #75, `row/QUANT-CT-MXFP4-FLASH-OCCUPANCY`, base `f7a1e322`). + +**The #69 premise is REFRAMED (and its "occupancy/L2-bound" line CORRECTED).** The +prior ncu pair was context-MISMATCHED: ours ran a ~5-token cli prompt +(`num_splits=1`, Split=false), vLLM ran lens=1024 (Split=true) - different kernel +instances. At the REAL c8 decode kernel both engines run the IDENTICAL grid +`1x3x64`. The matched-c8 ncu (both engines, full section set; vLLM `enforce_eager` +to dodge a FULL-cudagraph node-profiling driver-resource conflict; vLLM util +0.5->0.15 after ncu-replay + the 0.5 reservation OOM-rebooted the 119 GiB pool +once - the same #69 failure mode, now root-caused) shows: **occupancy IDENTICAL +8.33%** (both smem-limited to 1 CTA/SM by the byte-identical 81.92 KB smem - the +216-vs-241 reg delta is moot), **L2 ~1% on both** (KV streamed from DRAM, no +adjacency/warmth to win), identical short-scoreboard stall structure. The only +measured diff: ours +13% executed instructions. + +**W2 refuted every mirror-first lever, DECISIVELY via the exact vLLM recipe.** +Arch-mirror (compute_80 PTX driver-JIT) is neutral. Building vLLM's EXACT recipe +(compute_80 + `-use_fast_math`) reproduces vLLM's SASS profile EXACTLY - 241 reg, +17,008 instr vs vLLM 241/17,020 - and is STILL ~167 us, the same as our native +~165 us and ~10 us slower than vLLM's ~157 us. So the -13% instruction reduction +does NOT translate to speed (the kernel is not instruction-bound) and matching +vLLM's arch/fast-math/reg/instr does not close the gap. The residual is vLLM's +wheel-`ptxas` SASS-scheduling quality, un-reachable from nvcc 13.0. (native+fast- +math REGRESSES to 189.8 us per #68; compute_80+fast-math is neutral - the arch +changes how fast-math lands but neither wins.) CMake has no per-source +`CUDA_ARCHITECTURES` property, so build variants were made by swapping +`--generate-code` in the exact ninja command + manual ar/link, cuobjdump-verified. + +**Verdict.** MXFP4 stays BELOW-FLOOR at c2-c8 (binding unchanged +1.020/0.962/0.966/0.969). NO lever exists on our stack: occupancy, L2, num_splits, +reg count, instruction count, `__launch_bounds__`, arch-mirror and fast-math are +ALL refuted by measurement. The flash term (+450 us/step, ~40% of the c8 residual) +is an irreducible-for-us ptxas quality gap. NO default flip owed; no functional +code shipped (CMakeLists NOTE + benchmark-record #75 record the closed levers). +Box left clean (GPU idle, both locks free, worker down). Evidence: +`dgx:~/mxfp4-nsys/{ours,vllm,buildB,buildC}_flash_c8_ncu.ncu-rep`; PR #75. diff --git a/CMakeLists.txt b/CMakeLists.txt index c202820eb..fec00979e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1412,15 +1412,27 @@ if(VLLM_CPP_CUDA) src/vt/cuda/cuda_flash_attn_fa2.cu ${_FA2_KERNEL_SRCS}) target_compile_definitions(vllm PUBLIC VLLM_CPP_FLASH_ATTN) - # NOTE (QUANT-CT-MXFP4-FLASH-AUDIT #68): -use_fast_math was TRIED here to mirror - # vLLM's own flash-attn build (its _vllm_fa2_C is fast-math; it makes our decode - # kernel's SASS match vLLM's 4832/REG255 vs 4880/REG255). But a controlled - # same-build nsys A/B REFUTED it as a speed lever: this decode kernel is - # memory-latency-bound (ncu occ ~10.7%), so the +11.6% instruction reduction does - # not help, while fast-math's higher register count (246->255) LOWERS occupancy - # and REGRESSED flash 168.8 -> 189.8 us/call. So we do NOT pass -use_fast_math. - # The residual flash gap to vLLM (168.8 vs 156.3 us/call) is a runtime memory/L2 - # effect, not the SASS; see .agents/benchmark-record.md. + # NOTE (QUANT-CT-MXFP4-FLASH-AUDIT #68 + FLASH-OCCUPANCY #75): the c8 decode + # flash gap vs vLLM (ours ~168-174 vs vLLM ~156 us/call) is a COMPILER-CODEGEN + # gap, MEASURED to a verdict — NOT occupancy, NOT L2, NOT num_splits, and NOT a + # flag we can set. The owed matched-c8 ncu (both engines, IDENTICAL grid + # 1x3x64) shows occupancy is IDENTICAL 8.33% (both smem-limited to 1 CTA/SM by + # the byte-identical 81.92 KB kSmemSize — the 216-vs-241 reg delta is moot), + # L2 hit ~1% on BOTH (KV streamed from DRAM: no adjacency/warmth to win), and + # the same ~60% short-scoreboard stall structure. The ONLY difference: vLLM's + # SASS executes ~13% FEWER instructions (3.27M vs 3.70M) in this latency-bound + # ~1-warp/scheduler regime — but that instruction count is NOT the cause. #75 + # built the exact vLLM recipe: compiling the kernel TUs for compute_80 PTX + + # -use_fast_math REPRODUCES vLLM's SASS profile EXACTLY (241 reg, 17,008 instr + # vs vLLM 241/17,020) and is STILL ~167 us — the SAME as our native ~165 us and + # still ~10 us slower than vLLM's ~157 us. So matching vLLM's arch, fast-math, + # register AND instruction counts does NOT close the gap: the kernel is not + # instruction-bound, and the residual is vLLM's wheel-`ptxas` SASS-SCHEDULING + # quality, not reachable from our nvcc 13.0. (native+fast-math REGRESSES to + # 189.8 us, #68; compute_80 PTX driver-JIT alone is neutral, #75 Build B.) So + # we keep the plain native sm_121a compile and pass no -use_fast_math. The + # residual is an irreducible-for-us ptxas SASS-quality gap on a byte-identical + # kernel; see .agents/benchmark-record.md (#75). set_source_files_properties( src/vt/cuda/cuda_flash_attn_fa2.cu ${_FA2_KERNEL_SRCS} PROPERTIES diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 80480f7d5..39159e067 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -303,7 +303,7 @@ built on it rather than keeping the flattering one. | Kimi-Linear-48B-A3B (KDA+MLA+MoE) | Full-model GB10 e2e RUNS (bf16-resident §13), NEAR-TIE 106/128, pool math CLOSES; default OFF | Full model RUNS on GB10 (bf16-resident, RSS peak 1.7 GiB, min-avail 21 GiB, no OOM). Token NEAR-TIE 106/128 (6/8 prompts exact, numerics vs deterministic oracle). 1.59 tok/s. Detail: spec §13 | | vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin | | MiniMax-H3 FP4 speed (W-FP4a) | **Measured GB10 (`row/H3-FP4-GPU-E2E`).** Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) but frame is a non-scene patch-grid | Render coherence ROOT-CAUSED (#70): VAE fine, DiT latent spatially white. fp4 speed CLOSED. Detail: benchmark-record + spec §8 | -| MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us marlin). glue-into-marlin + `-use_fast_math` both REFUTED (#68). `FLASH-OCCUPANCY` spike: c8 decode-flash grid ours==vLLM 1×3×64, ours **216<241 reg** ⇒ codegen/L2, not occupancy (owed ncu diff) | +| MXFP4 Qwen3-8B (W4A16 Marlin) | **`KERNEL-MARLIN-DENSE-EXEC` x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less** (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | **VT_MARLIN_DENSE default-ON** (+951us). `FLASH-OCCUPANCY` #75: matched-c8 ncu, occupancy IDENTICAL 8.33%; built vLLM's exact flash recipe, matched reg+instr, STILL +10us, gap is ptxas SASS quality, no lever/flip | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | | cuBLAS invocation-parity guard | CI guard landed (CPU); `kGemvHeuristicAlgos` refactor build-verify owed | `nvcc` rebuild + SACRED gate on dgx | | Ampere consumer (`sm_86`, RTX 3090 class) | **No number owed; no such board here.** 2026-08-06 build-verify: 7/7 FA2 TUs 0-warn, real `sm_86` SASS. [Detail](../.agents/benchmark-record.md) | External RTX 3090 report. Floor is llama.cpp on that card (GGUF, not our Blackwell-only NVFP4 grid) | diff --git a/docs/STATUS.md b/docs/STATUS.md index ab6223cab..b3f670d85 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -58,7 +58,7 @@ token-for-token correctness against the pinned oracle. |---|---|---| | Qwen3.6-27B (NVFP4) text generation | Correctness-complete, at/above vLLM speed | Token-exact greedy on GB10; beats vLLM 0.25.0 total throughput at every concurrency (1.007-1.045x), effective parity 115/124 axes | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; 3-rep grid 0.93-1.03x. Async batch-1 token-0 degeneration FIXED: `VT_ASYNC_DEVICE_MIRROR` default ON | Token-exact SYNC+ASYNC (RED→GREEN); c16 0.93x; `VT_ASYNC_EXECUTOR` Option A (H2D out of capture) GREEN+RED but A/B NEUTRAL → OFF; c16 residual is prefill glue | -| Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-OCCUPANCY` spike: c8 flash grid ours==vLLM, 216<241 reg ⇒ codegen/L2. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | +| Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-OCCUPANCY` #75: occ IDENTICAL, gap=ptxas SASS quality, no lever. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | | Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete, speed-pending | Revalidated after merging current upstream: local throughput is unchanged at 0.99997x its prior run; against the freshly measured pinned oracle it is 0.9971x. TTFT 0.7719x and host PSS 0.3127x pass; TPOT/ITL 1.1244x and VRAM 1.0014x remain open. Direct ON/OFF outputs remain 128/128 identical | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 7c175f8f4..0ea2aa1a0 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -321,7 +321,7 @@ def features_errors(text: str) -> list[str]: # Lowering these numbers as the page is compacted is the gate closing. STATUS = ROOT / "docs/STATUS.md" STATUS_RATCHET = { - "chars": 284078, + "chars": 284073, "h2_sections": 11, "long_paragraphs": 89, "oversized_cells": 47,