perf(ops): warm L2 for the next projection from the MoE down tail, issuing the hint before the block barrier - #206
Merged
Neroued merged 1 commit intoSep 8, 2026
Conversation
…suing the hint before the block barrier
In the T=1 decode path sparse_moe_d4_nine_warp_kernel is the last producer before the next
layer's attention or GDN projection reads its weights from HBM. Those weights are a fixed,
known address range while D4 is still running, and L2 holds nothing of them: no earlier
kernel in the round has touched them. The kernel now receives that range and issues
prefetch.global.L2 over it, one 128-byte line per thread, clamped to 8 MiB against a 96 MiB
L2. The hint is issued before the block barrier rather than after the FP32 rank-order
epilogue, so it hides behind the slowest warp instead of extending the grid tail.
Ownership. The range is passed explicitly and nothing is stored. A new public struct
ops::SparseMoeHints names the span; a second overload of ops::sparse_moe() takes it, with
the existing six-argument signature kept and forwarding a default-constructed
SparseMoeHints{} to it - the idiom already at include/ninfer/ops/linear.h:107 and :121.
From there the value rides in the plan, as prefill and small-T already do: two fields on
SparseMoeDecodePlan, a third defaulted parameter on resolve_sparse_moe_decode_plan() so the
capacity call site is untouched, and sparse_moe_decode_launch() taking a
const SparseMoeDecodePlan&. The 8 MiB clamp became a pure function inside
resolve_sparse_moe_decode_plan(). There is no thread-local channel, no setter and no
module-global device state; the Op keeps no state between calls and is re-entrant, and CUDA
Graph capture is correct by construction because the pointer is an ordinary kernel argument
computed on the same path as the launch. No test and no bench file is touched.
This is a cache hint: it reads no value, writes nothing, and cannot change an addition
order. cuobjdump --dump-resource-usage over sparse_moe_decode_kernels.cu.o, matched by
demangled name with template arguments, shows the three Rows=1 bodies gaining two
parameters with registers 40 -> 40, shared 36 -> 36, stack 0 and spills 0, and every other
body unchanged; issuing the hint before rather than after the barrier changes 0 of 33
bodies. That census was taken on e3aeaf8 and has not been repeated on the current base;
the D4 kernel body it describes is byte-identical to the one here.
Claim level: schedule, cross-kernel producer-consumer.
Measured on one RTX 5090 (sm_120a, CUDA 13.1.115, driver 580.95.05, Release,
-DCMAKE_CUDA_ARCHITECTURES=120a), artifact Qwen3.6-35B-A3B, batch 1, every measurement
under a GPU-exclusive lock with the card-witness window opened inside the lock. All arms
built from one tree and run in one pass over one set of points; seven passes with pass 0
discarded and the median over the remaining six; arm order reversed on alternate passes;
one process per cell; two zero controls, a bit-identical copy of the master binary and a
rebuild of master.
ninfer_bench -pg 128,128;1000,128;4095,128;12000,128 -r 10 --warmup 2 --max-ctx 16384,
decode output t/s against master at mtp0, i.e. with neither speculative flag:
+1.944 / +1.873 / +1.808 / +1.720 %, pass spread +1.183 to +2.694 %, zero controls
-0.010 to +0.101 %. The claim is the interval, +1.7 to +1.9 %; the medians are printed
because that is what the raw says, but the zero-control band reaches +/-0.57 % at the
extremes on the short context. An earlier edition of this change, on ad0f3d3 and split
across two commits, read +1.74 / +1.70 / +1.53 / +1.54 %; both sets are left visible. The
effect is flat in context length because what is warmed is a weight block of fixed size,
not request data.
At --spec mtp --draft-tokens 3 the change measures -0.049 to +0.034 % against a zero
control of -0.139 to +0.039 %, and prefill medians are zero to within 0.06 % (a single run reaches +0.663 %). That zero is
structural, not weak: the prefetch parameters exist only on Rows=1 instantiations, T >= 2
goes to sparse_moe_small_t_launch which has no prefetch at all, and the MTP post-mixer
passes SparseMoeHints{} explicitly. Carried over from ad0f3d3 and not re-run: draft 1 and
draft 5 read +0.01 and -0.09 to -0.10 %, the CUDA Graph decode step at draft 3 read -0.03
to +0.04 %, and ninfer_sparse_moe_bench --sweep 2:12 read +0.00 % in all 22 cells - the
last a control rather than evidence, since the sweep starts at T=2 and never reaches the
T=1 kernel. There is no production-shape operator benchmark for this mechanism.
Output is bit-identical to master by MD5 and length over greedy generations with neither
speculative flag and at --spec mtp --draft-tokens 3, three prompts, two repeats, arms
alternated: 12 of 12, with a master-against-master control also at 12 of 12 and no reply
below the 40-byte floor. The gate is weaker than that score suggests and the weakness is
stated rather than left to be found: the repeats are the same deterministic run, and two of
the three prompts return identical output under both modes, so exactly one of the twelve
cells discriminates. The gate was also repaired mid-work - its first run returned 36
one-byte replies because a thinking model sent every permitted token to reasoning and
reasoning goes to stderr; the 40-byte floor caught it and --no-thinking fixed it.
ctest on the base this branch sits on, both arms from the same build tree: 107 of 114
registered targets executed, 106 passed, ninfer_attn_input_proj_test failed, 7 skipped as
artifact-gated - identical on the arm and on the bare base, so the single failure is an
upstream defect (issue Neroued#196) and not something this branch introduces. No round of
"114 of 114" is claimed. Real weights were not attached, so the five tests that exercise
the changed decode route did not run. clang-format 23.1.0 delta to the base is 0.
Base and provenance. This change targets upstream a16b644; it sits on a16b644's parent
487f897, which is the commit every measurement above was taken on, and the single commit
between the two touches documentation only and none of the eleven files changed here.
Earlier heads are named at the point of use for the figures taken on them. The bench flag
--mtp-draft-tokens no longer exists: the spelling is --spec mtp --draft-tokens N, and
--draft-tokens 0 is rejected for MTP (src/product/speculative_options.h:41-43), so the
zero-draft arm is spelled by passing neither flag - the identical engine state, since
SpeculativeOptions defaults to backend None with draft_tokens 0 and the validator accepts
None only when draft_tokens is 0.
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
This was referenced Sep 9, 2026
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
3 participants
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.
Ownership: the objection from #96, answered before anything else
You will ask this first, so it goes first. The form you rejected when you closed #96 is gone from
this branch. The earlier edition of this change carried the weight range to the op through a
thread_local WeightPrefetchSpaninsrc/ops/wrapper/sparse_moe.cppplus a publicset_next_weight_prefetch()setter called fromTextContext. That is, word for word, the patternyou named as unacceptable - "Hidden mutable channels such as thread-local setters or module-global
device state are also not acceptable ownership." That edition was never opened as a PR.
What is submitted here passes the parameter explicitly, down the call chain, and stores nothing:
struct ops::SparseMoeHints { const void* next_weight_prefetch; std::size_t next_weight_prefetch_bytes; }in the public op header;ops::sparse_moe()takingconst SparseMoeHints&, with the existingsix-argument signature kept and forwarding a default-constructed
SparseMoeHints{}to it. This isthe idiom already in the tree at
include/ninfer/ops/linear.h:107/:121- a policy-bearingoverload beside a plain one - not a new convention;
SparseMoeDecodePlangains two fields,resolve_sparse_moe_decode_plan()gains a third parameterdefaulted to
{}(so the capacity-computation call site is untouched), andsparse_moe_decode_launch()takesconst SparseMoeDecodePlan&instead of two loose arguments;resolve_sparse_moe_decode_plan(), where itis a pure function of its arguments.
Consequences worth stating rather than leaving to be inferred:
thread_localis gone, and the grep is in the package.git grep -n thread_localoversrc includeat the head returns four hits, all pre-existing random-number generators insrc/serve/*.cpp, none of them in a file this branch touches.WeightPrefetchSpan,set_next_weight_prefetch,set_next_projection_prefetchandprojection_prefetch_spanreturnzero hits anywhere, including
tests/andbench/. There are no mutable__device__globals.
computed on the same path as the launch. Under the old form the value was read from thread-local
storage, which made behaviour depend on which OS thread performed the capture and left the setter
dead on replay. That was a live defect, not only a style objection.
tests/andbench/are not touched by a single character:git diff --name-only 487f8977 HEAD | grep -E "^(tests|bench)/"is empty.ops::SparseMoeHints{}explicitly, so "zero under speculation" is aproperty of the structure rather than a coincidence; and the 27B target returns
{}from bothoverloads of
Variant::projection_prefetch_hintsand never learns about prefetch at all.The D4 kernel body, its arguments and its values are unchanged relative to the branch that was
measured. Only the host-side plumbing was rewritten. The measurements in this report were taken on
the submitted head, in this form.
Claim level
Schedule (cross-kernel producer-consumer). No arithmetic added, no numerics changed. The two
added kernel parameters are consumed only by
asm volatile("prefetch.global.L2 [%0];"), an instruction with no destination register and nomemory side effect, so bit-identity here is a property of the instruction rather than a result the
gate had to establish.
Linked issue and scope
Issue: none. Optimisation changes are opened directly in this repository, as #191, #193, #194 and
#195 were.
One commit.
48978ea0, parent487f8977, treeadbf1c77,git patch-id --stable00a1bbb5, 11 files changed, +140 -31, 30 hunks:Of those, the device-side change is one block of ten lines in
sparse_moe_d4_nine_warp_kernel; everything else is signature and plumbing.An earlier edition of this branch carried two commits - the warm-up, then a commit moving it
ahead of the barrier - and its evidence was a two-step ladder. That branch is not what is submitted.
The head here is a single commit that contains both, and the ladder is reported below as history,
explicitly, rather than as evidence for this head.
Environment
RTX 5090, sm_120a, driver 580.95.05, CUDA 13.1.115. Release,
-DCMAKE_CUDA_ARCHITECTURES=120a -DNINFER_BUILD_BENCHMARKS=ON -DBUILD_TESTING=ON.Artifact Qwen3.6-35B-A3B (21.2 GiB, revision
3d960f7b). Every measurement under a GPU-exclusivelock that refuses a card shared with a foreign process, with the card-witness window opened inside
the lock. Batch 1.
Protocol for the re-measured rows. The mtp0 session: five arms built from one tree, all arms in one
pass over one set of points, seven passes with pass 0 discarded and the median taken over the
remaining six. The mtp3 session below is smaller and is stated separately rather than folded in:
three arms, six passes, pass 0 discarded. Arm
order reversed on alternate passes, one process per cell. Two zero controls, both printed below:
M2is a bit-identical copy of the master binary (measures run-to-run noise) and
MRis masterrebuilt from the same tree (measures rebuild noise).
Observation
In the T=1 decode path
sparse_moe_d4_nine_warp_kernelis the last producer before the next layer'sattention or GDN projection reads its weights from HBM. That address range is known while D4 is still
running, and L2 holds nothing of it: no earlier kernel in the round has touched it. The consumer
therefore starts every layer cold for a range that was predictable one kernel earlier. The warmed
block is capped at 8 MiB against a 96 MiB L2, so it fits with room to spare and cannot evict what it
is meant to help.
Change
ops::SparseMoeHintsnames the next consumer's weight span; it reaches the decode path throughthe second
sparse_moe()overload and two fields ofSparseMoeDecodePlan(see Ownership).sparse_moe_d4_nine_warp_kernelissuesprefetch.global.L2over that span, one 128-byte line perthread, clamped to 8 MiB in
resolve_sparse_moe_decode_plan().__syncthreads()and the FP32 rank-order epilogue, so it hidesbehind the slowest warp instead of extending the grid tail.
Verification commands and results
End to end, decode output t/s against master, re-measured on
487f8977on the submitted headM2(bit-identical master)MR(master rebuilt)--spec mtp --draft-tokens 3M2at mtp3Prefill was zero in both sessions:
|delta| <= 0.06%on every median, against a zero control ofthe same order. Per single run the spread is wider - up to +0.663% in one cell - and since this
report defines a cell as one process run, the 0.06% figure is a statement about medians and is
labelled as such rather than left to be read as a per-run bound.
The three-digit medians are printed because they are what the raw says, but the claim this report
makes is the interval: +1.7 to +1.9%. The zero-control band on the short context reaches +/-0.57%
at the extremes, which is wider than the last digit of the median; the headline is stated as a range
for that reason, and the title says the range and not the medians.
The previous edition, on
ad0f3d38, read lower, and both sets are left visible:487f8977ad0f3d38(two commits)The mechanism therefore reads 0.18 to 0.28 percentage points higher than previously claimed, on
a newer base and after the ownership rewrite. Nothing was tuned to obtain that; the D4 kernel is
byte-identical to the one that produced the lower figures.
The effect is flat in context length because a weight block is of fixed size; request data scales
with context, weights do not.
Why the two-step ladder is not shown as evidence. The previous edition split its evidence into
"warm only" (+1.00 / +0.98 / +0.93 / +0.91%) and "warm + reorder" (+1.74 / +1.70 / +1.53 / +1.54%),
with the reorder worth +0.60 to +0.74% of its own. Those rows cannot be reproduced from what is
submitted here: the intermediate commit does not exist on this branch, so there is no arm to build
and no command that yields the split. They were measured on
e3aeaf8cwith-r 3 --warmup 1andfour passes, on a two-commit branch, and they are printed here as history only. If the split matters
for review, say so and I will re-cut the branch into two commits and re-measure it on
a16b6442rather than carry the old figures forward.
Controls carried over from
ad0f3d38, not re-run here: the CUDA Graph decode step at draft 3read -0.03 to +0.04% over contexts 128/4095/12000; the operator sweep
ninfer_sparse_moe_bench --sweep 2:12read +0.00% in all 22 cells of both codecs; end to end at--spec mtp --draft-tokens 1and5read +0.01% and -0.09 to -0.10%. Both zeros are structuraland the structure is unchanged on this head: the prefetch parameters exist only on
Rows=1instantiations,
T >= 2goes tosparse_moe_small_t_launch, which has no prefetch(
git grep prefetch -- src/ops/sparse_moe/small_t/is empty), and the MTP post-mixer passesSparseMoeHints{}explicitly. The re-measured mtp3 row above confirms the same zero on the new base.The edited route really executes (
nsys profile -t cuda --cuda-graph-trace=node,-pg "128,16" -r 1 --warmup 0 --max-ctx 2048, on487f8977):sparse_moe_d4_nine_warp_kernel<Q5Codec,1>(..., __nv_bfloat16*)...(..., __nv_bfloat16*, const char*, unsigned long long)d1Launch counts are equal, so the route is the same one; the signature carries the two new parameters,
so it is the edited kernel that runs; and D4 itself got 6.4% more expensive, which is how one
knows the
if (prefetch_data != nullptr)branch is actually taken - with a null pointer it would beskipped and the duration would match the base. At a standard deviation of 1688 ns over 777 launches,
a shift of 1181 ns is about twenty standard errors. The round as a whole got 1.8% cheaper while its
last kernel got 6.4% dearer; that is the trade the change makes, stated plainly.
Correctness: the bitwise gate, and why it is weaker than it looks
apps/ninfer,--greedy --seed 1 --max-new 160 --max-context 16384 --no-thinking, three prompts(instruction, code, arithmetic chain), two repeats, two modes (mtp0 and
--spec mtp --draft-tokens 3), arms alternated, output redirection performed inside theexclusive-lock wrapper, compared by MD5 and byte length:
Now the part that matters more than the score. Of those twelve cells, exactly one discriminates.
The two repeats are the same deterministic run and are always byte-for-byte equal, so they carry no
information. Prompts p1 and p3 produce the same output at mtp0 and at mtp3 - they do not
distinguish even two different execution routes. Only prompt p2 responds: it differs between mtp0
and mtp3, and it is also the only prompt that differs in a deliberately-broken control
(
--kv-dtype int8againstbf16on the bare base: 617 bytes32612429against 545 bytes678138cd, while p1 and p3 match). The instrument can tell outputs apart, but a battery ofthree short prompts leaves it with a single sensitive cell. The previous edition of this package ran
four prompts and 64 comparisons; this one runs twelve. That is a regression in the gate, and it is
named here rather than left for review to find.
And the gate nearly lied, once. Its first run returned 36 replies of one byte each. All MD5s
matched, and "12 of 12 IDENTICAL" would have printed as a success. The cause: the model is a
thinking model, all 160 permitted tokens went into reasoning, reasoning goes to stderr, and stdout
held a newline. What caught it was the 40-byte floor on reply size, plus the strength control
(mtp0 and mtp3 returning identical empty replies is not a plausible pass). The fix was
--no-thinking, and the floor and the strength control are now standing parts of the procedurerather than a one-off repair. This is the same class of mistake this package already documents -
measuring the wrapper's banner instead of the model's output - in a new disguise.
What the gate is and is not evidence for. The risk it guards is low by construction, not by
score:
prefetch.global.L2has no destination register, so the added values cannot enter arithmeticat all. The gate here checks wiring, not numerics. It should not be read as the numerical validation
a change with arithmetic in it would owe you.
ctest, re-run on487f8977on both arms from the same build tree:487f8977ninfer_attn_input_proj_testninfer_attn_input_proj_testThe formulation this package stands behind is the exact one: the arm's result matched the base,
and the single failure reproduces on bare
487f8977built from the same tree (upstream defect,issue #196). Two caveats belong with that count.
NINFER_QWEN3_6_35B_A3B_WEIGHTSwas not set, sothe seven artifact-gated tests were skipped - including the five that would exercise the changed
decode route, which means today's green
ctestverifies nothing on the edited path. And with realweights attached, this project has separately observed a second red test on a bare upstream base;
that observation is not from this package's raw and is mentioned only so the count above is not read
as more than it is.
clang-format(version 23.1.0, run in the repository tree with the repository's.clang-format): delta to the base is 0. Thirteen warnings exist insrc/targets/qwen3_6/impl/runtime/text_context_impl.hon the bare base and the same thirteen existon the head; they are master's own, at line numbers shifted by exactly +1 by this branch's single
added
#include. An earlier state of this rewrite carried +18 of its own; those are fixed.Resources, workspace, roofline
Workspace capacity is unchanged:
workspace_capacity_bytesreads 109453312 on every arm,re-measured on
487f8977. No CLI, graph-profile or artifact change. The span is a pointer and alength.
Register and shared-memory census - measured on
e3aeaf8c, not re-run on the current base;carried over because the D4 kernel body is byte-identical to the one censused.
cuobjdump --dump-resource-usageoversparse_moe_decode_kernels.cu.o, matched by demangled namewith template arguments: three
Rows=1bodies gain two parameters at registers 40 -> 40, shared36 -> 36, stack 0, spills 0; every other body unchanged. Placing the hint before rather than after
the barrier changed 0 of 33 bodies.
Roofline - measured on
ad0f3d38, not re-run. The operator table puts baseline MoE decode at715 to 867 GB/s of unique weight, i.e. 42.3 to 51.3% of the 1690.7 GB/s read ceiling measured on
this card. The 1792.1 GB/s constant printed by the benchmark header was not reproduced and is not
used. I did not decompose the T=1 step into per-kernel byte counts and therefore quote no
roofline figure for the step this change acts on - a check not run, not a number withheld.
Provenance of the ceilings: measured on this card with the in-tree probe
tools/hbm_bandwidth_probe.cu, two independent runs under the same exclusive lock. Sustained read1690.7 and 1690.6 GB/s (medians 1690.4 / 1690.3), write 1667.8, kernel copy 1496.8 GB/s of bus
traffic,
cudaMemcpyAsyncD2D 1530.8; the two runs agree to 0.01%. The probe also confirmsL2 = 96.0 MiB, 170 SMs, a 512-bit bus and an advertised peak of 1792.0 GB/s, so the sustained read
reaches 94.3% of the advertised figure. Percentages here are taken against the measured read rate,
not the advertised one.
Limitations and checks not run
zero. Read against the published 35B report, that means it acts on one of the five runs the report tabulates (
P0,M3,DS,DG,S3):P0("Long-context phases", backend
None) is non-speculative and is the arm this change moves;M3(MTP K=3),
DSandDG(DFlash K=7) are speculative and get nothing(
docs/performance/qwen3.6-35b-a3b.md, run table). That limit is in the title.this a sharper omission rather than a softer one: the public contract now does accept the span,
but
ninfer_sparse_moe_benchdoes not set it and its sweep starts atT=2, which never reachesthe T=1 kernel. So there is still no production-shape operator benchmark for this mechanism. The
evidence offered instead is the end-to-end delta, the structural controls and the
nsysrouteproof. If a benchmark is a condition of merge, say so and I will write one.
particular the
--max-contexttrap that this package itself documents - a prompt near 7500 tokenssilently producing empty output - is not exercised by this battery at all.
ctestdid not run on real weights, so the five tests covering the changed decode route didnot execute.
consistent with the timings - flat in context, D4 up 6.4%, round down 1.8% - but is not directly
demonstrated. Performance counters are closed by the driver default on the host used.
regression, which is weak evidence rather than proof.
a thread-local channel of unknown behaviour across requests - is gone; the plan-carried form is
re-entrant. What remains is simply that batched decode was not benchmarked.
projection_prefetch_hintsreturns{}.qwen3_6_35b_a3b/impl/variant.hwrites9216 * 2048and12288 * 2048. Both are correct today - the matrices areW8G32_F16S, one byte per element - butthey duplicate the weight's shape by hand, with no
static_assertand no query of theWeightobject. If the artifact ever changes shape, the literal will not complain; only the clamp keeps it
safe. I would take a suggestion on deriving the length from the weight instead.
Numerically irrelevant, but it is a loose end a reviewer will notice, and I would rather name it
than have it found.
dflash_round_benchdoes not run on current master (it bypassesEngineOptionsnormalization); root-only
context_cachefields were set locally to obtain the carried-over graphround-step control on
ad0f3d38.about an hour apart, with the same master binary in both. The noise band in the second is wider on
the short context (up to +/-0.57%) than in the first (+/-0.25%). Every number above is a paired
delta taken within a pass, so the sessions are not subtracted from one another.
Honest note on the origin of these numbers
Inside the 38-commit stack this mechanism was credited with +3.4% at mtp0, as two commits worth
+2.5% and +0.9%. Measured on its own against master it came out at +1.53 to +1.74% on
ad0f3d38, and at +1.72 to +1.94% on487f8977after the ownership rewrite. So the stackfigure was about twofold high, and it also failed to mention that the mechanism is inert under
speculation. The corrected figure is what the title claims.
Speculation counters: how much of this number is round count, and how much is round speed
ninfer_benchreports decode as generated tokens over elapsed time. A change that alters the numbersthe 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:
where
t = decode_seconds_mean / spec_rounds.spec_roundsis an exact integer printed by theengine and
decode_seconds_meanis the measured time, so the split is an identity rather than amodel.
At the operating point of the headline the question does not arise at all: at mtp0 there are no
speculative rounds, so the whole figure is round speed by construction.
For the speculative windows the check was run on
ad0f3d38and is carried over. Raw:data/all_raw.tgz, directorye2e/, one CSV per (arm, draft window, pass), carryingspec_rounds,spec_fallback_stepsandspec_acceptance_rate. Grouped by (draft window x point), 6 arms agreein all 16 groups, to the last digit:
Which proposal head this was taken at. Every cell cited above ran at the product default
ProposalHead::Full(include/ninfer/types.h:82). The same question was put separately to--lm-head-draft, the configuration the published 35B report uses(
docs/performance/qwen3.6-35b-a3b.md:29), 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 acceptance0.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