perf(sparse_moe): keep two Q4 group quads in flight in the routed gate/up dot product

[MichaelDementii]

Rebased onto 487f8977. Where the numbers below come from

Base. This package was written against ad0f3d38 and now targets upstream 487f8977
(perf(sparse_moe): one CTA per token in small-T S2, and a warp merge in place of eight dependent rounds), 83 commits further on. ad0f3d38 and the earlier heads still named
throughout this report are the commits the work was done and measured on. They are left
standing rather than rewritten, so that every number below stays attached to the tree it was
actually taken from.

Provenance of every number. Every timing, ratio, percentage, bandwidth, TFLOP/s figure,
register and shared-memory count and ctest line in this report was measured on ad0f3d38, or
on an earlier head named at the point of use. Nothing here has been re-measured on
487f8977. Two changes to the instrument itself bear on that, and are stated once here
instead of being repeated at every table:

• The engine context cache changed sides. On 487f8977 ninfer_bench switches it off
  itself - engine_options.context_cache.enabled = false,
  bench/targets/qwen3_6_27b/ninfer_bench.cpp:157, introduced upstream by 385b30ce. On
  ad0f3d38 that line does not exist at all, so the bench inherited the engine default
  ContextCacheOptions::enabled = true (include/ninfer/types.h:130 there, :132 here) and
  every end-to-end run behind this report was taken with the cache constructed. Two bounds
  on how far that reaches, both stated because the difference has not been measured. From
  above: a live context cache can serve a prefix and shorten a prefill, which for prefill and
  TTFT would be a first-order difference rather than a rounding one. From below: the bench also
  sets options.execution.allow_prefix_reuse = false on every request
  (bench/targets/qwen3_6_27b/ninfer_bench.cpp:65) - byte-identical on both bases - so no
  request in these runs was eligible for a prefix hit at all, and what the old configuration
  carried was a constructed-but-unused cache (host state slots and pinned host KV) rather than
  served prefixes. Which bound is nearer the truth is an open question here; until it is
  measured, read the prefill and TTFT figures as taken in a configuration the current bench no
  longer builds.
• The bench flag was renamed. --mtp-draft-tokens no longer exists anywhere in the tree.
  The current spelling is --spec mtp --draft-tokens N
  (bench/targets/qwen3_6_27b/ninfer_bench_support.cpp:356-359), the same pair the product CLI
  takes (apps/cli/options.cpp:141-145). Every reproduction line below has been rewritten to the
  current spelling; the runs behind the numbers used the old one. One trap in that rename is
  worth naming: --spec mtp --draft-tokens 0 is rejected - validate_speculative_cli_options
  requires [1,5] for MTP (src/product/speculative_options.h:41) - so the zero-draft arm,
  written --mtp-draft-tokens 0 throughout the old text and called mtp0 below, is spelled on
  the current bench by passing neither flag, which is already the default
  (SpeculativeBackend::None).

Registered tests. tests/CMakeLists.txt moved as well. The suite that registered 104
targets on ad0f3d38 registers 114 on 487f8977 (+10 targets), and one more of them is
artifact-gated: 20 named targets now carry an explicit SKIP_RETURN_CODE 77 against
19 on ad0f3d38, on top of the blanket rule inside ninfer_add_op_test that marks
every op test skippable. The addition is ninfer_qwen3_8_27b_dflash2_real_test, gated on
weights like the rest, so one further test skips on a host that carries none. Every ctest count printed below is consequently a
statement about ad0f3d38 or an earlier head, not about the current base; the counts are left
in place rather than restated, with the base each belongs to named beside it.

No ctest run in this report has been repeated on 487f8977, and no round "114 of 114" is
claimed for it. Two facts forbid that claim. ninfer_attn_input_proj_test is red on the
bare base - an upstream defect, filed as issue #196, not something this branch introduces.
And ctest and a direct run of the same test binary have been observed to disagree on this
host, so a green ctest is not by itself evidence of anything. The only formulation this
package will stand behind on the new base is the exact one: the arm's result matched the
base, and the single failure reproduces on bare 487f8977.

Line references re-checked against 487f8977. Every file:line citation in this report was
opened on the new base. The ones that moved are corrected in place; the number each one used to
carry is kept here so nothing is replaced silently.

citation	printed here before	on 487f8977	what stands at the old number now
include/ninfer/types.h, ProposalHead::Full	:80	:82	:80 is the comment // Startup-fixed K: MTP 1..5; DFlash and DFlash2 1..15 (query width K+1).

Linked Issue: #NNN (ISSUE_BODY.md in this package), scope confirmed there.
Base: ad0f3d38. One mechanism, one function body, one file.

What changes

dot_two_rows walks the routed gate/up K dimension in quads of four Q4 groups. One quad is one
128-byte warp transaction per row, and each lane owns eight consecutive K values, so a single
mantissa decode feeds eight FP32 FMAs. Today the loop issues one quad's code words and scales and
then decodes them immediately, so for the whole of that decode and its sixteen FMAs the load unit
has nothing outstanding.

This change issues both quads of a pair before either is decoded (kQuadsInFlight = 2), so
the decode and FMA chain of the first quad cover the memory latency of the second.

The groups are still visited in ascending order, so the order of additions into acc0 and
acc1 is unchanged and the output is bit-identical.

Diff: 50 insertions, 23 deletions, one file. No new template parameter, no new instantiation, no
launch-shape change, no workspace change.

About the static_assert. It pins kGroups % (4 * kQuadsInFlight) == 0, i.e. the
divisibility of the full K. That is not the quantity the loop depends on: the loop walks
[k_begin / kGroupK, k_end / kGroupK), so what the pairing needs is that the span is a
multiple of 4 * kQuadsInFlight. That is a property of the arguments and cannot be asserted at
compile time. The two coincide today because all four call sites pass k_begin = 0, k_end = K
(kGroups = 32, so 32 % 8 == 0); the assertion does not enforce that, and a future caller passing
a partial K range would walk past last_group without tripping it. The patch says exactly this in
a comment above the assertion rather than leaving it to be discovered.

Why this and not something else

Three measurements bound the claim. All three are in the Issue; the short form:

The path is not bandwidth-bound. Holding the kernel, the token count, the grid, the warp count
and the instruction count fixed and changing only how many distinct experts the routing touches
(--distribution same | trace-like | independent at --tokens 4 = 8, 22, 31 unique experts), the
mandatory weight traffic moves 3.30x while the kernel moves 1.62x. At the fixture's
trace-like point the kernel demands 930 GB/s = 55.1 % of the measured 1689.4 GB/s DRAM read
ceiling. Extrapolated to zero experts it still costs 16.8 us, i.e. 59 % of its own trace-like
time does not depend on the weights it reads at all.

The byte numerators there are computed from the routing geometry, not read from a counter; ncu
is unavailable on this machine. The comment inside the patch quotes the same 55 % and names the
same point.

It is not a tensor-core path. cuobjdump -sass over sparse_moe_decode_kernels.cu.o reports
zero HMMA, IMMA or OMMA. The kernel is SIMT fmaf throughout, so the MMA roofline does
not apply to it.

The grid is not the knob. The launch already issues exactly one CTA per unit of work
(6144 = 4 tokens x 3 path blocks x 512 columns), so a grid-stride form can only narrow it, and
every narrower grid measured is worse: the same work split rewritten as a strided loop costs
+2.22 %, half the columns +19.72 %, a quarter +65.59 %. Feed depth and grid together are worse
than feed depth alone (+53.63 %). These four figures, and the depth dose below, come from a
separate research build in which both knobs are environment-selected template arguments; they
are reported to bound the claim, and none of them is part of this change.

The same feed-depth change on the routed down kernel is -0.15 %, i.e. nothing, because its
k loop has two iterations rather than eight. It is not included.

Verification

Hardware and toolchain: one RTX 5090 (sm_120a), driver 580.159.03, 525 W cap, CUDA 13.1.115,
Release, -DCMAKE_CUDA_COMPILER=/usr/local/cuda-13.1/bin/nvcc. Artifact Qwen3.6-35B-A3B
(groupwise-int), codecs q4-q5, --spec mtp --draft-tokens 3 (T = 4).

Method: master and change are two separate binaries built from the same tree, differing only
by this diff; no environment variable selects behaviour. Every cell is its own GPU-exclusive
capture with its own card-witness window. Kernel medians come from nsys, not wall clock: the
bench's wall clock for this operation lands on a 2.048 us lattice, and the operation's own wall
median is 55.296 us at trace-like (47.104 at same, 63.488 at independent), so one lattice
division is 3.7 % of it — larger than the effect being measured.

How the witness window is taken, precisely. The timestamp that opens the window is written
before the run is handed to the lock wrapper, so the window includes any time spent waiting for
the card lock. A neighbour legitimately holding the card during that wait therefore dirties a
window that covers a clean run. The bias is one-sided — it can only reject a good cell, never
accept a bad one — and it did reject cells (below), but the method is not "the window covers the
run", and this body no longer says that it is. A later harness of mine stamps the window inside
the lock instead; that is the form to copy on any re-measurement.

Cells discarded. Thirteen attempts were discarded, on three cells: p1_same took five
attempts, p1_indep nine, k1_base two (raw/journal.txt, lines beginning DIRTY). Twelve of
those thirteen are in the p/k research series, whose numbers are not submitted. In the
submitted series — the two operator passes q1/q2 and the four end-to-end passes — every cell
was accepted on the first attempt and nothing was discarded (raw/journal.txt,
raw/prodround.witness.txt).

Operator

nsys profile -t cuda --cuda-graph-trace=node \
  ninfer_sparse_moe_bench --codec q4-q5 --tokens 4 --distribution <D> \
                          --execution graph --cache cold --warmup 3 --repeat 20

Two passes per arm, arm order rotated between passes; medians of the two passes:

distribution	unique experts	master (2 passes)	change (2 passes)	delta
same	8	21 472 / 21 376 ns	20 448 / 20 352 ns	-4.78 %
trace-like	22	28 768 / 28 736 ns	26 464 / 26 784 ns	-7.40 %
independent	31	34 624 / 34 592 ns	32 576 / 32 448 ns	-6.06 %

The depth itself was chosen by a dose rather than assumed. On a separate research build in which
the depth is an environment-selected template argument, trace-like gives depth 1 (master)
28 928 / 28 832 ns, depth 2 26 464 / 26 559 ns, depth 4 27 103 / 27 039 ns - the optimum is
interior. Two was then rebuilt as a constant and re-measured as the two binaries above.
Pass-to-pass spread inside a cell is 0.1-1.2 %.

Round

nsys profile -t cuda --cuda-graph-trace=node \
  ninfer_bench --weights qwen3_6_35b_a3b.ninfer -pg 1000,128 --prefill-chunk 8192 \
               -r 1 --warmup 1 --spec mtp --draft-tokens 3

Two arms x two passes. Routed gate/up kernel median 24 416 / 24 416 -> 23 744 / 23 744
(-2.75 %) over 5240 launches; kernel total 130 512 610 / 130 622 538 -> 126 452 934 /
126 528 751 ns (-3.11 / -3.13 %). The routed down kernel and the W8 draft instantiation of the
same kernel are the in-trace controls and do not move.

The operator fixture overstates this change by 2.7x, and that is the honest headline ratio.
The fixture's best case is -7.40 % on trace-like; the same kernel inside the real round moves
-2.75 %. The fixture is a cold-cache, fixed-distribution, T = 4 stand; the round is not. Anyone
reading -7.40 % as what the model gets is reading it wrong, and the number that belongs in a
decision is -2.75 % on the kernel and the end-to-end figure below.

An unreported regression on the second consumer of the same helper. dot_two_rows is shared
with sparse_moe_d3_nine_warp_kernel<Q4Codec>, the T = 1 decode path, which this change therefore
also alters. In the same two traces it gets slower:

	pass 1	pass 2
median, ns	11 136 -> 11 872 (+6.61 %)	11 200 -> 11 888 (+6.14 %)
minimum, ns	10 656 -> 11 392 (+6.91 %)	10 560 -> 11 456 (+8.48 %)
total over 40 launches, ns	452 610 -> 480 832 (+6.24 %)	451 263 -> 482 238 (+6.86 %)

The minima move together with the medians in both passes, so this is real and not a tail artefact.
In the round it costs 28 222 and 30 975 ns — 0.004 % of the trace's kernel time, and 0.7 % of
what the change saves on the routed path. It does not threaten the claim. It is reported because
it is a measured regression that the body would otherwise be silent about, and because it is a
direct answer to question 3 of the Issue: the two callers of dot_two_rows do want different
things, and separating them would remove this 6 % without costing anything on the routed path.

End to end

ninfer_bench --weights qwen3_6_35b_a3b.ninfer -pg 1000,128 --prefill-chunk 8192 \
             --spec mtp --draft-tokens 3 -r 4 --warmup 1

Four passes, arm order rotated, no profiler:

pass	master, decode tok/s	change, decode tok/s
1	414.25	416.40
2	413.36	415.80
3	413.73	415.93
4	413.11	415.82
mean	413.61	415.99

+0.57 %, every change cell above every master cell. Spread inside an arm is 0.28 % on
master and 0.14 % on the change — note that master's own spread is half the size of the effect,
which is why the separation of the cells, and not the spread, is what is claimed here. The prefill
column of the same runs, which this change cannot reach because prefill dispatches different
kernels, moves -0.17 %: the stand's systematic works against this claim, not for it.
Speculation acceptance (0.3567567568) and round counts (248/0) are identical across all eight
cells, so the arms did the same work.

An end-to-end number alone would not isolate an operator change; it is reported here as
corroboration of the operator measurement above, not as the claim.

Numerics

The claim is bit-identity, not a tolerance. ninfer_sparse_moe_test with
NINFER_OP_REPORT_STATS=1 over all sixteen registered cells (q4+q5 T = 1, 2, 46, 47, 768,
4097; q4+q6 T = 1, 2, 46, 47, 768; w8+w8 T = 1, 2, 19, 20, 768) produces byte-identical
statistics on both arms; both report OK sparse_moe correctness.

That the oracle can move was shown with a sabotage build (not part of this change) that rounds
each decoded Q4 weight to BF16 before the FMA:

cell	honest rel_l2	sabotage	factor
q4+q5 T=2	0.0015916	0.0082107	5.16
q4+q5 T=46	0.0016122	0.0086159	5.34
q4+q6 T=2	0.0016461	0.0083004	5.04
q4+q6 T=46	0.0016760	0.0086335	5.15

The other twelve cells are byte-identical under the sabotage. Stated plainly: that sabotage moves
the statistic without crossing the 1.2e-2 limit, so what is demonstrated is the sensitivity of the
number, not a failing verdict. That is sufficient here only because the claim is bit-identity,
which is strictly stronger than passing the tolerance. Rebuilding the same honest source into a
different binary reproduces all sixteen cells byte for byte, so the comparison is not a false
alarm.

Where that strength control does not reach, and it matters. The four cells the sabotage
moves are T = 2 and T = 46 on the two Q4 codec pairs. Both T = 1 cells are unmoved by it.
T = 1 is exactly the sparse_moe_d3_nine_warp_kernel path — the second consumer of
dot_two_rows, the one this change also alters and the one that regresses 6 % above. So the gate
covers that path only in the weak sense that it compares its output byte for byte; it has not been
shown able to detect damage there. The honest statement is: bit-identity on that path is asserted
by the same instrument, and the instrument's sensitivity has been demonstrated on the small-T
routed path and not on the T = 1 path.

Test suite. ctest was run in full on a clean ad0f3d38 and on this change from the same
build tree: 104 of 104 pass on both, with the same six opt-in real-artifact tests skipped on
both (raw/ctest/). ninfer_sparse_moe_test is #75 in that run.

Resources

cuobjdump --dump-resource-usage over sparse_moe_decode_kernels.cu.o, translation-unit hash
normalised away and bodies matched by demangled name with template arguments: 5 of 33 bodies
change register count, and shared memory, LOCAL and STACK change on none. Independently
re-taken for this submission in raw/fmt/d3_res_compare.txt.

instantiation	master	change
sparse_moe_d3_path_tiled_kernel<Q4Codec, 3, false> - the one selected at mtp3	40	40
sparse_moe_d3_path_tiled_kernel<Q4Codec, 3, true>	39	40
sparse_moe_d3_path_tiled_kernel<Q4Codec, 1, true>	39	40
sparse_moe_d3_path_tiled_kernel<Q4Codec, 9, false>	39	40
sparse_moe_d3_path_tiled_kernel<Q4Codec, 9, true>	37	39
sparse_moe_d3_nine_warp_kernel<Q4Codec>	37	40
every W8Codec and every sparse_moe_d4_* instantiation	unchanged	unchanged
shared memory, LOCAL, STACK, operation workspace	unchanged	unchanged

At 96 threads per block (PathsPerBlock * 32 with PathsPerBlock = 3), 40 registers allow 17
resident blocks per SM against the 16 the thread limit allows, so occupancy stays limited by threads and shared memory, not by registers.

Formatting, and what was re-checked after it

clang-format on the base file reports zero violations; on the first version of this change it
reported ten, all of them on lines this change adds (195-200 and 210-213). Both
clang-format 19.1.7 and 23.1.0 give the same ten. The file has been reformatted; both versions
now report zero (raw/fmt/cf_base_19.txt, cf_arm_19.txt, cf_final_19.txt, cf_all_23.txt,
and the reformat itself in cf_reformat.diff). Both sides of every comparison are real files in
two worktrees, per the folder's rule about reading .clang-format from disk rather than a pipe.

Every timing number and the ctest run above were taken on the unformatted arm. They carry
over because the reformatting, and the two comment edits made with it, compile to the same
machine code: recompiling the translation unit before and after gives SASS and -res-usage that
are identical once the translation-unit hash is normalised away (raw/fmt/object_identity_d3.txt,
raw/fmt/d3_res_reformat.txt). That is a measurement on the object, not an assumption about what
whitespace can do.

raw/fmt/hashes.txt carries the provenance chain: base ad0f3d38 / tree af48897f; the arm that
was measured, blob 9bef59fc / tree e23a981f — the blob hash being the index 12a68e0c..9bef59fc
value printed inside the earlier patch itself, and the tree matching the commit the ctest run was
taken on; and the submitted state after formatting, blob 831e626f / tree ad958599.

Affected behaviour and ownership

No public contract changes. No ownership or lifetime boundary is touched: the change is inside one
__device__ helper, reads the same buffers in the same order and writes nothing new. The helper is
shared with sparse_moe_d3_nine_warp_kernel (the T = 1 decode path, 0.057 % of the trace's
kernel time — nsys's own Time (%) column reads 0.1), which therefore also gets the change; that path is covered by the same bit-identity gate (the
T = 1 cells), with the sensitivity caveat stated under Numerics, and it is the path that
regresses 6 % under Round.

Checks not run, and the limitation each leaves

• The strength of the gate on the T = 1 path. The sabotage moves four cells, none of them
  T = 1. Bit-identity there is asserted, not demonstrated to be detectable.
• The 6 % regression on sparse_moe_d3_nine_warp_kernel<Q4Codec> is reported, not fixed.
  Splitting the two callers would remove it; that is question 3 of the Issue.
• No accuracy evaluation was run. Acceptable only because the output is bit-identical; if the
  bit-identity claim is doubted, that claim is the thing to re-check, not the accuracy.
• One artifact, one codec pair measured. q4-q6 is covered by the gate but not by the
  operator bench; the statement about it rests on it taking the same code path.
• No profiler counters. ncu is unavailable on this machine, so every attribution above rests
  on ablations and on byte budgets computed from geometry, not on measured DRAM traffic. The
  "% of the DRAM ceiling" figures are therefore demand, computed from mandatory weight bytes
  and time, not a counter reading.
• One instrument for the operator claim. Two passes of nsys plus an end-to-end run without a
  profiler; not two independent operator instruments.
• The witness window opens before the lock is taken (see Method). Rejection is conservative,
  the accepted cells stand, but a re-measurement should use the stamp-inside-the-lock form.
• The gate was not re-run on the reformatted source. It did not need to be: the object is
  byte-identical, which is shown rather than assumed.

DISPROVED

Kept verbatim, as the folder's rules require, with the correct statement beside each.

Was: "At the round's own operating point the kernel runs at 40.5 % of the measured
1689.4 GB/s DRAM read ceiling."

Correct: 40.5 % appears in no raw file of this package. It comes from the predecessor
campaign, which divides a geometry estimate of 16.72 MB (at an assumed E ~ 13) by the round
median 24 416 ns. The figure this package can support, and the one the patch's own comment
quotes, is the fixture's trace-like point: 26.74 MB / 28 752 ns = 930 GB/s = 55.1 %. The
comment in the submitted source and the body now name the same point.

Was: the Issue's table rows "trace-like … 28 800 ns … 928 … 55.0 %" and
"independent … 34 672 ns … 1060 … 62.8 %".

Correct: the two-pass means of the raw medians are 28 752 (28 768 / 28 736) and 34 608
(34 624 / 34 592), giving 930 GB/s / 55.1 % and 1062 GB/s / 62.9 %. The same row, 21 424 ns, was
already right. The conclusions (3.30x bytes against 1.62x time, 16.8 us extrapolated to zero
experts) are unchanged by the correction.

Was: "sparse_moe_d3_path_tiled_kernel<Q4Codec, 3, false> … 18.63 % of kernel time".

Correct: in the two attached round traces the kernel is 130 512 610 ns of 798 162 921 ns of
kernel time, 16.35 % — which is what nsys's own Time (%) column says (16.4). Excluding the
prefill kernels present in the same trace gives 17.71 %. 18.63 % is reproducible from neither,
and the denominator behind it was never stated.

Was: "any window containing a foreign process was discarded and re-taken (13 cells were)".

Correct: 13 attempts on 3 cells, twelve of them in the research series that is not
submitted. In the submitted series nothing was discarded at all.

Was: "0.16 % spread inside each arm".

Correct: 0.28 % on master and 0.14 % on the change. Master's spread is half the size of the
+0.57 % effect, which is why cell separation and not spread is what the claim rests on.

Was: "A static_assert pins the divisibility the pairing needs; both current callers
(kGroups = 32) satisfy it."

Correct: the assertion pins the divisibility of the full K, not of the span the loop walks.
They coincide only because every call site passes the full range. See "About the static_assert"
above; the patch now carries the same statement in a comment.

Was: the operator table and the round table quoted side by side with no relation between
them.

Correct: the ratio is 2.69 and it belongs in the body. The fixture overstates this change
by that factor.

Was: silence about sparse_moe_d3_nine_warp_kernel<Q4Codec>.

Correct: it regresses +6.1…+6.6 % on the median and +6.2…+6.9 % on the total, in both passes,
with the minima moving the same way. It costs 0.004 % of the trace's kernel time, and it is
reported above.

Was: "the bench's wall clock lands on a 2.048 us lattice, 3.4 % of the operation".

Correct: 3.7 %. 2.048 / 55.296 us at trace-like; 4.3 % at same and 3.2 % at independent.
The wall medians are in raw/journal.txt under the q1/q2 arms. The conclusion — that wall
clock cannot resolve this effect and the verdict must come from nsys kernel medians — is
unchanged.

Was: "the T = 1 decode path, 0.25 % of round kernel time".

Correct: 0.057 % — 452 610 ns and 451 263 ns of 798 162 921 ns and 798 871 366 ns of kernel
time in the two attached traces; nsys's own Time (%) column reads 0.1. Excluding the prefill
kernels of the same trace it is 0.061 %. The point the figure was making — that this path is
small — survives; the number did not.

Was: "Diff: 46 insertions, 23 deletions."

Correct: 50 insertions, 23 deletions, one file — 13 of the added lines are comments: a
nine-line block explaining the pairing and a four-line one above the static_assert, leaving 37
lines of code. The compiled object is unchanged (raw/fmt/object_identity_d3.txt).

Corrected. This line read "four added lines are the comment above the static_assert" and
was read by review as "4 of the 50 insertions are comment". Four is that one comment; the diff
carries 13 comment lines in total. Command:
grep "^+" patch.diff | grep -v "^+++" | grep -c "^+\s*//".

Speculation counters: how much of this number is round count, and how much is round speed

Added 2026-09-06 after a dedicated audit of every decode headline I have prepared for this
project. That audit was run with my own harness outside this repository and is not a file
of this package; it is named only to say where the section came from, and nothing below
leans on it — its own logs can be supplied on request. The section answers, from this
package's own raw, the one question a throughput figure in tokens per second cannot answer
by itself.

ninfer_bench reports decode as generated tokens over elapsed time. A change that alters the
numbers the model produces can alter draft acceptance, and with it the number of speculative
rounds spent on the same fixed output length. Such a change raises tokens per second without any
kernel running faster. The two effects separate exactly, because the bench prints the round count
itself:

decode_time(base) / decode_time(arm)  =  [ rounds(base) / rounds(arm) ]  *  [ t(base) / t(arm) ]

where t = decode_seconds_mean / spec_rounds is the time of one round. spec_rounds is an exact
integer printed by the engine and decode_seconds_mean is the measured time, so the split is an
identity rather than a model.

This package: the round count is identical across both arms, so the whole figure is round speed.
Raw: raw/q1_base.txt, q1_prod.txt, q2_base.txt, q2_prod.txt, q3_base.txt, q3_prod.txt,
q4_base.txt, q4_prod.txt - all eight at mtp_k=3 proposal_head=full, and all eight printing
the same spec acc and spec round/fb: 0.3567567568 and 248/0.

The end-to-end figure is taken under speculation, which is where a round-count substitution
could hide. It does not hide there: with the same round count in both arms the whole of the
end-to-end delta is the time of one round. That is what a bit-identical output requires, and it is
stated here because it was read out of the raw rather than assumed from bit-identity.

Which proposal head this was taken at. Every cell cited above ran at the product default
ProposalHead::Full (include/ninfer/types.h:82). The audit put the same question to
--lm-head-draft, the configuration docs/performance.md publishes in, on Qwen3.6-35B-A3B at
-pg 2048,384 --prefill-chunk 8192 --spec mtp --draft-tokens 3: a bit-exact arm and the base agree there
as well - 690 rounds and acceptance 0.7819767442 in both - so the split is the same under
either head. A non-bit-exact arm measured beside them in the same ladder does move the count
(714 rounds, acceptance 0.7422969188), which is what shows the instrument would have caught a
substitution had there been one.

🤖 Generated with Claude Code

[chatgpt-codex-connector]

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-09-07T08:22:35.497526Z	e52e1be	PR opened

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.