gqh: MMQ path, wide verify arms, and correctness gates for the wide and I8 arms - #35
Draft
DeanoC wants to merge 30 commits into
Draft
gqh: MMQ path, wide verify arms, and correctness gates for the wide and I8 arms#35DeanoC wants to merge 30 commits into
DeanoC wants to merge 30 commits into
Conversation
Widening the DFlash2 draft block past its trained 8 doubled acceptance on our artifact (avg_commit 7.22 to 10.98) but collapsed throughput 154 to 56 tok/s, because only 2..12 had exact-width arms and 13..16 fell to the generic runtime-guarded instantiation. Every width a verify actually dispatches now has one: GQH_MULTICOL_SPEC_MAX rises 12 to 16 (== GQH_MAX_COLS), with a separate wide arm from 13 (GQH_MULTICOL_WIDE_MIN) at GQH_WIDE_ROWS 2. The generic instantiation is kept only as the GGML_GQH_MULTICOL=0 A/B control. Measured on an R9700 (gfx1201), 10-prompt HumanEval under spec decode, block 16: 56.15 to 182.07 tok/s, a 3.23x speedup on the arm, reproduced three times at a 0.3% spread. The three control arms are unmoved -- GQH block 8 154.03, IQ4_XS block 16 160.54, IQ4_XS block 8 127.19 -- so the change is isolated to the widths it targets. That puts this artifact 13.4% ahead of the IQ4_XS arm it previously trailed, at 13.7% fewer bytes. Two premises this disproves, both of which had been used to argue the cliff was structural. It is not register pressure: occupancy was driven from 16 waves/SIMD to 3 with no movement. It is not instruction issue either: a 14% instruction cut bought 2.5%. The arm is memory-level-parallelism bound, which is why giving the wide widths their own exact-width shape is what moved it. Correctness: ctest -R gqh passes including the nvec 13 and 16 cases, and those cases drive x with one-hot basis vectors, so they exercise a single lane. A dense-activation probe was added alongside because that gate alone would not catch a lane-map error.
Section 8 concluded we were 2-5% behind upstream's best because GQH could not widen past ncols 8. The kernel change in this branch removes that constraint, so that conclusion is marked superseded rather than left to be quoted: GQH block 16 goes 56.40 to 182.07 tok/s and now leads IQ4_XS block 16 (160.54) by 13.4% at 13.7% fewer bytes, with all three control arms unmoved. Also records what section 8 got wrong. It blamed register pressure, then instruction issue; occupancy 16 to 3 waves/SIMD moved nothing and a 14% instruction cut bought 2.5%, so the arm is memory-level-parallelism bound. Both wrong diagnoses are kept visible because each one pointed at a fix that would not have worked. Finally, the correctness gate is documented as still incomplete: the nvec 13/16 ctest cases added earlier drive x with one-hot basis vectors and so exercise a single lane, blind past activation position 15 -- which is exactly the failure a wide lane map invites.
The wide arm from the previous commit used one ROWS for every width 13..16. Width 16 is the only one a DFlash2 verify dispatches at the tuned block size, and it wants a different shape from its neighbours: ROWS 2 -> 4. gqh_wide_rows() becomes the per-width table its NCOLS template parameter was always there for, and the kernel takes ROWS as a template argument instead of reading the macro, so every width that keeps ROWS 2 gets byte-identical codegen. Two static_asserts pin what the shape relies on: ROWS16 must be a power of two, and ROWS16 * GQH_MAX_COLS must stay a multiple of the warp so the widest arm keeps the reduce-scatter epilogue. R9700 (gfx1201), one build, one harness, 10 HumanEval prompts under spec decode, decode-only from the engine's own [spec-decode] timer: | target | block | tok/s | avg_commit | was | |--------|------:|------:|-----------:|----:| | GQH | 16 | 195.66| 10.9790 | 182.07 | | GQH | 8 | 154.25| 7.2250 | 154.03 | | IQ4_XS | 16 | 160.21| 12.1800 | 160.54 | | IQ4_XS | 8 | 126.80| 7.5740 | 127.19 | +7.5% on the targeted arm with all three controls inside 0.31%, and avg_commit identical to four decimals on every arm -- so this is per-step cost only, not an acceptance effect. ctest -R gqh passes on all four. Also lands GQH_WIRE_DEPTH2W as an opt-in A/B knob, default 0 and if constexpr-dead at ROWS16 4, for the 2-deep wide wire. Its ROWS16=2 arm read 180.54/180.31 against a scored 194.817 and was statistically identical to the same build's ROWS16=4 arm (194.60-195.71), so the knob stays off rather than shipping a shape that does not pay for itself.
DeanoC
changed the title
…ate can fail
The fused check drove x with one-hot basis vectors at position j < nvec. With
cols >= 256 that never places a nonzero activation past position 15, so the
k-reduction was only ever exercised in its first superblock and an activation
indexing error anywhere above that was invisible -- precisely the class of bug
the wide verify arms introduce.
Two passes are added, and neither weakens the existing bitwise claim:
A. offset-swept one-hot. Column j is e_(off+j) for offsets spanning the
reduction, so the dot product still has a SINGLE term and stays bit-exact,
but the nonzero now lands anywhere in k.
B. dense activations. Every k nonzero, so term count and accumulation order
are exercised -- which no single-term probe can see. Summation makes this
one tolerance-based (1e-4 relative); the observed margin is 8.04e-06 at
64x1024 and 1.3e-06 at 4x512, i.e. four orders of magnitude of headroom.
NEGATIVE CONTROL. A gate that has never failed proves nothing, so the new
checks were run against a deliberately broken kernel: gqh_load_x addressing
superblocks 2 and 3 swapped, chosen because it leaves superblock 0 -- where
every old one-hot nonzero lived -- untouched. On gqh3 64x1024 nvec 16:
FAIL gqh3 64x1024: dequant 0/65536 differ, fused 0/1024 differ,
offset-onehot 2044 differ over 3 offsets,
dense 1024 differ (max rel 55.6)
Both pre-existing checks report ZERO while both new checks catch it. The same
build passes 4x512, which has no superblock 2, so the coverage is specific and
not a blanket tightening. Kernel reverted; ctest -R gqh is 50/50 at 20.4 s.
A rocprofv3 kernel trace confirms these cases reach
gqh_matvec_kernel<gqh3, NCOLS=16, ROWS=4>, i.e. the wide arm the per-width ROWS
table introduced, so the new coverage lands on the shape that needed it.
KNOWN GAP, not addressed here: the trace shows no gqh_quant_x_* kernel in any
of these cases and the f32 x-load signature throughout, so the I8 activation
path -- what serving actually runs, and what the HE throughput numbers are
scored on -- still has NO ctest coverage. GGML_GQH_I8DOT=1 does not change the
dispatch in this harness: I8DOT=0 and I8DOT=1 both report an identical dense
max rel of 8.04e-06. Its only current gate is the AL >= 7.2 aggregate on HE.
… all
GGML_GQH_I8DOT defaults ON, so the int8 activation arm is what serving runs and
what every throughput number is scored on. Nothing verified it. A kernel trace
of the whole ctest suite shows only the f32 arm and no gqh_quant_x_* launch.
The env var was never the reason. gqh_i8_shape_ok() floors the arm at
gqh_i8_min_work() = 16 Mi of in*rows_total; the largest case in the suite is
64*1024 = 64 Ki, which is 256x under it. So the arm was unreachable from ctest
at ANY setting of GGML_GQH_I8DOT -- setting it to 1 changed nothing, and the f32
and I8 runs returned an identical dense max rel of 8.04e-06. The knob to lift
the floor already existed for exactly this: GGML_GQH_I8_MINWORK=0, which is
host-side only and moves no dispatch and no kernel.
The arm is not bit-exact -- both sides carry one symmetric int8 round, and the
two 1/127s fold into GQH_Q8_XSCALE -- so a new "i8" argument selects int8-sized
tolerances instead of memcmp:
* single-term (one-hot and offset-swept one-hot): absolute bound
max|w_row|/127, the weight LUT step, since the grid amax is exactly 1.0.
* dense: normalised by ||w_row|| * ||x_col||, bar 2e-2. The old max(1, |ref|)
denominator turns a quantisation-sized ABSOLUTE error on a near-zero
accumulator into a meaningless relative number, and it grew as sqrt(in):
0.039 at in=256, 0.085 at 512, 0.32 at 1024. Normalising by the dot's own
scale removes that -- measured 5.4e-4 / 6.6e-4 / 6.5e-4 across the same
three widths, i.e. shape-independent, with ~20x headroom.
36 new cases: rungs gqh3/gqh2_h/gqh4 x three shapes x nvec 8/13/16, plus a
GGML_GQH_I8_XSCALE=0 variant per rung/shape which also disables SB2 and is the
only coverage of the non-SB2 wide kernel. gqh2_c is excluded deliberately: it
returns before any I8 dispatch and has no I8 arm. 50 -> 86 tests, 20s -> 40s.
NEGATIVE CONTROL. Built with gqh_quant_x_wide16_kernel's activation group index
shifted -- a bug confined to the I8 wide-16 path, chosen so it never reads
group 0 -- in an isolated build directory so no shipped binary was affected:
I8 nvec 16: FAIL offset-onehot 2942 differ, dense 634 differ (max rel 0.123)
f32 nvec 16: OK bit-identical, dense max rel 8.04e-06
I8 nvec 8: OK dense max rel 0.000592
Caught on exactly the arm carrying the bug, clean on the f32 arm and on the
narrow I8 arm that uses a different quantiser: 0.123 against a 2e-2 bar and a
6.5e-4 clean margin. Note "fused 0/1024 differ" in the failing line -- the
pre-existing one-hot check reports ZERO with the bug present, because its
nonzeros all sit in group 0, so only the offset-swept pass sees it.
DeanoC
marked this pull request as draft
August 25, 2026 07:39
DeanoC
changed the title
…ong shape
Stage 1 of giving GQH a quantised matmul path so prefill stops dequantising the
weight set to fp16. GQH3 only, opt-in behind GGML_GQH_MMQ=1 and default OFF --
the dequant reference has to stay reachable on the same binary to be a control,
and the numbers below now independently justify keeping it that way.
Built on the registry-backed precedent (Q3_1_ROCMFP3_MIX), so the per-tensor
5-byte header reaches the kernel by value in the grid struct rather than through
tensor data. Uses the _16 vec_dot family, whose one-x-scale-per-16 against
one-y-scale-per-32 already matches GQH's ratio granularity, so no new activation
quantiser was needed.
THE RESULT IS NEGATIVE, AND IT IS THE POINT OF THE COMMIT. Target only,
--prefix-cache-slots 0, prefill from the server timer, arms interleaved 1/0/1
because the rig drifts upward:
prompt tokens dequant GQH3-MMQ delta
353 0.5 s 0.5 s under the 0.1 s timer resolution
2,033 2.4 s 2.5 s +4.2%
8,703 10.9 s 11.7 s +6.2%
Both arms reproduced to 0.1 s across the interleave, so the sign is not drift.
MMQ re-decodes the weight tile once per output column tile -- about 16 times at
2K and 68 at 8.7K -- where dequant unpacks once and cuBLAS streams fp16. GQH3
cannot widen out of it: it is already at 128x128 and mmq_x stops at 128. That
mechanism is a hypothesis; its prediction is that the regression grows with N,
and that is what the table shows. docs/QWEN38_R9700_REPRO.md section 9 predicted
the opposite and called this the highest-value remaining work, so it is
corrected here.
Two caveats that stop this being a refutation: GQH3 is only 139 of the artifact's
396 header-bearing tensors (the other 257 are GQH4 and still dequantise), and
this loader is deliberately naive. Neither explains a regression that grows
monotonically with N.
Gate: ctest -R gqh goes 86 -> 95, the 9 new cases covering the wider-than-16
widths that only MMQ can serve. NEGATIVE CONTROL: injecting w ^ 1 on the tile's
code-word index -- a pure index mis-map that leaves every magnitude and scale
intact -- fails 9 and passes the pre-existing 86, built in a separate directory
so no shipped binary ever held it, then restored md5-identical and re-verified.
A second control asserts MMQ actually dispatched (>= 2 launches): without it the
new cases pass happily on the dequant fallback, because the MMQ tolerances are
looser than the fallback's own error.
Grid codes were decoded from the artifact's own geoquant.gqh.headers rather than
assumed: GQH3 {3: 127, 4: 12}, GQH4 {2: 54, 3: 200, 4: 3}, no GQH2_H or GQH2_C.
Worst dynamic range 36.30:1, so the int8 range guard rejects nothing this
artifact contains -- it is currently unreachable code that goes live with GQH4.
NOT VERIFIED: the dp4a arm. On gfx1201 the WMMA branch executes, dp4a is dead
code, and should_use_mmq is gated to RDNA3.5/RDNA4 which both have WMMA -- so the
dp4a tile indices here are unreachable on every arch this enables and nothing on
this hardware tests them. A wrong dp4a tile size would pass the whole suite.
…y wins
Completes the MMQ rungs the shipping artifact uses (GQH3 + GQH4 = all 396
header-bearing tensors; it contains no GQH2_H or GQH2_C). Still opt-in behind
GGML_GQH_MMQ=1, default off. ctest -R gqh 95 -> 112.
PREFILL: MMQ IS THE WRONG SHAPE, now confirmed twice over. Same protocol, arms
interleaved 1/0/1, each reproducing to 0.1 s:
prompt tokens dequant MMQ GQH3 only MMQ both rungs
353 0.5 s 0.5 s 0.5 s
2,033 2.4 s 2.5 s (+4.2%) 2.6-2.7 s (+8.3..12.5%)
8,703 10.9 s 11.7 s (+6.2%) 12.2-12.5 s (+11.9..13.8%)
The regression grows with prompt length AND roughly doubled when MMQ's share of
the artifact went from 35% of tensors to 100%. Both are what per-output-column-
tile re-decode predicts: MMQ re-decodes the weight tile once per column tile
(~16x at 2K, ~68x at 8.7K) where dequant unpacks once and cuBLAS streams fp16.
BUT THE CODE WORKLOAD SAYS THE OPPOSITE, and that is the finding. Two readings
per arm, IQ4_XS as a machine control since GGML_GQH_MMQ cannot touch it:
arm metric off on delta
GQH b16 he_tok_s 194.02/195.98 191.25/191.25 -1.4..-2.4%
GQH b16 code_e2e 115.90/116.85 140.14/140.64 +20.6%
GQH b16 code_decode 138.06/139.33 152.74/153.42 +10.5%
GQH b8 code_e2e 112.26/112.42 131.27/130.53 +16.5%
IQ4_XS b16 he_tok_s 159.97/160.36 159.29/160.09 -0.3%
IQ4_XS b16 code_e2e 119.26/119.63 119.27/118.93 -0.3%
Control flat to 0.3% on every metric and both repeats, avg_commit identical to
four decimals, so this is not an acceptance effect. On the code workload GQH goes
from LOSING to IQ4_XS (116.85 vs 119.63) to BEATING it by 17.9% (140.64 vs
119.27) -- the end-to-end deficit that qualified every decode-only claim we have
made. code_decode excludes prefill and still gains 10.5%, so part of the win is
inside the decode loop, most plausibly the verify widths past GQH_MAX_COLS that
used to fall to the pathological dequant->GEMM. NOT INSTRUMENTED: recorded as a
hypothesis, not a claim.
So the shape is width-gated dispatch -- MMQ for the wide-verify band just past
GQH_MAX_COLS, dequant->cuBLAS kept for true prefill widths. The mix qtypes in
this file already do exactly that via mix_mmq_max_ne11. Finding the crossover is
the next measurement; until it exists, opt-in-off is correct.
DECODE MOVED, and it is a finding: he_tok_s on GQH b16 fell ~2% (191.25 against
194.02-195.98, whose own spread is 1.0% and the control's 0.3%). Small but real,
at or just above the ~1.3% floor. Nothing in the MMQ path touches a width the
matvec owns; the suspect is that should_use_mmq returning true for GQH also feeds
the graph-capture predicate and supports_op, so enabling it can change capture or
fusion on subgraphs whose arithmetic never changes. UNVERIFIED.
THE GUARD, and why building its test changed what it is justified by. Defeating
the int8 range guard did NOT breach the one-weight-step bound: grids 8 through 11
all passed it. int8 destroys the INNERMOST levels, so their absolute error stays
tiny -- grid 11's innermost level is 0.00195 of the superblock scale, a quarter
of one step. What dies is relative resolution, which no absolute bound can see.
Counting the code histogram of gqh4_64x1024: codes 7 and 8, the two levels grid
11 rounds to exactly zero, are 17.1% of all weights. That is the silent
corruption the numeric bar was blind to. The harness now also asserts no
non-zero reference weight returns exactly zero; it cannot false-positive on an
accepted grid, since within 127:1 every level quantises to |q| >= 1. Defeated,
it fails grids 10 and 11 while 8 and 9 pass -- so refusing 10/11 prevents hard
corruption and refusing 8/9 is a precaution about relative resolution. Guard
cases straddle the threshold: gqh4 grid 7 (112.84:1) must accept, grid 8
(164.69:1) must refuse, grid 11 (512:1) collapses, and gqh3 grid 11 (64:1) must
accept so the guard is shown not to over-refuse.
A HARNESS DEFECT this measurement exposed: ctest's ENVIRONMENT property ADDS to
the ambient environment rather than replacing it, so running the benchmark with
GGML_GQH_MMQ=1 silently moved 42 pre-existing cases onto MMQ while they still
demanded fp16-bitwise dequant output, and gqh_ctest_ok flipped 1 -> 0 on all four
rows. gqh_wide_measure.sh runs ctest -R gqh as the search loop's correctness gate
INSIDE the same environment as the benchmark it scores, so the loop would have
seen a fault that does not exist -- worse than seeing nothing, because it cannot
distinguish it from a real one. Fixed by pinning GGML_GQH_MMQ=0 on every case
that asserts the dequant result; verified 112/112 with the variable unset, =0
and =1.
NEGATIVE CONTROLS, both in a separate build directory, restored md5-identical:
GQH4 tile index mis-map (w ^ 1) fails exactly the 11 GQH4 MMQ cases while all 9
GQH3 MMQ cases, the 4 refuse cases and the other 88 pass. Guard defeated fails
exactly the 4 refuse cases. Guard defeated with the collapse check fails grids 10
and 11 (13007 non-zero weights collapsed to zero) and passes grid 9.
STILL UNVERIFIED: the dp4a arm in both loaders. On gfx1201 the WMMA branch
executes, dp4a is dead code, and the arch gate admits only RDNA3.5/RDNA4 which
both have WMMA -- so those tile indices are unreachable on every arch this
enables and nothing runnable here tests them.
MMQ and the dequant path pull in opposite directions for GQH: MMQ deletes the fp16 materialisation but re-decodes the weight tile once per output column tile, and the GQH unpack is expensive. So it wins narrow and loses wide, and a single all-or-nothing switch cannot serve both. gqh_mmq_max_ne11() makes dispatch width-bounded, defaulting to 160 and overridable by GGML_GQH_MMQ_MAX_NE11, mirroring mix_mmq_max_ne11. Widths 1..16 never reach it -- the matvec owns them and returns first -- so the gate is purely an upper bound. gqh-mmq-sweep is the standalone tool the threshold comes from: one mul_mat per width, arms selected by GGML_GQH_MMQ, each width reporting whether MMQ actually dispatched so a sweep cannot silently measure the same path twice. The harness no longer treats "did MMQ run" as a property of the run. A width gate can send the two wide nodes down different paths, so each numeric expectation now follows the launch count that actually happened rather than assuming one. 24 new cases pin both sides of the gate (in160/out161 for the compiled default, knob-in64/knob-out65 for the override) on every shape wide enough to straddle the bound. A gate that never declines is indistinguishable from no gate. ctest -R gqh: 112 -> 136.
… wins The width gate was justified by the kernel crossover curve alone. That is half a justification: it says where MMQ stops paying, not whether the workload ever goes there. Nobody had measured the second leg. GGML_GQH_NE11_LOG=1 turns on a census of the ne11 every GQH mul_mat dispatches, split by whether it got past ggml_cuda_gqh_mul_mat_vec -- the only calls the gate can affect, since the matvec owns 1..GQH_MAX_COLS and returns before the gate is read. It is off by default and dumps at exit. Over a DFlash2 verify at --draft-block-size 16 plus a 7.3k-token prefill the distribution is bimodal and has a hole in it: 40280 calls at ne11=16, none of which reach the gate; 4323 between 39 and 104; 5502 at exactly 512, the prefill chunk. Nothing at all in 105..511. So every threshold in that window dispatches identically and the exact value is not load-bearing -- which is worth knowing before anyone argues about 160 versus 256 again. What IS load-bearing is that the gate declines 512, where the curve says MMQ loses ~1.2x: lifting the bound costs about 7% of prefill throughput end to end. The crossover re-measured on a second R9700 box keeps its shape but sits further out, so 160 is conservative there rather than wrong. It stays, being the widest bound that is a win on both boxes and comfortably inside the window the workload leaves empty.
The harness prints its result as "pass@1 (greedy, raw)" but never put a temperature in the request, so the server fell back to the model card sampling defaults and the whole suite has been sampling. Measured on the GQH-shaped Qwen3.8-27B target, two runs of the SAME config agreed on only 27 of 164 replies (16.5%), and 30 items flipped verdict between them. The totals happened to land on 124 both times, but only because 15 items went pass->fail and 15 went fail->pass; that agreement was a coincidence of a churning set, not stability. The floor is wider than any effect this instrument gets pointed at, which is why an n=1 pass@1 from it carries much less information than it appears to. Pinning temperature=0 and top_k=1 makes it deterministic: same-config repeats now agree on 164/164 replies and 0 items flip. pass@1 on that target also rises from 0.756 to 0.884, greedy suiting code completion -- so the accidental sampling was costing real measured quality as well as reproducibility.
Dispatch has been width-gated since the gate landed, so MMQ now runs only where it measured faster and the dequant path keeps everything else. Ungated it costs ~7% of prefill; gated it does not. On-vs-off on an R9700, two readings per arm and order-balanced, because a run-order thermal effect worth ~1% otherwise reads as a result: prefill +1.0%, decode +1.1%. The IQ4_XS control is flat to 0.3% and provably inert rather than merely unaffected -- its ne11 census records no GQH mul_mat at all, so the flag cannot reach it. Enabling this CHANGES GENERATED TOKENS. MMQ is int8 where the dequant path is fp16, and the ragged sub-chunk prefill widths are exactly the ones the gate admits, so prompt logits shift slightly. That is a behaviour change, not a perf change, and it needs a quality number rather than an argument. HumanEval+, 164 tasks, four runs interleaved off/on/off/on: off 145/164 on 145/164 pass@1 0.884 both arms same-arm churn 0 items on - off = 0 items identical pass SET, item for item, in both repeats of both arms 12 of the 164 completions do differ textually between off and on; none of them change verdict (11 pass either way, 1 fails either way). So the tokens move and the correctness does not. That comparison is only worth quoting because the instrument was fixed first: it had been sampling, with a same-arm floor of 30 flipped items, which is wider than any effect worth arguing about. Deterministic it has a floor of 0, which is what makes a 0-item difference meaningful instead of merely small. ctest -R gqh is 136/136 with GGML_GQH_MMQ unset, =0 and =1. GGML_GQH_MMQ=0 still opts out, verified at runtime and not just in the source: unset now dispatches MMQ at ncols 128 and the gate still declines 320, and =0 is dequant at both.
Prior GQH-vs-vendor-quant numbers were spread across two boxes, two drafters and a quality harness that was silently sampling, so they could not be combined into anything a hardware partner could read. These four scripts produce the whole comparison on one box, one drafter, one binary, in one sitting. gqh_format_sweep_arm.py takes one reading for one arm. It launches a fresh server with every cache disabled (--prefix-cache-slots 0, --prefill-cache-slots 0, --prefill-compression off) so a repeat reading cannot be answered out of a cache instead of measured, then records, from that same server instance: decode tok/s with accept rate and avg_commit parsed out of the server own [spec-decode] line, prefill tok/s at a short and a long prompt from usage.timings.prefill_ms over prefilled_tokens, and end-to-end throughput from the in-tree bench_he_http.py. It also stores cache_hit and cached_prefix_tokens per request, which is what actually proves the caches were off rather than merely asked to be, plus rocm-smi edge temperature, VRAM and the KFD process list before and after every phase, so a reading taken under contention can be thrown away on evidence. gqh_format_sweep_timing.sh and gqh_format_sweep_quality.sh drive the arms in ABBA order rather than AABB. A previous sweep here found the second-in-pair about 1 percent slower purely from the card warming from 32 to 57 C, which was nearly written up as a property of the format. ABBA puts both arms at position-mean 2.5, so a linear thermal or position drift cancels in the arm means instead of being charged to whichever format ran second. Both wrappers gate on a cooldown back to 36 C between steps and abort outright if the temperature probe stops parsing, because a cooldown that silently does nothing is worse than no cooldown at all. Neither wrapper kills servers by name. quality_humaneval_plus.py already kills the process group it spawned, and this box has hosted a co-tenant dflash_server for days at a time, so a name-matched sweep could take somebody else multi-day job down. The quality wrapper instead waits for GPU 0 VRAM to fall back to baseline and stops with a report if it does not. gqh_format_sweep_report.py reduces the readings to one table. It reports a min-max spread for every figure rather than a point estimate, at two levels: within a reading across its 5 requests, and across the two readings. For quality it compares the two same-arm repeats reply by reply and item by item, so the same-arm floor comes out as an item count and as a verdict-flip count. That floor is the number that decides whether a cross-arm difference means anything, and it is exactly what was missing before.
The format sweep could not answer "does MMQ engage on the serving workload"
because it had no arm that turned MMQ off, and it was pointed at the wrong
drafter. Both are fixed here, and the drivers grew the instrumentation that
makes the answer measured rather than argued.
Drafter. gqh_format_sweep_arm.py hardcoded Qwen3.8-27B-DFlash2-Q8_0.gguf
(2,056,414,752 B). The canonical artifact is qwen38-dflash2-q8_0-canonical.gguf
(2,045,471,776 B). Accept rate is the figure most sensitive to which drafter is
loaded, so every accept number taken through the old default is uncomparable
with anything else. Now defaults to the canonical file, with GQH_SWEEP_DRAFTER
to move that axis deliberately.
Arms. An arm is now a (target, server-env) pair instead of just a target, so
holding the target fixed and moving only GGML_GQH_MMQ gives a same-binary
control, and both comparisons run through one driver with identical prompts,
cache flags and geometry. Adds gqh_mmqoff and gqh_mmqon.
ne11 census. GQH_SWEEP_CENSUS=1 collects the library's ne11 histogram, which
reports how many GQH mul_mat calls clear ggml_cuda_gqh_mul_mat_vec and at what
widths -- the half of the MMQ justification the kernel crossover curve cannot
supply. The table is only flushed by an atexit handler, so census mode shuts the
server down with SIGINT/SIGTERM and escalates only if it will not leave; a
timing run still ends with the immediate kill. It is a separate run because the
counters are not free.
Output paths. Each sweep writes under ~/bench-out/<GQH_SWEEP_RUN>/. Without this
the gqh arm of the format sweep and the gqh arm of the MMQ control share a
filename and the second silently overwrites the first.
Contention. This box is shared, and the sweep now behaves like it:
- timing.sh no longer runs `pkill -x dflash_server`, which would have killed a
co-tenant's server as readily as ours. It reaps by pid and only after
ps -o user= confirms the owner, and reports anything it declines to touch.
- smi() records absolute VRAM bytes for GPU 0 alongside VRAM%. Percent alone
cannot police a 34 GB card: it is an integer, so a 342 MB co-tenant rounds
to 0.
- a reading whose pre-launch VRAM is above the idle baseline is refused
outright, with the reason recorded. That case is not hypothetical -- it
happened mid-sequence during this run: a co-tenant took 13.7 GB, our target
would not fit, the load failed and the driver cheerfully reported 0.0 tok/s
and None prefill for every sample. Silent zeros averaged into an arm mean
are worse than a missing reading, so the guard makes it loud and skippable.
GQH_SWEEP_ALLOW_CONTENDED=1 overrides.
Cooldown. GPU 0 reports edge temperature as N/A whenever the card has gone to
sleep, which is exactly the cooled state cooldown waits for, so every gap burned
the full 600 s and logged "gave up waiting". It now falls back to the junction
sensor and, if no sensor reads at all, treats VRAM at the idle baseline as proof
the card is asleep and settles briefly instead.
smi() also no longer raises on an N/A sensor. float("N/A") would have killed a
measurement over a cosmetic sample.
report.py takes the run label and the arm pair as parameters, carries the drafter
and per-arm env into the table, and guards the cross-arm quality block so the MMQ
control -- which has no quality artifacts -- does not crash the timing report.
The GQH-vs-IQ4_XS accept-rate ordering reversed between an older local probe (--draft-block-size 16) and the shipping-default sweep (block 0, i.e. the drafter's own block_size 8). Same box and the same canonical drafter in both, so the drafter could not be the cause, and the two probes differed in more than one place at once. Arms carry an "args" list now, and gqh_b16 / iq4xs_b16 move ONLY the block width while holding prompts, cache flags, drafter and metric fixed. That turns the question into a one-variable experiment instead of an argument about which of three differences mattered. What it settled, measured here on the canonical drafter: block probe GQH avg_commit IQ4_XS avg_commit ordering 8 prose 4.56 4.33 GQH accepts more 8 code 7.34 7.08 GQH accepts more 16 prose 6.83 4.33 GQH accepts more 16 code 10.85 11.03 GQH accepts LESS The reversal needs block 16 AND code prompts together, which is exactly the older probe's geometry -- and it reproduces that probe's GQH figure (10.85 here vs 10.979 then). At the shipping default the ordering is the other way round, so the "GQH accepts less" reading describes a non-shipping configuration and should not be quoted against the shipping build. Two details worth keeping in view. IQ4_XS reads avg_commit 4.33 at both widths on the prose probe: it saturates and every one of the extra eight draft positions is wasted, so its accept percentage halves purely because the denominator doubled. And avg_commit is n_generated/n_draft_steps, which equals accept_pct * verify_cap here because n_generated == n_accept_sum -- worth remembering before comparing an avg_commit across two block widths, since the same commit count can sit under very different accept percentages.
…ts port Three instrumentation faults, all found by a run that died before producing a number rather than by a number that looked wrong. The quality driver hardcoded the WRONG drafter -- Qwen3.8-27B-DFlash2-Q8_0.gguf (2,056,414,752), not the canonical qwen38-dflash2-q8_0-canonical.gguf (2,045,471,776). Accept rate is the figure most sensitive to the drafter, and this is the second driver in this directory found pointing at that file. It now defaults to the canonical one AND hard-fails on an md5 mismatch, so a silent substitution cannot recur. Port 8765 is held by another user on the measurement box, bound to the tailscale address. A probe of 127.0.0.1:8765 returned nothing, so wait_ready failed instead of generating against a co-tenant's server -- correct behaviour by luck, not design. Added a PFLASH_PORT override and moved off the contended port. The cooldown gate called rocm-smi --showtemp alone, which on this card intermittently omits GPU[0] entirely (15/15 empty in a direct test; adding --showmemuse makes it appear), and the old code aborted the whole sequence on an unreadable temperature. The probe is now retried and report-only, and the co-tenant gate keys on exact VRAM bytes against the idle baseline instead -- which matched on every pre- and post-reading check of four runs. Thermals cannot change greedy token output, so an unreadable temperature must not abort a quality run.
The reproduction doc had grown into an archaeology of its own revisions: a reader had to assemble the answer from a sanity-values section, a head-to-head table whose own footnote said not to quote it, a block-width section, and a prefill section whose title still claimed GQH had no MMQ path after it had one. Three different accept-rate figures appeared with no statement of which configuration each came from. Anyone outside this work would have drawn the wrong conclusion, and anyone inside it had to remember the order the numbers arrived in. Now it opens with Results: one configuration, stated once, and every number in that section from that configuration alone. R9700 on the canonical drafter with its md5, shipping defaults, two readings per arm, order-balanced, co-tenant checked against an exact VRAM baseline. bytes 13,440,110,432 vs 15,567,824,480 -13.67% HumanEval+ 145/164 both arms identical, McNemar p=1.0000 decode 95.42-95.46 vs 81.60-82.46 tok/s +16% decode ms/step 47.56-47.78 vs 52.25-53.50 -10% prefill 119 tok 526.8-530.2 vs 875.8-891.8 -40% prefill 6850 tok 771.4-773.3 vs 1073.2-1074.5 -28% end-to-end 118.83-118.88 vs 115.84-116.55 +2.6% Smaller, indistinguishable on quality, faster decode, slower prefill, net slightly ahead end-to-end. The prefill loss is stated as a loss rather than buried, because anyone benchmarking a prefill-heavy workload will find it. The quality claim is backed rather than asserted: determinism was established first (164/164 byte-identical replies within each arm, so the same-arm floor is 0 items and a one-item gap would have been real), and the arms are then shown to be discordant on 10 items, 5 each way, with only 54 of 164 replies byte-identical between them. So it is "no measurable difference on this benchmark", not "the same model". Also corrected, because both were load-bearing and wrong: * Section 8's "why GQH cannot widen" described the ncols-8 cliff as a property. It was a defect and it is fixed -- that arm went 56.4 to 195.8 tok/s. Retitled so the table reads as the diagnosis it was, not as current behaviour. * Section 9 predicted one MMQ path would fix prefill AND the wide widths. Half right. MMQ loses at prefill widths and loses harder with length, because it re-decodes the tile once per output column tile. Replaced the prediction with the measurement, the census that localises where MMQ does win, and the reason the width gate's exact threshold turns out not to matter. Section 7's superseded table is deleted rather than annotated; it tuned only the ddtree budget, which left both arms at the drafter's trained block width. Adds the two traps that cost the most time here: use the canonical drafter (three DFlash2-shaped drafters exist, all load, two drivers in server/scripts were found pointing at the wrong one), and avg_commit is accept% x block width, so figures taken at different widths are not comparable -- an apparent accept-rate reversal in the earlier notes was entirely this.
DeanoC
changed the title
The dp4a arm in both GQH tile loaders has never executed. On every arch the GQH MMQ gate admits, the wave32 WMMA branch is what compiles and runs; dp4a is dead code, and there is no build switch that forces it, so its tile indices have no test coverage at all. The only thing keeping it unreachable was the two-entry arch allowlist, which means adding a validated arch to that list would have silently armed an untested loader instead. amd_wmma_available(cc) goes in as an extra conjunct. It is not a replacement for the allowlist: GGML_CUDA_CC_IS_RDNA3 covers RDNA 3.0 as well, and swapping the list for the predicate would WIDEN the gate to tile shapes that were never measured on RDNA 3.0. The allowlist stays the validated-arch list and the predicate enforces the invariant the comment already asserted, so the two cannot drift apart. No behaviour change on gfx1201: amd_wmma_available admits RDNA 3.0, 3.5 and 4, a strict superset of the allowlist, so the conjunction is exactly the allowlist today. ctest -R gqh is 136/136 with GGML_GQH_MMQ unset, =0 and =1, and the mmq cases in that suite fail if MMQ declines, so dispatch is proven rather than assumed.
gqh_q8_level baked every rung's level grid to int8 against a flat
denominator of 127. That is the WIDEST lattice, not the best-fitting one:
every grid's amax is exactly 1.0, so any N <= 127 lands the extreme level on
+-N with no clamping, and the levels in between land wherever k/N happens to
fall. On a curved grid that fit is poor, and the cost is pure loss --
ggml_gqh_q8_denom now picks the N in 1..127 that minimises the level error
instead, and the compensating 127/N folds into the per-tensor weight scale,
so a better N costs one host-side multiply per dispatch and not one extra
instruction in any kernel.
THE OBJECTIVE IS OCCUPANCY-WEIGHTED RMS ABSOLUTE ERROR, not worst-case
relative error. A dot product accumulates sum_k w_k x_k, so what reaches the
output is the absolute perturbation of each weight; a level's relative error
appears nowhere in that sum. An earlier pass minimised relative error, which
chases the innermost level of a curved grid -- the level that contributes
least to any dot -- and reported a 20.7x gain that does not exist.
The prior is UNIFORM, not measured occupancy. Occupancy-optimal N was worth
about 3% more and would bake one model's histogram into a table every GQH
tensor of every model shares. Re-derived from the grid tables here rather
than trusted: the shipping artifact's grids come out at gqh3 {3: N=120,
4: N=120} and gqh4 {2: N=114, 3: N=122, 4: N=110}, and the rms level error
against N=127 improves 3.96x, 2.32x, 1.35x, 1.49x and 1.68x respectively.
That holds under a Gaussian occupancy prior too (1.30x to 3.37x on the same
five grids), so the uniform choice is not an artifact of the prior.
WHERE THE COMPENSATING FACTOR GOES IS THE WHOLE CORRECTNESS ARGUMENT. It is
NOT folded into GQH_Q8_XSCALE. That constant rides the activation pre-pass,
and the paired arm shares ONE quantised activation between the gate and up
tensors, which carry different grid codes -- and so different N -- for 26% of
this model's pairs. One activation cannot carry two weight-side factors, so
folding there would silently rescale one half of every fused pair. It goes on
the per-tensor weight scale, which is already selected per half by blockIdx.y,
under if constexpr (I8DOT) and #if !GQH_FP8DOT. The registry's tensor_scale is
left pristine, so the f32 arm, the dequant kernels and the fp16 converters
never see 1/N.
N rides the by-value grid struct that already travels as a kernarg, so the
LUT fill and the scale read one field sourced once. Since amax is exactly
1.0, gqh_q8_level(1.0f, N) == N exactly, and pair index 0 is the (-1, -1)
entry and SIZE-1 the (+1, +1) entry for all three rungs -- so the fill checks
both against +-N and traps. Verified reachable: baking the LUT at n+1 while
the scale keeps 127/n aborts with SIGABRT and an HSA hardware exception
instead of quietly rescaling every weight in the tensor.
The one-hot tolerance in test-gqh-backend stays LITERALLY max|w_row|/127.
Reparameterising it to /N would loosen it and make the tolerance move with
the thing it measures; it stays a fixed bar that the measurement now clears
by a wide margin, which is what turns it into a real measurement.
What certifies the change is a SECOND, tight, per-grid bound: a one-hot dot
on the int8 arm has one term and an exact activation, so the error is the
weight LUT error alone, and row_max * ggml_gqh_q8_maxe bounds it. Measured
slack used is 1.000 or below on all nine vector shapes and attains the bound
exactly on five of them, so 1.05 is a cushion for fp32 rounding rather than a
fudge factor.
NEGATIVE CONTROL, nine new cases. GGML_GQH_Q8N=127 puts the old flat
denominator back in the kernel while the bound stays where the derivation
puts it, and every one of the nine FAILS as it must -- slack used 1.18 to
49.56 against the 1.05 allowed. That control is also the only thing in the
suite that proves the int8 arm dispatched at all: had it declined, the f32
arm would answer bit-exactly and every tolerance in the file would pass.
ctest -R gqh is 145/145 (136 + 9 controls) with GGML_GQH_MMQ unset, =0 and
=1. GPU 0 returned to its idle VRAM baseline before and after.
Decode throughput does not move, which is the expected result: the kernel
instruction stream is identical either side, only one host-side float
changes. Order-balanced ABBA, block 16, 10 HumanEval prompts, decode-only
[spec-decode] timer, GPU 0 at its exact idle VRAM baseline before and after:
derived N 193.76 191.21 tok/s avg_commit 10.928 (both)
N=127 194.33 195.13 tok/s avg_commit 10.979 (both)
The 1.15% arm gap is smaller than the derived arm's own 1.33% spread and
sits under the rig's ~1.3% floor, so it is not a reading. avg_commit is NOT
literally invariant, though: it lands on 10.928 against 10.979, exactly
repeatable within each arm. A more accurate weight LUT moves the logits, so
it moves which draft tokens are accepted; -0.46% on a decode path is a
different token sequence, not worse acceptance.
Section 10 still prescribed GQH_HE_GOAL_TPS=163.0. The block-16 arm is at 195.84 tok/s (R9700, canonical drafter, 10 HumanEval prompts, decode-only [spec-decode] timer) -- the 3.47x section 8 already records -- so that goal is beaten by 20% and a search aimed at it passes every candidate on its first iteration and ranks on nothing. The bar is now 205.0, and it is labelled as what it is: a threshold derived from the current best plus clearance over the rig's own ~1.3% noise floor, not a measurement. It says to recalibrate after any move that shifts the best, and to do it on the box that will run the search. The AL floor goes with it. GQH_HE_AL_MIN=10.5 existed because throughput ranking alone preferred abandoning wide verify -- block-8 was 2.75x faster and rejected anyway. The wide arm now wins on throughput too (195.84 against 154.19), so the floor no longer buys the search anything it would not choose on its own; both stale overrides are called out explicitly rather than quietly deleted, so anyone holding the old recipe can see why it changed. Also fixed the closing note, which still calibrated against a 157-163 bar that no longer appears anywhere above it.
The per-grid denominators are a real fidelity gain -- weight rms 1.35x to 3.96x better across the grids this artifact uses, at zero runtime cost -- but they were gated on HumanEval+ and did not earn the default. Four readings, order-balanced ABBA, canonical drafter, R9700, greedy and byte-identical within each arm: derived N 144/164, 144/164 flat 127 145/164, 145/164 Same-arm floor is 0 items in BOTH arms, so the 1-item deficit is reproducible rather than noise. It is NOT statistically separable: the arms are discordant on 5 items, 3 to 2, McNemar exact two-sided p ~ 1.0. So this is not evidence of harm, and the derivation stands. But the measured throughput gain was also nil -- 193.76/191.21 tok/s derived against 194.33/195.13 flat, inside the arm own 1.33% spread and under the rig ~1.3% floor -- and a change that alters generated tokens for no measurable benefit does not get to be the default. GGML_GQH_Q8N opts in. Reconsider the default if a workload appears where the extra weight precision shows up in an output metric.
The format string has four %d -- "at ncols %d and %d on grid %d ... over %d offsets" -- and only three values were passed. Every conversion after the third %d was therefore shifted by one: %s was handed a double, and %.9g read past the end of the varargs. Undefined behaviour, on the SUCCESS path of the two guard-refuse cases, which is exactly why nothing ever noticed -- the cases pass, so nobody reads the line, and the numbers in it were wrong. Measured, gqh4 4x512 forced to grid 8 at mmq33, before: MMQ DECLINED (0 launches) at ncols 33 and 8 on grid 5, and the dequant path answered fp16-exact over -347811548 offsets(null) (scale 4.7186799) Every field after the launch count is wrong: the widths are 512 and 33, the grid is 8, and the offset count is 5. The refusal covers both widths the graph holds -- cols for the wide node and nwide for the probe -- and the sentence names both, so the missing value is cols. After: MMQ DECLINED (0 launches) at ncols 512 and 33 on grid 8, and the dequant path answered fp16-exact over 5 offsets (scale 4.7186799) Found while making the per-grid bound self-consistent; unrelated to that work, so it is separate. -Wformat had been reporting it all along.
Flipping the int8 LUT default to the flat 127 turned 35 of 145 gqh cases red
without anything being wrong with any of them. The cause was a coupling, not a
regression: test-gqh-backend built its per-grid max|e| bound from
ggml_gqh_q8_denom -- the DERIVED table -- and deliberately did not read
GGML_GQH_Q8N. That was how the nine WILL_FAIL cases failed. It also meant the
bound was a statement about a quantiser that was not necessarily the one running,
so the moment the default moved, every ordinary i8 case ran loose against a tight
bar.
A BOUND MUST DESCRIBE THE QUANTISER THAT RAN. Two knobs now, and they are
separate:
ggml_gqh_q8_denom_eff(type, code) the denominator ACTUALLY IN FORCE. Moved out
of gqh.cu into ggml/src/gqh.cpp so there is exactly one definition of the
GGML_GQH_Q8N policy and of the shipping default, read by the int8 arms and by
the harness. The harness cannot hold a stale copy of a default it does not
own. GGML_GQH_Q8N also accepts opt / derived / auto now, because "set it to
something outside 1..127" was a poor way to spell the opt-in.
ggml_gqh_q8_maxe_at(type, code, n) maxe at an ARBITRARY n, not only at the
argmin of rms. Same gqh_q8_round chain as the search, so host and device still
cannot drift. ggml_gqh_q8_maxe(t, c) is this at ggml_gqh_q8_denom(t, c).
GQH_TEST_BOUND_N harness-only: build the bound at this
denominator whatever the kernel uses. Unset means follow the kernel.
Unset, an ordinary i8 case is SELF-CONSISTENT and passes with the default either
way. The negative control gets its mismatch by SAYING SO -- kernel pinned to 127,
bound pinned to derived -- instead of exploiting what the default happens to be.
Both sides pinned, so the control is invariant to the default; the ordinary cases
pin neither, so ctest exercises whatever policy the caller set.
NO BOUND WAS WEAKENED. GRID_MAXE_SLACK is still 1.05 and the slack USED is still
at most 1.000 on all 36 ordinary i8 cases in every state measured -- attained
exactly at 1.000 on 19 of them under the flat default, so the cushion covers fp32
rounding and nothing else. Verified green at 145/145 with GGML_GQH_Q8N unset, at
127, at derived, at 0 and at 110.
The nine controls fail identically in all three states, kernel N=127 against a
derived-N bound:
gqh3 1x256 grid 5 bound N=71 slack 1.289
gqh3 4x512 grid 6 bound N=124 slack 2.368
gqh3 64x1024 grid 6 bound N=124 slack 2.368
gqh2_h 1x256 grid 4 bound N=81 slack 11.147
gqh2_h 4x512 grid 5 bound N=84 slack 49.562
gqh2_h 64x1024 grid 5 bound N=84 slack 49.562
gqh4 1x256 grid 4 bound N=110 slack 1.178
gqh4 4x512 grid 2 bound N=114 slack 1.223
gqh4 64x1024 grid 2 bound N=114 slack 1.311
The control now also proves GQH_TEST_BOUND_N is HONOURED: drop it and the case is
self-consistent and PASSES (measured: gqh3 1x256 slack 0.667, gqh4 1x256 0.933,
gqh2_h 4x512 1.000), which WILL_FAIL reports as a failure. Checked in all three
states.
The OK and FAIL lines now print the worst level error itself, max|e|/row_max, not
only the slack. Slack is measured against the bound and therefore moves with the
denominator; the level error is the physical quantity and is what makes two
denominators comparable. maxe is constant within a run so the argmax is the same
item and the product is exact -- no extra tracking.
The flat default is byte-for-byte the OLD arithmetic. The pre-optimal-s binaries
in build-ctl and the current build agree on i8 dense max rel to every printed
digit on all 36 ordinary cases (4.57e-4 to 9.55e-4), so the ~6.5e-4 neighbourhood
is exactly where it was before optimal-s landed.
… floor Two scorer faults, found only because a gate result turned on a single item. The extractor required a CLOSING fence: `re.findall(r"```(?:python)?\s*\n(.*?)```")`. A reply that runs to its token cap mid-block therefore matched nothing, `code` fell back to the entire reply, and the literal ```python line reached the grader as source -- SyntaxError on line 1, scored as a failure. 12 of 164 items take that path. HumanEval/28 was a complete, correct concatenate marked failed purely for a missing closing fence. The pattern now lets the last block run to end-of-reply and strips any residual fence line, so a properly closed block behaves exactly as before. Verified by re-grading stored replies: the old extractor reproduces the harness's own counts to the item, and the fix moves every reading by +1. The second fault is not fixable, only quotable: THE GRADER IS NOT DETERMINISTIC. Re-grading byte-identical replies flips a verdict -- HumanEval/39 (prime_fib) goes pass, fail, pass across grading passes of the same file. That is an execution timeout, not sampling. Every "same-arm floor = 0 items" claim we have made, including in the results table, measured REPLY determinism; verdict determinism is weaker and the floor is about +-1 item. So the optimal-s justification is corrected in place. It flipped to opt-in on the strength of "1-item deficit against a 0-item floor", and that reasoning does not survive: the gap is inside the floor. The default still stands, on the argument that does survive -- the throughput gain was nil, and a change that alters 38 of 164 replies for no measurable benefit does not get to be the default. "No better, and not free of consequence", not "worse". The published comparison is unaffected in direction: both arms shift together under the fence fix, and a 1-item difference is inside the floor either way, so GQH and IQ4_XS remain indistinguishable on this benchmark. The absolute pass@1 in the results table is about one item low, and is left as measured with the floor stated beside it rather than silently restated -- re-running under the fixed scorer is the honest way to raise it, not re-grading and quoting the better number. Resolving a one-item effect here would need on the order of 3,700 items at the observed discordance rate, and that assumes a deterministic scorer.
…lects
Prefill -23% on the R9700, decode flat. Found by the evolutionary search over
mmq.cuh, iterations 1-5, and this is the cumulative validated state of that run:
iter 1 small MMQ tile for the GQH rungs 1.0114 prefill +8.5%
iter 3 half-k tile second half in registers 1.0205 prefill +3.9%
iter 4 that tile was in SCRATCH, not registers NEUTRAL (diagnosis, no ship)
iter 5 v_perm_b32 for the LUT decode 1.0921 prefill -23%
Cumulatively prefill goes 405.9 -> 599.2 tok/s on the search's own harness at
block 16, with avg_commit identical at 10.9790 throughout.
MEASUREMENT. Paired A/B, arms swapped, three rounds, libggml-hip.so md5 verified
distinct per arm and repeating within an arm:
round base prefill cand prefill d prefill d ranked
1 127.87 ms 99.07 ms -22.52% -4.49%
2 130.40 100.05 -23.28% -5.05%
3 132.49 100.35 -24.26% -5.41%
Same sign in all three rounds and the magnitude is ~5x the base arm's own 5.2%
within-arm spread (cand's spread is 1.4%). Both arms degrade monotonically round
to round and cand is always measured LATER within a round, so the thermal drift
works against the result rather than producing it.
DO NOT QUOTE THE 23% ALONE. Prefill is only ~22% of the ranked quantity
(prefill_s ~1.28 s against ~4.43 s of decode), so the honest pair is prefill -23%,
end-to-end -5%. The committed CSV reads LATENCY_SPEEDUP 1.0921 rested and ~1.07
mid-burst; those are the same change on a cool and a hot box, not a regression.
The mechanism accounts for the measurement. Weighting the estimated dynamic VALU
cuts by each rung's measured trace time (GQH4 164.78 ms, GQH3 113.03 ms, so
59.3%/40.7%) predicts -25.7% VALU against a measured -23% prefill.
WHAT THIS IS NOT: iteration 6 built on this and was cut off mid-A/B, so its
candidate is unmeasured and is NOT included here. The search clone's working tree
carries it (1126 lines dirty against this patch's 1007) and the handoff's "the
dirty state IS the committed state" was true at the end of iteration 5 and stale
by the end of iteration 6. This commit is the patch at the iteration-5 record,
not the tree.
Prefill was the one axis where this artifact lost to IQ4_XS (526.8-530.2 tok/s
against 875.8-891.8 at a 119-token prompt). This narrows it. The published
results table is NOT updated here -- these figures come off the search harness at
block 16, and the table's configuration has to be re-measured before it moves.
Both arms re-measured on ONE box at ONE commit, rather than carrying figures forward from the box the previous table was taken on. Two readings per arm, order-balanced, canonical drafter, shipping defaults, caches off and cache_hit=true never observed across 120 readings. The table's largest loss more than halves and the end-to-end win roughly quadruples: row was now prefill @119 tok -40.2% -18.5% prefill @6850 tok -28.1% -30.6% decode tok/s +16.4% +15.1% decode ms/step -9.8% -8.7% end-to-end +2.3% +9.0% Short-prompt prefill moved 38% WITHIN this branch, measured same-box one commit apart: 507.7-515.4 tok/s at the parent against 701.2-711.3 at the head, non-overlapping ranges. A same-box parent comparison was used deliberately -- a cross-box reference cannot separate "the change moved it" from "the box differs", and it turns out two GQH rows do carry a 3-4% box offset that the untouched IQ4_XS arm bounds at <=1.9% on every row. Long-prompt prefill did not move (+0.02%) and structurally cannot: gqh_mmq_max_ne11 is 160, so a 512-wide prefill chunk never enters the GQH MMQ kernel where this code lives. That is where the remaining 30.6% gap sits, and it is a weight materialisation cost rather than a kernel-tuning one -- so the section now says plainly that prefill is still the axis this artifact loses on, and that a prefill-dominated workload will not see the end-to-end figure. Decode, accept % and avg_commit are unchanged to the digit, which is what a bit-identical decode path should give, and the accept/avg_commit rows are identical to the previous table's -- an independent check that harness and drafter match the earlier setup. The stale "Where the speed comes from" subsection is deleted rather than annotated; its 47.6 ms/step figure is superseded by the re-measure.
The 160 bound was the widest all-shapes-win crossover when it was measured, but it predates the v_perm_b32 weight-LUT decode. Re-sweeping gqh-mmq-sweep with the gate lifted leaves every artifact shape ahead at 640, and makes the 512-wide prefill chunk a 7-27% kernel win rather than the ~1.2x loss the old comment recorded. The ne11 census over a shipping-default serve shows why that matters: of the 36549 GQH mul_mat calls that reach the gate, 27510 are the 194 and 512 buckets, and 6850 = 13*512 + 194 exactly, with 25545/1965 = 13.0. Those two buckets are the long prefill, chunked, and the gate was sending all of it to the dequant path. Measured end to end on lucebox5, every cache off, ABBA order-balanced, two readings per arm and five samples per reading, cooled to 36C between: arm decode tok/s prefill @119 prefill @6850 gate 160 95.00-95.30 702.5-709.6 723.3-751.3 gate 512 95.00-95.30 702.9-708.8 852.5-866.5 +15.6% Prefill at 119 tokens is the negative control and does not move: ne11 119 was already under the old bound, so the gate cannot touch it. Decode does not move either, since it never reaches the gate. The whole gain is long-prompt prefill, which is the worst row of the published GQH-vs-IQ4_XS table. 512 rather than the 640 the kernel curve would allow, because the shipping prefill chunk is min(DFLASH_QWEN35MOE_PREFILL_CHUNK, prompt_len) and that env defaults to 512, so no prefill can present a wider ne11. A 640 bound admits nothing extra for any prompt and spends GQH3's whole margin: 0.999 at 640 is break-even, 0.927 at 512 is not. A 640 arm measured identically end to end, which doubles as the internal control on the noise floor. This retracts the previous claim in the gate comment that lifting it cost ~7% of prefill throughput end to end. That held before v_perm_b32; at this commit lifting it to 512 is a 15.6% prefill gain. Tests: gate-in160/out161 pinned the old default, and gate-out161 would now dispatch. They are retargeted to gate-in512/out513, keeping the property that MMQ dispatches at the threshold and declines one above it. Both need a vector that straddles the default (the probe requires nwide <= cols), so they now register only on 64x1024. The knob pair moves to its own width condition so its three-shape coverage of the gate mechanism is not tied to whatever the compiled default happens to be. ctest -R gqh is 137/137 with GGML_GQH_MMQ unset, =0 and =1. The small row tile in mmq.cuh was chosen under the old bound, where a gate-admitted multiply fitted one output column tile. At 512 it no longer does, so the weight superblock is decoded more than once. The change was measured end to end with that tile in place, but the comment there is corrected and the tile is flagged as unswept at the new width rather than left resting on the stale ntx == 1 argument.
prefill @6,850 goes 719.8-750.6 -> 852.5-866.5 tok/s, so that row's delta against IQ4_XS moves -30.6% -> -19.5%. Together with the earlier v_perm_b32 change the prefill axis roughly halves at both ends: -40.2% -> -18.5% short, -30.6% -> -19.5% long. The previous text said long-prompt prefill "did not move, and structurally cannot", because gqh_mmq_max_ne11 was 160 and a 512-wide chunk therefore never reached the MMQ kernel. The structural claim was right and the conclusion was wrong: the bound itself was the problem. It had been measured BEFORE v_perm_b32 made the kernel ~2.3x faster, so the crossover had moved out from ~256 to ~640 and nobody had re-read it. MMQ is 7-27% faster than dequant at width 512. Raised to 512 rather than the crossover's 640, and the reason is better than the crossover: qwen35moe_prefill_chunk_limit is min(DFLASH_QWEN35MOE_PREFILL_CHUNK, prompt_len) with an env default of 512, so no prefill can ever present ne11 > 512. 640 admits nothing extra for any prompt while spending gqh3's entire remaining margin -- 0.999 at 640 is break-even against 0.927 at 512. Causation is measured, not argued. The MMQ launch counter reads 0 at gate 511 and 1 at gate 512 for an ne11=512 node; an ne11 census over a shipping serve puts 27,510 of 36,549 gate-visible calls in exactly the 194 and 512 buckets, and 6,850 = 13x512 + 194 with a measured 13.0 ratio, so those buckets are the long prefill chunked. Controls: prefill @119 does not move (already inside the old bound), decode does not move (never reaches the gate), and gates 512 and 640 agree to 0.14pp because nothing lands between 195 and 511. Both of this branch's largest prefill wins came from a stale assumption rather than from more search: the gate was set by a measurement that predated the kernel it gated, and the earlier comment claiming a lift "costs ~7% of prefill throughput end to end" was from the same era. One further assumption of that kind is now flagged in mmq.cuh rather than trusted: the small-tile choice rested on ntx == 1, which only held while dispatched widths were <= 160.
…x == 1 The small-tile comment justified itself partly with "the multiply fitted ONE output column tile and no superblock was decoded twice", which held only while gqh_mmq_max_ne11 capped dispatched widths at 160. At the 512 default a multiply spans several column tiles and that half of the argument is void. The previous commit corrected the text and left the tile choice explicitly open. This closes it by measurement, and the answer is that the tile does not change. Both arms are this commit differing only in the mmq_gqh_small_tile constexpr, selected by swapping libggml-hip.so so ONE sweep binary and one server binary measure both and the comparison cannot drift in harness, widths or iterations. GGML_GQH_MMQ_MAX_NE11 is lifted for the sweep so MMQ dispatches at every width in BOTH arms -- without that everything above the bound falls to dequant on both sides and the sweep measures nothing -- and the per-width mmq/iter column is the proof it dispatched. Default tile passes ctest -R gqh 137/137 with GGML_GQH_MMQ unset, 0 and 1, so this is a throughput question and not a correctness one. gqh-mmq-sweep, iters 400, eight readings per cell over two order-balanced passes. Ratio is default-tile time over small-tile time, so above 1.0 the small tile is faster: ncols 119 194 256 384 512 640 gqh3 1.055 1.000 1.015 0.982 0.982 0.996 gqh4 K17408 1.196 1.071 1.090 0.983 0.981 1.028 gqh4 K6144 1.178 1.064 1.094 0.976 0.979 1.026 The crossover is between 256 and 384 and the two widths this workload dispatches straddle it, so neither tile dominates: small is 0-7% ahead at 194, default ~2% ahead at 512. Weighting by the real chunking of a 6850-token prefill, 13*512 + 194 from the ne11 census, predicts the default tile winning that one prefill by 1.8-2.0% and nothing else. That prediction is what the real workload then returns, which is why the sweep above is worth trusting as the map of this trade. ABBA order-balanced, two readings per arm, five samples per reading, prefix and prefill caches off and no cache_hit observed, shipping gate: arm prefill @119 prefill @6850 e2e HE-10 decode small tile 688.7-709.6 848.9-865.6 126.57 / 126.78 94.9-95.2 default tile 643.6-648.1 863.1-876.9 124.34 / 124.76 95.0-95.2 The default tile takes 1.5-1.7% of the long prefill and gives back 8.8% of the short one and 1.6-1.9% of end-to-end. Decode is flat in both arms, as it must be: the matvec owns widths 1..16 and returns before the gate ever sees them. The small-tile arm is also the pessimistic one -- its first reading started at 43 C against 36 C for every other reading, which if anything flatters the default arm -- and it wins regardless. So the rung keeps the small tile, and the comment now rests on this sweep rather than on a premise the gate change retired. What would overturn it is a workload whose prefill widths sit above the 256-384 crossover with the short buckets gone, which the 512 chunk limit and the census say this one is not. Comment only; no code changes. The rebuilt libggml-hip.so differs from the pre-change one in 1248 of 139052904 bytes, all line-constant metadata shifted by the 42 inserted comment lines: ctest is 137/137 on all three MMQ arms and every sweep width reproduces its recorded range.
The results table's authority rests on a stated method -- both arms, one box, one commit -- and two cells no longer honour it. The GQH prefill @6,850 row was patched in from a GQH-only run after the width gate default moved, while the IQ4_XS column still comes from the earlier both-arms run. A GQH-only gate change cannot move IQ4_XS, so the figure is not wrong; the problem is that patching a single cell quietly converts a self-consistent table into a composite one, and the method line is the reason a reader trusts it at all. The end-to-end row has the same weakness from a different direction: it did not reproduce to its stated 126.60-126.68 band in a later A/B, which read 126.79-126.88 on both arms. That is inside noise, but it is not the same measurement, and a later A/B also showed the apparent e2e movement between gate arms was within-reading phase ordering -- the pinned-160 arm spends 46.1 s in the long prefills against 39.9 s and enters e2e 2-4 C hotter -- not the gate. The ten e2e prompts are 39-93 tokens, below both bounds, so the gate provably cannot touch them. Caveat rather than silent patch, and rather than reverting a correct number: both arms are being re-measured at the head, and this note comes out when the table is self-consistent again.
…thod supports Re-measured both arms at this commit so the table honours its own method again. The provenance caveat comes out; the deltas that were patched in from a GQH-only run are now from the same run as the column they are compared against. row GQH IQ4_XS delta decode tok/s 95.00-95.20 82.50-82.70 +15.1% decode ms/step 47.89-48.00 52.42-52.50 -8.6% prefill @119 703.3-712.6 858.6-868.6 -18.0% prefill @6,850 851.7-866.3 1064.0-1074.7 -19.7% end-to-end (isolated) 127.18-127.23 117.89-118.02 +7.8% accept / avg_commit 56.9 / 4.56 54.2 / 4.33 Everything reproduced except end-to-end, and that row is corrected downward: +9.0% -> +7.8%. Two reasons, both method rather than kernel. First, e2e is thermally coupled to the phase before it: run in sequence after the long prefills, GQH enters e2e 2-5 C hotter than IQ4_XS because it spends ~8 s longer in that phase, and its within-arm spread inflates to 0.87 against IQ4_XS's 0.14. Measured on an isolated cold server -- all readings entering at 38-39 C -- the spreads collapse to 0.05 and 0.13 and the delta lands at +7.8%. The in-sequence figure flattered us by ~0.4 points of delta, so the isolated number is the one quoted. Second, the previous 0.08-wide band was finer than the instrument: e2e measures 763 tokens over ~6.0-6.5 s, so 0.05 s of wall clock moves it ~0.8%. It is now quoted to about +-0.3 rather than to two decimals it cannot support. The quality row is labelled *carried over* rather than absorbed into the one-commit claim: it was not re-measured in this run, and the table's authority comes from that claim being exactly true. Also corrected two stale annotations that both still said the width gate was 160 after the default moved to 512 -- the defaults bullet in this doc, and the env_note that gqh_format_sweep_arm.py copies into every JSON artifact it writes. The second was cosmetic but it labels the provenance of measurements, which is the last place a wrong constant should sit.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Adds a GQH MMQ path, widens the GQH verify arms, and puts correctness gates on the
two arms that had none. Measured on an R9700 (gfx1201).
Results
Both arms measured on the same box at the same commit: R9700/gfx1201, canonical
drafter
qwen38-dflash2-q8_0-canonical.gguf(2,045,471,776, md5a98fb401578886f082315c7031f419a2), shipping defaults (MMQ on, width gate 160,draft block 8,
GGML_GQH_Q8Nunset), prefix and prefill caches off withcache_hit=truenever observed, two readings per arm, order-balanced, GPU at itsexact idle VRAM baseline before and after every reading.
13.7% smaller, indistinguishable quality, decodes ~15% faster, still prefills
slower, ~9% ahead end-to-end.
Prefill is still the axis this loses on, and that is stated as a loss. It was
-40.2% at the short prompt before this branch and is -18.5% now; the long prompt
is unchanged at -30.6%.
gqh_mmq_max_ne11is 160, so a 512-wide prefill chunknever enters the GQH MMQ kernel -- the remaining long-prompt gap is a weight
materialisation cost, not a kernel-tuning one. A prefill-dominated workload will
not see the end-to-end figure above.
The short-prompt gain was measured same-box, one commit apart (507.7-515.4 ->
701.2-711.3 tok/s, non-overlapping) rather than against a cross-box reference,
because a cross-box comparison cannot separate the change from the box. Two GQH
rows do carry a 3-4% box offset, which the untouched IQ4_XS arm bounds at <=1.9%.
The quality claim, and what backs it
Determinism first: with the harness pinned to
temperature: 0, top_k: 1, bothreadings of each arm returned 164/164 byte-identical replies. Item-by-item the
arms are not the same model -- 140 pass in both, 14 fail in both, 10 discordant,
5 each way -- and only 54 of 164 replies are byte-identical between arms.
McNemar exact, two-sided: p = 1.0000.
The instrument has a floor of about one item. Two scorer faults were found
here: an unterminated code fence used to fail a correct answer (12 of 164 items
take that path; fixed in this branch), and the grader itself is not
deterministic -- re-grading byte-identical replies flips
HumanEval/39(
prime_fib) on an execution timeout. So a one-item difference on this benchmarkis not a result, in either direction.
What the change is
Wide verify arms. Every dispatched verify width gets its own exact-width arm
(
GQH_MULTICOL_SPEC_MAX12 -> 16), then width 16 gets its own ROWS via a per-widthtable. That arm went 56.4 -> 195.8 tok/s (3.47x) with all control arms unmoved
and
avg_commitidentical to four decimals, so it is per-step cost only.MMQ for GQH3 and GQH4, width-gated by
gqh_mmq_max_ne11()(default 160) andon by default. MMQ re-decodes the weight tile once per output column tile, so it
loses at prefill widths and loses harder with length; it wins the narrow chunks,
1.68x on a 119-token prefill. Decode never reaches it -- the matvec owns widths
1..16 and returns first. So MMQ is creditable for prefill, never for decode. An
ne11census puts 9,039 of 85,839 GQHmul_matcalls inside the gate; theworkload's widths are bimodal with a hole (8, 39-119, 512), so no width lands
between 105 and 511 and the exact threshold is not load-bearing.
An int8 range guard, because int8 cannot represent every GQH grid: GQH4 codes
8-11 exceed 127:1 and codes 10/11 round two levels to exactly zero -- 17.1% of
weights in the grid we sampled. The shipping artifact uses only GQH3 {3,4} and
GQH4 {2,3,4} (worst 36.30:1), so the guard rejects nothing it contains, but the
failure it prevents is silent.
Correctness gates: 50 -> 136 tests
Two arms had no coverage at all, and neither gap was where it looked.
Wide activation positions. The fused check drove
xwith one-hot vectors atposition
j < nvec, so withcols >= 256it never placed a nonzero activationpast position 15 -- the k-reduction was only ever exercised in its first
superblock. Added offset-swept one-hot (still single-term, so still bit-exact) and
dense activations.
The I8 activation arm -- what serving actually runs. It had no coverage, and
the env var was never the reason:
gqh_i8_shape_ok()floors the arm at 16 Mi ofin*rowsand the largest suite case is 64 Ki, 256x under it. So the arm wasunreachable from ctest at any setting of
GGML_GQH_I8DOT, and both settingsreturned an identical dense max rel of 8.04e-06. Lifted host-side with
GGML_GQH_I8_MINWORK=0; tolerances are int8-sized and normalised by||w_row||*||x_col||rather thanmax(1,|ref|), which removes a sqrt(in) artefact(0.039/0.085/0.32 at in=256/512/1024 became 5.4e-4/6.6e-4/6.5e-4).
Every gate has a negative control. Each was built in a separate directory so no
shipped library held a bugged build, then restored and md5-verified:
zero differences while the new ones catch 2044 and 1024 -- the old gate was
provably blind to it
a 2e-2 bar), passes the f32 arm and the narrow I8 arm
w ^ 1): fails exactly the MMQ cases<=-><, and threshold 160 -> 192): fails exactly the boundarycases
numeric bound did not catch it, because int8 destroys the innermost levels
whose absolute error is tiny. A collapse check (no non-zero weight may return
exactly zero) is what catches it.
dequant fallback, since MMQ's tolerances are looser than the fallback's error
Two harness faults fixed here, both of which corrupted results
The quality harness was silently sampling.
quality_humaneval_plus.pyprinted"pass@1 (greedy, raw)" and never sent
temperature, so the server used the modelcard's sampling defaults. Two runs of one config agreed on 16.5% of replies and
30 of 164 items flipped verdict; two totals matched only because 15 items went
each way. Any n=1 pass@1 from this script before the fix is unreliable. Pinned to
temperature: 0, top_k: 1, which also raised pass@1 from 0.756 to 0.884.ctest's
ENVIRONMENTproperty adds to the ambient environment rather thanreplacing it, so benchmarking with
GGML_GQH_MMQ=1silently moved 42 existingcases onto MMQ while they still demanded bitwise dequant output. The search loop
runs
ctest -R gqhas its correctness gate inside the benchmark's ownenvironment, so it would have seen a fault that does not exist. Pinned
GGML_GQH_MMQ=0on every case that asserts the dequant result.Reviewer notes
code, and
should_use_mmqadmits only RDNA3.5/RDNA4 -- both of which have WMMA.So those tile indices are unreachable on every arch this enables and nothing
runnable on this hardware tests them. A wrong dp4a tile size would pass the whole
suite. This is the one place I would want a second opinion, or a decision to make
the gate refuse non-WMMA arches outright.
acceptance shifts). The quality table above is the evidence they do not get
worse; it is not a claim that output is unchanged.
two drivers in
server/scripts/were found pointing at the wrong one.qwen38-dflash2-q8_0.gguf(1,838,540,000) came from a converter that silentlydrops DFlash2's 23 conv and candidate-selector tensors.
avg_commitisaccept% x draft block width, so figures at different widths arenot comparable. This produced an apparent accept-rate reversal in earlier notes
that was entirely the metric.